__shfl_down_sync para reduções
Faça reduções de warp sem barreiras.
__shfl_down_sync para reduções é 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.
O embaralhamento move registradores
Um embaralhamento permite que uma faixa leia diretamente o valor do registrador de outra. Os dados passam entre threads sem memória compartilhada nem barreira intermediária.
Conheça o shfl_down_sync
O recurso principal para reduções é o shfl_down_sync. Cada faixa obtém um valor de uma faixa que está a um número fixo de posições acima no warp.
float v = __shfl_down_sync(mask, val, offset);Lendo os argumentos
A chamada recebe uma máscara ativa, o valor a ser compartilhado e um deslocamento. A faixa i recebe o valor mantido pela faixa i mais o deslocamento.
Some o que você recebe
Uma soma do warp simplesmente adiciona novamente o valor embaralhado. A faixa i pega o número da vizinha e o incorpora ao seu próprio total acumulado.
val += __shfl_down_sync(mask, val, offset);Reduza o deslocamento pela metade
Como em uma redução em árvore, o deslocamento começa em 16 e é reduzido pela metade a cada etapa: 16, 8, 4, 2, 1. Após cinco etapas, todo o warp está somado.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);Por que cinco etapas
Um warp tem 32 faixas, e 2 elevado à quinta potência é 32. Portanto, cinco etapas de redução pela metade são exatamente suficientes para combinar os 32 valores em um só.
A resposta fica na faixa 0
Após o laço, o total completo do warp fica na faixa 0. As outras faixas contêm resultados parciais sem utilidade, portanto somente a faixa 0 deve gravar o resultado.
Não são necessárias barreiras
Como a troca ocorre por meio dos registradores em sincronia, você pode ignorar completamente syncthreads. O sufixo sync já coordena as faixas mascaradas.
Mais rápido que a memória compartilhada
Uma redução por embaralhamento do warp supera a versão com memória compartilhada: menos instruções, nenhum conflito de bancos e nenhuma espera em barreira. Esse é o caminho rápido para as últimas 32 faixas.
Passe a máscara correta
Se algumas faixas já tiverem saído, uma máscara completa estará errada. Capture o conjunto ativo com activemask para que cada embaralhamento toque apenas nas faixas presentes.
Reduções em dois níveis
Reduções grandes geralmente reduzem cada warp com embaralhamentos e depois combinam os resultados dos warps na memória compartilhada. Os embaralhamentos lidam muito bem com o nível interno, de baixo custo.
Verificação rápida
Pense em qual faixa contém a resposta quando o laço termina.
Recapitulação
Você somou um warp com shfl_down_sync: reduza o deslocamento pela metade cinco vezes, e a faixa 0 conterá o total, sem necessidade de barreiras. A seguir: votação e ballot. ✨
Aprenda C++ com um tutor de IA — grátis
Escreva e execute código real no seu navegador, obtenha ajuda instantânea de um tutor de IA 24/7 e continue de onde parou na web ou no app.
- Cursos
- 30
- Aulas
- 120
Perguntas Frequentes
A aula “__shfl_down_sync para reduções” é grátis?
Sim — o texto completo de “__shfl_down_sync para reduções” é 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 “__shfl_down_sync para reduções”?
Faça reduções de warp sem barreiras. 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 “__shfl_down_sync para reduções”?
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
- Warps, faixas e máscaras
- __shfl_down_sync para reduções
- Funções de votação e cédula
- Grupos cooperativos