cub::GridMappingStrategy enumerates alternative strategies for mapping constant-sized tiles of device-wide data onto a grid of CUDA thread blocks.
| Enumerator |
|---|
| GRID_MAPPING_RAKE | An a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.
- Overview
- The input is evenly partitioned into
p segments, where p is constant and corresponds loosely to the number of thread blocks that may actively reside on the target device. Each segment is comprised of consecutive tiles, where a tile is a small, constant-sized unit of input to be processed to completion before the thread block terminates or obtains more work. The kernel invokes p thread blocks, each of which iteratively consumes a segment of n/p elements in tile-size increments.
|
| GRID_MAPPING_STRIP_MINE | An a "strip mining" access pattern in which the input tiles assigned to each thread block are separated by a stride equal to the the extent of the grid.
- Overview
- The input is evenly partitioned into
p sets, where p is constant and corresponds loosely to the number of thread blocks that may actively reside on the target device. Each set is comprised of data tiles separated by stride tiles, where a tile is a small, constant-sized unit of input to be processed to completion before the thread block terminates or obtains more work. The kernel invokes p thread blocks, each of which iteratively consumes a segment of n/p elements in tile-size increments.
|
| GRID_MAPPING_DYNAMIC | A dynamic "queue-based" strategy for assigning input tiles to thread blocks.
- Overview
- The input is treated as a queue to be dynamically consumed by a grid of thread blocks. Work is atomically dequeued in tiles, where a tile is a unit of input to be processed to completion before the thread block terminates or obtains more work. The grid size
p is constant, loosely corresponding to the number of thread blocks that may actively reside on the target device.
|