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

BlockRadixRankMatch&lt; BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH &gt; Class Template Reference#

hipCUB: hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH > Class Template Reference
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>

Inheritance diagram for hipcub::BlockRadixRankMatch< BLOCK_DIM_X, RADIX_BITS, IS_DESCENDING, INNER_SCAN_ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, ARCH >:

Public Types

enum  { BINS_TRACKED_PER_THREAD = base_type::digits_per_thread }
 
using TempStorage = typename base_type::storage_type
 

Public Member Functions

Collective constructors
__device__ BlockRadixRankMatch ()
 Collective constructor using a private static allocation of shared memory as temporary storage.
 
__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_storageReference 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]keysKeys for this tile
[out]ranksFor each key, the local rank within the tile
[in]digit_extractorThe 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]keysKeys for this tile
[out]ranksFor each key, the local rank within the tile (out parameter)
[in]digit_extractorThe digit extractor
[out]exclusive_digit_prefixThe 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.7.0/hipcub/include/hipcub/backend/rocprim/block/block_radix_rank.hpp