CUDA Academy · Les

__shfl_down_sync voor reducties

Warp-reducties zonder barrières.

Les 2 van 413 stappen

__shfl_down_sync voor reducties is een gratis CUDA Academy-les op CoddyKit. Dit is les 2 van 4. Je kunt de volledige les hieronder gratis lezen en daarna in de browser praktisch oefenen met een ingebouwde code-editor en een AI-begeleider die 24/7 beschikbaar is. Deze les maakt deel uit van het leertraject CUDA Academy. Je voortgang wordt gesynchroniseerd op het web en in de CoddyKit-app. De cursus CUDA Academy bevat in totaal 4 lessen.

Shuffle verplaatst registers

Met een shuffle kan één baan rechtstreeks de registerwaarde van een andere baan lezen. Gegevens springen tussen threads zonder gedeeld geheugen en zonder tussenliggende barrière.

Maak kennis met shfl_down_sync

Het belangrijkste hulpmiddel voor reducties is shfl_down_sync. Elke baan haalt een waarde op uit een baan die een vast aantal posities hoger in de warp staat.

float v = __shfl_down_sync(mask, val, offset);

De argumenten lezen

De aanroep krijgt een actief masker, de te delen waarde en een offset. Baan i ontvangt de waarde die wordt vastgehouden door baan i plus de offset.

Tel op wat u ontvangt

Een warpsom telt de geshuffelde waarde er eenvoudig weer bij op. Baan i neemt het getal van zijn buur en voegt dat toe aan zijn eigen lopende totaal.

val += __shfl_down_sync(mask, val, offset);

Halveer de offset

Net als bij een boomreductie begint de offset bij 16 en halveert die bij elke stap: 16, 8, 4, 2, 1. Na vijf stappen is de hele warp opgeteld.

for (int o = 16; o > 0; o >>= 1)
  val += __shfl_down_sync(mask, val, o);

Waarom vijf stappen

Een warp heeft 32 banen en 2 tot de vijfde macht is 32. Daarom zijn vijf halveringsstappen precies genoeg om alle 32 waarden tot één waarde te combineren.

Het antwoord komt terecht in baan 0

Na de lus staat het volledige warp-totaal in lane 0. De andere lanes bevatten gedeeltelijke rommel, dus alleen lane 0 mag het resultaat wegschrijven.

Geen barrières nodig

Omdat de uitwisseling synchroon via registers gebeurt, kun je syncthreads volledig overslaan. Het achtervoegsel _sync coördineert de gemaskeerde lanes al.

Sneller dan gedeeld geheugen

Een reductie met warp-shuffles is sneller dan de versie met gedeeld geheugen: minder instructies, geen bankconflicten en geen vertraging door barrières. Dit is het snelle pad voor de laatste 32 elementen.

Geef het juiste masker door

Als sommige lanes al zijn gestopt, is een volledig masker onjuist. Leg de actieve set vast met activemask, zodat elke shuffle alleen aanwezige lanes aanraakt.

Reducties in twee niveaus

Bij grote reducties wordt elke warp vaak eerst met shuffles gereduceerd, waarna de warpresultaten in shared memory worden gecombineerd. Shuffles verwerken het goedkope binnenste niveau uitstekend.

Korte controle

Denk na over welke lane het antwoord bevat wanneer de lus klaar is.

Samenvatting

Je hebt een warp opgeteld met shfl_down_sync: halveer de offset vijf keer en lane 0 bevat het totaal, zonder barrières. Hierna: ballot en vote. ✨

Gratis beginnen

Leer C++ met een AI-tutor — gratis

Schrijf echte code en voer die uit in je browser, krijg direct hulp van een AI-tutor die 24/7 beschikbaar is en ga verder waar je gebleven bent op het web of in de app.

Cursussen
30
Lessen
120

Veelgestelde vragen

Is de les “__shfl_down_sync voor reducties” gratis?

Ja — de volledige tekst van “__shfl_down_sync voor reducties” kun je hier gratis op het web lezen. Als je interactief wilt oefenen met een ingebouwde code-editor en een AI-begeleider die 24/7 beschikbaar is, en de rest van de cursus CUDA Academy wilt ontgrendelen, kun je upgraden naar CoddyKit PRO. De cursus CUDA Academy bevat in totaal 4 lessen.

Wat leer ik in “__shfl_down_sync voor reducties”?

Warp-reducties zonder barrières. Je oefent met CUDA Academy door code rechtstreeks in de browser uit te voeren. Een AI-begeleider die 24/7 beschikbaar is beantwoordt je vragen terwijl je de les doorwerkt.

Heb ik ervaring nodig om met CUDA Academy te beginnen?

Ervaring vooraf is niet nodig. CUDA Academy op CoddyKit is opgebouwd voor beginners tot gevorderden, zodat je hier of bij het begin kunt starten en in je eigen tempo kunt leren. Dit is les 2 van 4.

Hoe lang duurt de les “__shfl_down_sync voor reducties”?

De meeste lessen van CoddyKit duren ongeveer 5–10 minuten. Elke les is kort en interactief, zodat je gestaag vooruitgaat en op het web en in de app precies verdergaat waar je was gebleven.

Kan ik code schrijven en uitvoeren in deze les over CUDA Academy?

Ja. Elke les over CUDA Academy bevat een ingebouwde code-editor, zodat je rechtstreeks in je browser echte code kunt schrijven en uitvoeren en direct feedback van AI krijgt — lokale installatie is niet nodig.

Alle lessen in deze cursus

  1. Warps, lanes en masks
  2. __shfl_down_sync voor reducties
  3. Ballot- en votefuncties
  4. Cooperative Groups
← Terug naar CUDA Academy