0Pricing
CUDA Academy · Leçon

Construire un histogramme

Utilisez des opérations atomiques avec privatisation de la mémoire partagée.

Construire un histogramme est une leçon CUDA Academy gratuite sur CoddyKit. Ceci est la leçon 3 sur 4. Tu peux lire la leçon complète ci-dessous gratuitement — puis la pratiquer en direct dans le navigateur avec un éditeur de code intégré et un tuteur IA 24/7. Elle fait partie du parcours d'apprentissage CUDA Academy, et ta progression se synchronise sur le web et l'application CoddyKit. Le cours CUDA Academy comprend 4 leçons au total.

Certaines parties de cette leçon n'ont pas encore été traduites et s'affichent en anglais.

What a Histogram Counts

A histogram counts how many inputs fall into each bin. Many threads will want to increment the same bin, so this is an atomics problem. 📊

The Naive Approach

Each thread reads one element, finds its bin, and increments that bin. Without protection, popular bins lose counts to races.

The Global Atomic Version

The simplest correct fix is one atomicAdd per element straight into global memory. It works, but hot bins serialize threads.

atomicAdd(&hist[bin], 1);

The Contention Problem

When data clusters into a few bins, thousands of threads pile onto the same address. That contention can make global atomics painfully slow.

Privatization to the Rescue

Privatization gives each block its own private histogram in fast shared memory. Threads collide only inside their block, not across the whole grid.

Declare the Shared Histogram

Each block declares a shared array sized to the number of bins. It lives on-chip, so atomics there are far cheaper than global ones.

__shared__ int local[NBINS];

Step 1: Clear the Bins

Threads cooperatively zero the shared histogram, then call __syncthreads so no one counts before clearing is done.

local[tid] = 0;
__syncthreads();

Step 2: Count Locally

Now each thread atomically bumps its bin in shared memory. Same atomicAdd, but on the fast on-chip copy instead of global.

atomicAdd(&local[bin], 1);

Step 3: Merge to Global

After a barrier, threads add each shared bin into the global histogram with one atomicAdd per bin. Far fewer global atomics than before.

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

Why This Is Faster

Shared-memory atomics are quick, and the costly global atomics now fire once per bin per block instead of once per element.

Watch the Bin Count

The private histogram must fit in shared memory. With too many bins, split them into passes or fall back to global atomics.

Quick Check

One question on the histogram strategy.

Recap: Building a Histogram

You built a histogram with global atomics, then sped it up using shared-memory privatization: clear, count locally, merge. ✅

Questions Fréquemment Posées

La leçon « Construire un histogramme » est-elle gratuite ?

Oui — le texte complet de « Construire un histogramme » est gratuit à lire ici sur le web. Pour la pratiquer de manière interactive (un éditeur de code intégré et un tuteur IA 24/7) et déverrouiller le reste du cours CUDA Academy, passe à CoddyKit PRO. Le cours CUDA Academy comprend 4 leçons au total.

Qu'est-ce que j'apprendrai dans « Construire un histogramme » ?

Utilisez des opérations atomiques avec privatisation de la mémoire partagée. Tu pratiques CUDA Academy avec du code pratique que tu exécutes directement dans le navigateur, et un tuteur IA 24/7 répond à tes questions au fur et à mesure que tu avances dans la leçon.

Dois-je avoir de l'expérience pour commencer CUDA Academy ?

Aucune expérience préalable n'est requise. CUDA Academy sur CoddyKit est structuré pour les débutants jusqu'aux apprenants avancés, donc tu peux commencer ici ou depuis le début et avancer à ton rythme. Ceci est la leçon 3 sur 4.

Combien de temps prend la leçon « Construire un histogramme » ?

La plupart des leçons CoddyKit prennent environ 5–10 minutes. Chacune est courte et interactive, tu progresses régulièrement et tu repiques exactement où tu t'es arrêté sur le web et l'app.

Peux-tu écrire et exécuter du code dans cette leçon CUDA Academy ?

Oui. Chaque leçon CUDA Academy inclut un éditeur de code intégré, tu écris et exécutes du vrai code directement dans ton navigateur et tu reçois des retours IA instantanés — aucune configuration locale requise.

Toutes les leçons de ce cours

  1. Conditions de concurrence sur le GPU
  2. atomicAdd et fonctions associées
  3. Construire un histogramme
  4. Opérations atomiques personnalisées avec atomicCAS
← Retour à CUDA Academy