Implémentation d'une section critique dans CUDA

Ok, j'ai compris, et c'est encore un autre-un-de-la-cuda-paradigme-douleurs.

Comme tout bon cuda programmeur sait (avis que je ne me souviens pas de ce qui fait de moi une mauvaise cuda programmeur, je crois), tous les threads d'un warp doivent exécuter le même code. Le code que j'ai écrit fonctionne parfaitement si ce n'est pour ce fait. Comme il est, toutefois, il est fort probable que les deux fils dans la même courbure accès à la même serrure. Si l'un d'entre eux acquiert le verrou, il a juste oublie de l'exécution de la boucle, mais il ne peut pas continuer au-delà de la boucle jusqu'à ce que tous les autres threads dans sa courbure ont terminé la boucle. Malheureusement, l'autre thread ne sera jamais complète, car il est en attente pour le premier à débloquer.

Ici est un noyau qui va faire le tour sans erreur:

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    bool leaveLoop = false;
    while (!leaveLoop) {
        if (atomicExch(&(locks[id]), 1u) == 0u) {
            //critical section
            leaveLoop = true;
            atomicExch(&(locks[id]),0u);
        }
    } 
}

L'affiche a déjà trouvé une réponse à sa propre question. Néanmoins, dans le code ci-dessous, je suis en fournissant un cadre général pour mettre en œuvre un section critique en CUDA. Plus en détail, le code exécute un bloc de comptage, mais il est facilement modifiable pour accueillir d'autres opérations qui doivent être effectuées dans un section critique. Ci-dessous, je suis également état d'une explication du code, avec certains, "typique" des erreurs dans la mise en œuvre de sections critiques dans CUDA.

LE CODE

#include <stdio.h>
#include "Utilities.cuh"
#define NUMBLOCKS  512
#define NUMTHREADS 512 * 2
/***************/
/* LOCK STRUCT */
/***************/
struct Lock {
int *d_state;
//--- Constructor
Lock(void) {
int h_state = 0;                                        //--- Host side lock state initializer
gpuErrchk(cudaMalloc((void **)&d_state, sizeof(int)));  //--- Allocate device side lock state
gpuErrchk(cudaMemcpy(d_state, &h_state, sizeof(int), cudaMemcpyHostToDevice)); //--- Initialize device side lock state
}
//--- Destructor
__host__ __device__ ~Lock(void) { 
#if !defined(__CUDACC__)
gpuErrchk(cudaFree(d_state)); 
#else
#endif  
}
//--- Lock function
__device__ void lock(void) { while (atomicCAS(d_state, 0, 1) != 0); }
//--- Unlock function
__device__ void unlock(void) { atomicExch(d_state, 0); }
};
/*************************************/
/* BLOCK COUNTER KERNEL WITHOUT LOCK */
/*************************************/
__global__ void blockCountingKernelNoLock(int *numBlocks) {
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
}
/**********************************/
/* BLOCK COUNTER KERNEL WITH LOCK */
/**********************************/
__global__ void blockCountingKernelLock(Lock lock, int *numBlocks) {
if (threadIdx.x == 0) {
lock.lock();
numBlocks[0] = numBlocks[0] + 1;
lock.unlock();
}
}
/****************************************/
/* BLOCK COUNTER KERNEL WITH WRONG LOCK */
/****************************************/
__global__ void blockCountingKernelDeadlock(Lock lock, int *numBlocks) {
lock.lock();
if (threadIdx.x == 0) { numBlocks[0] = numBlocks[0] + 1; }
lock.unlock();
}
/********/
/* MAIN */
/********/
int main(){
int h_counting, *d_counting;
Lock lock;
gpuErrchk(cudaMalloc(&d_counting, sizeof(int)));
//--- Unlocked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelNoLock << <NUMBLOCKS, NUMTHREADS >> >(d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the unlocked case: %i\n", h_counting);
//--- Locked case
h_counting = 0;
gpuErrchk(cudaMemcpy(d_counting, &h_counting, sizeof(int), cudaMemcpyHostToDevice));
blockCountingKernelLock << <NUMBLOCKS, NUMTHREADS >> >(lock, d_counting);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(&h_counting, d_counting, sizeof(int), cudaMemcpyDeviceToHost));
printf("Counting in the locked case: %i\n", h_counting);
gpuErrchk(cudaFree(d_counting));
}

EXPLICATION DU CODE

Les sections critiques sont des séquences d'opérations qui doivent être exécutées de manière séquentielle par le CUDA de threads.

Suppose de construire un noyau qui a la tâche de calcul du nombre de thread se bloque d'un fil de grille. Une idée possible est de laisser chaque fil dans chaque bloc ayant threadIdx.x == 0 augmenter un compteur global. Pour éviter des conditions de course, toutes les augmentations se produisent de manière séquentielle, de sorte qu'ils doivent être intégrés dans une section critique.

Le code ci-dessus a deux fonctions du noyau: blockCountingKernelNoLock et blockCountingKernelLock. L'ancien ne pas utiliser une section critique pour augmenter le compteur et, comme on peut le voir, les retours de mauvais résultats. Ce dernier incarne le compteur augmenter à l'intérieur d'une section critique et produit des résultats corrects. Mais comment fonctionne la section critique du travail?

La section critique est régi par un état global d_state. D'abord, l'état est 0. En outre, deux __device__ méthodes, lock et unlockpeut modifier cet état. Le lock et unlock méthodes peuvent être invoquées que par un seul thread à l'intérieur de chaque bloc et, en particulier, par le thread ayant thread local de l'indice de threadIdx.x == 0.

Au hasard lors de l'exécution, l'un des fils ayant thread local de l'indice de threadIdx.x == 0 et global index de threads, disons, t sera la première invocation de la lock méthode. En particulier, il lancera atomicCAS(d_state, 0, 1). Car d'abord d_state == 0puis d_state sera mis à jour pour 1atomicCAS sera de retour 0 et le fil de sortie du lock fonction, en passant à la mise à jour de l'instruction. En attendant un tel thread effectue les operations mentionnees, tous les autres threads de tous les autres blocs ayant threadIdx.x == 0 va exécuter la lock méthode. Ils devront cependant trouver une valeur de d_state égal à 1de sorte que atomicCAS(d_state, 0, 1) effectuera pas de mise à jour et sera de retour 1afin de laisser ces threads d'exécution de la boucle while. Après ce thread t effectue la mise à jour, puis il exécute la unlock fonction, à savoir atomicExch(d_state, 0)rétablissant ainsi l' d_state à 0. À ce stade, au hasard, un autre des threads avec threadIdx.x == 0 se verrouille de nouveau à l'état.

Le code ci-dessus contient également une troisième fonction noyau, à savoir blockCountingKernelDeadlock. Cependant, c'est une autre mauvaise exécution de la section critique, conduisant à des blocages. En effet, nous rappelons que les croisements de fonctionner au même rythme et qu'ils se synchronisent après chaque instruction. Donc, quand nous exécutons blockCountingKernelDeadlockil y a la possibilité que l'un des threads d'un warp, dire un fil avec fil local de l'indice de t≠0verrouiller l'état. Dans ce cas, les autres threads du même courbure de tnotamment avec threadIdx.x == 0exécute les mêmes, alors que le rapport de boucle de fil tl'exécution de threads dans la même courbure effectuée au même rythme. En conséquence, tous les threads en attente pour quelqu'un pour déverrouiller l'état, mais aucun autre thread ne sera en mesure de le faire, et le code sera coincé dans une impasse.