Template Class BlockMergeSortStrategy

Nested Relationships

Nested Types

Class Documentation

template<typename KeyT, typename ValueT, int NUM_THREADS, int ITEMS_PER_THREAD, typename SynchronizationPolicy>
class hipcub::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 <hipcub/hipcub.hpp> // or equivalently <hipcub/block/block_merge_sort.hpp>

constexpr int BLOCK_THREADS = 256;
constexpr int ITEMS_PER_THREAD = 9;

class BlockMergeSort : public BlockMergeSortStrategy<int,
                                                     hipcub::NullType,
                                                     BLOCK_THREADS,
                                                     ITEMS_PER_THREAD,
                                                     BlockMergeSort>
{
  using BlockMergeSortStrategyT =
    BlockMergeSortStrategy<int,
                           hipcub::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. hipcub::NullType indicates a keys-only sort

  • SynchronizationPolicy – Provides a way of synchronizing threads. Should be derived from BlockMergeSortStrategy.

Public Functions

BlockMergeSortStrategy() = delete
__device__ __forceinline__ inline explicit BlockMergeSortStrategy(unsigned int linear_tid)
__device__ __forceinline__ inline BlockMergeSortStrategy(TempStorage &temp_storage, unsigned int linear_tid)
__device__ __forceinline__ inline unsigned int get_linear_tid() const
template<typename CompareOp>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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>
__device__ __forceinline__ 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 hipcub::Uninitialized<_TempStorage>

{BlockMergeSort}

Public Types

enum [anonymous]

Values:

typedef UnitWord<_TempStorage>::DeviceWord DeviceWord

Biggest memory-access word that T is a whole multiple of and is not larger than the alignment of T.

Public Functions

__host__ __device__ __forceinline__ inline _TempStorage &Alias()

Alias.

Public Members

DeviceWord storage[WORDS]

Backing storage.