워프, 레인과 마스크
32개 스레드 단위와 활성 마스크를 알아봅니다.
워프, 레인과 마스크은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 1번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 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/7 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
“워프, 레인과 마스크”에서 뭘 배우나요?
32개 스레드 단위와 활성 마스크를 알아봅니다. 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.
CUDA Academy을(를) 시작하는 데 경험이 필요한가요?
사전 경험은 필요하지 않습니다. CoddyKit의 CUDA Academy은(는) 초급자부터 고급 학습자까지를 위해 구성되어 있으므로, 여기서 시작하거나 처음부터 시작할 수 있으며 자신의 속도대로 진행할 수 있습니다. 이것은 4개 중 1번째 강의입니다.
“워프, 레인과 마스크” 강의는 얼마나 걸리나요?
대부분의 CoddyKit 강의는 약 5~10분이 소요됩니다. 각 강의는 간결하고 인터랙티브하여 꾸준한 진행이 가능하며, 웹과 앱에서 중단한 부분부터 바로 시작할 수 있습니다.
이 CUDA Academy 강의에서 코드를 작성하고 실행할 수 있나요?
네. 모든 CUDA Academy 강의에는 내장 코드 에디터가 포함되어 있으므로, 브라우저에서 바로 실제 코드를 작성하고 실행한 후 즉시 AI 피드백을 받을 수 있습니다 — 로컬 설정이 필요 없습니다.
이 강의의 모든 강의
- 워프, 레인과 마스크
- 축소를 위한 __shfl_down_sync
- 투표와 선택 함수
- 협력 그룹