Aula 25: Scan Paralelo — o Kernel de Duplicação
Na aula passada vimos o que uma soma de prefixos calcula e a ideia da duplicação: a cada rodada, todo elemento soma o vizinho a uma distância d, e d dobra (1, 2, 4...). Aqui vamos construir isso como um kernel de CUDA de verdade. Vamos trabalhar dentro de um único block, com um array em shared memor
Cada thread é dona de uma posição no array compartilhado. Em cada passo, ela primeiro lê o valor d lugares à sua esquerda e guarda esse valor de lado, espera todo mundo terminar de ler (__syncthreads), só então soma esse valor a si mesma, e espera de novo até todo mundo terminar de escrever antes do próximo passo. As esperas são o que mantém a ordem e evita sobrescrita.
- Hillis-Steele
- O kernel de scan por duplicação em log2(n) passos: no passo d, cada thread com tid>=d soma o valor de tid-d. d dobra a cada passo.
- tile em shared memory
- Um array __shared__ para o block inteiro. O scan carrega a entrada nele, trabalha no próprio lugar ao longo de todos os passos, e por fim escreve em out.
- barreiras de leitura e escrita
- __syncthreads() depois da leitura (para que todas leiam antes de qualquer escrita) e depois da escrita (para que as escritas fiquem visíveis para o próximo passo).
- número de passos log2(n)
- O loop for(d=1; d<n; d*=2) roda log2(n) vezes, porque d dobra a cada passo até chegar a n.
