CUDA et Programmation Générale sur GPU

Comment Programmer ?

Architecture de la solution

Le kit de développement de CUDA est composé de plusieurs éléments :

Au sommet, le code source du programme utilisant le CPU et le GPU, et des bibliothèques optimisées fournies par nVIDIA facilitant le travail des développeurs En dessous, le compilateur (nvcc) prends ces fichiers en entrée et sépare le code GPU du code CPU. Il compile code du GPU en assembleur propriétaire (PTX) et compile le code du CPU avec GCC. Le code en assembleur GPU (PTX) passe, lors de l'exécution, par le pilote CUDA pour la carte graphique ou par le débugger pour etre exécuté sur le GPU Le code du CPU est exécuté en parallèle sur le CPU. C'est lui qui, cependant, déclenche l'exécution du code sur le GPU.

Comment Programmer sur GPU ?

Pour programmer l'exécution de code sur un GPU avec CUDA, il faut définir le bout de code à exécuter (qui sera exécuté en une multitude de threads), et lancer son exécution depuis le thread principal s'exécutant sur le CPU.

Ainsi, le bout de code à exécuter sur le GPU est défini par le programmeur en une fonction en C respectant certaines contraintes, et est appelé "kernel" selon la terminologie CUDA.

Également, l'exécution d'un kernel sur un GPU réponds à une certaine syntaxe, où le programmeur indique le nombre de threads exécutant son kernel, sous quelle organisation, avec une certaine quantité de mémoire partagée, etc.).

Les Kernels CUDA

Un kernel CUDA est une fonction C soumise à certaines contraintes :

• Identifidateurs précédant la déclaration de la fonction :
  • __device__ pour un kernel appelable uniquement depuis le GPU
  • __global__ pour un kernel appelable uniquement depuis le CPU
  • ces mots clés peuvent être combinés
• Pas d'accès direct à la mémoire centrale de l'ordinateur (RAM)
• retour void obligatoire
• Pas d'arguments variables (varargs)
• Pas de récursivité
• Pas de variables statique (à ne pas confondre avec l'espace mémoire statique, cf. partie sur la mémoire)

 

Le lancement d'un kernel s'effectue de la façon suivante :

myKernel<<< dimGrid, dimBlock[, dimMem ]>>>(params); Où : 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 ...

Comme exemple simple, voici deux bouts de code (un pour CPU et un pour GPU avec CUDA) qui ont le même but : incrémenter de 1 chaque case d'un tableau :

Code pour CPU : Dans ce cas, le thread principal appelle la fonction qui, dans une boucle, incrémente chaque case du tableau.

Code pour GPU : Ici, on déclare un kernel qui, en premier lieu, calcule son indice de thread dans sa grille (indice du bloc dans la grille x taille des blocs + indice du thread dans son bloc). Ensuite, il se sert de cet indice pour incrémenter la case correspondante du tableau. Ainsi, dans le thread principal, on lance l'exécution du kernel avec autant de threads que de case du tableau à incrémenter. À noter, les variables blockIdx, blockDim et threadIdx sont fournies par CUDA au kernel, et dépendent des valeurs données lors de son lancement. Ces variables disposent d'attributs x, y et z pour faciliter leur repérage dans l'espace, mais ici nous n'utilisons que l'attribut x car nous avons spécifié qu'une seule dimension lors de la déclaration des variables dimBlock et dimGrid de type dim3 dans le main().