BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH > Class Template Reference

• hipcub
• BlockRadixRankMatch

Classes | Public Types | List of all members

hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH > Class Template Reference

#include <block_radix_rank.hpp>

Classes
struct	TempStorage
	\smemstorage{BlockScan} More...

Public Types
enum	{ BINS_TRACKED_PER_THREAD = rocprim::maximum<int>()(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS) }

Public Member Functions
Collective constructors
__device__	BlockRadixRankMatch (TempStorage &temp_storage)
	Collective constructor using the specified memory allocation as temporary storage. More...
Raking
template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >
__device__ __forceinline__ void	RankKeys (UnsignedBits(&keys)[KEYS_PER_THREAD], int(&ranks)[KEYS_PER_THREAD], DigitExtractorT digit_extractor)
	Rank keys. More...
template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >
__device__ __forceinline__ void	RankKeys (UnsignedBits(&keys)[KEYS_PER_THREAD], int(&ranks)[KEYS_PER_THREAD], DigitExtractorT digit_extractor, int(&exclusive_digit_prefix)[BINS_TRACKED_PER_THREAD])
	Rank keys. For the lower RADIX_DIGITS threads, digit counts for each digit are provided for the corresponding thread. More...

Detailed Description

template<int BLOCK_DIM_X, int RADIX_BITS, bool IS_DESCENDING, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int ARCH = 1>
class hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH >

Radix-rank using match.any

Member Enumeration Documentation

◆ anonymous enum

template<int BLOCK_DIM_X, int RADIX_BITS, bool IS_DESCENDING, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int ARCH = 1>

anonymous enum

Enumerator
BINS_TRACKED_PER_THREAD	Number of bin-starting offsets tracked per thread.

Constructor & Destructor Documentation

◆ BlockRadixRankMatch()

template<int BLOCK_DIM_X, int RADIX_BITS, bool IS_DESCENDING, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int ARCH = 1>

__device__ hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH >::BlockRadixRankMatch ( TempStorage &  temp_storage )

inline

Collective constructor using the specified memory allocation as temporary storage.

Parameters

[in]

temp_storage

Reference to memory allocation having layout type TempStorage

Member Function Documentation

◆ RankKeys() [1/2]

template<int BLOCK_DIM_X, int RADIX_BITS, bool IS_DESCENDING, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int ARCH = 1>

template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >

__device__ __forceinline__ void hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH >::RankKeys ( UnsignedBits(&)  keys[KEYS_PER_THREAD], int(&)  ranks[KEYS_PER_THREAD], DigitExtractorT  digit_extractor  )

inline

Rank keys.

Parameters

[in]	keys	Keys for this tile
[out]	ranks	For each key, the local rank within the tile
[in]	digit_extractor	The digit extractor

◆ RankKeys() [2/2]

template<int BLOCK_DIM_X, int RADIX_BITS, bool IS_DESCENDING, BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int ARCH = 1>

template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >

__device__ __forceinline__ void hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH >::RankKeys ( UnsignedBits(&)  keys[KEYS_PER_THREAD], int(&)  ranks[KEYS_PER_THREAD], DigitExtractorT  digit_extractor, int(&)  exclusive_digit_prefix[BINS_TRACKED_PER_THREAD]  )

inline

Rank keys. For the lower RADIX_DIGITS threads, digit counts for each digit are provided for the corresponding thread.

Parameters

[in]	keys	Keys for this tile
[out]	ranks	For each key, the local rank within the tile (out parameter)
[in]	digit_extractor	The digit extractor
[out]	exclusive_digit_prefix	The exclusive prefix sum for the digits [(threadIdx.x * BINS_TRACKED_PER_THREAD) ... (threadIdx.x * BINS_TRACKED_PER_THREAD) + BINS_TRACKED_PER_THREAD - 1]

---

The documentation for this class was generated from the following file:

• /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.4.3/hipcub/include/hipcub/backend/rocprim/block/block_radix_rank.hpp