Redução final em vários blocos

Uma redução dentro de um bloco é fácil, mas os blocos executam de forma independente e não podem sincronizar entre si durante o núcleo. Portanto, um único lançamento não pode somar tudo.

Assim, cada bloco reduz sua própria parte a um número, uma soma parcial, e a grava em um pequeno array de saída indexado por blockIdx.

Com 1000 blocos, você passa de um milhão de entradas para 1000 somas parciais. A parte difícil terminou; resta combinar apenas um array pequeno.

A forma mais simples de concluir é fazer um segundo lançamento do mesmo núcleo sobre as somas parciais. Repita até restar apenas um valor.

Cada passagem reduz o array pelo tamanho do bloco. Alguns lançamentos recursivos reduzem milhões de valores a uma única soma final.

Como alternativa, a thread 0 de cada bloco pode adicionar sua soma parcial diretamente a um único total global usando atomicAdd, evitando um segundo núcleo.

As operações atômicas são simples e exigem apenas um lançamento, mas muitos blocos competindo pelo mesmo endereço podem ser serializados. Com poucas somas parciais, geralmente funcionam bem.

Um laço com passo da grade permite que cada thread some primeiro muitos elementos em um registrador, reduzindo bastante o número de blocos necessários antes da etapa final.

Fazer mais trabalho por thread no início significa menos somas parciais e menos lançamentos. Isso costuma ser melhor que criar uma thread por elemento.

Se usar operações atômicas, lembre-se de zerar o total global antes de iniciar, ou a soma começará com lixo deixado nessa memória.

Entradas pequenas se beneficiam da simplicidade das operações atômicas; entradas enormes favorecem um projeto de duas passagens ou com avanço pela grade. Faça medições com seus dados para escolher.

Cada bloco produz uma soma parcial, e depois você as combina com uma segunda inicialização, operações atômicas ou avanço pela grade. Agora você pode reduzir vetores de qualquer tamanho. 🏁