Comprendre les dimensions de la grille CUDA, les dimensions des blocs et l'organisation des fils (explication simple) [fermé]

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

Matériel

Si un périphérique GPU a, par exemple, 4 unités multiprocesseurs, et qu'ils peuvent exécuter 768 threads chacun: alors à un moment donné, pas plus de 4 * 768 threads ne fonctionneront vraiment en parallèle (si vous avez planifié plus de threads, ils attendront leur tour).

Logiciel

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

Évidemment, si vous avez besoin de plus de 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 file d'attente de blocs en attente pour entrer dans le GPU (car, dans notre exemple, le GPU dispose de 4 multiprocesseurs et seuls 4 blocs sont exécutés simultanément).

Maintenant, un cas simple: traiter une image 512x512

Supposons que nous voulons 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 (donc pour avoir 512x512 threads = 4096 * 64)

Il est courant d'organiser (pour faciliter l'indexation de l'image) les threads en 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 4096 blocs nécessaires). Je préfère l'appeler numBlocks.

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

Le noyau est lancé comme ceci:

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

Enfin: il y aura quelque chose comme "une file d'attente de 4096 blocs", où un bloc attend d'être affecté à l'un des multiprocesseurs du GPU pour que ses 64 threads soient exécutés.

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

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

Supposons un GPU 9800GT:

• il dispose de 14 multiprocesseurs (SM)
• chaque SM dispose de 8 processeurs de thread (processeurs de flux AKA, SP ou cœurs)
• autorise jusqu'à 512 threads par bloc
• warpsize est de 32 (ce qui signifie que chacun des 14x8 = 112 processeurs de threads peut planifier jusqu'à 32 threads)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Un bloc ne peut pas avoir plus de threads actifs que 512 par conséquent __syncthreadsne peut synchroniser qu'un nombre limité de threads. ie Si vous exécutez ce qui suit avec 600 threads:

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

alors le noyau doit s'exécuter deux fois et l'ordre d'exécution sera:

1. func1 est exécuté pour les 512 premiers threads
2. func2 est exécuté pour les 512 premiers threads
3. func1 est exécuté pour les threads restants
4. func2 est exécuté pour les threads restants

Remarque:

Le point principal __syncthreadsest une opération à l'échelle du bloc et elle ne synchronise pas tous les threads.

Je ne suis pas sûr du nombre exact de threads qui __syncthreadspeuvent se synchroniser, car vous pouvez créer un bloc avec plus de 512 threads et laisser le warp gérer la planification. À ma connaissance, il est plus précis de dire: func1 est exécuté au moins pour les 512 premiers threads.

Avant de modifier cette réponse (en 2010), j'ai mesuré que les threads 14x8x32 étaient synchronisés à l'aide de __syncthreads.

J'apprécierais beaucoup que quelqu'un teste à nouveau cela pour obtenir une information plus précise.