워프 분기 없애기
워프가 계속 작업하도록 인덱스를 다시 지정합니다.
워프 분기 없애기은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 2번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 AI 튜터와 함께 브라우저에서 직접 실습할 수 있습니다. 이 강의는 CUDA Academy 학습 경로의 일부이며, 진행 상황이 웹과 CoddyKit 앱에 동기화됩니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
이 강의의 일부는 아직 번역되지 않았으며 영어로 표시됩니다.
Warps Run in Lockstep
A warp is 32 threads that execute the same instruction together. When their paths agree, the hardware runs at full speed.
What Divergence Costs
If threads in a warp take different branches, that is divergence. The hardware runs each path serially, leaving some lanes idle and wasting cycles.
The Naive Reduction Diverges
The simple version uses tid % (2*s) to pick active threads. Active and idle threads interleave inside every warp, so each warp diverges hard.
if (tid % (2 * s) == 0)
data[tid] += data[tid + s];Idle Lanes Still Cost
Even though half the threads do nothing, they still occupy the warp. The warp cannot finish until both the active and idle paths are handled.
Reindex by Thread ID
The fix is to map active work to the lowest thread IDs instead of scattered ones. Compute an index from tid and the stride.
int index = 2 * s * tid;
if (index < blockDim.x)
data[index] += data[index + s];Why That Helps
Now the busy threads are contiguous: tid 0,1,2,... all work, the rest all rest. Whole warps are either fully active or fully idle.
Fully Idle Warps Are Free
A warp where every lane is idle just retires with no work. There is no per-lane serialization, so the cost of divergence largely disappears.
The Modulo Trap
The hidden villain was the modulo condition. It scattered active threads across each warp, which is exactly what creates divergence.
Same Work, Better Mapping
You did not change the math or the number of additions. You only remapped which thread does each add, and the warps thank you for it.
It Compounds at Scale
Across thousands of blocks and many steps, removing divergence is a real speedup, often a couple of times faster than the naive kernel.
Still One Snag Left
This version reads neighbors that are interleaved in shared memory, which can cause bank conflicts. The next lesson fixes that too.
Quick Check
Think about what causes warp divergence in the naive reduction.
Recap
You killed divergence by giving work to the lowest thread IDs, so warps are all-active or all-idle. Same math, faster reduction. Next: bank conflicts. 🚀
자주 묻는 질문
“워프 분기 없애기” 강의는 무료인가요?
네 — “워프 분기 없애기” 전체 내용을 이 웹사이트에서 무료로 읽을 수 있습니다. 인터랙티브하게 실습하려면(내장 코드 에디터와 24/7 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
“워프 분기 없애기”에서 뭘 배우나요?
워프가 계속 작업하도록 인덱스를 다시 지정합니다. 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.
CUDA Academy을(를) 시작하는 데 경험이 필요한가요?
사전 경험은 필요하지 않습니다. CoddyKit의 CUDA Academy은(는) 초급자부터 고급 학습자까지를 위해 구성되어 있으므로, 여기서 시작하거나 처음부터 시작할 수 있으며 자신의 속도대로 진행할 수 있습니다. 이것은 4개 중 2번째 강의입니다.
“워프 분기 없애기” 강의는 얼마나 걸리나요?
대부분의 CoddyKit 강의는 약 5~10분이 소요됩니다. 각 강의는 간결하고 인터랙티브하여 꾸준한 진행이 가능하며, 웹과 앱에서 중단한 부분부터 바로 시작할 수 있습니다.
이 CUDA Academy 강의에서 코드를 작성하고 실행할 수 있나요?
네. 모든 CUDA Academy 강의에는 내장 코드 에디터가 포함되어 있으므로, 브라우저에서 바로 실제 코드를 작성하고 실행한 후 즉시 AI 피드백을 받을 수 있습니다 — 로컬 설정이 필요 없습니다.
이 강의의 모든 강의
- 축소 트리 개념
- 워프 분기 없애기
- 순차 주소 지정
- 여러 블록의 최종 축소