ワープ、レーン、マスク
32スレッド単位とアクティブマスクを学びます。
「ワープ、レーン、マスク」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン1/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
このレッスンの一部はまだ翻訳されておらず、英語で表示されています。
Threads Run in Warps
The GPU does not schedule threads one by one. It groups them into a warp of 32 threads that march together, executing the same instruction in lockstep.
Why 32 Matters
On NVIDIA hardware a warp is always 32 threads. It is the real unit of execution, so good kernels think in groups of 32, not single threads.
Each Thread Is a Lane
Inside a warp, every thread has a position from 0 to 31 called its lane. The lane id is how warp primitives know who is talking to whom.
int lane = threadIdx.x % 32;Finding the Lane Id
You can read the lane directly from a special register instead of computing it. The laneid register always holds 0 through 31 for the current thread.
Lockstep Has a Catch
Threads share one program counter per warp. When an if sends lanes down different paths, the warp diverges and runs each path in turn, hurting speed.
The Active Mask
Not every lane is always running. A 32-bit mask marks which lanes are active right now, with one bit per lane set to 1 when that lane participates.
Why Masks Exist
After divergence only some lanes are live. Warp primitives need to know exactly who is present, so you pass them an active mask to stay correct.
The Full Mask
When you are sure all 32 lanes are active, the mask is 0xffffffff, every bit set. This full mask is the most common value you will pass.
unsigned mask = 0xffffffff;Building a Mask Safely
Inside a branch, do not guess the mask. Call activemask to capture exactly which lanes reached this point right now.
unsigned mask = __activemask();Lanes Talk Without Memory
The big win is that lanes in one warp can swap data directly through registers. No shared memory and no barriers are needed for warp-local exchange.
Sync Means the Warp
The sync suffix on these intrinsics is a warp-level handshake, not a block barrier. It only coordinates the lanes named in the mask you give it.
Quick Check
Recall what a warp is and how big it is on NVIDIA GPUs.
Recap
A warp is 32 lanes running in lockstep, and a mask tracks who is active. That foundation lets lanes share data fast. Next: shuffles for reductions. ✨
よくある質問
「ワープ、レーン、マスク」レッスンは無料ですか？
はい。「ワープ、レーン、マスク」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「ワープ、レーン、マスク」で何を学びますか？
32スレッド単位とアクティブマスクを学びます。 ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン1/4です。
「ワープ、レーン、マスク」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。