cub::GridEvenShare< OffsetT > Struct Template Reference

GridEvenShare is a descriptor utility for distributing input among CUDA thread blocks in an "even-share" fashion. Each thread block gets roughly the same number of input tiles. More...

Detailed Description

template<typename OffsetT>
struct cub::GridEvenShare< OffsetT >

GridEvenShare is a descriptor utility for distributing input among CUDA thread blocks in an "even-share" fashion. Each thread block gets roughly the same number of input tiles.

Overview

Each thread block is assigned a consecutive sequence of input tiles. To help preserve alignment and eliminate the overhead of guarded loads for all but the last thread block, to GridEvenShare assigns one of three different amounts of work to a given thread block: "big", "normal", or "last". The "big" workloads are one scheduling grain larger than "normal". The "last" work unit for the last thread block may be partially-full if the input is not an even multiple of the scheduling grain size.

Before invoking a child grid, a parent thread will typically construct an instance of GridEvenShare. The instance can be passed to child thread blocks which can initialize their per-thread block offsets using BlockInit().

Definition at line 74 of file grid_even_share.cuh.

Public Member Functions
__host__ __device__ __forceinline__	GridEvenShare ()
	Constructor.
__host__ __device__ __forceinline__ void	DispatchInit (OffsetT num_items, int max_grid_size, int tile_items)
	Dispatch initializer. To be called prior prior to kernel launch.
template<int TILE_ITEMS>
__device__ __forceinline__ void	BlockInit (int block_id, Int2Type< GRID_MAPPING_RAKE >)
	Initializes ranges for the specified thread block index. Specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.
template<int TILE_ITEMS>
__device__ __forceinline__ void	BlockInit (int block_id, Int2Type< GRID_MAPPING_STRIP_MINE >)
	Block-initialization, specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.
template<int TILE_ITEMS, GridMappingStrategy STRATEGY>
__device__ __forceinline__ void	BlockInit ()
	Block-initialization, specialized for "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.
template<int TILE_ITEMS>
__device__ __forceinline__ void	BlockInit (OffsetT block_offset, OffsetT block_end)
	Block-initialization, specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.

Data Fields
OffsetT	num_items
	Total number of input items.
int	grid_size
	Grid size in thread blocks.
OffsetT	block_offset
	OffsetT into input marking the beginning of the owning thread block's segment of input tiles.
OffsetT	block_end
	OffsetT into input of marking the end (one-past) of the owning thread block's segment of input tiles.
OffsetT	block_stride
	Stride between input tiles.

Private Attributes
OffsetT	total_tiles
int	big_shares
OffsetT	big_share_items
OffsetT	normal_share_items
OffsetT	normal_base_offset

Constructor & Destructor Documentation

GridEvenShare()

template<typename OffsetT >

__host__ __device__ __forceinline__ cub::GridEvenShare< OffsetT >::GridEvenShare ( )

inline

Constructor.

Definition at line 105 of file grid_even_share.cuh.

Member Function Documentation

BlockInit() [1/4]

template<typename OffsetT >

template<int TILE_ITEMS, GridMappingStrategy STRATEGY>

__device__ __forceinline__ void cub::GridEvenShare< OffsetT >::BlockInit ( )

inline

Block-initialization, specialized for "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.

Definition at line 191 of file grid_even_share.cuh.

BlockInit() [2/4]

template<typename OffsetT >

template<int TILE_ITEMS>

__device__ __forceinline__ void cub::GridEvenShare< OffsetT >::BlockInit ( int  block_id, Int2Type< GRID_MAPPING_RAKE >    )

inline

Initializes ranges for the specified thread block index. Specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.

Definition at line 146 of file grid_even_share.cuh.

BlockInit() [3/4]

template<typename OffsetT >

template<int TILE_ITEMS>

__device__ __forceinline__ void cub::GridEvenShare< OffsetT >::BlockInit ( int  block_id, Int2Type< GRID_MAPPING_STRIP_MINE >    )

inline

Block-initialization, specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.

Definition at line 173 of file grid_even_share.cuh.

BlockInit() [4/4]

template<typename OffsetT >

template<int TILE_ITEMS>

__device__ __forceinline__ void cub::GridEvenShare< OffsetT >::BlockInit ( OffsetT  block_offset, OffsetT  block_end  )

inline

Block-initialization, specialized for a "raking" access pattern in which each thread block is assigned a consecutive sequence of input tiles.

Parameters

[in]	block_offset	Threadblock begin offset (inclusive)
[in]	block_end	Threadblock end offset (exclusive)

Definition at line 203 of file grid_even_share.cuh.

DispatchInit()

template<typename OffsetT >

__host__ __device__ __forceinline__ void cub::GridEvenShare< OffsetT >::DispatchInit ( OffsetT  num_items, int  max_grid_size, int  tile_items  )

inline

Dispatch initializer. To be called prior prior to kernel launch.

Parameters

num_items	Total number of input items
max_grid_size	Maximum grid size allowable (actual grid size may be less if not warranted by the the number of input items)
tile_items	Number of data items per input tile

Definition at line 122 of file grid_even_share.cuh.

Field Documentation

big_share_items

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::big_share_items

private

Definition at line 80 of file grid_even_share.cuh.

big_shares

template<typename OffsetT >

int cub::GridEvenShare< OffsetT >::big_shares

private

Definition at line 79 of file grid_even_share.cuh.

block_end

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::block_end

OffsetT into input of marking the end (one-past) of the owning thread block's segment of input tiles.

Definition at line 96 of file grid_even_share.cuh.

block_offset

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::block_offset

OffsetT into input marking the beginning of the owning thread block's segment of input tiles.

Definition at line 93 of file grid_even_share.cuh.

block_stride

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::block_stride

Stride between input tiles.

Definition at line 99 of file grid_even_share.cuh.

grid_size

template<typename OffsetT >

int cub::GridEvenShare< OffsetT >::grid_size

Grid size in thread blocks.

Definition at line 90 of file grid_even_share.cuh.

normal_base_offset

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::normal_base_offset

private

Definition at line 82 of file grid_even_share.cuh.

normal_share_items

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::normal_share_items

private

Definition at line 81 of file grid_even_share.cuh.

num_items

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::num_items

Total number of input items.

Definition at line 87 of file grid_even_share.cuh.

total_tiles

template<typename OffsetT >

OffsetT cub::GridEvenShare< OffsetT >::total_tiles

private

Definition at line 78 of file grid_even_share.cuh.

---

The documentation for this struct was generated from the following file:

• openfpm_data/src/util/cuda/cub_old/grid/grid_even_share.cuh