CUDA Academy · Aula

__shfl_down_sync para reduções

Faça reduções de warp sem barreiras.

Aula 2 de 413 etapas

__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. ✨

Grátis para começar

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

  1. Warps, faixas e máscaras
  2. __shfl_down_sync para reduções
  3. Funções de votação e cédula
  4. Grupos cooperativos
← Voltar para CUDA Academy