__shfl_down_sync für Reduktionen
Warp-Reduktionen ohne Barrieren.
__shfl_down_sync für Reduktionen ist eine kostenlose CUDA Academy-Lektion auf CoddyKit. Dies ist Lektion 2 von 4. Du kannst die komplette Lektion unten kostenlos lesen – dann übst du sie direkt im Browser mit einem integrierten Code-Editor und einem KI-Tutor rund um die Uhr. Sie ist Teil des CUDA Academy-Lernpfads, und dein Fortschritt wird über Web und CoddyKit-App synchronisiert. Der CUDA Academy-Kurs umfasst insgesamt 4 Lektionen.
Teile dieser Lektion wurden noch nicht übersetzt und werden auf Englisch angezeigt.
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. ✨
Häufig gestellte Fragen
Ist die Lektion „__shfl_down_sync für Reduktionen“ kostenlos?
Ja — der vollständige Text von „__shfl_down_sync für Reduktionen“ ist hier im Web kostenlos zu lesen. Um sie interaktiv zu üben (integrierter Code-Editor und 24/7 KI-Tutor) und den Rest des CUDA Academy-Kurses freizuschalten, upgrade auf CoddyKit PRO. Der CUDA Academy-Kurs umfasst insgesamt 4 Lektionen.
Was lerne ich in „__shfl_down_sync für Reduktionen“?
Warp-Reduktionen ohne Barrieren. Du übst CUDA Academy mit praktischem Code, den du direkt im Browser ausführst, und ein 24/7 KI-Tutor beantwortet deine Fragen während du die Lektion bearbeitest.
Brauche ich Erfahrung, um CUDA Academy zu starten?
Keine Vorkenntnisse erforderlich. CUDA Academy auf CoddyKit ist für Anfänger bis fortgeschrittene Lernende strukturiert, sodass du hier starten oder von Anfang an beginnen und in deinem eigenen Tempo voranschreiten kannst. Dies ist Lektion 2 von 4.
Wie lange dauert die Lektion „__shfl_down_sync für Reduktionen“?
Die meisten CoddyKit-Lektionen dauern etwa 5–10 Minuten. Jede ist kompakt und interaktiv, sodass du stetig Fortschritte machst und genau dort weitermachst, wo du aufgehört hast – im Web und in der App.
Kann ich in dieser CUDA Academy-Lektion Code schreiben und ausführen?
Ja. Jede CUDA Academy-Lektion enthält einen integrierten Code-Editor, sodass du echten Code direkt in deinem Browser schreibst und ausführst und sofort KI-Feedback erhältst — ohne lokale Einrichtung erforderlich.
Alle Lektionen in diesem Kurs
- Warps, Lanes und Masken
- __shfl_down_sync für Reduktionen
- Ballot- und Vote-Funktionen
- Cooperative Groups