cub::LoadDirectStripedVectorized#

template<int ThreadsPerBlock, ::cuda::std::size_t VecSize, CacheLoadModifier Modifier = LOAD_DEFAULT, typename InputT, typename T, int ItemsPerThread>
void cub::LoadDirectStripedVectorized(
int linear_tid,
const InputT *block_src_ptr,
T (&dst_items)[ItemsPerThread]
)#

Load a linear segment of items into a striped arrangement of vectors across the thread block.

Added in version 3.6.0.

Assumes a striped arrangement of vectors of VecSize consecutive items across the thread block, where threadi owns the vectors (i), (i + block-threads), …, (i + (block-threads * (items-per-thread / VecSize - 1))). With VecSize == 1, this is the arrangement of LoadDirectStriped.

Each vector is loaded as a CubVector<InputT, VecSize>::Type. The following conditions will prevent vectorization and the same arrangement is loaded item by item:

  • block_src_ptr is not aligned to the vector access width

  • InputT is volatile or not trivially copyable

Template Parameters:

  • ThreadsPerBlock – The thread block size in threads

  • VecSize – The number of consecutive items in each vector, from 1 to 4. Must evenly divide ItemsPerThread.

  • Modifier – [optional] Cache load modifier for reading the input (default: LOAD_DEFAULT)

  • InputT – [inferred] The value type of the input, which is loaded with vectors and converted to T.

  • T – [inferred] The data type to load.

  • ItemsPerThread – [inferred] The number of items partitioned onto each thread.

Parameters:

  • linear_tid – [in] A suitable 1D thread-identifier for the calling thread (e.g., (threadIdx.y * blockDim.x) + linear_tid for 2D thread blocks)

  • block_src_ptr – [in] The thread block’s base pointer for loading from

  • dst_items – [out] Destination to load data into