Class to represent a CUDA grid. More...

#include <GridTraits.hpp>

Static Public Member Functions
template<int I>
constexpr static __device__ int	blockDim ()
	Get the I'th block dimension, which may be a compile-time constant. More...

template<int I>
constexpr static __device__ int	gridDim ()
	Get the I'th grid dimension, which may be a compile-time constant. More...

constexpr static __device__ sp::Vec< int, BlockRank >	blockDims ()
	Get a `Vec` representing the block size. More...

constexpr static __device__ sp::Vec< int, GridRank >	gridDims ()
	Get a `Vec` representing the block size. More...

static __device__ sp::Vec< int, GridRank >	blockIdx ()
	Get the block index as a `Vec`. More...

template<int I>
static __device__ int	blockIdx ()
	Get the I'th block idx. More...

static __device__ sp::Vec< int, BlockRank >	threadIdx ()
	Get the thread index as a `Vec`. More...

template<int I>
static __device__ int	threadIdx ()
	Get the I'th thread idx. More...

static __device__ int	blockId ()
	Get the block ID. This can be thought of as the 1D block index. More...

static __device__ int	threadId ()
	Get the thread ID. This can be thought of as the 1D thread index within the block. More...

constexpr static __device__ int	numBlocks ()
	Total number of blocks. More...

constexpr static __device__ int	threadsPerBlock ()
	Number of threads per block. More...

constexpr static __device__ int	warpsPerBlock ()
	Warps per block. More...

constexpr static __device__ int	warpsInGrid ()
	Total number of warps in the grid. More...

constexpr static __device__ int	threadsInGrid ()
	Total number of threads in the grid. More...

static __device__ int	warpIdx ()
	Get the warp index, given the block dimensionality. More...

static __device__ int	warpGridIdx ()
	Index of this warp in grid. More...

static __device__ int	threadGridIdx ()
	Index of this thread in grid. More...

static __device__ int	laneIdx ()
	Get the lane index of this thread. More...

Static Public Attributes
constexpr static int	GridRank = GridSizes::Size

constexpr static int	BlockRank = BlockSizes::Size

Detailed Description

template<typename BlockSizes, typename GridSizes>
struct sp::Grid< BlockSizes, GridSizes >

Class to represent a CUDA grid.

Provides convenient accessors to blockIdx/blockDim/etc. using numerical indices.
Allows accessing blockIdx/blockDim/etc. as sp::Vec.
Allows you to specify assumptions about the range of values blockIdx/blockDim/etc. can take, allowing the compiler to constant-propagate them away.
Provides APIs for higher-level grid queries, like warpIdx().

When you declare your kernel's Grid type, you specify the dimensionality and any fixed size components of the block or grid dimensions. Those values can then be used in constexpr contexts when accessed through blockDim<I>() and friends. This also improves optimisation, since the compiler can't typically determine when some of these are constants.

Example

// Blocks are always 1D, and exactly 512 threads.
// Grids are always 1D, but the number of blocks is not known.
// The compiler will be able to constant-propagate gridDims.y/z as 1, blockDims.y/z as 1, and blockDims.x as 512.
// Any of those can be used as template parameters when accessed through `MyGrid`, as they will be constexpr.
using MyGrid = sp::Grid<sp::BlockDims<512>, sp::GridDims<-1>>;

Template Parameters

BlockSizes	An `sp::int_sequence<>` specifying the block size. Use `-1` for entries that are non-constant.
GridSizes	An `sp::int_sequence<>` specifying the grid size. Use `-1` for entries that are non-constant.