축소를 위한 __shfl_down_sync
장벽 없이 워프를 축소합니다.
축소를 위한 __shfl_down_sync은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 2번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 AI 튜터와 함께 브라우저에서 직접 실습할 수 있습니다. 이 강의는 CUDA Academy 학습 경로의 일부이며, 진행 상황이 웹과 CoddyKit 앱에 동기화됩니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
이 강의의 일부는 아직 번역되지 않았으며 영어로 표시됩니다.
Shuffle Moves Registers
A shuffle lets one lane read another lane's register value directly. Data hops between threads with no shared memory and no barrier in between.
Meet shfl_down_sync
The workhorse for reductions is shfl_down_sync. Each lane grabs a value from a lane a fixed number of positions higher in the warp.
float v = __shfl_down_sync(mask, val, offset);Reading the Arguments
The call takes an active mask, the value to share, and an offset. Lane i receives the value held by lane i plus offset.
Add What You Receive
A warp sum just adds the shuffled value back in. Lane i takes its neighbor's number and folds it into its own running total.
val += __shfl_down_sync(mask, val, offset);Halve the Offset
Like a tree reduction, the offset starts at 16 and halves each step: 16, 8, 4, 2, 1. After five steps the whole warp is summed.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Why Five Steps
A warp has 32 lanes, and 2 to the fifth power is 32. So five halving steps are exactly enough to combine all 32 values into one.
Answer Lands in Lane 0
After the loop, the complete warp total sits in lane 0. The other lanes hold partial junk, so only lane 0 should write the result out.
No Barriers Needed
Because exchange happens through registers in lockstep, you skip syncthreads entirely. The sync suffix already coordinates the masked lanes.
Faster Than Shared
A warp shuffle reduction beats the shared-memory version: fewer instructions, no bank conflicts, and no barrier stalls. It is the fast path for the last 32.
Pass the Right Mask
If some lanes already exited, a full mask is wrong. Capture the live set with activemask so each shuffle only touches present lanes.
Two-Level Reductions
Big reductions often reduce each warp with shuffles, then combine warp results in shared memory. Shuffles handle the cheap inner level beautifully.
Quick Check
Think about which lane holds the answer when the loop finishes.
Recap
You summed a warp with shfl_down_sync: halve the offset five times, and lane 0 holds the total, no barriers needed. Next: ballot and vote. ✨
자주 묻는 질문
“축소를 위한 __shfl_down_sync” 강의는 무료인가요?
네 — “축소를 위한 __shfl_down_sync” 전체 내용을 이 웹사이트에서 무료로 읽을 수 있습니다. 인터랙티브하게 실습하려면(내장 코드 에디터와 24/7 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
“축소를 위한 __shfl_down_sync”에서 뭘 배우나요?
장벽 없이 워프를 축소합니다. 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.
CUDA Academy을(를) 시작하는 데 경험이 필요한가요?
사전 경험은 필요하지 않습니다. CoddyKit의 CUDA Academy은(는) 초급자부터 고급 학습자까지를 위해 구성되어 있으므로, 여기서 시작하거나 처음부터 시작할 수 있으며 자신의 속도대로 진행할 수 있습니다. 이것은 4개 중 2번째 강의입니다.
“축소를 위한 __shfl_down_sync” 강의는 얼마나 걸리나요?
대부분의 CoddyKit 강의는 약 5~10분이 소요됩니다. 각 강의는 간결하고 인터랙티브하여 꾸준한 진행이 가능하며, 웹과 앱에서 중단한 부분부터 바로 시작할 수 있습니다.
이 CUDA Academy 강의에서 코드를 작성하고 실행할 수 있나요?
네. 모든 CUDA Academy 강의에는 내장 코드 에디터가 포함되어 있으므로, 브라우저에서 바로 실제 코드를 작성하고 실행한 후 즉시 AI 피드백을 받을 수 있습니다 — 로컬 설정이 필요 없습니다.
이 강의의 모든 강의
- 워프, 레인과 마스크
- 축소를 위한 __shfl_down_sync
- 투표와 선택 함수
- 협력 그룹