Macro CONFLICT_FREE_OFFSET utilisée dans l'algorithme de préfixe parallèle de GPU Gems 3

Question

Tout d'abord, voici le lien vers l'algorithme:

Gemmes GPU 3, Chapitre 39: Somme des préfixes parallèles (Scan) avec CUDA.

Afin d'éviter les conflits de banques, un remplissage est ajouté au tableau de mémoire partagée tous les NUM_BANKS (c'est-à-dire 32 pour les périphériques de calculabilité 2.x) éléments.Ceci est fait par (comme dans la figure 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]

Je ne comprends pas en quoi ai / NUM_BANKS est équivalent à la macro:

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

N'est-ce pas égal à

n >> LOG_NUM_BANKS

Toute aide est appréciée.Je vous remercie

Answer 1

J'ai écrit ce code et co-écrit l'article, et je vous demande d'utiliser l'article uniquement pour en apprendre davantage sur les algorithmes de balayage, et de ne pas utiliser le code qu'il contient.Il a été écrit quand CUDA était nouveau, et j'étais nouveau sur CUDA.Si vous utilisez une implémentation moderne de scan dans CUDA, vous n'avez pas besoin d'éviter les conflits bancaires.

Si vous souhaitez effectuer des analyses de manière simple, utilisez thrust::inclusive_scan ou thrust::exclusive_scan.

Si vous voulez vraiment implémenter une analyse, reportez-vous à des articles plus récents comme celui-ci [1].Ou pour un vrai opus avec un code plus rapide mais qui demandera un peu plus d'étude, celui-ci [2].Ou lire Tutoriel de Sean Baxter (bien que ce dernier n'inclue pas de citations des travaux fondateurs sur l'algorithme de balayage).

[1] Shubhabrata Sengupta, Mark Harris, Michael Garland et John D.Owens."Algorithmes d'analyse parallèle efficaces pour les GPU à plusieurs cœurs".Dans le livre de Jakub Kurzak, David A.Bader et Jack Dongarra, éditeurs, Calcul scientifique avec multicœur et accélérateurs, Chapman & Hall/CRC Computational Science, chapitre 19, pages 413-442.Taylor et Francis, janvier 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1041

[2] Merrill, auteur principal.et Grimshaw, A.Balayage parallèle pour les architectures de flux.Rapport technique CS2009-14, Département d'informatique, Université de Virginie.Déc.2009.