Aula 15: Diagnosticando e Corrigindo Acesso com Stride
Na aula anterior vimos o que é coalescing: quando threads consecutivas em um warp tocam endereços consecutivos, uma transação atende a todas elas, e quando pulam em um stride grande o warp consome muitas transações desperdiçadas. Agora vamos aprender a habilidade prática — identificar um acesso unco
Imagine 4 trabalhadores colhendo maçãs em fileiras de árvores. No início, cada trabalhador corre para uma coluna diferente em uma extremidade distante do pomar — 4 corridas longas. Se em vez disso cada trabalhador pega a próxima maçã em uma fileira bem próxima, todos trabalham lado a lado em uma única passada. Não mudamos o pomar — só quem-colhe-o-quê.
- strided access
- Um padrão em que threads consecutivas pulam em um stride grande (por exemplo data[threadIdx.x * N]), de modo que seus endereços ficam bem afastados e cada um cai em uma transação de memória diferente.
- coalesced vs uncoalesced
- Coalesced: os endereços de threads vizinhas são contíguos, uma transação para o warp inteiro. Uncoalesced: os endereços estão espalhados, o warp consome muitas transações, a maioria desperdiçada.
- index-remap fix
- A correção para acesso espalhado: fazer com que threads vizinhas toquem endereços vizinhos — dar à thread t o elemento base + t em vez de base + t*N, tudo em global memory.
- ordem row-major
- Um layout em que elementos da mesma linha são vizinhos na memória. Percorrer por linha faz com que threads vizinhas toquem endereços vizinhos (coalesced); percorrer por coluna pula width (espalhado).
