Grid-Stride Loops
Handling arrays larger than the grid.
Grid-Stride Loops is a free CUDA Academy lesson on CoddyKit — lesson 4 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
When the Array Is Huge
Sometimes the data is far bigger than the threads you launch. One thread per element no longer fits, so you need each thread to handle several elements.
The Grid Has a Size
The total number of threads is the grid size: blocks times threads per block. This stride is how far apart each thread's elements sit.
int stride = blockDim.x * gridDim.x;Start, Then Stride
Each thread begins at its usual global index, then jumps forward by the grid size again and again until it runs off the array.
The Grid-Stride Loop
This loop is the whole pattern: start at i, step by stride, stop at n. Any array size is covered no matter how many threads you launch.
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < n;
i += stride) {
out[i] = a[i] + b[i];
}The Built-in Guard
Notice the loop condition i < n is itself the bounds check. Threads that start past the end simply never enter the loop.
How the Work Splits
With a stride of 4 threads, thread 0 does elements 0, 4, 8, while thread 1 does 1, 5, 9. The work is interleaved, not chunked.
Interleaving Helps Coalescing
Because neighboring threads still touch neighboring addresses each step, the access stays coalesced and memory bandwidth stays high.
Decouple Threads From Data
Now your launch size no longer depends on n. You can pick a thread count that fits the GPU and let the loop absorb any workload.
Tune for the Hardware
A common choice is enough blocks to fill every multiprocessor, then let each thread loop. This keeps the GPU busy without overlaunching.
It Also Works When Tiny
If n is smaller than the grid, each thread runs the loop body at most once. The pattern degrades gracefully to the simple case.
A Robust Default
Many CUDA pros write every elementwise kernel as a grid-stride loop. It is flexible, safe, and rarely the wrong choice. 🚀
Quick Check
Identify the stride.
Recap
You learned the grid-stride loop: start at the global index and step by blockDim.x * gridDim.x until i reaches n. One kernel now handles any array size. 🎉
Frequently asked questions
Is the “Grid-Stride Loops” lesson free?
Yes — the full text of “Grid-Stride Loops” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Grid-Stride Loops”?
Handling arrays larger than the grid. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 4 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Grid-Stride Loops” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.