0Pricing
CUDA Academy · Aula

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

  1. A ideia da árvore de redução
  2. Elimine a divergência de warps
  3. Endereçamento sequencial
  4. Redução final em vários blocos
← Voltar para CUDA Academy