Endereçamento sequencial
Passos sem conflitos na memória compartilhada.
Endereçamento sequencial é uma aula grátis de CUDA Academy no CoddyKit. Esta é a aula 3 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.
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. ✨
Perguntas Frequentes
A aula “Endereçamento sequencial” é grátis?
Sim — o texto completo de “Endereçamento sequencial” é 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 “Endereçamento sequencial”?
Passos sem conflitos na memória compartilhada. 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 3 de 4.
Quanto tempo leva a aula “Endereçamento sequencial”?
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