GPU Computing CUDA

Architecture Logicielle

Présentation

La Runtime CUDA est une extension du langage C++, et permet de dialoguer avec la partie matérielle. Cette extension inclut les éléments suivants:

• Mots clefs : ils permettent de définir par quel composant la fonction sera appélée, ainsi où celle-ci sera appelée.
• __device__ : kernel exécuté sur le GPU mais appelé par le CPU
• __global__ : kernel exécuté et appelé par le GPU
• __host__ : fonction exécutée et appelée par le CPU.


• Variables : elles permettent d'identifier le thread en cours d'exécution. Elles sont de type dim3.
• blockIdx : index du bloc dans la grille
• threadIdx : index du thread dans le bloc
• blockDim : nombre de threads par bloc

Les kernels

Les kernels sont des fonctions que nous allons executer sur le GPU. Chaque instance de ce code est appelé un thread . Un seul kernel est exécuté sur le GPU à la fois, mais ne posséde pas la même localisation dans la grille. Ainsi, nous pouvons définir un traitement pour chaque thread en fonction de son indice.

Ces codes possèdent certaines propriétés :

• Pas d'accès direct à la mémoire centrale de l'ordinateur (RAM)
• Ne retourne aucune donnée, le type de retour des fonctions est void
• Pas d'arguments variables (varargs)
• Pas de récursivité

Un kernel ne peut accéder à la mémoire centrale. Si nous souhaitons accéder/modifier des variables situées dans la mémoire centrale, par le GPU, nous devons dans un premier temps, le charger dans la VRAM (Video RAM), grâce à la fonction cudaMemcpyHostToDevice(). Ensuite, le recharger dans la mémoire centrale grâce la fonction cudaMemcpyDeviceToHost().

Il est possible de synchroniser les threads d'un bloc avec la fonction __syncthreads(). Cette fonction peut être particulièrement utile lorsque l'on souhaite attendre que tous les threads du bloc aient, à la suite de calculs, finis d'inscrire leurs résultats dans la mémoire partagée.

Lancement d'un Kernel

Le lancement d'un kernel s'effectue de la maniére suivante :

								
							kernel<<< dimGrid, dimBlock[, dimMem ]>>>(params);
							

• kernel : Nom de la fonction
• dimGrid : taille de la grille (en nombre de blocs)
• dimBlock : taille de chaque bloc (en nombre de threads)
• dimMem (optionel) : taille de la mémoire partagée allouée par bloc

Le lancement d'un kernel est non bloquant pour le thread qui l'exécute.

Un exemple

Dans cet exemple, nous allons incrémenter la valeur d'un vecteur d'une constante passée en paramétre. Le tableau représente le vecteur de dimension n.

				
					__global__ void incrementVectorGPU(int *a, int b, int n)
					{
						/* Calcul de la position dans la grille */
						int idx = blockIdx.x * blockDim.x + threadIdx.x;
						
						/* Test pour éviter le débordement de la taille du vecteur */
						if (idx < n )
							a[idx] = a[idx] + b;
					}

						/* Exemple d'éxécution de Kernel */
					int main(void)
					{						
						dim3 dimBlock(blockSize);
						dim3 dimGrid(N/(int)blockSize);
						
						incrementVectorGPU<< dimGrid, dimBlock>>>(vector1, scalar, n);
					}
				

Explication

Dans notre exemple, nous avons définis que notre vecteur possède une taille de 16, et que la taille d'un block est de 4. Ainsi, pour chaque block, les informations sont les suivantes :

Grâce aux différents indices calculés 'idx', nous pouvons voir que chaque case du tableau est bien identifiée par un Thread qui exécutera le calcul sur un seul indice du tableau. La condition permet simplement de ne pas déborder du tableau, si la taille de la grille, ou du block est plus important que la taille du vecteur.

					Int idx = blockDim.x * blockId.x+ threadIdx.x 
				

Schema récapitulatif


Chaque thread s'occupant d'un indice du tableau, exécutés simultanement, il ne faut qu'un cylce GPU pour additionner un scalaire à un vecteur, alors que pour le CPU il aurait fallu 'n' cycles pour effectuer ce calcul.

Exemple code complet

Présentation du code de lancement d'un Kernel. Comme je l'ai précisé, il est impossible à partir d'un Kernel d'accéder à la mémoire centrale, donc nous devons copier les données dans la VRAM.

					// allocation de la mémoire du vecteur
					unsigned int numBytes = N * sizeof(float)
					float* h_A = (float*) malloc (numBytes);

					// allocation de la mémoire sur la carte graphique
					float* d_A = 0;
					cudaMalloc ((void**)&d_A, numbytes);

					// copie des données de la mémoire centrale vers VRAM
					cudaMemcpy(d_A, h_A, numBytes, cudaMemcpyHostToDevice);

					// exécution du programme
					increment_gpu<<< N/blockSize, blockSize>>>(d_A, b);

					// rapatriement des données dans la mémoire centrale
					cudaMemcpy(h_A, d_A, numBytes, cudaMemcpyDeviceToHost);

					// libération de la VRAM
					cudaFree(d_A);