The Load-Sync-Compute Pattern
Staging tiles before computing on them.
The Load-Sync-Compute Pattern 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.
Three Simple Phases
Tiling follows one rhythm in every kernel: load a tile into shared memory, sync, then compute from the fast copy.
Phase One: Load
In the load phase, each thread reads one element from global memory and stores it into a shared-memory tile its whole block can see.
tile[threadIdx.x] = in[globalIndex];A Cooperative Effort
Loading is a team job. Every thread fetches its slice, so together the block stages one full tile with a single coalesced pass.
Phase Two: Sync
Before anyone reads a neighbor s value, all threads must finish loading. __syncthreads is the barrier that guarantees the tile is ready.
__syncthreads();Why Sync Is Mandatory
Skip the barrier and a thread may read a tile slot that its neighbor has not written yet. That is a race and gives wrong results. 💥
Phase Three: Compute
Now every value lives on chip. In the compute phase threads read tile entries freely, since shared memory is orders faster than global.
Fast Reuse Pays Off
Because the tile sits in shared memory, a value loaded once is reused by many threads with almost no extra cost. That is the whole win.
Sometimes Sync Again
If the compute phase writes back into the same tile for a next step, add a second __syncthreads before reusing those slots.
All Threads or None
Every thread in the block must reach __syncthreads. If some skip it inside a branch, the kernel deadlocks or misbehaves.
A Reusable Skeleton
Load, sync, compute is a template you will reuse for convolutions, matrix multiply, and reductions throughout this category.
Mind the Tile Size
The shared tile must fit the block. Tile width usually equals blockDim, so each thread owns exactly one slot to load and reuse.
Quick Check
Why is __syncthreads needed between load and compute?
Recap
Tiling is load, sync, compute: stage a tile cooperatively, barrier so it is complete, then reuse it fast on chip. ✅
Frequently asked questions
Is the “The Load-Sync-Compute Pattern” lesson free?
Yes — the full text of “The Load-Sync-Compute Pattern” 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 “The Load-Sync-Compute Pattern”?
Staging tiles before computing on them. 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 “The Load-Sync-Compute Pattern” 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
- The Data Reuse Problem
- The Load-Sync-Compute Pattern
- Stencil and Sliding Windows
- Handling Edge Tiles