CUDA Academy · 课时

使用 __syncthreads 同步

让线程保持同步的屏障

使用 __syncthreads 同步 是 CoddyKit 上的免费 CUDA Academy 课时。这是第 2 节课，共 4 节。你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

本课时的部分内容尚未翻译，以英文显示。

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. 🎯

常见问题解答

「使用 __syncthreads 同步」课时是免费的吗？

是的 — 「使用 __syncthreads 同步」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。

「使用 __syncthreads 同步」这节课中我会学到什么？

让线程保持同步的屏障你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗？

无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。这是第 2 节课，共 4 节。

「使用 __syncthreads 同步」课时需要多长时间？

大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。

我能在这节 CUDA Academy 课中编写并运行代码吗？

能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。

此课程中的所有课时

← 返回 CUDA Academy