Catching Sync Errors with synccheck
Divergent __syncthreads and hazards.
Catching Sync Errors with synccheck 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.
The Barrier Contract
A __syncthreads barrier only works if every thread in the block reaches it. Skip it for some threads and behavior turns undefined. ⚠️
Divergent Barriers
Calling a barrier inside a branch some threads skip is a divergent sync, the exact hazard the synccheck tool is built to catch.
Meet synccheck
Run the synccheck tool of compute-sanitizer to detect illegal or divergent uses of __syncthreads and related barriers.
compute-sanitizer --tool synccheck ./tileA Buggy Conditional Sync
Here only some threads call the barrier, leaving the rest behind. synccheck flags this as a divergent __syncthreads.
if (tid < half) {
__syncthreads();
}The Correct Pattern
Pull the barrier out of the branch so all threads reach it together, while only the guarded work stays conditional.
if (tid < half) {
sum += tile[tid];
}
__syncthreads();Barriers in Loops
Every thread must run the same number of barriers, so a loop with a thread-dependent trip count plus an inner sync is a trap.
Early Returns Are Risky
A thread that hits an early return before a barrier never arrives, so synccheck reports the remaining threads stuck waiting.
Reading the Report
synccheck names the barrier and the threads that diverged, so you can trace which branch or return broke the contract.
Warp-Level Syncs Too
synccheck also checks __syncwarp masks, flagging when a thread participates with a mask that does not match its lane.
Add Line Info
Compile with -lineinfo so synccheck can point at the exact barrier call instead of just naming the kernel.
nvcc -lineinfo tile.cu -o tileOne Tool per Run
Run synccheck on its own pass, separate from memcheck and racecheck, since each sanitizer tool targets a different class of bug.
Quick Check
What problem does the synccheck tool specifically detect?
Recap
You learned the barrier contract, ran synccheck, and fixed divergent syncs by moving barriers out of branches. Your block stays in step. 🙌
Frequently asked questions
Is the “Catching Sync Errors with synccheck” lesson free?
Yes — the full text of “Catching Sync Errors with synccheck” 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 “Catching Sync Errors with synccheck”?
Divergent __syncthreads and hazards. 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 “Catching Sync Errors with synccheck” 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
- Stepping Kernels in cuda-gdb
- Finding Leaks with memcheck
- Hunting Races with racecheck
- Catching Sync Errors with synccheck