cub::GridMappingStrategy

enum cub::GridMappingStrategy

cub::GridMappingStrategy enumerates alternative strategies for mapping constant-sized tiles of device-wide data onto a grid of CUDA thread blocks.

Values:

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.

enumerator 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.

enumerator 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.