Coalesced vs Strided Reads
How thread-to-address mapping matters.
Coalesced vs Strided Reads is a free CUDA Academy lesson on CoddyKit — lesson 2 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.
It Is All About Mapping
Coalescing depends on which address each thread touches. The map from thread to address decides whether one transaction serves the whole warp.
The Coalesced Pattern
When neighboring threads read neighboring addresses, the warp covers one contiguous run. That tidy layout is a coalesced access.
int i = blockIdx.x * blockDim.x + threadIdx.x;
float v = data[i];Why It Wins
Thread 0 hits index 0, thread 1 hits index 1, and so on. All 32 addresses fall in one aligned line, so the GPU needs a single transaction.
Enter the Stride
A stride is a fixed gap between the addresses successive threads touch. The moment that gap grows past one element, coalescing starts to break.
float v = data[i * stride];Strided Reads Spread Out
With a stride of 2, thread 0 reads 0, thread 1 reads 2, thread 2 reads 4. The warp now spans twice the memory, so it needs more transactions.
The Cost Scales Up
Double the stride and you roughly double the lines touched. Big strides can push a warp toward many transactions while using only a sliver of each line.
A Common Trap
Giving each thread a whole column of a row-major matrix creates a huge stride. It looks neat in code but quietly scatters every warp.
float v = matrix[threadIdx.x * width + row];Transpose the Mapping
Often the fix is just swapping which index varies fastest with the thread. Let consecutive threads walk consecutive memory and the reads coalesce again.
float v = matrix[row * width + threadIdx.x];Offsets Hurt Too
Even a coalesced pattern suffers if it starts mid-line. A misaligned offset shifts the warp across a boundary, splitting one read into two.
Think in Warps
To judge a pattern, do not picture one thread. Picture all 32 lanes at once and ask how many lines their addresses cover together. 🔍
The Rule of Thumb
Make the fastest-changing index follow threadIdx.x. That single habit keeps most of your global reads coalesced for free.
Quick Check
Pick the access pattern that coalesces best.
Recap
You saw that coalesced reads keep neighbors together while a stride scatters them across lines. Keep threadIdx.x driving the fastest index. 🎉
Frequently asked questions
Is the “Coalesced vs Strided Reads” lesson free?
Yes — the full text of “Coalesced vs Strided Reads” 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 “Coalesced vs Strided Reads”?
How thread-to-address mapping matters. 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 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Coalesced vs Strided Reads” 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.
All lessons in this course
- What a Memory Transaction Is
- Coalesced vs Strided Reads
- Structure of Arrays vs Array of Structs
- Measuring Effective Bandwidth