cub::AgentWarpReducePolicy#

template<int _BLOCK_THREADS, int _WARP_THREADS, int NOMINAL_ITEMS_PER_THREAD_4B, typename ComputeT, int _VECTOR_LOAD_LENGTH, CacheLoadModifier _LOAD_MODIFIER>
struct AgentWarpReducePolicy#

Parameterizable tuning policy type for AgentReduce.

Template Parameters:

  • _BLOCK_THREADS – Threads per thread block

  • _WARP_THREADS – Threads per warp

  • NOMINAL_ITEMS_PER_THREAD_4B – Items per thread (per tile of input)

  • ComputeT – Dominant compute type

  • _VECTOR_LOAD_LENGTH – Number of items per vectorized load

  • _LOAD_MODIFIER – Cache load modifier for reading input elements

Public Static Attributes

static constexpr int WARP_THREADS = _WARP_THREADS#

Number of threads per warp.

static constexpr int VECTOR_LOAD_LENGTH = _VECTOR_LOAD_LENGTH#

Number of items per vectorized load.

static constexpr int BLOCK_THREADS = _BLOCK_THREADS#

Number of threads per block.

static constexpr int ITEMS_PER_THREAD = detail::MemBoundScaling<0, NOMINAL_ITEMS_PER_THREAD_4B, ComputeT>::ITEMS_PER_THREAD#

Number of items per thread.

static constexpr CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER#

Cache load modifier for reading input elements.

static constexpr int ITEMS_PER_TILE = ITEMS_PER_THREAD * WARP_THREADS#

Number of items per tile.

static constexpr int SEGMENTS_PER_BLOCK = BLOCK_THREADS / WARP_THREADS#

Number of segments per block.