Conflitto_free_offset macro utilizzato nell'algoritmo prefisso parallelo da GPU Gems 3

Question

Prima di tutto, ecco il link all'algoritmo:

GPU Gems 3, Capitolo 39: Parallel Prefisx Sum (Scansione) con CUDA .

Per evitare conflitti bancari, l'imbottitura viene aggiunta all'array di memoria condivisa ogni numero num_banks (I.e., 32 per dispositivi di elementi di Computability 2.x).Questo è fatto da (come nella 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]

.

Non capisco come AI / NUM_BANKS è equivalente alla macro:

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

.

non è uguale a

n >> LOG_NUM_BANKS

.

Qualsiasi aiuto è apprezzato.Grazie

Answer 1

Ho scritto quel codice e ho scritto l'articolo, e chiedo di utilizzare l'articolo solo per l'apprendimento degli algoritmi di scansione e non utilizzare il codice in esso. È stato scritto quando Cuda era nuovo, ed ero nuovo a Cuda. Se si utilizza una moderna implementazione della scansione in CUDA non hai bisogno di alcun conflitto di conflitto bancario.

Se vuoi eseguire la scansione del modo semplice, utilizzare thrust::inclusive_scan o thrust::exclusive_scan .

Se vuoi veramente implementare una scansione, fare riferimento a articoli più recenti come questo [1] . O per un vero opus con un codice più rapido, ma questo richiederà un po 'più di studio, questo [2] . O leggere tutorial di Sean Baxter (anche se quest'ultimo non include citazioni del lavoro seminale sull'algoritmo di scansione).

[1] Shubhabrata Sengupta, Mark Harris, Michael Garland e John D. Owens. "Algoritmi di scansione parallela efficienti per GPU a molti core". A Jakub Kurzak, David A. Bader, e Jack Dongarra, redattori, calcolo scientifico con multicore e acceleratori, Chapman & Hall / CRC Computational Science, Capitolo 19, pagine 413-442. Taylor & Francis, gennaio 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id= 1041

[2] Merrill, D. e Grimshaw, A. Scansione parallela per architetture in streaming. Rapporto tecnico CS2009-14, Dipartimento di Informatica, Università della Virginia. 29 dicembre 2009.