Elimine a divergência de warps
Reindexe para manter os warps ocupados.
Elimine a divergência de warps é uma aula grátis de CUDA Academy no CoddyKit. Esta é a aula 2 de 4. Você pode ler a aula completa abaixo gratuitamente — depois pratica ao vivo no navegador com um editor de código integrado e um tutor de IA 24/7. Faz parte do caminho de aprendizado de CUDA Academy, e seu progresso é sincronizado entre a web e o app CoddyKit. O curso de CUDA Academy inclui 4 aulas no total.
Partes desta aula ainda não foram traduzidas e aparecem em inglês.
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. 🚀
Perguntas Frequentes
A aula “Elimine a divergência de warps” é grátis?
Sim — o texto completo de “Elimine a divergência de warps” é grátis para ler aqui na web. Para praticá-la interativamente (um editor de código integrado e um tutor de IA 24/7) e desbloquear o restante do curso de CUDA Academy, atualize para CoddyKit PRO. O curso de CUDA Academy inclui 4 aulas no total.
O que vou aprender em “Elimine a divergência de warps”?
Reindexe para manter os warps ocupados. Você pratica CUDA Academy com código prático que executa diretamente no navegador, e um tutor de IA 24/7 responde suas dúvidas enquanto trabalha na aula.
Preciso ter experiência prévia para começar CUDA Academy?
Nenhuma experiência prévia é necessária. CUDA Academy no CoddyKit é estruturado para alunos iniciantes até avançados, então você pode começar aqui ou desde o início e aprender no seu ritmo. Esta é a aula 2 de 4.
Quanto tempo leva a aula “Elimine a divergência de warps”?
A maioria das aulas CoddyKit leva cerca de 5–10 minutos. Cada uma é compacta e interativa, então você faz progresso constante e retoma exatamente de onde parou entre web e app.
Posso escrever e executar código nesta aula de CUDA Academy?
Sim. Cada aula de CUDA Academy inclui um editor de código integrado, então você escreve e executa código real direto no navegador e recebe feedback de IA instantaneamente — nenhuma configuração local necessária.
Todas as aulas deste curso
- A ideia da árvore de redução
- Elimine a divergência de warps
- Endereçamento sequencial
- Redução final em vários blocos