1 /******************************************************************************
2 * Copyright (c) 2011, Duane Merrill. All rights reserved.
3 * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4 *
5 * Redistribution and use in source and binary forms, with or without
6 * modification, are permitted provided that the following conditions are met:
7 * * Redistributions of source code must retain the above copyright
8 * notice, this list of conditions and the following disclaimer.
9 * * Redistributions in binary form must reproduce the above copyright
10 * notice, this list of conditions and the following disclaimer in the
11 * documentation and/or other materials provided with the distribution.
12 * * Neither the name of the NVIDIA CORPORATION nor the
13 * names of its contributors may be used to endorse or promote products
14 * derived from this software without specific prior written permission.
15 *
16 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26 *
27 ******************************************************************************/
28
29 /******************************************************************************
30 * Simple example of DeviceRadixSort::SortPairs().
31 *
32 * Sorts an array of float keys paired with a corresponding array of int values.
33 *
34 * To compile using the command line:
35 * nvcc -arch=sm_XX example_device_radix_sort.cu -I../.. -lcudart -O3
36 *
37 ******************************************************************************/
38
39 // Ensure printing of CUDA runtime errors to console
40 #define CUB_STDERR
41
42 #include <stdio.h>
43 #include <algorithm>
44
45 #include <cub/util_allocator.cuh>
46 #include <cub/device/device_radix_sort.cuh>
47
48 #include "../../test/test_util.h"
49
50 using namespace cub;
51
52
53 //---------------------------------------------------------------------
54 // Globals, constants and typedefs
55 //---------------------------------------------------------------------
56
57 bool g_verbose = false; // Whether to display input/output to console
58 CachingDeviceAllocator g_allocator(true); // Caching allocator for device memory
59
60
61 //---------------------------------------------------------------------
62 // Test generation
63 //---------------------------------------------------------------------
64
65 /**
66 * Simple key-value pairing for floating point types. Distinguishes
67 * between positive and negative zero.
68 */
69 struct Pair
70 {
71 float key;
72 int value;
73
operator <Pair74 bool operator<(const Pair &b) const
75 {
76 if (key < b.key)
77 return true;
78
79 if (key > b.key)
80 return false;
81
82 // Return true if key is negative zero and b.key is positive zero
83 unsigned int key_bits = SafeBitCast<unsigned int>(key);
84 unsigned int b_key_bits = SafeBitCast<unsigned int>(b.key);
85 unsigned int HIGH_BIT = 1u << 31;
86
87 return ((key_bits & HIGH_BIT) != 0) && ((b_key_bits & HIGH_BIT) == 0);
88 }
89 };
90
91
92 /**
93 * Initialize key-value sorting problem.
94 */
Initialize(float * h_keys,int * h_values,float * h_reference_keys,int * h_reference_values,int num_items)95 void Initialize(
96 float *h_keys,
97 int *h_values,
98 float *h_reference_keys,
99 int *h_reference_values,
100 int num_items)
101 {
102 Pair *h_pairs = new Pair[num_items];
103
104 for (int i = 0; i < num_items; ++i)
105 {
106 RandomBits(h_keys[i]);
107 RandomBits(h_values[i]);
108 h_pairs[i].key = h_keys[i];
109 h_pairs[i].value = h_values[i];
110 }
111
112 if (g_verbose)
113 {
114 printf("Input keys:\n");
115 DisplayResults(h_keys, num_items);
116 printf("\n\n");
117
118 printf("Input values:\n");
119 DisplayResults(h_values, num_items);
120 printf("\n\n");
121 }
122
123 std::stable_sort(h_pairs, h_pairs + num_items);
124
125 for (int i = 0; i < num_items; ++i)
126 {
127 h_reference_keys[i] = h_pairs[i].key;
128 h_reference_values[i] = h_pairs[i].value;
129 }
130
131 delete[] h_pairs;
132 }
133
134
135 //---------------------------------------------------------------------
136 // Main
137 //---------------------------------------------------------------------
138
139 /**
140 * Main
141 */
main(int argc,char ** argv)142 int main(int argc, char** argv)
143 {
144 int num_items = 150;
145
146 // Initialize command line
147 CommandLineArgs args(argc, argv);
148 g_verbose = args.CheckCmdLineFlag("v");
149 args.GetCmdLineArgument("n", num_items);
150
151 // Print usage
152 if (args.CheckCmdLineFlag("help"))
153 {
154 printf("%s "
155 "[--n=<input items> "
156 "[--device=<device-id>] "
157 "[--v] "
158 "\n", argv[0]);
159 exit(0);
160 }
161
162 // Initialize device
163 CubDebugExit(args.DeviceInit());
164
165 printf("cub::DeviceRadixSort::SortPairs() %d items (%d-byte keys %d-byte values)\n",
166 num_items, int(sizeof(float)), int(sizeof(int)));
167 fflush(stdout);
168
169 // Allocate host arrays
170 float *h_keys = new float[num_items];
171 float *h_reference_keys = new float[num_items];
172 int *h_values = new int[num_items];
173 int *h_reference_values = new int[num_items];
174
175 // Initialize problem and solution on host
176 Initialize(h_keys, h_values, h_reference_keys, h_reference_values, num_items);
177
178 // Allocate device arrays
179 DoubleBuffer<float> d_keys;
180 DoubleBuffer<int> d_values;
181 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_keys.d_buffers[0], sizeof(float) * num_items));
182 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_keys.d_buffers[1], sizeof(float) * num_items));
183 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_values.d_buffers[0], sizeof(int) * num_items));
184 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_values.d_buffers[1], sizeof(int) * num_items));
185
186 // Allocate temporary storage
187 size_t temp_storage_bytes = 0;
188 void *d_temp_storage = NULL;
189
190 CubDebugExit(DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items));
191 CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes));
192
193 // Initialize device arrays
194 CubDebugExit(cudaMemcpy(d_keys.d_buffers[d_keys.selector], h_keys, sizeof(float) * num_items, cudaMemcpyHostToDevice));
195 CubDebugExit(cudaMemcpy(d_values.d_buffers[d_values.selector], h_values, sizeof(int) * num_items, cudaMemcpyHostToDevice));
196
197 // Run
198 CubDebugExit(DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items));
199
200 // Check for correctness (and display results, if specified)
201 int compare = CompareDeviceResults(h_reference_keys, d_keys.Current(), num_items, true, g_verbose);
202 printf("\t Compare keys (selector %d): %s\n", d_keys.selector, compare ? "FAIL" : "PASS");
203 AssertEquals(0, compare);
204 compare = CompareDeviceResults(h_reference_values, d_values.Current(), num_items, true, g_verbose);
205 printf("\t Compare values (selector %d): %s\n", d_values.selector, compare ? "FAIL" : "PASS");
206 AssertEquals(0, compare);
207
208 // Cleanup
209 if (h_keys) delete[] h_keys;
210 if (h_reference_keys) delete[] h_reference_keys;
211 if (h_values) delete[] h_values;
212 if (h_reference_values) delete[] h_reference_values;
213
214 if (d_keys.d_buffers[0]) CubDebugExit(g_allocator.DeviceFree(d_keys.d_buffers[0]));
215 if (d_keys.d_buffers[1]) CubDebugExit(g_allocator.DeviceFree(d_keys.d_buffers[1]));
216 if (d_values.d_buffers[0]) CubDebugExit(g_allocator.DeviceFree(d_values.d_buffers[0]));
217 if (d_values.d_buffers[1]) CubDebugExit(g_allocator.DeviceFree(d_values.d_buffers[1]));
218 if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage));
219
220 printf("\n\n");
221
222 return 0;
223 }
224
225
226
227