cub::BlockMergeSortStrategy#

template<typename KeyT, typename ValueT, int NumThreads, int ItemsPerThread, typename SynchronizationPolicy, bool _Unroll = true> class BlockMergeSortStrategy#

Generalized merge sort algorithm.

This class is used to reduce code duplication. Warp and Block merge sort differ only in how they compute thread index and how they synchronize threads. Since synchronization might require access to custom data (like member mask), CRTP is used.

The code snippet below illustrates the way this class can be used.

#include <cub/cub.cuh> // or equivalently <cub/block/block_merge_sort.cuh>

constexpr int BLOCK_THREADS = 256;
constexpr int ItemsPerThread = 9;

class BlockMergeSort : public BlockMergeSortStrategy<int,
                                                     cub::NullType,
                                                     BLOCK_THREADS,
                                                     ItemsPerThread,
                                                     BlockMergeSort>
{
  using BlockMergeSortStrategyT =
    BlockMergeSortStrategy<int,
                           cub::NullType,
                           BLOCK_THREADS,
                           ItemsPerThread,
                           BlockMergeSort>;
public:
  __device__ __forceinline__ explicit BlockMergeSort(
    typename BlockMergeSortStrategyT::TempStorage &temp_storage)
      : BlockMergeSortStrategyT(temp_storage, threadIdx.x)
  {}

  __device__ __forceinline__ void SyncImplementation() const
  {
    __syncthreads();
  }
};

Template Parameters:

KeyT – KeyT type
ValueT – ValueT type. cub::NullType indicates a keys-only sort
SynchronizationPolicy – Provides a way of synchronizing threads. Should be derived from BlockMergeSortStrategy.

Subclassed by cub::BlockMergeSort< KeyT, BlockDimX, ItemsPerThread, ValueT, BlockDimY, BlockDimZ, _Unroll >, cub::WarpMergeSort< KeyT, ITEMS_PER_THREAD, LOGICAL_WARP_THREADS, ValueT >

Public Functions

BlockMergeSortStrategy() = delete#

inline explicit BlockMergeSortStrategy(unsigned int linear_tid)#

inline BlockMergeSortStrategy( TempStorage &temp_storage, unsigned int linear_tid )#

inline unsigned int get_linear_tid() const#

template<typename CompareOp> inline void Sort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

: Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second

template<typename CompareOp> inline void Sort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op, int valid_items, KeyT oob_default )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.
It is required that oob_default is ordered after any value in the valid_items boundaries and that all threads provide the same oob_default and valid_items. The algorithm always sorts a fixed amount of elements, which is equal to ItemsPerThread * BLOCK_THREADS. If there is a value that is ordered after oob_default, it won’t be placed within valid_items boundaries.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort
oob_default – [in] Value that must be ordered after any value within the valid_items boundaries

template<typename CompareOp> inline void Sort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

: Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second

template<typename CompareOp, bool IS_LAST_TILE = true> inline void Sort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op, int valid_items, KeyT oob_default )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.
It is required that oob_default is ordered after any value in the valid_items boundaries and that all threads provide the same oob_default and valid_items. The algorithm always sorts a fixed amount of elements, which is equal to ItemsPerThread * BLOCK_THREADS. If there is a value that is ordered after oob_default, it won’t be placed within valid_items boundaries.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs) CompareOp is a model of Strict Weak Ordering.
IS_LAST_TILE – True if valid_items isn’t equal to the ITEMS_PER_TILE

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort
oob_default – [in] Value that must be ordered after any value within the valid_items boundaries

template<typename CompareOp> inline void Sort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op, int valid_items )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Only the first valid_items elements are sorted; no out-of-bounds default value is required.

Added in version 3.6.

Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.
On output, the first valid_items positions of the tile hold the sorted valid keys; keys beyond the valid_items boundary have unspecified values. Unlike the overload taking oob_default, no sentinel value that orders after all valid keys is required, which supports comparators and types without one.
All threads in the thread block must provide the same valid_items.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort

template<typename CompareOp> inline void Sort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op, int valid_items )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Only the first valid_items elements are sorted; no out-of-bounds default value is required.

Added in version 3.6.

Sort is not guaranteed to be stable. That is, suppose that i and j are equivalent: neither one is less than the other. It is not guaranteed that the relative order of these two elements will be preserved by sort.
On output, the first valid_items positions of the tile hold the sorted valid keys; keys beyond the valid_items boundary have unspecified values. Unlike the overload taking oob_default, no sentinel value that orders after all valid keys is required, which supports comparators and types without one.
All threads in the thread block must provide the same valid_items.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort

template<typename CompareOp> inline void StableSort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

: StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second

template<typename CompareOp> inline void StableSort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

: StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second

template<typename CompareOp> inline void StableSort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op, int valid_items, KeyT oob_default )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.
It is required that oob_default is ordered after any value in the valid_items boundaries and that all threads provide the same oob_default and valid_items. The algorithm always sorts a fixed amount of elements, which is equal to ItemsPerThread * BLOCK_THREADS. If there is a value that is ordered after oob_default, it won’t be placed within valid_items boundaries.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort
oob_default – [in] Value that must be ordered after any value within the valid_items boundaries

template<typename CompareOp, bool IS_LAST_TILE = true> inline void StableSort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op, int valid_items, KeyT oob_default )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Added in version 2.2.0: First appears in CUDA Toolkit 12.3.

StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.
It is required that oob_default is ordered after any value in the valid_items boundaries and that all threads provide the same oob_default and valid_items. The algorithm always sorts a fixed amount of elements, which is equal to ItemsPerThread * BLOCK_THREADS. If there is a value that is ordered after oob_default, it won’t be placed within valid_items boundaries.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.
IS_LAST_TILE – True if valid_items isn’t equal to the ITEMS_PER_TILE

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort
oob_default – [in] Value that must be ordered after any value within the valid_items boundaries

template<typename CompareOp> inline void StableSort( KeyT (&keys)[ItemsPerThread], CompareOp compare_op, int valid_items )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Only the first valid_items elements are sorted; no out-of-bounds default value is required.

Added in version 3.6.

StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.
On output, the first valid_items positions of the tile hold the sorted valid keys; keys beyond the valid_items boundary have unspecified values. Unlike the overload taking oob_default, no sentinel value that orders after all valid keys is required, which supports comparators and types without one.
All threads in the thread block must provide the same valid_items.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort

template<typename CompareOp> inline void StableSort( KeyT (&keys)[ItemsPerThread], ValueT (&items)[ItemsPerThread], CompareOp compare_op, int valid_items )#

Sorts items partitioned across a CUDA thread block using a merge sorting method.

Only the first valid_items elements are sorted; no out-of-bounds default value is required.

Added in version 3.6.

StableSort is stable: it preserves the relative ordering of equivalent elements. That is, if x and y are elements such that x precedes y, and if the two elements are equivalent (neither x < y nor y < x) then a postcondition of StableSort is that x still precedes y.
On output, the first valid_items positions of the tile hold the sorted valid keys; keys beyond the valid_items boundary have unspecified values. Unlike the overload taking oob_default, no sentinel value that orders after all valid keys is required, which supports comparators and types without one.
All threads in the thread block must provide the same valid_items.

Template Parameters:

CompareOp – functor type having member bool operator()(KeyT lhs, KeyT rhs). CompareOp is a model of Strict Weak Ordering.

Parameters:

keys – [inout] Keys to sort
items – [inout] Values to sort
compare_op – [in] Comparison function object which returns true if the first argument is ordered before the second
valid_items – [in] Number of valid items to sort

struct TempStorage : public Uninitialized<_TempStorage>#: The operations exposed by BlockMergeSort require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the __shared__ keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or union’d with other storage allocation types to facilitate memory reuse.