cub::BlockLoadAlgorithm#

enum cub::BlockLoadAlgorithm#

cub::BlockLoadAlgorithm enumerates alternative algorithms for cub::BlockLoad to read a linear segment of data from memory into a blocked arrangement across a CUDA thread block.

Values:

enumerator BLOCK_LOAD_DIRECT#

Overview#

A blocked arrangement of data is read directly from memory.

Performance Considerations#

The utilization of memory transactions (coalescing) decreases as the access stride between threads increases (i.e., the number items per thread).

enumerator BLOCK_LOAD_STRIPED#

Overview#

A striped arrangement of data is read directly from memory.

Performance Considerations#

The utilization of memory transactions (coalescing) doesn’t depend on the number of items per thread.

enumerator BLOCK_LOAD_VECTORIZE#

Overview#

A blocked arrangement of data is read from memory using CUDA’s built-in vectorized loads as a coalescing optimization. For example, ld.global.v4.s32 instructions will be generated when T = int and ItemsPerThread % 4 == 0.

Performance Considerations#

  • The utilization of memory transactions (coalescing) remains high until the the access stride between threads (i.e., the number items per thread) exceeds the maximum vector load width (typically 4 items or 64B, whichever is lower).

  • The following conditions will prevent vectorization and loading will fall back to cub::BLOCK_LOAD_DIRECT:

    • The RandomAccessIterator is not a contiguous iterator (cub::CacheModifiedInputIterator is an exception and is still vectorized)

    • The value type of the RandomAccessIterator is not T

    • The block input offset is not aligned to the vector access width

    • The input is volatile or T is not trivially copyable

enumerator BLOCK_LOAD_TRANSPOSE#

Overview#

A striped arrangement of data is read efficiently from memory and then locally transposed into a blocked arrangement.

Performance Considerations#

  • The utilization of memory transactions (coalescing) remains high regardless of items loaded per thread.

  • The local reordering incurs slightly longer latencies and throughput than the direct cub::BLOCK_LOAD_DIRECT and cub::BLOCK_LOAD_VECTORIZE alternatives.

enumerator BLOCK_LOAD_WARP_TRANSPOSE#

Overview#

A warp-striped arrangement of data is read efficiently from memory and then locally transposed into a blocked arrangement.

Usage Considerations#

  • ThreadsPerBlock must be a multiple of WARP_THREADS

Performance Considerations#

  • The utilization of memory transactions (coalescing) remains high regardless of items loaded per thread.

  • The local reordering incurs slightly larger latencies than the direct cub::BLOCK_LOAD_DIRECT and cub::BLOCK_LOAD_VECTORIZE alternatives.

  • Provisions more shared storage, but incurs smaller latencies than the BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED alternative.

enumerator BLOCK_LOAD_WARP_TRANSPOSE_TIMESLICED#

Overview#

Like BLOCK_LOAD_WARP_TRANSPOSE, a warp-striped arrangement of data is read directly from memory and then is locally transposed into a blocked arrangement. To reduce the shared memory requirement, only one warp’s worth of shared memory is provisioned and is subsequently time-sliced among warps.

Usage Considerations#

  • ThreadsPerBlock must be a multiple of WARP_THREADS

Performance Considerations#

  • The utilization of memory transactions (coalescing) remains high regardless of items loaded per thread.

  • Provisions less shared memory temporary storage, but incurs larger latencies than the BLOCK_LOAD_WARP_TRANSPOSE alternative.

enumerator BLOCK_LOAD_STRIPED_VECTORIZED#

Overview#

A striped arrangement of vectors is read from memory using CUDA’s built-in vectorized loads. Each vector holds consecutive items, and thread t owns the vectors at positions t, t + ThreadsPerBlock, and so on. Each vector holds as many items as the vector with which cub::BLOCK_LOAD_VECTORIZE loads the items of a thread, and is loaded as a CubVector of these items. If the vector of cub::BLOCK_LOAD_VECTORIZE does not hold a whole number of items, each vector is a single item, which is the arrangement of cub::BLOCK_LOAD_STRIPED.

Performance Considerations#

  • Consecutive threads load consecutive vectors, so the accesses stay coalesced regardless of the number of items per thread.

  • The following conditions will prevent vectorization and the same arrangement is loaded item by item:

    • The RandomAccessIterator is not a contiguous iterator (cub::CacheModifiedInputIterator is an exception and is still vectorized)

    • The value type of the RandomAccessIterator is not T

    • The block input offset is not aligned to the vector access width

    • The input is volatile or T is not trivially copyable

  • Guarded loads use the arrangement of cub::BLOCK_LOAD_STRIPED.