What is a grid in CUDA?

All the blocks one kernel launch creates, in one, two or three dimensions.

The grid is the only part of a launch that has to know how big your data is, and it does not know. <<<blocks, threads>>> starts exactly blocks * threads threads and has never seen your array. Whether those threads cover n elements is arithmetic you did on the host beforehand. That is the difference between a grid and a for loop, and it is where the first real CUDA bug usually lives: day 4 sized a grid with n / 256 for n = 2000, got 7 blocks, covered 1792 elements, and left 208 with no owner. Exit code zero, no error string, three quarters of the array correct.

There are two honest ways to size it. Size to the data, (n + threads - 1) / threads, and pair it with an if (i < n) guard because ceiling division always over-provides. Or size to the machine, pick a grid that fills the GPU once and let each thread walk the array in a grid-stride loop. Day 8 built the machine-sized grid on this card as 160 blocks of 256 threads, four blocks per SM across 40 SMs, and then showed what that costs without the loop: it is correct up to 40,960 elements and wrong from 40,961, first failure at element 40,960. A fixed grid is not a covering strategy on its own.

The published limits are generous enough that they are almost never what stops you. Table 30 gives a maximum x-dimension of 2^31 - 1 blocks and 65,535 in y and z, on every compute capability from 7.5 through 12.0, so the seven reference GPUs this course names all carry the same grid maxima and a per-GPU table would be seven identical rows. The ceiling you actually hit is a block's, not the grid's: 1024 threads. Zero is the other end and it is not free. A grid sized to the data with n = 0 asks for zero blocks, and day 8's run had the driver refuse it with invalid configuration argument, which is cudaErrorInvalidConfiguration, code 9, reported synchronously at the launch.

Two dimensions change the arithmetic and nothing else. A 2D grid gives each block a blockIdx.y as well as an .x, which is convenient for images, and the built-in index variables still give you nothing but coordinates. Blocks in a grid run in no defined order and cannot wait for each other, which is why a grid-wide barrier needs cooperative groups and a launch API that checks the grid fits on the device first.

Measured

One card behind both tables: a Tesla T4 on driver 595.84 with CUDA 12.6 (V12.6.85), compiled nvcc -O3 -arch=sm_75, run 2026-08-30 on the project's verification node.

Day 4, 2000 elements, three grids over the same kernel:

Launch	Threads started	Unwritten	Idle
`<<<7, 256>>>`, truncating division	1792	208	0
`<<<8, 256>>>`, ceiling division	2048	0	48
`<<<2, 1024>>>`, ceiling division	2048	0	48

Day 8 compared three grid strategies across ten sizes on the same card, with a device grid of 160 blocks x 256 threads = 40,960 threads:

n	Grid sized to n	Grid sized to the GPU, no loop	Grid-stride loop
0	refused	ok	ok
611	ok	ok	ok
40,960	ok	ok	ok
40,961	ok	wrong at element 40,960	ok
4,194,304	ok	wrong at element 40,960	ok

The stride-loop column also held under <<<1, 1>>>, <<<1, 256>>>, <<<160, 256>>> and <<<10000, 32>>> at n = 611, which is the property that makes it worth the extra line: the launch configuration stops being part of the correctness argument.

Code

From code/day04-indexing/indexing.cu, the two grid sizes that produced the first table. One + kThreadsPerBlock - 1 between them, 208 elements between their results.

    // Integer division truncates, so this grid is one block short whenever
    // kElems is not a multiple of the block size. Nothing warns you. The
    // launch succeeds, the kernel returns, and the tail of the array keeps
    // whatever was in it before.
    const int shortBlocks = static_cast<int>(kElems / kThreadsPerBlock);

    // Ceiling division. One more block than you strictly need, and the guard
    // inside the kernel switches off the threads that block over-provides.
    const int blocks =
        static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);

Diagram

Original SVG: a row of eight block boxes over a 2000-cell strip, with the eighth box drawn half outside the strip. The seven-block grid is drawn above it, stopping short, with the 208 uncovered cells shaded.

Alt text: "Two grids over the same 2000-element array. Seven blocks of 256 threads stop at element 1791 and leave 208 cells with no block above them. Eight blocks cover every cell and hang 48 threads off the end, which the bounds check switches off."

Related terms

Where you meet this

Day 4, grid, block and thread indexing, the owner of this term
Day 7, two-dimensional grids, where blockIdx.y starts carrying a row number
Day 8, bounds checks and grid-stride loops, which produced the second table
Day 10, choosing threads per block, for the other half of the launch
invalid configuration argument, what a grid of zero blocks returns

Sources

CUDA Programming Guide, Table 30: maximum grid x-dimension 2^31 - 1, maximum y and z 65,535, maximum dimensionality 3, identical across compute capabilities 7.x to 12.x: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
CUDA Programming Guide, on block-to-SM assignment being outside the application's control: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)
CUDA Programming Guide, built-in variables, where gridDim is documented as dim3: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
CUDA Runtime API, the cudaError enum, where cudaErrorInvalidConfiguration = 9: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html (checked 2026-08-29)

Byline

Written by: unassigned. Reviewed by: unassigned. Written on: not set. Last checked: not set. Numbers captured 2026-08-30 on the project's verification node. This entry stays a draft until a named author and a different named reviewer sign it.