CUDA Academy · Les

Een histogram bouwen

Gebruik atomics met privatisering in shared memory.

Les 3 van 413 stappen

Een histogram bouwen is een gratis CUDA Academy-les op CoddyKit. Dit is les 3 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.

Wat een histogram telt

Een histogram telt hoeveel invoerwaarden in elke bak vallen. Veel threads zullen dezelfde bak willen verhogen, dus dit is een probleem voor atomics. 📊

De naïeve aanpak

Elke thread leest één element, bepaalt de bijbehorende bak en verhoogt die bak. Zonder bescherming verliezen populaire bakken tellingen door races.

De globale atomische versie

De eenvoudigste correcte oplossing is één atomicAdd per element, rechtstreeks naar het globale geheugen. Dit werkt, maar drukke bakken voeren threads na elkaar uit.

atomicAdd(&hist[bin], 1);

Het probleem van concurrentie

Wanneer gegevens zich ophopen in enkele bakken, richten duizenden threads zich op hetzelfde adres. Die concurrentie kan globale atomics pijnlijk traag maken.

Privatisering als oplossing

Privatisering geeft elk blok een eigen privéhistogram in snel gedeeld geheugen. Threads botsen alleen binnen hun blok, niet in het hele grid.

Het gedeelde histogram declareren

Elk blok declareert een gedeelde array met een grootte die overeenkomt met het aantal bakken. Deze staat op de chip, waardoor atomics daar veel goedkoper zijn dan globale atomics.

__shared__ int local[NBINS];

Stap 1: maak de bakken leeg

Threads zetten gezamenlijk het gedeelde histogram op nul en roepen daarna __syncthreads aan, zodat niemand begint te tellen voordat het leegmaken klaar is.

local[tid] = 0;
__syncthreads();

Stap 2: tel lokaal

Nu verhoogt elke thread atomisch zijn bak in het gedeelde geheugen. Hetzelfde atomicAdd, maar op de snelle kopie op de chip in plaats van in het globale geheugen.

atomicAdd(&local[bin], 1);

Stap 3: voeg samen in het globale geheugen

Na een barrière tellen threads elke gedeelde bak met één atomicAdd op bij het globale histogram. Veel minder globale atomics dan voorheen.

atomicAdd(&hist[i], local[i]);

Waarom dit sneller is

Atomics in gedeeld geheugen zijn snel, en de dure globale atomics worden nu één keer per bak per blok uitgevoerd in plaats van één keer per element.

Let op het aantal bakken

Het privéhistogram moet in het gedeelde geheugen passen. Als er te veel bakken zijn, verwerk je ze in meerdere doorgangen of val je terug op globale atomics.

Snelle controle

Eén vraag over de histogramstrategie.

Samenvatting: een histogram bouwen

Je hebt een histogram gebouwd met globale atomics en het daarna versneld met privatisering in gedeeld geheugen: leegmaken, lokaal tellen, samenvoegen. ✅

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 “Een histogram bouwen” gratis?

Ja — de volledige tekst van “Een histogram bouwen” 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 “Een histogram bouwen”?

Gebruik atomics met privatisering in shared memory. 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 3 van 4.

Hoe lang duurt de les “Een histogram bouwen”?

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. Racecondities op de GPU
  2. atomicAdd en verwante functies
  3. Een histogram bouwen
  4. Aangepaste atomics met atomicCAS
← Terug naar CUDA Academy