Launch Geometry and the Arithmetic of Thread Indexing

A kernel launch passes three integers per dimension and nothing else that describes the work. The device derives every identifier the kernel will use from those integers by arithmetic, and the shape of that arithmetic constrains what sizes are legal, how the compiler generates the guard for out-of-range work-items, and how much of the machine a single block can occupy.

From three extents to a linear identifier

CUDA’s launch configuration is <<<gridDim, blockDim>>>, both of type dim3 and both three-dimensional. The device exposes the values to the kernel as special registers, which PTX reads as %tid.x and friends for the position within a block, %ntid for the block’s extents, %ctaid for the block’s position in the grid, and %nctaid for the grid’s extents. CUDA C wraps these in threadIdx, blockDim, blockIdx, and gridDim.

The conventional linearisation is row-major over the declared extents:

__global__ void scale(float *data, float k, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) data[i] *= k;
}

Each of the four special-register components is a 32-bit unsigned quantity as far as the arithmetic is concerned, but the declared extents are much smaller for the y and z dimensions than for x. blockDim.x and blockDim.y are each capped at 1024 and blockDim.z at 64, with the product of all three capped at 1024 as well, so a block of 1024 threads may be shaped (1024,1,1), (32,32,1), or (16,16,4) but not (1024,2,1). The grid dimensions are capped per component: gridDim.x may reach 2^31 - 1 while y and z stop at 65535, which makes a one-dimensional grid the only choice when the work-item count exceeds roughly two billion.

The mapping from a multi-dimensional index to the linearised form is left to the programmer, and this is deliberate. Nothing in the contract requires y to be the second-fastest varying dimension. A kernel working on a two-dimensional tile of a matrix can linearise however it likes, and the only constraint is that every work-item computes a distinct, correct address. The device does not care; the memory system very much does, because the order in which identifiers map onto addresses decides how many memory transactions a warp issues. The layout choice is where the index space meets the memory hierarchy, and the API offers no help with it.

Why the guard predicate exists

The guard if (i < n) appears in almost every real kernel, and its cost is not zero. Where n is a multiple of the block size the predicate is always true and the compiler removes it. Where n is not, the last block in the grid contains active lanes and inactive lanes, and a block of 256 threads processing 1000 elements leaves 24 lanes idle in its final block.

The idle lanes are harmless when the kernel is uniform. They become expensive when the kernel contains a barrier, because a barrier requires every work-item in the group to arrive and the inactive lanes must still reach the barrier instruction, which is why the guard has to be written so that the return path does not skip a __syncthreads(). Writing if (i < n) { ... __syncthreads(); ... } deadlocks or produces undefined behaviour depending on the hardware. The standard idiom hoists the barrier out of the conditional, or makes the out-of-range lanes participate without touching memory.

__global__ void reduce_partial(const float *in, float *out, int n) {
    __shared__ float tile[256];        // launch with exactly 256 threads per block
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // Out-of-range lanes contribute the identity, then take part in the barrier.
    tile[threadIdx.x] = (i < n) ? in[i] : 0.0f;
    __syncthreads();
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (threadIdx.x < s) tile[threadIdx.x] += tile[threadIdx.x + s];
        __syncthreads();
    }
    if (threadIdx.x == 0) out[blockIdx.x] = tile[0];
}

The alternative to a guard is to pad the grid, which shifts the problem to the allocation: the buffer must have gridDim.x * blockDim.x elements of real storage, and the extra reads must not fault. Padding is sometimes viable for read-only inputs; it is not for writes, which would clobber memory beyond the buffer.

Block size as a scheduling constraint

The block size is not a hint. It fixes the granularity at which the hardware allocates a multiprocessor’s resources, and that determines how many blocks can be resident at once.

A multiprocessor has a finite pool of registers and a finite partition of scratch memory, and the number of resident blocks is bounded by how much of each pool a single block consumes. A block of 1024 threads on a machine with 64K registers per multiprocessor cannot be resident alongside anything else if the kernel needs more than 64 registers per thread, because 1024 threads times 64 registers exhausts the pool exactly. A block of 128 threads with the same per-thread register count allows several blocks to share the multiprocessor.

The block size also sets the floor on wasted lanes. Threads are grouped into warps of 32 on NVIDIA hardware, and a block whose size is not a multiple of 32 leaves the trailing warp partially populated for the block’s entire lifetime. A block of 100 threads occupies four warps, three full and one with four active lanes: 28 of 128 lanes idle permanently, a 21 percent loss of issue capacity before any work is done. Sizes that are multiples of 32, and preferably powers of two, avoid the effect.

A launch whose block size exceeds these limits does not degrade; it fails, and the error is returned by the launch call rather than raised later. The set of limits a given device reports is one entry in the broader property enumeration discussed in Capability Negotiation Across Architectures. The hardware does not split a block across SMs, so a block that exceeds the per-SM limit in any resource cannot be scheduled at all and the launch fails rather than degrading. OpenCL exposes the corresponding bound as CL_DEVICE_MAX_WORK_GROUP_SIZE and the per-dimension limits as CL_DEVICE_MAX_WORK_ITEM_SIZES, both queried at runtime; the specification only guarantees a small floor, so a portable program that assumes 1024 is relying on a value the standard does not promise.

Persistent blocks and the grid-stride loop

Because the grid may be far larger than the number of resident blocks, and because launching a grid costs a driver round trip, the alternative pattern is to launch exactly enough blocks to fill the machine and let each one iterate.

__global__ void fill_stride(float *data, float k, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n;
         i += gridDim.x * blockDim.x) {
        data[i] = k;
    }
}

The launch is asynchronous in the usual way, so the loop’s completion is observed through the queue mechanisms covered in Command Queues, Events, and Cross-Stream Ordering. The stride between iterations equals the total number of threads launched, so consecutive lanes still touch consecutive addresses on every iteration and coalescing is preserved. The launch can be sized by a query of the device’s occupancy limits for this kernel rather than by the problem size, which decouples the two: the same launch configuration handles any n.

The cost is that the loop adds an induction variable and a second branch per trip, and that every lane executes the same number of iterations only when n divides evenly. Neither is significant beside the elimination of one launch per tile of work, which matters most when the kernels are short.

Limits that differ between APIs and change between generations

The numbers above are the CUDA limits from compute capability 2.0 onward, and they have been stable for a long time: 1024 threads per block, 2^31 - 1 blocks in x, 65535 in y and z. Devices of compute capability 1.x allowed 512 threads per block, so a launch sized for 1024 is rejected there, while a kernel sized for 512 remains valid on everything newer.

Vulkan and SPIR-V express the same structure with different vocabulary and looser bounds. The local size is declared per-dimension in the shader’s LocalSize execution mode or overridden at dispatch, and the total is bounded by the device’s maxComputeWorkGroupInvocations, which on desktop hardware is commonly 1024 but is a queried property, not a constant.

The shared lesson across all of them: block size and grid size are encoded in 32-bit unsigned fields, the individual caps are small enough to be reached by ordinary problem sizes, and the only safe sizing arithmetic multiplies in a 64-bit type before comparing. A grid computed as (n + 255) / 256 in int overflows for n above about 2^31, and the resulting negative block count is either rejected by the runtime or, worse, silently wrapped.