/
device_scan.hpp
686 lines (639 loc) · 30.7 KB
/
device_scan.hpp
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
// Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
#ifndef ROCPRIM_DEVICE_DEVICE_SCAN_HPP_
#define ROCPRIM_DEVICE_DEVICE_SCAN_HPP_
#include <iostream>
#include <iterator>
#include <type_traits>
#include "../config.hpp"
#include "../detail/temp_storage.hpp"
#include "../detail/various.hpp"
#include "../functional.hpp"
#include "../type_traits.hpp"
#include "../types/future_value.hpp"
#include "detail/config/device_scan.hpp"
#include "detail/device_scan.hpp"
#include "detail/device_scan_common.hpp"
#include "device_scan_config.hpp"
#include "device_transform.hpp"
BEGIN_ROCPRIM_NAMESPACE
/// \addtogroup devicemodule
/// @{
namespace detail
{
// Single kernel scan (performs scan on one thread block only)
template<bool Exclusive,
class Config,
class InputIterator,
class OutputIterator,
class BinaryFunction,
class AccType>
ROCPRIM_DEVICE ROCPRIM_FORCE_INLINE void single_scan_kernel_impl(InputIterator input,
const size_t input_size,
AccType initial_value,
OutputIterator output,
BinaryFunction scan_op)
{
static constexpr scan_config_params params = device_params<Config>();
constexpr unsigned int block_size = params.kernel_config.block_size;
constexpr unsigned int items_per_thread = params.kernel_config.items_per_thread;
using block_load_type
= ::rocprim::block_load<AccType, block_size, items_per_thread, params.block_load_method>;
using block_store_type
= ::rocprim::block_store<AccType, block_size, items_per_thread, params.block_store_method>;
using block_scan_type = ::rocprim::block_scan<AccType, block_size, params.block_scan_method>;
ROCPRIM_SHARED_MEMORY union
{
typename block_load_type::storage_type load;
typename block_store_type::storage_type store;
typename block_scan_type::storage_type scan;
} storage;
AccType values[items_per_thread];
// load input values into values
block_load_type().load(input, values, input_size, *(input), storage.load);
::rocprim::syncthreads(); // sync threads to reuse shared memory
single_scan_block_scan<Exclusive, block_scan_type>(values, // input
values, // output
initial_value,
storage.scan,
scan_op);
::rocprim::syncthreads(); // sync threads to reuse shared memory
// Save values into output array
block_store_type().store(output, values, input_size, storage.store);
}
template<bool Exclusive,
class Config,
class InputIterator,
class OutputIterator,
class BinaryFunction,
class InitValueType,
class AccType>
ROCPRIM_KERNEL
__launch_bounds__(device_params<Config>().kernel_config.block_size) void single_scan_kernel(
InputIterator input,
const size_t size,
const InitValueType initial_value,
OutputIterator output,
BinaryFunction scan_op)
{
single_scan_kernel_impl<Exclusive, Config>(input,
size,
static_cast<AccType>(get_input_value(initial_value)),
output,
scan_op);
}
// Single pass (look-back kernels)
template<bool Exclusive,
class Config,
class InputIterator,
class OutputIterator,
class BinaryFunction,
class InitValueType,
class AccType,
class LookBackScanState>
ROCPRIM_KERNEL
__launch_bounds__(device_params<Config>().kernel_config.block_size) void lookback_scan_kernel(
InputIterator input,
OutputIterator output,
const size_t size,
const InitValueType initial_value,
BinaryFunction scan_op,
LookBackScanState lookback_scan_state,
const unsigned int number_of_blocks,
AccType* previous_last_element = nullptr,
AccType* new_last_element = nullptr,
bool override_first_value = false,
bool save_last_value = false)
{
lookback_scan_kernel_impl<Exclusive, Config>(
input,
output,
size,
static_cast<AccType>(get_input_value(initial_value)),
scan_op,
lookback_scan_state,
number_of_blocks,
previous_last_element,
new_last_element,
override_first_value,
save_last_value);
}
#define ROCPRIM_DETAIL_HIP_SYNC(name, size, start) \
if(debug_synchronous) \
{ \
std::cout << name << "(" << size << ")"; \
auto error = hipStreamSynchronize(stream); \
if(error != hipSuccess) return error; \
auto end = std::chrono::high_resolution_clock::now(); \
auto d = std::chrono::duration_cast<std::chrono::duration<double>>(end - start); \
std::cout << " " << d.count() * 1000 << " ms" << '\n'; \
}
#define ROCPRIM_DETAIL_HIP_SYNC_AND_RETURN_ON_ERROR(name, size, start) \
{ \
auto _error = hipGetLastError(); \
if(_error != hipSuccess) return _error; \
if(debug_synchronous) \
{ \
std::cout << name << "(" << size << ")"; \
auto __error = hipStreamSynchronize(stream); \
if(__error != hipSuccess) return __error; \
auto _end = std::chrono::high_resolution_clock::now(); \
auto _d = std::chrono::duration_cast<std::chrono::duration<double>>(_end - start); \
std::cout << " " << _d.count() * 1000 << " ms" << '\n'; \
} \
}
template<bool Exclusive,
class Config,
class InputIterator,
class OutputIterator,
class InitValueType,
class BinaryFunction,
class AccType>
inline auto scan_impl(void* temporary_storage,
size_t& storage_size,
InputIterator input,
OutputIterator output,
const InitValueType initial_value,
const size_t size,
BinaryFunction scan_op,
const hipStream_t stream,
bool debug_synchronous)
{
using config = wrapped_scan_config<Config, AccType>;
detail::target_arch target_arch;
hipError_t result = host_target_arch(stream, target_arch);
if(result != hipSuccess)
{
return result;
}
const scan_config_params params = dispatch_target_arch<config>(target_arch);
using scan_state_type = detail::lookback_scan_state<AccType>;
using scan_state_with_sleep_type = detail::lookback_scan_state<AccType, true>;
const unsigned int block_size = params.kernel_config.block_size;
const unsigned int items_per_thread = params.kernel_config.items_per_thread;
const auto items_per_block = block_size * items_per_thread;
const size_t size_limit = params.kernel_config.size_limit;
const size_t aligned_size_limit
= ::rocprim::max<size_t>(size_limit - size_limit % items_per_block, items_per_block);
size_t limited_size = std::min<size_t>(size, aligned_size_limit);
const bool use_limited_size = limited_size == aligned_size_limit;
unsigned int number_of_blocks = (limited_size + items_per_block - 1)/items_per_block;
// Pointer to array with block_prefixes
void* scan_state_storage;
AccType* previous_last_element;
AccType* new_last_element;
detail::temp_storage::layout layout{};
hipError_t layout_result
= scan_state_type::get_temp_storage_layout(number_of_blocks, stream, layout);
if(layout_result != hipSuccess)
{
return layout_result;
}
const hipError_t partition_result = detail::temp_storage::partition(
temporary_storage,
storage_size,
detail::temp_storage::make_linear_partition(
// This is valid even with offset_scan_state_with_sleep_type
detail::temp_storage::make_partition(&scan_state_storage, layout),
detail::temp_storage::ptr_aligned_array(&previous_last_element,
use_limited_size ? 1 : 0),
detail::temp_storage::ptr_aligned_array(&new_last_element, use_limited_size ? 1 : 0)));
if(partition_result != hipSuccess || temporary_storage == nullptr)
{
return partition_result;
}
// Start point for time measurements
std::chrono::high_resolution_clock::time_point start;
if( number_of_blocks == 0u )
return hipSuccess;
if(number_of_blocks > 1 || use_limited_size)
{
// Create and initialize lookback_scan_state obj
scan_state_type scan_state{};
hipError_t result
= scan_state_type::create(scan_state, scan_state_storage, number_of_blocks, stream);
scan_state_with_sleep_type scan_state_with_sleep{};
result = scan_state_with_sleep_type::create(scan_state_with_sleep,
scan_state_storage,
number_of_blocks,
stream);
if(result != hipSuccess)
{
return result;
}
hipDeviceProp_t prop;
int deviceId;
static_cast<void>(hipGetDevice(&deviceId));
static_cast<void>(hipGetDeviceProperties(&prop, deviceId));
if(debug_synchronous) start = std::chrono::high_resolution_clock::now();
#if HIP_VERSION >= 307
int asicRevision = prop.asicRevision;
#else
int asicRevision = 0;
#endif
size_t number_of_launch = (size + limited_size - 1)/limited_size;
for (size_t i = 0, offset = 0; i < number_of_launch; i++, offset+=limited_size )
{
size_t current_size = std::min<size_t>(size - offset, limited_size);
number_of_blocks = (current_size + items_per_block - 1)/items_per_block;
auto grid_size = (number_of_blocks + block_size - 1)/block_size;
if(debug_synchronous)
{
std::cout << "use_limited_size " << use_limited_size << '\n';
std::cout << "aligned_size_limit " << aligned_size_limit << '\n';
std::cout << "number_of_launch " << number_of_launch << '\n';
std::cout << "index " << i << '\n';
std::cout << "size " << current_size << '\n';
std::cout << "block_size " << block_size << '\n';
std::cout << "number of blocks " << number_of_blocks << '\n';
std::cout << "items_per_block " << items_per_block << '\n';
}
if(std::string(prop.gcnArchName).find("908") != std::string::npos && asicRevision < 2)
{
init_lookback_scan_state_kernel<scan_state_with_sleep_type>
<<<dim3(grid_size), dim3(block_size), 0, stream>>>(scan_state_with_sleep,
number_of_blocks);
} else
{
init_lookback_scan_state_kernel<scan_state_type>
<<<dim3(grid_size), dim3(block_size), 0, stream>>>(scan_state,
number_of_blocks);
}
ROCPRIM_DETAIL_HIP_SYNC_AND_RETURN_ON_ERROR("init_lookback_scan_state_kernel", number_of_blocks, start)
if(debug_synchronous) start = std::chrono::high_resolution_clock::now();
grid_size = number_of_blocks;
if(std::string(prop.gcnArchName).find("908") != std::string::npos && asicRevision < 2)
{
lookback_scan_kernel<Exclusive, // flag for exclusive scan operation
config,
InputIterator,
OutputIterator,
BinaryFunction,
InitValueType,
AccType,
scan_state_with_sleep_type>
<<<dim3(grid_size), dim3(block_size), 0, stream>>>(input + offset,
output + offset,
current_size,
initial_value,
scan_op,
scan_state_with_sleep,
number_of_blocks,
previous_last_element,
new_last_element,
i != size_t(0),
number_of_launch > 1);
}
else
{
if(debug_synchronous)
{
std::cout << "use_limited_size " << use_limited_size << '\n';
std::cout << "aligned_size_limit " << aligned_size_limit << '\n';
std::cout << "size " << current_size << '\n';
std::cout << "block_size " << block_size << '\n';
std::cout << "number of blocks " << number_of_blocks << '\n';
std::cout << "items_per_block " << items_per_block << '\n';
}
lookback_scan_kernel<Exclusive, // flag for exclusive scan operation
config,
InputIterator,
OutputIterator,
BinaryFunction,
InitValueType,
AccType,
scan_state_type>
<<<dim3(grid_size), dim3(block_size), 0, stream>>>(input + offset,
output + offset,
current_size,
initial_value,
scan_op,
scan_state,
number_of_blocks,
previous_last_element,
new_last_element,
i != size_t(0),
number_of_launch > 1);
}
ROCPRIM_DETAIL_HIP_SYNC_AND_RETURN_ON_ERROR("lookback_scan_kernel", current_size, start)
// Swap the last_elements
if(number_of_launch > 1)
{
hipError_t error = ::rocprim::transform(new_last_element,
previous_last_element,
1,
::rocprim::identity<AccType>(),
stream,
debug_synchronous);
if(error != hipSuccess) return error;
}
}
}
else
{
if(debug_synchronous)
{
std::cout << "size " << size << '\n';
std::cout << "block_size " << block_size << '\n';
std::cout << "number of blocks " << number_of_blocks << '\n';
std::cout << "items_per_block " << items_per_block << '\n';
start = std::chrono::high_resolution_clock::now();
}
single_scan_kernel<Exclusive, // flag for exclusive scan operation
config,
InputIterator,
OutputIterator,
BinaryFunction,
InitValueType,
AccType>
<<<dim3(1), dim3(block_size), 0, stream>>>(input, size, initial_value, output, scan_op);
ROCPRIM_DETAIL_HIP_SYNC_AND_RETURN_ON_ERROR("single_scan_kernel", size, start);
}
return hipSuccess;
}
#undef ROCPRIM_DETAIL_HIP_SYNC_AND_RETURN_ON_ERROR
#undef ROCPRIM_DETAIL_HIP_SYNC
} // end of detail namespace
/// \brief Parallel inclusive scan primitive for device level.
///
/// inclusive_scan function performs a device-wide inclusive prefix scan operation
/// using binary \p scan_op operator.
///
/// \par Overview
/// * Supports non-commutative scan operators. However, a scan operator should be
/// associative.
/// * When used with non-associative functions (e.g. floating point arithmetic operations):
/// - the results may be non-deterministic and/or vary in precision,
/// - and bit-wise reproducibility is not guaranteed, that is, results from multiple runs
/// using the same input values on the same device may not be bit-wise identical.
/// * Returns the required size of \p temporary_storage in \p storage_size
/// if \p temporary_storage in a null pointer.
/// * Ranges specified by \p input and \p output must have at least \p size elements.
/// * By default, the input type is used for accumulation. A custom type
/// can be specified using the \p AccType type parameter, see the example below.
///
/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config or a class derived from it.
/// \tparam InputIterator - random-access iterator type of the input range. Must meet the
/// requirements of a C++ InputIterator concept. It can be a simple pointer type.
/// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
/// requirements of a C++ OutputIterator concept. It can be a simple pointer type.
/// \tparam BinaryFunction - type of binary function used for scan. Default type
/// is \p rocprim::plus<T>, where \p T is a \p value_type of \p InputIterator.
/// \tparam AccType - accumulator type used to propagate the scanned values. Default type
/// is value type of the input iterator.
///
/// \param [in] temporary_storage - pointer to a device-accessible temporary storage. When
/// a null pointer is passed, the required allocation size (in bytes) is written to
/// \p storage_size and function returns without performing the scan operation.
/// \param [in,out] storage_size - reference to a size (in bytes) of \p temporary_storage.
/// \param [in] input - iterator to the first element in the range to scan.
/// \param [out] output - iterator to the first element in the output range. It can be
/// same as \p input.
/// \param [in] size - number of element in the input range.
/// \param [in] scan_op - binary operation function object that will be used for scan.
/// The signature of the function should be equivalent to the following:
/// <tt>T f(const T &a, const T &b);</tt>. The signature does not need to have
/// <tt>const &</tt>, but function object must not modify the objects passed to it.
/// Default is BinaryFunction().
/// \param [in] stream - [optional] HIP stream object. Default is \p 0 (default stream).
/// \param [in] debug_synchronous - [optional] If true, synchronization after every kernel
/// launch is forced in order to check for errors. Default value is \p false.
///
/// \returns \p hipSuccess (\p 0) after successful scan; otherwise a HIP runtime error of
/// type \p hipError_t.
///
/// \par Example
/// \parblock
/// In this example a device-level inclusive sum operation is performed on an array of
/// integer values (<tt>short</tt>s are scanned into <tt>int</tt>s).
///
/// \code{.cpp}
/// #include <rocprim/rocprim.hpp>
///
/// // Prepare input and output (declare pointers, allocate device memory etc.)
/// size_t input_size; // e.g., 8
/// short * input; // e.g., [1, 2, 3, 4, 5, 6, 7, 8]
/// int * output; // empty array of 8 elements
///
/// size_t temporary_storage_size_bytes;
/// void * temporary_storage_ptr = nullptr;
/// // Get required size of the temporary storage
/// rocprim::inclusive_scan(
/// temporary_storage_ptr, temporary_storage_size_bytes,
/// input, output, input_size, rocprim::plus<int>()
/// );
///
/// // allocate temporary storage
/// hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);
///
/// // perform scan
/// rocprim::inclusive_scan(
/// temporary_storage_ptr, temporary_storage_size_bytes,
/// input, output, input_size, rocprim::plus<int>()
/// );
/// // output: [1, 3, 6, 10, 15, 21, 28, 36]
/// \endcode
///
/// The same example as above, but now a custom accumulator type is specified.
///
/// \code{.cpp}
/// #include <rocprim/rocprim.hpp>
///
/// size_t input_size;
/// short * input;
/// int * output;
///
/// size_t temporary_storage_size_bytes;
/// void * temporary_storage_ptr = nullptr;
///
/// rocprim::inclusive_scan(
/// temporary_storage_ptr, temporary_storage_size_bytes,
/// input, output, input_size, rocprim::plus<int>()
/// );
///
/// hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);
///
/// // Use type parameter to set custom accumulator type
/// rocprim::inclusive_scan<rocprim::default_config,
/// short*,
/// int*,
/// rocprim::plus<int>,
/// int>(temporary_storage_ptr,
/// temporary_storage_size_bytes,
/// input_iterator,
/// output,
/// input_size,
/// rocprim::plus<int>());
/// \endcode
/// \endparblock
template<class Config = default_config,
class InputIterator,
class OutputIterator,
class BinaryFunction
= ::rocprim::plus<typename std::iterator_traits<InputIterator>::value_type>,
class AccType = typename std::iterator_traits<InputIterator>::value_type>
inline hipError_t inclusive_scan(void* temporary_storage,
size_t& storage_size,
InputIterator input,
OutputIterator output,
const size_t size,
BinaryFunction scan_op = BinaryFunction(),
const hipStream_t stream = 0,
bool debug_synchronous = false)
{
// input_type() is a dummy initial value (not used)
return detail::
scan_impl<false, Config, InputIterator, OutputIterator, AccType, BinaryFunction, AccType>(
temporary_storage,
storage_size,
input,
output,
AccType{},
size,
scan_op,
stream,
debug_synchronous);
}
/// \brief Parallel exclusive scan primitive for device level.
///
/// exclusive_scan function performs a device-wide exclusive prefix scan operation
/// using binary \p scan_op operator.
///
/// \par Overview
/// * Supports non-commutative scan operators. However, a scan operator should be
/// associative.
/// * When used with non-associative functions (e.g. floating point arithmetic operations):
/// - the results may be non-deterministic and/or vary in precision,
/// - and bit-wise reproducibility is not guaranteed, that is, results from multiple runs
/// using the same input values on the same device may not be bit-wise identical.
/// * Returns the required size of \p temporary_storage in \p storage_size
/// if \p temporary_storage in a null pointer.
/// * Ranges specified by \p input and \p output must have at least \p size elements.
///
/// \tparam Config - [optional] configuration of the primitive, has to be \p scan_config or a class derived from it.
/// \tparam InputIterator - random-access iterator type of the input range. Must meet the
/// requirements of a C++ InputIterator concept. It can be a simple pointer type.
/// \tparam OutputIterator - random-access iterator type of the output range. Must meet the
/// requirements of a C++ OutputIterator concept. It can be a simple pointer type.
/// \tparam InitValueType - type of the initial value.
/// \tparam BinaryFunction - type of binary function used for scan. Default type
/// is \p rocprim::plus<T>, where \p T is a \p value_type of \p InputIterator.
/// \tparam AccType - accumulator type used to propagate the scanned values. Default type
/// is 'InitValueType', unless it's 'rocprim::future_value'. Then it will be the wrapped input type.
///
/// \param [in] temporary_storage - pointer to a device-accessible temporary storage. When
/// a null pointer is passed, the required allocation size (in bytes) is written to
/// \p storage_size and function returns without performing the scan operation.
/// \param [in,out] storage_size - reference to a size (in bytes) of \p temporary_storage.
/// \param [in] input - iterator to the first element in the range to scan.
/// \param [out] output - iterator to the first element in the output range. It can be
/// same as \p input.
/// \param [in] initial_value - initial value to start the scan.
/// A rocpim::future_value may be passed to use a value that will be later computed.
/// \param [in] size - number of element in the input range.
/// \param [in] scan_op - binary operation function object that will be used for scan.
/// The signature of the function should be equivalent to the following:
/// <tt>T f(const T &a, const T &b);</tt>. The signature does not need to have
/// <tt>const &</tt>, but function object must not modify the objects passed to it.
/// The default value is \p BinaryFunction().
/// \param [in] stream - [optional] HIP stream object. The default is \p 0 (default stream).
/// \param [in] debug_synchronous - [optional] If true, synchronization after every kernel
/// launch is forced in order to check for errors. The default value is \p false.
///
/// \returns \p hipSuccess (\p 0) after successful scan; otherwise a HIP runtime error of
/// type \p hipError_t.
///
/// \par Example
/// \parblock
/// In this example a device-level exclusive min-scan operation is performed on an array of
/// integer values (<tt>short</tt>s are scanned into <tt>int</tt>s) using custom operator.
///
/// \code{.cpp}
/// #include <rocprim/rocprim.hpp>
///
/// // custom scan function
/// auto min_op =
/// [] __device__ (int a, int b) -> int
/// {
/// return a < b ? a : b;
/// };
///
/// // Prepare input and output (declare pointers, allocate device memory etc.)
/// size_t input_size; // e.g., 8
/// short * input; // e.g., [4, 7, 6, 2, 5, 1, 3, 8]
/// int * output; // empty array of 8 elements
/// int start_value; // e.g., 9
///
/// size_t temporary_storage_size_bytes;
/// void * temporary_storage_ptr = nullptr;
/// // Get required size of the temporary storage
/// rocprim::exclusive_scan(
/// temporary_storage_ptr, temporary_storage_size_bytes,
/// input, output, start_value, input_size, min_op
/// );
///
/// // allocate temporary storage
/// hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes);
///
/// // perform scan
/// rocprim::exclusive_scan(
/// temporary_storage_ptr, temporary_storage_size_bytes,
/// input, output, start_value, input_size, min_op
/// );
/// // output: [9, 4, 4, 4, 2, 2, 1, 1]
/// \endcode
/// \endparblock
template<class Config = default_config,
class InputIterator,
class OutputIterator,
class InitValueType,
class BinaryFunction
= ::rocprim::plus<typename std::iterator_traits<InputIterator>::value_type>,
class AccType = detail::input_type_t<InitValueType>>
inline hipError_t exclusive_scan(void* temporary_storage,
size_t& storage_size,
InputIterator input,
OutputIterator output,
const InitValueType initial_value,
const size_t size,
BinaryFunction scan_op = BinaryFunction(),
const hipStream_t stream = 0,
bool debug_synchronous = false)
{
return detail::scan_impl<true,
Config,
InputIterator,
OutputIterator,
InitValueType,
BinaryFunction,
AccType>(temporary_storage,
storage_size,
input,
output,
initial_value,
size,
scan_op,
stream,
debug_synchronous);
}
/// @}
// end of group devicemodule
END_ROCPRIM_NAMESPACE
#endif // ROCPRIM_DEVICE_DEVICE_SCAN_HPP_