What is the CUDA execution configuration?

The <<<blocks, threads, sharedBytes, stream>>> syntax between a kernel's name and its arguments.

Four parameters, and most code writes two. The first is the grid shape, the second is the block shape, and both are dim3 values that accept a plain integer for the one-dimensional case. The third is dynamic shared memory in bytes per block, which defaults to 0 and which day 13 is the first to use. The fourth is the stream the launch is queued on, which defaults to the null stream and which day 51 uses to overlap two kernels. The guide writes the full form as <<<grid_dim, block_dim, dynamic_smem_bytes, stream>>> in section 5.4.3.

Picking the first two is arithmetic, not taste. Threads per block should be a multiple of 32, because the SM issues in warps of 32 and a partial warp still costs a whole warp slot. Blocks come from rounding up: (n + threadsPerBlock - 1) / threadsPerBlock, never n / threadsPerBlock, which rounds a small n to zero blocks and launches nothing. Rounding up always launches more threads than you have elements, so every kernel written this way needs a bounds check on its global thread index before it touches memory.

Two failures come out of this line and both are loud. Threads per block above the device maximum of 1024, usually from a 2D block like dim3(32, 32, 2), is rejected at the launch with invalid configuration argument; so is a zero from that integer division. A launch returns void, so the error surfaces at the next cudaGetLastError() rather than at the launch itself, which is why the check goes on the line after. See invalid configuration argument for the full list of causes.

Measured

On a Tesla T4 (driver 595.84, CUDA 12.6, built with nvcc -O3 -arch=sm_75), day 5 added two vectors of n = 611 with 256 threads per block. The ceiling division gives 3 blocks, so 768 threads start, 157 of them find no element and return at the bounds check, and all 611 results match the CPU reference. The 157 is the point: an execution configuration that fits n exactly is the exception, not the rule.

Block shape decides how much of a warp you waste, and day 2 measured that on the same card. The same 128 threads as 8 blocks of 16 take 8 warp slots and leave 128 of those warps' 256 lanes idle; as 4 blocks of 32 they take 4 warp slots with no idle lanes. Day 1 launches <<<2, 4>>> for readability alone, and pays 28 idle lanes per block for it. Captured 2026-08-30; transcripts in code/day05-vector-add/evidence/run-2026-08-30.txt and code/day02-gpu-vs-cpu/evidence/run-2026-08-30.txt.

Code

The whole configuration, three slices of code/day05-vector-add/vector_add.cu.

constexpr int kThreadsPerBlock = 256;  // 8 warps
// ...
    const int blocks =
        static_cast<int>((kElems + kThreadsPerBlock - 1) / kThreadsPerBlock);
    const size_t launched = static_cast<size_t>(blocks) * kThreadsPerBlock;
// ...
    vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);

launched exists so the program can print how many threads have nothing to do. A configuration that hides that number is where off-by-one bugs live.

Related terms

• kernel
• grid
• thread block
• CUDA stream
• shared memory
• global thread index

Where you meet this

• Day 1, your first CUDA kernel, which owns this term and launches <<<2, 4>>>
• Day 4, grid, block and thread indexing, where the numbers get chosen properly
• Day 5, vector addition, which owns the measurement above
• Day 10, choosing threads per block, the sweep that says which block size to pick
• invalid configuration argument, when the launch is refused

Sources

• CUDA Programming Guide, C++ language extensions, section 5.4.3 on the execution configuration and its four parameters: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-29)
• CUDA Programming Guide, compute capabilities appendix, for the per-architecture maxima a configuration has to stay under: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-29)

Byline

Written by: unassigned. Reviewed by: unassigned. This entry is a draft and cannot publish until both are named people, and two different ones. Written on: not set. Last checked: not set. Numbers captured 2026-08-30 on the project's verification node.