线程束、通道与掩码
了解由 32 个线程组成的执行单元与活跃掩码
线程束、通道与掩码 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 1 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 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. ✨
常见问题解答
「线程束、通道与掩码」课时是免费的吗?
是的 — 「线程束、通道与掩码」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「线程束、通道与掩码」这节课中我会学到什么?
了解由 32 个线程组成的执行单元与活跃掩码 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗?
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 1 节课,共 4 节。
「线程束、通道与掩码」课时需要多长时间?
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗?
能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。
此课程中的所有课时
- 线程束、通道与掩码
- 使用 __shfl_down_sync 进行归约
- 投票与表决函数
- 协作组