MUDA
Loading...
Searching...
No Matches
device_histogram.h
1#pragma once
2#include <muda/cub/device/cub_wrapper.h>
3#include "details/cub_wrapper_macro_def.inl"
4#ifndef __INTELLISENSE__
5#include <cub/device/device_histogram.cuh>
6#endif
7
8namespace muda
9{
10//ref: https://nvlabs.github.io/cub/structcub_1_1_device_histogram.html
11
12class DeviceHistogram : public CubWrapper<DeviceHistogram>
13{
14 using Base = CubWrapper<DeviceHistogram>;
15
16 public:
17 using Base::Base;
18
19 // DeviceVector:
20
21 // HistogramEven (single channel, 1D input)
22 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
23 DeviceHistogram& HistogramEven(SampleIteratorT d_samples,
24 CounterT* d_histogram,
25 int num_levels,
26 LevelT lower_level,
27 LevelT upper_level,
28 OffsetT num_samples)
29 {
30 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::HistogramEven(d_temp_storage,
31 temp_storage_bytes,
32 d_samples,
33 d_histogram,
34 num_levels,
35 lower_level,
36 upper_level,
37 num_samples,
38 _stream,
39 false));
40 }
41
42 // HistogramEven (single channel, 2D input)
43 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
44 DeviceHistogram& HistogramEven(SampleIteratorT d_samples,
45 CounterT* d_histogram,
46 int num_levels,
47 LevelT lower_level,
48 LevelT upper_level,
49 OffsetT num_row_samples,
50 OffsetT num_rows,
51 size_t row_stride_bytes)
52 {
53 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::HistogramEven(d_temp_storage,
54 temp_storage_bytes,
55 d_samples,
56 d_histogram,
57 num_levels,
58 lower_level,
59 upper_level,
60 num_row_samples,
61 num_rows,
62 row_stride_bytes,
63 _stream,
64 false));
65 }
66
67 // MultiHistogramEven (multiple channels, 1D input)
68 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
69 DeviceHistogram& MultiHistogramEven(SampleIteratorT d_samples,
70 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
71 int num_levels[NUM_ACTIVE_CHANNELS],
72 LevelT lower_level[NUM_ACTIVE_CHANNELS],
73 LevelT upper_level[NUM_ACTIVE_CHANNELS],
74 OffsetT num_pixels)
75 {
76 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::MultiHistogramEven(d_temp_storage,
77 temp_storage_bytes,
78 d_samples,
79 d_histogram,
80 num_levels,
81 lower_level,
82 upper_level,
83 num_pixels,
84 _stream,
85 false));
86 }
87
88 // MultiHistogramEven (multiple channels, 2D input)
89 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
90 DeviceHistogram& MultiHistogramEven(SampleIteratorT d_samples,
91 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
92 int num_levels[NUM_ACTIVE_CHANNELS],
93 LevelT lower_level[NUM_ACTIVE_CHANNELS],
94 LevelT upper_level[NUM_ACTIVE_CHANNELS],
95 OffsetT num_row_pixels,
96 OffsetT num_rows,
97 size_t row_stride_bytes)
98 {
99 MUDA_CUB_WRAPPER_IMPL(
100 cub::DeviceHistogram::MultiHistogramEven(d_temp_storage,
101 temp_storage_bytes,
102 d_samples,
103 d_histogram,
104 num_levels,
105 lower_level,
106 upper_level,
107 num_row_pixels,
108 num_rows,
109 row_stride_bytes,
110 _stream,
111 false));
112 }
113
114 // HistogramRange (single channel, 1D input)
115 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
116 DeviceHistogram& HistogramRange(SampleIteratorT d_samples,
117 CounterT* d_histogram,
118 int num_levels,
119 LevelT* d_levels,
120 OffsetT num_samples)
121 {
122 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::HistogramRange(
123 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels, num_samples, _stream, false));
124 }
125
126 // HistogramRange (single channel, 2D input)
127 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
128 DeviceHistogram& HistogramRange(SampleIteratorT d_samples,
129 CounterT* d_histogram,
130 int num_levels,
131 LevelT* d_levels,
132 OffsetT num_row_samples,
133 OffsetT num_rows,
134 size_t row_stride_bytes)
135 {
136 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::HistogramRange(d_temp_storage,
137 temp_storage_bytes,
138 d_samples,
139 d_histogram,
140 num_levels,
141 d_levels,
142 num_row_samples,
143 num_rows,
144 row_stride_bytes,
145 _stream,
146 false));
147 }
148
149 // MultiHistogramRange (multiple channels, 1D input)
150 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
151 DeviceHistogram& MultiHistogramRange(SampleIteratorT d_samples,
152 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
153 int num_levels[NUM_ACTIVE_CHANNELS],
154 LevelT* d_levels[NUM_ACTIVE_CHANNELS],
155 OffsetT num_pixels)
156 {
157 MUDA_CUB_WRAPPER_IMPL(cub::DeviceHistogram::MultiHistogramRange(
158 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels, num_pixels, _stream, false));
159 }
160
161 // MultiHistogramRange (multiple channels, 2D input)
162 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
163 DeviceHistogram& MultiHistogramRange(SampleIteratorT d_samples,
164 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
165 int num_levels[NUM_ACTIVE_CHANNELS],
166 LevelT* d_levels[NUM_ACTIVE_CHANNELS],
167 OffsetT num_row_pixels,
168 OffsetT num_rows,
169 size_t row_stride_bytes)
170 {
171 MUDA_CUB_WRAPPER_IMPL(
172 cub::DeviceHistogram::MultiHistogramRange(d_temp_storage,
173 temp_storage_bytes,
174 d_samples,
175 d_histogram,
176 num_levels,
177 d_levels,
178 num_row_pixels,
179 num_rows,
180 row_stride_bytes,
181 _stream,
182 false));
183 }
184
185 // Origin:
186
187 // HistogramEven (single channel, 1D input)
188 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
189 DeviceHistogram& HistogramEven(void* d_temp_storage,
190 size_t& temp_storage_bytes,
191 SampleIteratorT d_samples,
192 CounterT* d_histogram,
193 int num_levels,
194 LevelT lower_level,
195 LevelT upper_level,
196 OffsetT num_samples)
197 {
198 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(cub::DeviceHistogram::HistogramEven(
199 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, lower_level, upper_level, num_samples, _stream, false));
200 }
201
202 // HistogramEven (single channel, 2D input)
203 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
204 DeviceHistogram& HistogramEven(void* d_temp_storage,
205 size_t& temp_storage_bytes,
206 SampleIteratorT d_samples,
207 CounterT* d_histogram,
208 int num_levels,
209 LevelT lower_level,
210 LevelT upper_level,
211 OffsetT num_row_samples,
212 OffsetT num_rows,
213 size_t row_stride_bytes)
214 {
215 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(
216 cub::DeviceHistogram::HistogramEven(d_temp_storage,
217 temp_storage_bytes,
218 d_samples,
219 d_histogram,
220 num_levels,
221 lower_level,
222 upper_level,
223 num_row_samples,
224 num_rows,
225 row_stride_bytes,
226 _stream,
227 false));
228 }
229
230 // MultiHistogramEven (multiple channels, 1D input)
231 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
232 DeviceHistogram& MultiHistogramEven(void* d_temp_storage,
233 size_t& temp_storage_bytes,
234 SampleIteratorT d_samples,
235 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
236 int num_levels[NUM_ACTIVE_CHANNELS],
237 LevelT lower_level[NUM_ACTIVE_CHANNELS],
238 LevelT upper_level[NUM_ACTIVE_CHANNELS],
239 OffsetT num_pixels)
240 {
241 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(cub::DeviceHistogram::MultiHistogramEven(
242 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, lower_level, upper_level, num_pixels, _stream, false));
243 }
244
245 // MultiHistogramEven (multiple channels, 2D input)
246 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
247 DeviceHistogram& MultiHistogramEven(void* d_temp_storage,
248 size_t& temp_storage_bytes,
249 SampleIteratorT d_samples,
250 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
251 int num_levels[NUM_ACTIVE_CHANNELS],
252 LevelT lower_level[NUM_ACTIVE_CHANNELS],
253 LevelT upper_level[NUM_ACTIVE_CHANNELS],
254 OffsetT num_row_pixels,
255 OffsetT num_rows,
256 size_t row_stride_bytes)
257 {
258 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(
259 cub::DeviceHistogram::MultiHistogramEven(d_temp_storage,
260 temp_storage_bytes,
261 d_samples,
262 d_histogram,
263 num_levels,
264 lower_level,
265 upper_level,
266 num_row_pixels,
267 num_rows,
268 row_stride_bytes,
269 _stream,
270 false));
271 }
272
273 // HistogramRange (single channel, 1D input)
274 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
275 DeviceHistogram& HistogramRange(void* d_temp_storage,
276 size_t& temp_storage_bytes,
277 SampleIteratorT d_samples,
278 CounterT* d_histogram,
279 int num_levels,
280 LevelT* d_levels,
281 OffsetT num_samples)
282 {
283 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(cub::DeviceHistogram::HistogramRange(
284 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels, num_samples, _stream, false));
285 }
286
287 // HistogramRange (single channel, 2D input)
288 template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
289 DeviceHistogram& HistogramRange(void* d_temp_storage,
290 size_t& temp_storage_bytes,
291 SampleIteratorT d_samples,
292 CounterT* d_histogram,
293 int num_levels,
294 LevelT* d_levels,
295 OffsetT num_row_samples,
296 OffsetT num_rows,
297 size_t row_stride_bytes)
298 {
299 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(
300 cub::DeviceHistogram::HistogramRange(d_temp_storage,
301 temp_storage_bytes,
302 d_samples,
303 d_histogram,
304 num_levels,
305 d_levels,
306 num_row_samples,
307 num_rows,
308 row_stride_bytes,
309 _stream,
310 false));
311 }
312
313 // MultiHistogramRange (multiple channels, 1D input)
314 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
315 DeviceHistogram& MultiHistogramRange(void* d_temp_storage,
316 size_t& temp_storage_bytes,
317 SampleIteratorT d_samples,
318 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
319 int num_levels[NUM_ACTIVE_CHANNELS],
320 LevelT* d_levels[NUM_ACTIVE_CHANNELS],
321 OffsetT num_pixels)
322 {
323 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(cub::DeviceHistogram::MultiHistogramRange(
324 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels, num_pixels, _stream, false));
325 }
326
327 // MultiHistogramRange (multiple channels, 2D input)
328 template <int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
329 DeviceHistogram& MultiHistogramRange(void* d_temp_storage,
330 size_t& temp_storage_bytes,
331 SampleIteratorT d_samples,
332 CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
333 int num_levels[NUM_ACTIVE_CHANNELS],
334 LevelT* d_levels[NUM_ACTIVE_CHANNELS],
335 OffsetT num_row_pixels,
336 OffsetT num_rows,
337 size_t row_stride_bytes)
338 {
339 MUDA_CUB_WRAPPER_FOR_COMPUTE_GRAPH_IMPL(
340 cub::DeviceHistogram::MultiHistogramRange(d_temp_storage,
341 temp_storage_bytes,
342 d_samples,
343 d_histogram,
344 num_levels,
345 d_levels,
346 num_row_pixels,
347 num_rows,
348 row_stride_bytes,
349 _stream,
350 false));
351 }
352};
353} // namespace muda
354
355#include "details/cub_wrapper_macro_undef.inl"
muda::CubWrapper
Definition cub_wrapper.h:14
muda::DeviceHistogram
Definition device_histogram.h:13
muda::LaunchCore
Definition launch_base.h:42