Grid Assumptions

A system for telling the compiler about constant properties of your kernel grid. More...

Classes
struct	sp::Grid< BlockSizes, GridSizes >
	Class to represent a CUDA grid. More...
struct	sp::BlockTraits< Ds >
	Represents the dimensions of a CUDA block, in cases where this is constexpr. More...
struct	sp::GridTraits< B, Ds >
	Represents the dimensions of a CUDA grid, in cases where this is constexpr. More...

Typedefs
template<int... Ints>
using	sp::GridDims = sp::int_sequence< Ints... >
template<int... Ints>
using	sp::BlockDims = sp::int_sequence< Ints... >

Detailed Description

A system for telling the compiler about constant properties of your kernel grid.

Often, the programmer knows something about the thread or block dimensions that the compiler does not.

For example, you may always launch a 2D grid, so the third dimension is unused. Or perhaps one or more of the block or grid dimensions is fixed.

sp::Grid allows you to explicitly encode these assumptions. For example, if our kernel always launches blocks of size 32x32, and launches an unknown but always 1D quantity of blocks, we can write:

using MyGrid = sp::Grid<
    // -1 represents "this value is unknown". The number of elements given represents the number
    // of dimensions that are used. So the below means "blocks are always 32x32, and the grid is
    // always 1D".
    sp::BlockDims<32, 32>,
    sp::GridDims<-1>
>;

In our kernel, we now use MyGrid instead of accessing CUDA builtins like blockDim or threadIdx directly. Our grid object will use the assumptions we specified to optimise for us. Any dimensions that we specified as constants will be constexpr and may be used as template parameters!

__global__ void myKernel() {
    constexpr int BlockWidth = MyGrid::blockDim<0>();
    constexpr int BlockHeight = MyGrid::blockDim<1>();
 
    // Can also get it as a constexpr vec because it's all constant:
    sp::Vec<int, 2> BlockDims = MyGrid::blockDims();
 
    sp::Vec<float, BlockWidth> foo; // One element for each row of the grid.
 
    // A static tensor with one element for each thread
    sp::StaticSquareMatrix<float, BlockHeight, BlockWidth> bar;
 
    // Get the total number of warps in the grid.
    // This isn't a compile-time constant, but it will use the assumptions we specified about grid dimensionality
    // (and the constant block size) to optimise the calculation to be as efficient as possible.
    int totalWarps = MyGrid::warpsInGrid();
}

This allows you to write high-level code that uses thread and grid indexes generically without worrying about chewing up unnecessary math using dimensions that are zero.

See also

Grid