0Pricing
CUDA Academy · Leçon

Éliminer la divergence des warps

Réindexez pour maintenir les warps actifs.

Éliminer la divergence des warps est une leçon CUDA Academy gratuite sur CoddyKit. Ceci est la leçon 2 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.

Warps Run in Lockstep

A warp is 32 threads that execute the same instruction together. When their paths agree, the hardware runs at full speed.

What Divergence Costs

If threads in a warp take different branches, that is divergence. The hardware runs each path serially, leaving some lanes idle and wasting cycles.

The Naive Reduction Diverges

The simple version uses tid % (2*s) to pick active threads. Active and idle threads interleave inside every warp, so each warp diverges hard.

if (tid % (2 * s) == 0)
  data[tid] += data[tid + s];

Idle Lanes Still Cost

Even though half the threads do nothing, they still occupy the warp. The warp cannot finish until both the active and idle paths are handled.

Reindex by Thread ID

The fix is to map active work to the lowest thread IDs instead of scattered ones. Compute an index from tid and the stride.

int index = 2 * s * tid;
if (index < blockDim.x)
  data[index] += data[index + s];

Why That Helps

Now the busy threads are contiguous: tid 0,1,2,... all work, the rest all rest. Whole warps are either fully active or fully idle.

Fully Idle Warps Are Free

A warp where every lane is idle just retires with no work. There is no per-lane serialization, so the cost of divergence largely disappears.

The Modulo Trap

The hidden villain was the modulo condition. It scattered active threads across each warp, which is exactly what creates divergence.

Same Work, Better Mapping

You did not change the math or the number of additions. You only remapped which thread does each add, and the warps thank you for it.

It Compounds at Scale

Across thousands of blocks and many steps, removing divergence is a real speedup, often a couple of times faster than the naive kernel.

Still One Snag Left

This version reads neighbors that are interleaved in shared memory, which can cause bank conflicts. The next lesson fixes that too.

Quick Check

Think about what causes warp divergence in the naive reduction.

Recap

You killed divergence by giving work to the lowest thread IDs, so warps are all-active or all-idle. Same math, faster reduction. Next: bank conflicts. 🚀

Questions Fréquemment Posées

La leçon « Éliminer la divergence des warps » est-elle gratuite ?

Oui — le texte complet de « Éliminer la divergence des warps » 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 « Éliminer la divergence des warps » ?

Réindexez pour maintenir les warps actifs. 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 2 sur 4.

Combien de temps prend la leçon « Éliminer la divergence des warps » ?

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. L’idée de l’arbre de réduction
  2. Éliminer la divergence des warps
  3. Adressage séquentiel
  4. Réduction finale sur plusieurs blocs
← Retour à CUDA Academy