Barrière entre blocs sur CUDA

Je souhaite implémenter une barrière interbloc sur CUDA, mais je rencontre un problème sérieux.

Je ne peux pas comprendre pourquoi cela ne fonctionne pas.

#include  #include  #include  #define SIZE 10000000 #define BLOCKS 100 using namespace std; struct Barrier { int *count; __device__ void wait() { atomicSub(count, 1); while(*count) ; } Barrier() { int blocks = BLOCKS; cudaMalloc((void**) &count, sizeof(int)); cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice); } ~Barrier() { cudaFree(count); } }; __global__ void sum(int* vec, int* cache, int *sum, Barrier barrier) { int tid = blockIdx.x; int temp = 0; while(tid < SIZE) { temp += vec[tid]; tid += gridDim.x; } cache[blockIdx.x] = temp; barrier.wait(); if(blockIdx.x == 0) { for(int i = 0 ; i < BLOCKS; ++i) *sum += cache[i]; } } int main() { int* vec_host = (int *) malloc(SIZE * sizeof(int)); for(int i = 0; i < SIZE; ++i) vec_host[i] = 1; int *vec_dev; int *sum_dev; int *cache; int sum_gpu = 0; cudaMalloc((void**) &vec_dev, SIZE * sizeof(int)); cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice); cudaMalloc((void**) &sum_dev, sizeof(int)); cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice); cudaMalloc((void**) &cache, BLOCKS * sizeof(int)); cudaMemset(cache, 0, BLOCKS * sizeof(int)); Barrier barrier; sum<<>>(vec_dev, cache, sum_dev, barrier); cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(vec_dev); cudaFree(sum_dev); cudaFree(cache); free(vec_host); return 0; } 

En fait, même si je réécris wait () comme suit

  __device__ void wait() { while(*count != 234124) ; } 

Le programme se termine normalement. Mais je m’attends à obtenir une boucle infinie dans ce cas.

Malheureusement, ce que vous souhaitez réaliser (communication / synchronisation interbloc) n’est pas ssortingctement possible dans CUDA. Le guide de programmation CUDA indique que “les blocs de threads doivent être exécutés indépendamment: ils doivent pouvoir être exécutés dans n’importe quel ordre, en parallèle ou en série”. La raison de cette ressortingction est de permettre une flexibilité dans le planificateur de blocs de thread et de permettre au code de s’adapter de manière agnostique au nombre de cœurs. La seule méthode de synchronisation inter-blocs prise en charge consiste à lancer un autre kernel: les lancements de kernel (au sein du même stream) sont des points de synchronisation implicites.

Votre code enfreint la règle d’indépendance des blocs car il suppose implicitement que les blocs de threads de votre kernel s’exécutent simultanément (cf. en parallèle). Mais il n’y a aucune garantie qu’ils font. Pour voir pourquoi cela est important pour votre code, considérons un GPU hypothétique avec un seul cœur. Nous supposerons également que vous ne voulez lancer que deux blocs de thread. Votre kernel spinloop bloquera réellement dans cette situation. Si le bloc de threads zéro est programmé en premier sur le kernel, il sera bouclé à jamais lorsqu’il atteindra la barrière, car le bloc de thread 1 n’a jamais la possibilité de mettre à jour le compteur. Etant donné que le bloc de threads zéro n’est jamais échangé (les blocs de threads sont exécutés jusqu’à leur achèvement), il empêche le bloc de threads de se centrer pendant son rotation.

Certaines personnes ont essayé des schémas tels que le vôtre et ont connu du succès, car le planificateur a programmé des blocs de manière fortuite, de manière à ce que les hypothèses soient vérifiées. Par exemple, il y a eu un moment où lancer autant de blocs de threads qu’un GPU a de SM signifie que les blocs sont réellement exécutés simultanément. Mais ils ont été déçus quand un changement de pilote, de runtime CUDA ou de GPU a invalidé cette hypothèse, cassant ainsi leur code.

Pour votre application, essayez de trouver une solution qui ne dépend pas de la synchronisation entre blocs, car (sauf modification significative du modèle de programmation CUDA), cela n’est tout simplement pas possible.

Bloquer pour bloquer la synchronisation est possible. Voir ce papier .
Le document n’entre pas dans les détails sur son fonctionnement, mais il repose sur l’opération de __syncthreads (); pour créer la barrière de pause pour le bloc actuel, … en attendant que les autres blocs atteignent le sharepoint synchronisation.

Un élément qui n’a pas été noté dans le document est que la synchronisation n’est possible que si le nombre de blocs est suffisamment petit ou si le nombre de MS est suffisamment grand pour la tâche à accomplir. Par exemple, si vous avez 4 SM et que vous essayez de synchroniser 5 blocs, le kernel se bloquera.

Grâce à leur approche, j’ai été en mesure de répartir une longue tâche en série sur plusieurs blocs, ce qui représente un gain de temps considérable de 30% sur une approche mono-bloc. c’est-à-dire que la synchronisation des blocs a fonctionné pour moi.

Ressemble à un problème d’optimisation du compilateur. Je ne suis pas doué pour lire le code PTX, mais il semblerait que le compilateur ait omis la boucle while (du tout) (même lorsqu’il a été compilé avec -O0 ):

 .loc 3 41 0 cvt.u64.u32 %rd7, %ctaid.x; // Save blockIdx.x to rd7 ld.param.u64 %rd8, [__cudaparm__Z3sumPiS_S_7Barrier_cache]; mov.s32 %r8, %ctaid.x; // Now calculate ouput address mul.wide.u32 %rd9, %r8, 4; add.u64 %rd10, %rd8, %rd9; st.global.s32 [%rd10+0], %r5; // Store result to cache[blockIdx.x] .loc 17 128 0 ld.param.u64 %rd11, [__cudaparm__Z3sumPiS_S_7Barrier_barrier+0]; // Get *count to rd11 mov.s32 %r9, -1; // put -1 to r9 atom.global.add.s32 %r10, [%rd11], %r9; // Do AtomicSub, storing the result to r10 (will be unused) cvt.u32.u64 %r11, %rd7; // Put blockIdx.x saved in rd7 to r11 mov.u32 %r12, 0; // Put 0 to r12 setp.ne.u32 %p3, %r11, %r12; // if(blockIdx.x == 0) @%p3 bra $Lt_0_5122; ld.param.u64 %rd12, [__cudaparm__Z3sumPiS_S_7Barrier_sum]; ld.global.s32 %r13, [%rd12+0]; mov.s64 %rd13, %rd8; mov.s32 %r14, 0; 

En cas de code CPU, un tel comportement est empêché en déclarant la variable avec un préfixe volatile . Mais même si nous déclarons count tant que int __device__ count (et modifions le code de manière appropriée), l’ajout d’un spécificateur volatile interrompt simplement la compilation (avec un argument of type "volatile int *" is incompatible with parameter of type "void *" loke d’erreur argument of type "volatile int *" is incompatible with parameter of type "void *" )

Je suggère de regarder exemple threadFenceReduction de CUDA SDK. Là, ils font à peu près la même chose que vous, mais le bloc pour effectuer la sommation finale est choisi au moment de l’exécution, plutôt que prédéfini, et la boucle while est éliminée, car le locking en rotation d’une variable globale doit être très lent.