Bounds checks and CUDA grid-stride loops
Day 5's kernel requires gridDim.x * blockDim.x >= n. C does not express or check that requirement.
NVIDIA's post on grid-stride loops launches its example like this (https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/ , checked 2026-08-30):
saxpy<<<32*numSMs, 256>>>(1 << 20, 2.0, x, y);
That grid is sized to the GPU, not the array. It works because the kernel loops over more than one element per thread.
Use the same launch with a one-thread-per-element kernel, and only the first 32 * numSMs * 256 elements of the 1,048,576-element array get results. Both cudaGetLastError() and cudaDeviceSynchronize() report success because every executed operation was valid. The process exits 0 with the rest of the output unchanged.
This page rewrites that kernel to work with any valid launch configuration, including <<<1, 1>>>, and tests n = 0.
The launch configuration is an argument your kernel never declared
Day 4 gave you the global thread index and day 5 put a guard under it:
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = a[i] + b[i];
}
That guard stops extra threads from writing past the buffer. For 611 elements and 256 threads per block, the launch starts 768 threads. The last 157 find i >= n and do not write.
The guard cannot detect elements that no thread owns. If a grid has fewer threads than the array has elements, each running thread makes a valid write while later elements remain untouched.
CUDA has no fault to report.
The grid-stride loop turns the precondition into a loop condition:
for (size_t i = blockIdx.x * blockDim.x + threadIdx.x;
i < n; i += blockDim.x * gridDim.x) {
Let T be the total number of threads in the grid. Thread t handles elements t, t + T, t + 2T, and so on until the index reaches n.
Thread i % T writes element i on pass i / T, exactly once for any T >= 1. The loop condition is the bounds check.
The loop also keeps the warp's memory access pattern. On each pass, thread t and thread t + 1 access adjacent elements, so a warp's 32 lanes cover 128 consecutive bytes.
Each pass repeats this coalesced access. Day 11 measures it.
Where the ceiling actually is
One common reason for the loop is that N can exceed the largest valid grid. In one dimension, that limit is 2^31 - 1 blocks in x on each architecture this course targets (Table 30, https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-30).
At 256 threads per block, that grid covers about 5.5 * 10^11 elements, or 2.2 TB of floats. Current GPUs cannot hold that array in device memory.
The y and z grid limits are easier to reach. Both stop at 65,535, so a 2D launch with one block per output row fails at 65,536 rows. Day 7 introduces 2D grids.
The arithmetic in front of the launch has a ceiling of its own. The form every tutorial prints, int blocks = (n + 255) / 256; with an int n, overflows at about 2.1 * 10^9 elements, which is 8.6 GB of floats and fits on an A100 80GB. The count goes negative before the launch ever sees it.
The loop's main benefit is that kernel correctness no longer depends on one grid size. Days 10 and 11 use that property to test launch choices and memory access.
The tidy fix that is the wrong one
Another approach gives each thread a consecutive chunk. Thread t handles elements [t * C, (t + 1) * C), so the chunks cover the array.
const size_t base = t * chunkSize;
for (int k = 0; k < chunkSize; ++k) {
out[base + k] = in[base + k];
}
This pattern suits a CPU because each core works through one region in its private cache. On a GPU, the 32 lanes of a warp execute the same k at once, so their addresses are chunkSize elements apart. That creates a strided access.
Day 11's recorded run on driver 595.84 measured 232.9 GB/s for coalesced access and 32.6 GB/s for one consecutive chunk per thread, with the same byte count. A grid-stride loop uses the coalesced pattern on each pass.
The chunked version gives the right answer but moved the bytes about eight times slower in that run.
One kernel, three columns, and a prediction
The full program is in code/day08-grid-stride/grid_stride.cu. It runs the same vector add three ways at ten sizes and checks each result.
The two kernels differ by a loop and nothing else. Same arithmetic, same buffers, same block size, so nothing but the mapping from threads to elements can move a result.
__global__ void addOnePerElement(const float* a, const float* b, float* out,
size_t n) {
const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
if (i < n) {
out[i] = a[i] + b[i];
}
}
__global__ void addGridStride(const float* a, const float* b, float* out,
size_t n) {
const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
i < n; i += step) {
out[i] = a[i] + b[i];
}
}
The grid comes from the driver, not from n. Ask how many blocks of this kernel fit on one SM, multiply by the SM count, and you have a grid that covers the machine once. n appears nowhere in it, which is the point: one grid serves all ten rows.
int blocksPerSm = 0;
CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&blocksPerSm, addGridStride, kThreadsPerBlock, 0));
const int deviceBlocks = blocksPerSm * prop.multiProcessorCount;
const size_t totalThreads =
static_cast<size_t>(deviceBlocks) * kThreadsPerBlock;
The output is filled with NaN before every launch, not zero. An element no thread wrote has to fail the comparison. Fill with zero and a kernel that skips element 0 still matches, because the reference value there is also zero.
Every check is a branch that returns EXIT_FAILURE, and one of them is a prediction. One thread per element on a device-sized grid should be right exactly while n <= totalThreads and wrong after, and when it is wrong the first wrong element should be totalThreads itself, because that is the first element no thread owns. The program fails if the hardware disagrees in either direction.
This program tests correctness under several launch configurations but measures no time. The bandwidth figures above come from day 11's run.
Results
Re-verified on a Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88)
on 2026-09-02. The original CUDA 12.6 transcript remains in the page's
evidence array. All correctness behavior held; only the zero-block error
GPU: Tesla T4 (compute capability 7.5), 40 SMs
device grid: 160 blocks x 256 threads = 40960 threads (4 blocks per SM x 40 SMs)
A = one thread per element, grid sized to n (day 5)
B = one thread per element, grid sized to the GPU
C = grid-stride loop, grid sized to the GPU
n blocks(A) A B C
0 0 refused ok ok
1 1 ok ok ok
31 1 ok ok ok
32 1 ok ok ok
611 3 ok ok ok
1024 4 ok ok ok
1025 5 ok ok ok
40960 160 ok ok ok
40961 161 ok wrong@40960 ok
4194304 16384 ok wrong@40960 ok
A at n = 0 asks for a grid of 0 blocks. This driver refused it: invalid argument
C alone, n = 611, four launch configurations:
<<< 1, 1>>> ok
<<< 1, 256>>> ok
<<< 160, 256>>> ok
<<<10000, 32>>> ok
all 10 sizes and 4 launch configurations agree with the CPU reference where the model says they should
The prediction held on all three counts.
C, the grid-stride loop, is correct in all ten sizes, including n = 0 and
n = 1. It also passes all four launch configurations, from <<<1, 1>>> to
<<<10000, 32>>>. Day 10 relies on this result
when it sweeps block sizes without changing the kernel.
A, one thread per element with the grid sized to n, refuses to launch at
n = 0. A grid of zero blocks is not a legal launch. CUDA 13.0 returned
invalid argument; the preserved CUDA 12.6 run returned invalid configuration argument.
The behavior did not change, but code must not rely on one version's error string. Tests often miss the empty case.
B, one thread per element with the grid sized to the GPU, goes wrong at n = 40961. This device holds 40,960 threads in its 160-block grid, so the first element past that is never written. The first wrong element is at index 40960, exactly where the model says it should be.
Run it yourself
Compiler Explorer needs no local GPU, account, or install. Target sm_75 or lower for its current runner; a higher target can compile but fail on the device.
The largest case uses three 16 MiB device buffers and the full run launches 34 kernels. The one-file program fits within the runner's 20-second compile and run limits.
A Colab CUDA runtime or a local CUDA GPU works the same way:
nvcc -std=c++17 -O3 -arch=sm_75 -o grid_stride grid_stride.cu
Exercise
Rewrite day 5's solve() so it is correct for any n under any grid, then run it at the odd sizes and say in one sentence which element the monolithic version stopped writing and why that index and not another.
Time: 25 to 35 minutes. Submit: your grid_stride.cu, plus that one sentence.
Check: the harness owns the buffers, stream, and main(). It runs your function on nine sizes, from 0 and 1 through 611, 1,024, 1,025, and 4,194,304.
It reports the smallest failing case and the first ten wrong indices with their block, thread, and grid-stride pass. If the kernel stops after T elements, the report says you processed the first T elements only. This needs a grid-stride loop, or more blocks. Full contract at /reference/harness.
Hint 1
Your kernel is handed a grid it did not choose and cannot see the size of in advance. When there are fewer threads than elements, which elements does one thread own, and what decides?
Hint 2
Thread t owns element t. Which element does it own second? Pick the smallest step that still leaves lane 0 and lane 1 of a warp one element apart when they take it.
Solution
The solution needs three lines:
const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
i < n; i += step) {
out[i] = a[i] + b[i];
}
The diff against the starter deletes if (i < n) { and its closing brace and replaces them with the for, because the loop condition is that bounds check. Nothing else moves. The full file is code/day08-grid-stride/grid_stride.cu; its measured transcript on a named GPU lands in the Results section above when the program runs, and this exercise is graded on correctness rather than on a number.
The monolithic version stops writing at index gridDim.x * blockDim.x, the total thread count, because that is the first element with no thread behind it. Every element below it had one, so every write that happened was right, which is why nothing complains.
A monolithic kernel depends on its launch configuration for coverage, while a grid-stride kernel does not. Day 10 can therefore test several block sizes without changing the kernel.
Pitfalls
Your kernel passes at 1,024 elements and fails at 4,194,304. The grid was sized once, or hard coded, and it covers the small case. The report calls it a short grid: every wrong index is at or above the total thread count. Size the grid to the device and loop, or recompute the grid from n on every call, but do not do half of each.
Stepping by blockDim.x instead of blockDim.x * gridDim.x. Every block then walks the whole array from its own starting point, so element i is written by as many threads as there are blocks below it. A plain map survives that, because every one of those threads writes the same value, and your test passes while the kernel does the work several times over. Anything that accumulates into out[i] is wrong, and it is wrong by a factor that changes with the grid.
i += blockDim.x * gridDim.x with no cast. Both are unsigned int (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html#built-in-variables , checked 2026-08-30), so the product is computed in 32 bits and wraps above 2^32 threads. It is the same trap as the index line on day 5, it is silent, and the wrong answer lands in the middle of the array.
Writing the guard as if (i >= n) return; at all. In the monolithic kernel, it behaves like if (i < n) { ... } until a later edit adds a barrier. On day 13, threads that return before __syncthreads() never reach the barrier, so the block hangs.
The grid-stride loop has no separate early return. Never return above a barrier.
n = 0 with a grid sized to n. (0 + 255) / 256 is 0 blocks, and a grid with nothing in it is not a launch. Guard it in host code, or launch a grid that does not depend on n and let the loop condition be false on its first test. The second option is why n = 0 is a row in the table above and not a special case in the source.
Assuming the loop is free. It adds a register and a compare per pass, and it hides the trip count from the compiler. When the grid does cover n the loop runs once and the cost is small, but "small" is a number nobody on this page has measured. Day 45 is where launch configuration stops being a correctness question and becomes a measured one.
Go deeper
- CUDA Programming Guide, compute-capability appendix, Table 30, for the grid dimension caps: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
- Mark Harris, "CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops", the post the launch line at the top of this page comes from: https://developer.nvidia.com/blog/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/ (checked 2026-08-30)
cuda-samples,cpp/0_Introduction/vectorAdd, NVIDIA's own monolithic version of the kernel this page rewrites: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Introduction/vectorAdd (checked 2026-08-30)- Programming Massively Parallel Processors, 4th edition, chapter 2, on the launch configuration and the boundary check for a length that is not a multiple of the block size: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0
Next
Day 9 times one kernel five ways. Day 10 then sweeps block sizes without changing the kernel and measures occupancy. Both finish module 1.
Day 61 returns to bounds checks with Compute Sanitizer, which finds overruns that normal test sizes miss.