Thread-Level Operations (Utilities)#

Scan#

exclusive#

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_exclusive(T (&input)[LENGTH], T (&output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix = true)#

Perform a sequential exclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • prefix – [in] Prefix to seed scan with

  • apply_prefix – [in] Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

  • prefix – [in] - Prefix to seed scan with

  • apply_prefix – [in] - Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Template Parameters

  • LENGTH – - Length of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_exclusive(T *input, T *output, ScanOp scan_op, T prefix, bool apply_prefix = true)#

Perform a sequential exclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • prefix – [in] Prefix to seed scan with

  • apply_prefix – [in] Whether or not the calling thread should apply its prefix. If not, the first output element is undefined. (Handy for preventing thread-0 from applying a prefix.)

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

  • prefix – [in] - Prefix to seed scan with

  • apply_prefix – [in] - Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Template Parameters

  • LENGTH – - Length of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_exclusive(T inclusive, T exclusive, T *input, T *output, ScanOp scan_op, Int2Type<LENGTH>)#

Perform a sequential exclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • inclusive – [in] Initial value for inclusive aggregate

  • exclusive – [in] Initial value for exclusive aggregate

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • inclusive – [in] - Initial value for inclusive aggregate

  • exclusive – [in] - Initial value for exclusive aggregate

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

Template Parameters

  • LENGTH – - Length of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

inclusive#

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_inclusive(T inclusive, T *input, T *output, ScanOp scan_op, Int2Type<LENGTH>)#

Perform a sequential exclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • inclusive – [in] Initial value for inclusive aggregate

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • inclusive – [in] - Initial value for inclusive aggregate

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

Template Parameters

  • LENGTH – - Length of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_inclusive(T *input, T *output, ScanOp scan_op)#

Perform a sequential inclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Template Parameters

  • LENGTH – - LengthT of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Parameters

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_inclusive(T (&input)[LENGTH], T (&output)[LENGTH], ScanOp scan_op)#

Perform a sequential inclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

Template Parameters

  • LENGTH – - LengthT of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_inclusive(T *input, T *output, ScanOp scan_op, T prefix, bool apply_prefix = true)#

Perform a sequential inclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Parameters

  • input – [in] Input array

  • output – [out] Output array (may be aliased to input)

  • scan_op – [in] Binary scan operator

  • prefix – [in] Prefix to seed scan with

  • apply_prefix – [in] Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

  • prefix – [in] - Prefix to seed scan with

  • apply_prefix – [in] - Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Template Parameters

  • LENGTH – - LengthT of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Returns

- Aggregate of the scan

template<int LENGTH, typename T, typename ScanOp>
__device__ inline T rocprim::thread_scan_inclusive(T (&input)[LENGTH], T (&output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix = true)#

Perform a sequential inclusive prefix scan over LENGTH elements of the input array. The aggregate is returned.

Template Parameters

  • LENGTH – - LengthT of input and output arrays

  • T – - [inferred] The data type to be scanned.

  • ScanOp – - [inferred] Binary scan operator type having member T operator()(const T &a, const T &b)

Parameters

  • input – [in] - Input array

  • output – [out] - Output array (may be aliased to input)

  • scan_op – [in] - Binary scan operator

  • prefix – [in] - Prefix to seed scan with

  • apply_prefix – [in] - Whether or not the calling thread should apply its prefix. (Handy for preventing thread-0 from applying a prefix.)

Returns

- Aggregate of the scan