WarpLoadAlgorithm#

enum cub::WarpLoadAlgorithm#

cub::WarpLoadAlgorithm enumerates alternative algorithms for cub::WarpLoad to read a linear segment of data from memory into a CUDA warp.

Values:

enumerator WARP_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 WARP_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 WARP_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 ITEMS_PER_THREAD % 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::WARP_LOAD_DIRECT:

    • ITEMS_PER_THREAD is odd

    • The InputIteratorT is not a simple pointer type

    • The block input offset is not quadword-aligned

    • The data type T is not a built-in primitive or CUDA vector type (e.g., short, int2, double, float2, etc.)

enumerator WARP_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::WARP_LOAD_DIRECT and cub::WARP_LOAD_VECTORIZE alternatives.