__shfl_down_sync for reduksjoner
Utfør warp-reduksjoner uten barrierer
__shfl_down_sync for reduksjoner er en gratis leksjon i CUDA Academy på CoddyKit. Dette er leksjon 2 av 4. Du kan lese hele leksjonen gratis nedenfor – og deretter øve praktisk i nettleseren med en innebygd kodeeditor og en AI-veileder som er tilgjengelig døgnet rundt. Den er en del av læringsløpet i CUDA Academy, og fremdriften din synkroniseres mellom nettet og CoddyKit-appen. Kurset i CUDA Academy inneholder totalt 4 leksjoner.
Shuffle flytter registre
En shuffle lar én lane lese registerverdien til en annen lane direkte. Data flyttes mellom tråder uten delt minne og uten en barriere imellom.
Møt shfl_down_sync
Arbeidshesten for reduksjoner er shfl_down_sync. Hver lane henter en verdi fra en lane et fast antall posisjoner høyere i warpen.
float v = __shfl_down_sync(mask, val, offset);Les argumentene
Kallet tar en aktiv maske, verdien som skal deles, og et offset. Lane i mottar verdien som holdes av lane i pluss offset.
Legg sammen det De mottar
En warp-sum legger ganske enkelt den shuflede verdien tilbake. Lane i tar naboens tall og legger det inn i sin egen løpende total.
val += __shfl_down_sync(mask, val, offset);Halver offsetet
Som i en trærreduksjon starter offsetet på 16 og halveres for hvert trinn: 16, 8, 4, 2, 1. Etter fem trinn er hele warpen summert.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Hvorfor fem trinn
En warp har 32 laner, og 2 opphøyd i femte er 32. Derfor er fem halveringstrinn akkurat nok til å kombinere alle de 32 verdiene til én.
Svaret havner i lane 0
Etter løkken ligger den fullstendige warp-summen i lane 0. De andre lanene inneholder delvise søppelverdier, så det er bare lane 0 som skal skrive ut resultatet.
Ingen barrierer er nødvendige
Fordi utvekslingen skjer gjennom registre i låstakt, kan De hoppe over syncthreads helt. Sync-suffikset koordinerer allerede de maskerte lanene.
Raskere enn shared memory
En reduksjon med warp shuffle slår varianten med shared memory: færre instruksjoner, ingen bankkonflikter og ingen barrier-stans. Dette er den raske veien for de siste 32.
Send riktig maske
Hvis noen laner allerede har avsluttet, er en full maske feil. Fang opp de aktive lanene med activemask, slik at hver shuffle bare berører laner som fortsatt er til stede.
Reduksjoner på to nivåer
Store reduksjoner reduserer ofte hver warp med shuffles og kombinerer deretter warp-resultatene i shared memory. Shuffles håndterer det kostnadseffektive indre nivået svært godt.
Rask kontroll
Tenk over hvilken lane som inneholder svaret når løkken er ferdig.
Oppsummering
De summerte en warp med shfl_down_sync: halver offseten fem ganger, så inneholder lane 0 totalen, uten behov for barrierer. Neste: ballot og vote. ✨
Lær deg C++ med en AI-veileder – gratis
Skriv og kjør ekte kode i nettleseren, få umiddelbar hjelp fra en AI-veileder som er tilgjengelig døgnet rundt, og fortsett der du slapp – på nettet eller i appen.
- Kurs
- 30
- Leksjoner
- 120
Ofte stilte spørsmål
Er leksjonen «__shfl_down_sync for reduksjoner» gratis?
Ja – hele teksten i «__shfl_down_sync for reduksjoner» er gratis å lese her på nettet. For å øve interaktivt med en innebygd kodeeditor og en AI-veileder som er tilgjengelig døgnet rundt, og for å låse opp resten av CUDA Academy-kurset, kan du oppgradere til CoddyKit PRO. Kurset i CUDA Academy inneholder totalt 4 leksjoner.
Hva lærer jeg i «__shfl_down_sync for reduksjoner»?
Utfør warp-reduksjoner uten barrierer Du øver på CUDA Academy med praktisk kode som du kjører direkte i nettleseren, mens en AI-veileder som er tilgjengelig døgnet rundt, svarer på spørsmålene dine mens du jobber deg gjennom leksjonen.
Trenger jeg erfaring for å begynne med CUDA Academy?
Ingen tidligere erfaring er nødvendig. CUDA Academy på CoddyKit er lagt opp for både nybegynnere og viderekomne, så De kan begynne her eller helt fra start og lære i Deres eget tempo. Dette er leksjon 2 av 4.
Hvor lang tid tar leksjonen «__shfl_down_sync for reduksjoner»?
De fleste CoddyKit-leksjoner tar omtrent 5–10 minutter. Hver leksjon er kort og interaktiv, slik at De gjør jevne fremskritt og kan fortsette akkurat der De slapp – både på nettet og i appen.
Kan jeg skrive og kjøre kode i denne CUDA Academy-leksjonen?
Ja. Alle CUDA Academy-leksjoner har en innebygd kodeeditor, slik at De kan skrive og kjøre ekte kode direkte i nettleseren og få umiddelbar tilbakemelding fra AI – uten lokal konfigurering.
Alle leksjonene i dette kurset
- Warps, lanes og masker
- __shfl_down_sync for reduksjoner
- Ballot- og vote-funksjoner
- Kooperative grupper