Comprendre les dimensions de la grid CUDA, les dimensions des blocs et l’organisation des threads (explication simple)

Comment les threads sont-ils organisés pour être exécutés par un GPU?

Matériel

Si un périphérique GPU dispose, par exemple, de 4 unités de multitraitement, et qu’il peut exécuter 768 threads chacun: alors, à un moment donné, pas plus de 4 * 768 threads ne fonctionneront réellement en parallèle (si vous planifiez plus de threads, ils attendent leur tour).

Logiciel

les discussions sont organisées en blocs. Un bloc est exécuté par une unité de multitraitement. Les threads d’un bloc peuvent être identifiés (indexés) en utilisant 1Dimension (x), 2Dimensions (x, y) ou les index 3Dim (x, y, z) mais en tout cas x y z <= 768 pour notre exemple (d'autres restrictions s'appliquent) à x, y, z, voir le guide et la capacité de votre appareil).

De toute évidence, si vous avez besoin de plus que ces 4 * 768 threads, vous avez besoin de plus de 4 blocs. Les blocs peuvent également être indexés 1D, 2D ou 3D. Il y a une queue de blocs en attente d’entrer dans le GPU (car, dans notre exemple, le GPU possède 4 multiprocesseurs et 4 blocs seulement sont exécutés simultanément).

Maintenant, un cas simple: traiter une image 512×512

Supposons que nous voulions qu’un thread traite un pixel (i, j).

Nous pouvons utiliser des blocs de 64 threads chacun. Ensuite, nous avons besoin de 512 * 512/64 = 4096 blocs (pour avoir 512×512 threads = 4096 * 64)

Il est courant d’organiser (pour faciliter l’indexation de l’image) les threads dans les blocs 2D ayant blockDim = 8 x 8 (les 64 threads par bloc). Je préfère l’appeler threadsPerBlock.

dim3 threadsPerBlock(8, 8); // 64 threads 

et 2D gridDim = 64 x 64 blocs (les blocs 4096 nécessaires). Je préfère l’appeler numBlocks.

 dim3 numBlocks(imageWidth/threadsPerBlock.x, /* for instance 512/8 = 64*/ imageHeight/threadsPerBlock.y); 

Le kernel est lancé comme ceci:

 myKernel <<>>( /* params for the kernel function */ ); 

Enfin: il y aura quelque chose comme “une queue de 4096 blocs”, où un bloc attend d’être assigné à l’un des multiprocesseurs du GPU pour exécuter ses 64 threads.

Dans le kernel, le pixel (i, j) à traiter par un thread est calculé de la manière suivante:

 uint i = (blockIdx.x * blockDim.x) + threadIdx.x; uint j = (blockIdx.y * blockDim.y) + threadIdx.y; 

supposons un GPU 9800GT: 14 multiprocesseurs, chacun ayant 8 processeurs de threads et warpsize 32, ce qui signifie que chaque processeur de thread gère jusqu’à 32 threads. 14 * 8 * 32 = 3584 est le nombre maximum de filetages à courant continu.

si vous exécutez ce kernel avec plus de 3584 threads (disons 4000 threads et que ce n’est pas important de définir le bloc et la grid), gpu les traitera de la même manière:

 func1(); __syncthreads(); func2(); __syncthreads(); 

alors l’ordre d’exécution de ces deux fonctions est le suivant:

1.func1 est exécuté pour les premiers 3584 threads

2.func2 est exécuté pour les premiers 3584 threads

3.func1 est exécuté pour les threads restants

4.func2 est exécuté pour les threads restants