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 DeviceSelect::Unique().
31 *
32 * Selects the first element from each run of identical values from a sequence
33 * of int keys.
34 *
35 * To compile using the command line:
36 * nvcc -arch=sm_XX example_device_select_unique.cu -I../.. -lcudart -O3
37 *
38 ******************************************************************************/
39
40 // Ensure printing of CUDA runtime errors to console
41 #define CUB_STDERR
42
43 #include <stdio.h>
44
45 #include <cub/util_allocator.cuh>
46 #include <cub/device/device_select.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 /**
67 * Initialize problem, setting runs of random length chosen from [1..max_segment]
68 */
Initialize(int * h_in,int num_items,int max_segment)69 void Initialize(
70 int *h_in,
71 int num_items,
72 int max_segment)
73 {
74 int key = 0;
75 int i = 0;
76 while (i < num_items)
77 {
78 // Randomly select number of repeating occurrences uniformly from [1..max_segment]
79 unsigned short max_short = (unsigned short) -1;
80 unsigned short repeat;
81 RandomBits(repeat);
82 repeat = (unsigned short) ((float(repeat) * (float(max_segment) / float(max_short))));
83 repeat = CUB_MAX(1, repeat);
84
85 int j = i;
86 while (j < CUB_MIN(i + repeat, num_items))
87 {
88 h_in[j] = key;
89 j++;
90 }
91
92 i = j;
93 key++;
94 }
95
96 if (g_verbose)
97 {
98 printf("Input:\n");
99 DisplayResults(h_in, num_items);
100 printf("\n\n");
101 }
102 }
103
104
105 /**
106 * Solve unique problem
107 */
Solve(int * h_in,int * h_reference,int num_items)108 int Solve(
109 int *h_in,
110 int *h_reference,
111 int num_items)
112 {
113 int num_selected = 0;
114 if (num_items > 0)
115 {
116 h_reference[num_selected] = h_in[0];
117 num_selected++;
118 }
119
120 for (int i = 1; i < num_items; ++i)
121 {
122 if (h_in[i] != h_in[i - 1])
123 {
124 h_reference[num_selected] = h_in[i];
125 num_selected++;
126 }
127 }
128
129 return num_selected;
130 }
131
132
133 //---------------------------------------------------------------------
134 // Main
135 //---------------------------------------------------------------------
136
137 /**
138 * Main
139 */
main(int argc,char ** argv)140 int main(int argc, char** argv)
141 {
142 int num_items = 150;
143 int max_segment = 40; // Maximum segment length
144
145 // Initialize command line
146 CommandLineArgs args(argc, argv);
147 g_verbose = args.CheckCmdLineFlag("v");
148 args.GetCmdLineArgument("n", num_items);
149 args.GetCmdLineArgument("maxseg", max_segment);
150
151 // Print usage
152 if (args.CheckCmdLineFlag("help"))
153 {
154 printf("%s "
155 "[--n=<input items> "
156 "[--device=<device-id>] "
157 "[--maxseg=<max segment length>]"
158 "[--v] "
159 "\n", argv[0]);
160 exit(0);
161 }
162
163 // Initialize device
164 CubDebugExit(args.DeviceInit());
165
166 // Allocate host arrays
167 int* h_in = new int[num_items];
168 int* h_reference = new int[num_items];
169
170 // Initialize problem and solution
171 Initialize(h_in, num_items, max_segment);
172 int num_selected = Solve(h_in, h_reference, num_items);
173
174 printf("cub::DeviceSelect::Unique %d items (%d-byte elements), %d selected (avg run length %d)\n",
175 num_items, (int) sizeof(int), num_selected, num_items / num_selected);
176 fflush(stdout);
177
178 // Allocate problem device arrays
179 int *d_in = NULL;
180 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(int) * num_items));
181
182 // Initialize device input
183 CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice));
184
185 // Allocate device output array and num selected
186 int *d_out = NULL;
187 int *d_num_selected_out = NULL;
188 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_out, sizeof(int) * num_items));
189 CubDebugExit(g_allocator.DeviceAllocate((void**)&d_num_selected_out, sizeof(int)));
190
191 // Allocate temporary storage
192 void *d_temp_storage = NULL;
193 size_t temp_storage_bytes = 0;
194 CubDebugExit(DeviceSelect::Unique(d_temp_storage, temp_storage_bytes, d_in, d_out, d_num_selected_out, num_items));
195 CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes));
196
197 // Run
198 CubDebugExit(DeviceSelect::Unique(d_temp_storage, temp_storage_bytes, d_in, d_out, d_num_selected_out, num_items));
199
200 // Check for correctness (and display results, if specified)
201 int compare = CompareDeviceResults(h_reference, d_out, num_selected, true, g_verbose);
202 printf("\t Data %s ", compare ? "FAIL" : "PASS");
203 compare = compare | CompareDeviceResults(&num_selected, d_num_selected_out, 1, true, g_verbose);
204 printf("\t Count %s ", compare ? "FAIL" : "PASS");
205 AssertEquals(0, compare);
206
207 // Cleanup
208 if (h_in) delete[] h_in;
209 if (h_reference) delete[] h_reference;
210 if (d_in) CubDebugExit(g_allocator.DeviceFree(d_in));
211 if (d_out) CubDebugExit(g_allocator.DeviceFree(d_out));
212 if (d_num_selected_out) CubDebugExit(g_allocator.DeviceFree(d_num_selected_out));
213 if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage));
214
215 printf("\n\n");
216
217 return 0;
218 }
219
220
221
222