Accès à la mémoire CUDA atomique et non atomique

Question

J'ai deux fonctions CUDA qui manipulent des listes liées dans la mémoire globale. La fonction pmalloc supprime l'élément head de l'une des listes. Il choisit d'abord une liste, puis appelle pmallocBucket qui supprime réellement l'élément principal. Si la liste choisie est vide, pmalloc va essayer d'autres listes. D'autre part, la fonction pfree va insérer un nouvel élément head dans une liste.

L'exclusion mutuelle est réalisée par des sémaphores, un pour chaque liste liée. L'implémentation pour les sémaphores provient du livre CUDA Par l'exemple. Dans un autre code de test, le sémaphore fonctionne parfaitement.

Le problème que j'ai avec le code est le suivant: Parfois, plusieurs threads vont essayer d'accéder simultanément à la même liste chaînée. Ces accès sont séquencés avec succès par le sémaphore, mais parfois, un thread enlève le même élément head de la liste qu'un thread précédent. Cela peut arriver immédiatement consécutivement, ou il peut y avoir un ou plusieurs autres threads entre les deux. Le thread va alors free une zone de mémoire non allouée et mon programme se bloque.

Voici les fonctions mentionnées. mmd est une structure dans la mémoire globale qui est initialisée à partir d'une autre fonction.

extern __device__ void wait(int* s) { 
    while(atomicCAS(s, 0, 1) != 0); 
} 

extern __device__ void signal(int* s) { 
    atomicExch(s, 0); 
} 

__device__ void pfree(Expression* node) { 
    LinkedList* l = (LinkedList*) malloc(sizeof(LinkedList)); 
    l->cell = node; 
    node->type = EMPTY; 
    node->funcidx = 0; 
    node->name = NULL; 
    node->len = 0; 
    node->value = 0; 
    node->numParams = 0; 
    free(node->params); 

    int targetBin = (blockIdx.x * mmd.bucketSize + threadIdx.x)/BINSIZE; 
    /* 
    * The for loop and subsequent if are necessary to make sure that only one 
    * thread in a warp is actively waiting for the lock on the semaphore. 
    * Leaving this out will result in massive headaches. 
    * See "CUDA by example", p. 273 
    */ 

    for(int i = 0; i < WARPSIZE; i++) { 
    if(((threadIdx.x + blockIdx.x * blockDim.x) % WARPSIZE) == i) { 
     wait(&mmd.bucketSemaphores[targetBin]); 
     l->next = mmd.freeCells[targetBin]; 
     mmd.freeCells[targetBin] = l; 
     signal(&mmd.bucketSemaphores[targetBin]); 
    } 
    } 
} 

__device__ Expression* pmalloc() { 
    Expression* retval = NULL; 
    int i = 0; 

    int bucket = (blockIdx.x * mmd.bucketSize + threadIdx.x)/BINSIZE; 

    while(retval == NULL && i < mmd.numCellBins) { 
    retval = pmallocBucket((i + bucket) % mmd.numCellBins); 
    i++; 
    } 

    if(retval == NULL) { 
    printf("(%u, %u) Out of memory\n", blockIdx.x, threadIdx.x); 
    } 

    return retval; 
} 

__device__ Expression* pmallocBucket(int bucket) { 
    Expression* retval = NULL; 

    if(bucket < mmd.numCellBins) { 
    LinkedList* l = NULL; 

    for(int i = 0; i < WARPSIZE; i++) { 
     if(((threadIdx.x + blockIdx.x * blockDim.x) % WARPSIZE) == i) { 
     wait(&mmd.bucketSemaphores[bucket]); 
      l = mmd.freeCells[bucket]; 

      if(l != NULL) { 
      retval = l->cell; 
      mmd.freeCells[bucket] = l->next; 
      } 
     signal(&mmd.bucketSemaphores[bucket]); 
     free(l); 
     } 
    } 
    } 

    return retval; 
} 

Je suis tout à fait à court de ressources. Je ne sais pas ce qui se passe réellement et toutes mes tentatives jusqu'ici pour l'éclaircir ont été infructueuses. Toute aide est grandement appréciée. P.: Oui, je réalise que l'utilisation des opérations atomiques et des sémaphores est loin d'être idéale pour les applications CUDA. Mais dans ce cas, pour l'instant je n'ai aucune idée de comment cela pourrait être mis en œuvre différemment et mon projet est sur un délai absolument fixe qui approche très rapidement, donc cela devra faire.

Answer 1

Vous devez vous assurer que la manipulation de liste est entièrement effectuée dans la section critique protégée par le sémaphore, sans utiliser de données obsolètes avant l'acquisition du sémaphore.

Déclarez l->next et mmd.freeCells comme volatils, ou manipulez-les par des fonctions atomiques (atomicExch()).

Vous pouvez également utiliser un assemblage en ligne avec des opérateurs de cache appropriés. L'utilisation de mov.cg pour la charge devrait suffire à garantir qu'aucune valeur localement mise en cache n'est utilisée, avec un __threadfence() juste avant le signal() pour s'assurer que l'écriture a atteint la mémoire globale avant la publication du sémaphore. Assurez-vous d'utiliser asm volatile(...), ou encore le compilateur serait libre de déplacer l'entier en ligne asm hors de la section critique.