Macro CONFLICT_FREE_OFFSET usada no algoritmo de prefixo paralelo do GPU Gems 3

Question

Em primeiro lugar, aqui está o link para o algoritmo:

Gemas da GPU 3, Capítulo 39: Soma de prefixo paralelo (varredura) com CUDA.

Para evitar conflitos de bancos, o preenchimento é adicionado ao array de memória compartilhada a cada NUM_BANKS (ou seja, 32 para dispositivos de computabilidade 2.x) elementos.Isso é feito (como na Figura 39-5):

int ai = offset*(2*thid+1)-1
int bi = offset*(2*thid+2)-1
ai += ai/NUM_BANKS
bi += ai/NUM_BANKS
temp[bi] += temp[ai]

Não entendo como ai/NUM_BANKS é equivalente à macro:

   #define NUM_BANKS 16  
   #define LOG_NUM_BANKS 4  
   #define CONFLICT_FREE_OFFSET(n) \  
          ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS))  

Não é igual a

n >> LOG_NUM_BANKS

Qualquer ajuda é apreciada.Obrigado

Answer 1

Eu escrevi esse código e co-escrevi o artigo, e solicito que você use o artigo apenas para aprender sobre algoritmos de varredura e não use o código nele contido.Foi escrito quando o CUDA era novo e eu era novo no CUDA.Se você usar uma implementação moderna de varredura em CUDA, não precisará evitar conflitos bancários.

Se você quiser fazer varreduras da maneira mais fácil, use thrust::inclusive_scan ou thrust::exclusive_scan.

Se você realmente deseja implementar uma verificação, consulte artigos mais recentes como este [1].Ou para uma obra real com código mais rápido, mas que exigirá um pouco mais de estudo, esta [2].Ou leia Tutorial de Sean Baxter (embora este último não inclua citações do trabalho seminal sobre o algoritmo de varredura).

[1] Shubhabrata Sengupta, Mark Harris, Michael Garland e John D.Owens."Algoritmos de varredura paralela eficientes para GPUs de vários núcleos" .Em Jakub Kurzak, David A.Bader e Jack Dongarra, editores, Scientific Computing with Multicore and Accelerators, Chapman & Hall/CRC Computational Science, capítulo 19, páginas 413–442.Taylor e Francis, janeiro de 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1041

[2] Merril, D.e Grimshaw, A.Varredura paralela para arquiteturas de fluxo.Relatório Técnico CS2009-14, Departamento de Ciência da Computação, Universidade da Virgínia.Dez.2009.