CUDA Academy · レッスン

__syncthreadsで同期する

スレッドの足並みをそろえるバリアを学びます。

「__syncthreadsで同期する」はCoddyKit上の無料CUDA Academyレッスンです。これはレッスン2/4です。下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応の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で同期する」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。

「__syncthreadsで同期する」で何を学びますか？

スレッドの足並みをそろえるバリアを学びます。ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。

CUDA Academyを始めるのに経験は必要ですか？

事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。これはレッスン2/4です。

「__syncthreadsで同期する」レッスンにはどのくらい時間がかかりますか？

ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。

このCUDA Academyレッスンでコードを書いて実行できますか？

はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。

このコースのすべてのレッスン

← CUDA Academyに戻る