cub::BlockMergeSortStrategy

Defined in cub/block/block_merge_sort.cuh

template<typename KeyT, typename ValueT, int NUM_THREADS, int ITEMS_PER_THREAD, typename SynchronizationPolicy>
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 ITEMS_PER_THREAD = 9;

class BlockMergeSort : public BlockMergeSortStrategy<int,
                                                     cub::NullType,
                                                     BLOCK_THREADS,
                                                     ITEMS_PER_THREAD,
                                                     BlockMergeSort>
{
  using BlockMergeSortStrategyT =
    BlockMergeSortStrategy<int,
                           cub::NullType,
                           BLOCK_THREADS,
                           ITEMS_PER_THREAD,
                           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.

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)[ITEMS_PER_THREAD], CompareOp compare_op)

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

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)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items, KeyT oob_default)

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

  • 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.

  • The value of oob_default is assigned to all elements that are out of valid_items boundaries. It’s expected that oob_default is ordered after any value in the valid_items boundaries. The algorithm always sorts a fixed amount of elements, which is equal to ITEMS_PER_THREAD * 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] Default value to assign out-of-bound items

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

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

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)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items, KeyT oob_default)

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

  • 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.

  • The value of oob_default is assigned to all elements that are out of valid_items boundaries. It’s expected that oob_default is ordered after any value in the valid_items boundaries. The algorithm always sorts a fixed amount of elements, which is equal to ITEMS_PER_THREAD * 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] Default value to assign out-of-bound items

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

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

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)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THREAD], CompareOp compare_op)

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

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)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items, KeyT oob_default)

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

  • 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.

  • The value of oob_default is assigned to all elements that are out of valid_items boundaries. It’s expected that oob_default is ordered after any value in the valid_items boundaries. The algorithm always sorts a fixed amount of elements, which is equal to ITEMS_PER_THREAD * 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] Default value to assign out-of-bound items

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

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

  • 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.

  • The value of oob_default is assigned to all elements that are out of valid_items boundaries. It’s expected that oob_default is ordered after any value in the valid_items boundaries. The algorithm always sorts a fixed amount of elements, which is equal to ITEMS_PER_THREAD * 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] Default value to assign out-of-bound items

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.