Synchronizing with __syncthreads
Barriers that keep threads in step.
Synchronizing with __syncthreads 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.
Threads Run Out of Step
Threads in a block do not march in lockstep. One may finish writing while another is still loading, so you cannot assume the data is ready yet.
Meet the Barrier
A barrier is a line in the code where every thread must wait until all of them arrive. Only then does the block move on together.
The __syncthreads Call
CUDA gives you __syncthreads() as the block-wide barrier. Call it and every thread in the block pauses until the last one shows up. 🛑
__syncthreads();The Load-Sync-Use Order
The golden pattern is simple: every thread loads its piece into shared memory, you sync, then everyone safely reads what others wrote.
tile[threadIdx.x] = input[i];
__syncthreads();
float left = tile[threadIdx.x - 1];Skip It and You Get Garbage
Forget the barrier and a thread may read a slot before its neighbor wrote it. That is a race condition, and the result is silent, wrong data.
It Is a Memory Fence Too
The barrier also makes shared writes visible to all threads after it. So once you pass __syncthreads, everyone sees the freshly written values.
All or None Must Reach It
The strict rule: every thread in the block must hit the same __syncthreads. If some skip it, the block can hang forever.
The Divergent Branch Trap
Never put a barrier inside an if that only some threads enter. Threads taking the other path never arrive, and the block deadlocks. ⚠️
if (threadIdx.x < 64) {
__syncthreads();
}Block Scope, Not Grid Scope
One important limit: __syncthreads only synchronizes one block. It cannot coordinate threads across different blocks of the grid.
Syncing Inside Loops
In tiled algorithms you often sync twice per phase: once after loading a tile and once after computing, before loading the next.
Cheap but Not Free
A barrier costs a little time while threads wait. Use it where correctness needs it, but avoid extra calls that just stall fast threads.
Quick Check
Let us test your grasp of the barrier.
Recap
You learned that __syncthreads() is a block-wide barrier that keeps threads in step, and that placing it inside divergent branches can deadlock. Next: bank conflicts. 🎯
Frequently asked questions
Is the “Synchronizing with __syncthreads” lesson free?
Yes — the full text of “Synchronizing with __syncthreads” 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 “Synchronizing with __syncthreads”?
Barriers that keep threads in step. 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 “Synchronizing with __syncthreads” 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
- Declaring __shared__ Arrays
- Synchronizing with __syncthreads
- Avoiding Bank Conflicts
- Dynamic Shared Memory