Kernels, grids, and indexing
Write and test a bounds-safe grid-stride kernel whose ownership proof covers zero, tail, and very large sizes
Before you begin
Follow a linked prerequisite when you cannot explain it from memory. Background skills without links are stated as abilities rather than hidden terminology.
- C++ functions, pointers, loops, and array bounds
- Integer division and remainder
- Chapter 4 definitions of CUDA thread, block, grid, and warp
A kernel does not receive an array element automatically. Every thread sees its block and thread indices, and the program must use those coordinates to select logical work. This mapping is part of the algorithm. If two threads believe they own the same output, they can race. If no thread owns an output, it remains unchanged. If neighboring lanes map to distant addresses, the program can be correct and still transfer far more data than necessary.
Begin with a one-dimensional vector. The expression blockIdx.x * blockDim.x + threadIdx.x gives a unique starting index for every thread in a one-dimensional grid. A bounds check protects the final partial block. This direct mapping is sufficient when the grid can contain at least one thread for every element.
A grid-stride loop extends the same idea. After processing its first element, a thread advances by blockDim.x * gridDim.x, which is the total number of launched threads. The launch can now be chosen for the machine while the loop covers any input size. This pattern also makes it easy to test the ownership rule: each valid index has one starting residue modulo the grid stride.
Prove the ownership mapping
For an output with n elements, a correct mapping must satisfy coverage and uniqueness. Coverage means that every index from 0 through n - 1 is processed. Uniqueness means that exactly one thread writes each output unless the algorithm explicitly combines concurrent updates. Write these properties before considering block size or occupancy.
For the grid-stride loop, the starting indices are 0 through S - 1, where S is the total number of launched threads. Thread t visits t, t + S, t + 2S, and so on. Every nonnegative index has exactly one remainder modulo S, so it belongs to exactly one thread. The condition i < n removes values outside the array. This short proof is more reliable than checking a few convenient sizes.
Multidimensional mappings need the same treatment. If a row-major tensor has shape R × C, then logical coordinate (r, c) maps to r * row_stride + c * column_stride. Contiguous storage usually has column_stride = 1. A transposed view can have different strides without moving data. Never flatten a view unless the stride contract permits it.
Map warp lanes to addresses
After correctness, write the addresses used by lanes 0 through 31 for one load or store instruction. When threadIdx.x changes across the contiguous dimension, those addresses are usually adjacent. The memory system can then service them with a small number of aligned transactions. If threadIdx.x selects rows with a large row stride, each lane may require a different transaction.
The block should own the region whose threads must cooperate. A matrix tile that is staged in shared memory belongs naturally to one block because its threads can synchronize. Independent output tiles can become separate blocks. This rule connects indexing to the synchronization scope instead of treating grid coordinates as arbitrary dimensions.
Separate block geometry from grid size
Block shape determines warp composition, shared-memory cooperation, per-block resource reservation, and which threads can use a block barrier. Grid size determines how many blocks are available to distribute across the device and how much work each block repeats. A block with 256 threads and a grid with 256 blocks describes two different scales; combining them into one occupancy intuition causes errors.
A very small grid can leave SMs without blocks. A large grid is normally processed in waves, but blocks that do very little work can make launch and setup overhead significant. Grid-stride loops allow a moderate grid to process a larger vector, but too few blocks can still reduce parallelism. Treat the number of blocks and the work per block as separate experimental variables.
Tails also require a deliberate policy. A bounds predicate around one final load is often inexpensive. Deeply tiled kernels can accumulate predicates in inner loops, so padding or a specialized aligned path may be faster for common shapes. Use wide enough integer types for byte offsets; multiplying a valid 32-bit element index by a large stride can overflow before pointer arithmetic occurs.
Worked example: Mapping a row-wise operation
A tensor has R rows and C contiguous columns. Each output element applies an independent transform, and C is not necessarily a multiple of the block width.
- Map threadIdx.x across columns so neighboring lanes read neighboring values. Let blockIdx.x identify a row or a tile of a long row, depending on C.
- Use a bounds predicate for the final partial tile. Test C smaller than a warp, equal to a block, one greater than a block, and a large non-multiple.
- If R is small and C is huge, give each row several blocks or use a two-dimensional grid; otherwise the number of rows may not expose enough blocks.
- Inspect memory transactions and compare alternative block widths. The best width balances coalescing, launch parallelism, and any per-block resource use rather than obeying a universal constant.
Conclusion: The row and column mapping first establishes unique ownership, then assigns adjacent lanes to adjacent columns. Boundary tests validate coverage and uniqueness. Transaction and residency measurements determine whether the chosen block and grid shapes also use the hardware effectively.
Common errors
- Flattening a non-contiguous tensor as though it were contiguous. Logical adjacency and physical stride must not be confused.
- Choosing block size from occupancy alone. Memory pattern, cooperation, instruction structure, and shape determine the useful geometry.
- Testing only dimensions divisible by the tile. Most indexing failures and slow cleanup paths live at the boundaries.
Reference summary
__global__ void saxpy(int n, float a,
const float* x, float* y) {
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < n;
i += blockDim.x * gridDim.x) {
y[i] = a * x[i] + y[i];
}
}The first index gives each thread a unique starting element. The increment equals the total number of launched threads, so later loop iterations cover the remaining elements without overlap. The condition i < n protects both a partial final block and later grid-stride iterations.
Correctness requires coverage and unique ownership. Performance usually requires adjacent lanes to access adjacent elements and enough blocks to distribute work across the SMs. For multidimensional tensors, calculate addresses from actual strides rather than assuming contiguous storage.
- Start with a block size that is a multiple of the warp width; benchmark plausible alternatives.
- Launch enough blocks to occupy the device, but do not equate maximum occupancy with maximum performance.
- Use 64-bit indexing when tensor sizes or byte offsets can exceed 32-bit range.
- Test n=0, n<blockDim, non-multiple tails, and very large n.
Knowledge check
Why does mapping consecutive threads to a large stride often increase global-memory transactions?
Show the answer guide
- Consecutive warp lanes then request addresses separated by the large stride instead of adjacent words, so their requested bytes occupy many distinct memory sectors or cache lines.
- The memory system transfers the containing sectors, including bytes that no lane requested, which increases physical traffic relative to useful payload.
- A correct diagnosis compares requested bytes, transferred bytes or sectors, and elapsed time for the same mapping; the source-level element count alone cannot show transaction efficiency.
Verify indexing and access stride
Run a grid-stride copy or SAXPY over tail sizes and compare contiguous lane access with strides of 2, 8, and 32 elements. Validate every output before profiling.
Evidence to keep: Correctness results for zero and tail sizes, address formulas, raw timings, and transferred-sector or global-load-efficiency evidence across strides.
A kernel launch defines the available work. Register capacity, shared-memory capacity, and architectural limits determine the resident work on each SM.
Sources and further reading
The lesson explanations synthesize primary papers and official references. Follow the originals for proofs, exact APIs, architecture details, and version-sensitive behavior.