Exemple d'interblocage réaliste dans CUDA/OpenCL

Question

Pour un tutoriel que j'écris, je cherche un exemple "réaliste" et simple d'un blocage causé par l'ignorance de SIMT/SIMD.

Je suis venu avec cet extrait, qui semble être un bon exemple.

Toute contribution serait appréciée.

… 
int x = threadID/2; 
if (threadID > x) { 
    value[threadID] = 42; 
    barrier(); 
    } 
else { 
    value2[threadID/2] = 13 
    barrier(); 
} 
result = value[threadID/2] + value2[threadID/2]; 

Je sais, ce n'est ni bon CUDA C, ni OpenCL C.

Answer 1

Une impasse simple qui est en fait facile à attraper par le programmeur CUDA novice est quand on essaie de mettre en œuvre une section critique pour un seul thread, qui devrait finalement être effectué par tous les threads. Il va plus ou moins comme ceci:

__global__ kernel() { 
    __shared__ int semaphore; 
    semaphore=0; 
    __syncthreads(); 
    while (true) { 
    int prev=atomicCAS(&semaphore,0,1); 
    if (prev==0) { 
     //critical section 
     semaphore=0; 
     break; 
    } 
    } 
} 

L'instruction atomicCAS assure que exaclty un thread obtient 0 affecté à prev, tandis que les autres obtiennent 1. Lorsque qu'un thread a fini sa section critique, il définit la sémaphores retour à 0 pour que les autres threads aient une chance d'entrer dans la section critique. Le problème est que, alors que 1 thread obtient prev = 0, 31 threads, appartenant à la même unité SIMD, reçoivent une valeur 1. A l'instruction if CUDA scheduler place ce thread unique en attente (le masque) et laissez les autres 31 fils continuer leur travail. Dans des circonstances normales, c'est une bonne stratégie, mais dans ce cas particulier, vous vous retrouvez avec 1 thread de section critique qui n'est jamais exécuté et 31 threads qui attendent l'infini. Impasse.

Notez également, l'existence de break qui conduit le flux de contrôle en dehors de la boucle while. Si vous passez l'instruction break et que vous avez d'autres opérations après le if-block qui sont supposées être exécutées par tous les threads, cela peut aider le planificateur à éviter le blocage. En ce qui concerne votre exemple donné dans la question: dans CUDA il est explicitement interdit de mettre __syncthreads() dans un code SIMD divergent. Le compilateur ne l'attrape pas mais le manuel parle de "comportement indéfini". En pratique, sur les dispositifs pré-Fermi, tous les __syncthreads() sont considérés comme les mêmes barrières. Avec cette supposition, votre code se terminerait réellement sans erreur. On devrait cependant s'appuyer sur ce comportement.