Radix-rank using match.any. More...

#include <block_radix_rank.hpp>

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

[legend]

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

[legend]

Public Types
enum	{ BINS_TRACKED_PER_THREAD = base_type::digits_per_thread }

using	TempStorage = typename base_type::storage_type

Public Member Functions
Collective constructors

HIPCUB_DEVICE	BlockRadixRankMatch ()
	Collective constructor using a private static allocation of shared memory as temporary storage.

HIPCUB_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. 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 = HIPCUB_ARCH>
class 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 = HIPCUB_ARCH>

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 = HIPCUB_ARCH>

HIPCUB_DEVICE 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 = HIPCUB_ARCH>

template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >

__device__ __forceinline__ void 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 = HIPCUB_ARCH>

template<typename UnsignedBits , int KEYS_PER_THREAD, typename DigitExtractorT >

__device__ __forceinline__ void 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:

hipcub/include/hipcub/backend/rocprim/block/block_radix_rank.hpp

Public Types

Public Member Functions

Detailed Description

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ BlockRadixRankMatch()

Member Function Documentation

◆ RankKeys() [1/2]

◆ RankKeys() [2/2]