Adressage séquentiel
Des pas sans conflit dans la mémoire partagée.
Adressage séquentiel 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.
Shared Memory Has Banks
Shared memory is split into 32 banks, one per warp lane. When 32 threads hit 32 different banks, all reads happen in a single fast cycle.
Bank Conflicts Slow You Down
If two threads in a warp touch the same bank, that is a bank conflict. The hardware serializes those accesses, costing extra cycles.
Interleaved Addressing
The previous reduction used interleaved addressing: stride starts small and doubles, so partners are close together in shared memory.
int index = 2 * s * tid;
data[index] += data[index + s];Why Interleaving Conflicts
With small, doubling strides, several lanes in a warp map onto the same bank. Those accesses can no longer happen in one cycle.
Flip the Stride Order
Sequential addressing starts the stride large and halves it each step, the reverse of interleaving. This single change removes the conflicts.
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s)
data[tid] += data[tid + s];
__syncthreads();
}Big Stride, Clean Banks
A large stride spreads partner addresses far apart, so each lane lands on its own bank. The warp reads conflict-free in one cycle.
The tid < s Guard
Only the lower half of threads work each step, written as tid < s. That keeps active threads contiguous, so warps stay non-divergent too.
Two Wins at Once
Sequential addressing fixes bank conflicts and avoids warp divergence in the same kernel. One layout change, two performance problems solved.
Still Sync Each Step
You still need a __syncthreads after each step. Threads must see the previous level's writes before they read for the next level.
Result Lands at Index 0
As the stride halves toward zero, all partial sums fold into data[0]. Thread 0 then writes that block's result back to global memory.
A Classic Optimization
This pattern comes straight from NVIDIA's famous reduction guide. Sequential addressing is a textbook step toward a conflict-free kernel.
Quick Check
Think about why a large, halving stride beats a small, doubling one.
Recap
You swapped interleaved for sequential addressing: stride starts large and halves, killing bank conflicts and divergence at once. Up next: multi-block sums. ✨
Questions Fréquemment Posées
La leçon « Adressage séquentiel » est-elle gratuite ?
Oui — le texte complet de « Adressage séquentiel » 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 « Adressage séquentiel » ?
Des pas sans conflit dans 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 « Adressage séquentiel » ?
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
- L’idée de l’arbre de réduction
- Éliminer la divergence des warps
- Adressage séquentiel
- Réduction finale sur plusieurs blocs