Conflict_free_offset macro utilizado en el algoritmo de prefijo paralelo de GPU GEMS 3

Question

En primer lugar, aquí está el enlace al algoritmo:

GPU GEMS 3, Capítulo 39: Resumen de prefijo paralelo (escaneo) con CUDA .

Para evitar los conflictos bancarios, se agrega el relleno a la matriz de memoria compartida cada num_banks (es decir, 32 para dispositivos de computabilidad 2.x) elementos.Esto se hace (como en la 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]

No entiendo cómo AI / NUM_BANKS ES EQUIVALENTE A LA MACRO:

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

¿No es igual a

n >> LOG_NUM_BANKS

Se aprecia cualquier ayuda.Gracias

Answer 1

Escribí ese código y co-escribió el artículo, y solicito que utilice el artículo solo para aprender sobre algoritmos de escaneo, y no utilice el código en él. Fue escrito cuando Cudá era nuevo, y yo era nuevo en Cudá. Si utiliza una implementación moderna de escaneo en CUDA, no necesita ningún evitación de conflictos bancarios.

Si desea obtener escanear la forma fácil, use thrust::inclusive_scan o thrust::exclusive_scan .

Si realmente desea implementar un escaneo, consulte los artículos más recientes, como este [1] . O para un opus real con un código más rápido, pero eso requerirá un poco más de estudio, este [2] . O lea Tutorial de Sean Baxter (aunque este último no incluye las citas del trabajo seminal en el algoritmo de escaneo).

[1] Shubhabrata Sengupta, Mark Harris, Michael Garland y John D. Owens. "Algoritmos de escaneo paralelo eficientes para gpus de muchos núcleos". En Jakub Kurzak, David A. Bader, y Jack Dongarra, editores, computación científica con multinúcleo y aceleradores, Chapman y Hall / CRC Ciencia computacional, Capítulo 19, Páginas 413-442. Taylor & Francis, enero de 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id= 1041

[2] Merrill, D. y Grimshaw, A. Exploración paralela para arquitecturas de flujo. Informe técnico CS2009-14, Departamento de Informática, Universidad de Virginia. Diciembre de 2009.