Tutoriel CUDA

Introduction

Dans ce tutoriel, nous présentons la composante de calcul hautement parallèle qu'est le processeur graphique (ou GPU pour graphics processing unit); le langage de programmation parallèle CUDA; et quelques-unes des librairies numériques CUDA utilisées en calcul de haute performance.

Prérequis

Ce tutoriel montre comment utiliser CUDA pour accélérer des programmes en C ou en C++; une bonne connaissance d'un de ces langages vous permettra d’en tirer le meilleur profit. Si CUDA sert aussi aux programmes en Fortran, nous nous limiterons ici à CUDA pour C/C++ et utiliserons le terme CUDA C. Il s'agit essentiellement de produire des fonctions en C/C++ pouvant être exécutées par les CPUs et les GPUs.

Objectifs d'apprentissage

• Comprendre l'architecture d'un GPU
• Comprendre le déroulement d'un programme CUDA
• Comprendre et gérer les différents types de mémoires GPU
• Écrire et compiler un exemple de code CUDA

Qu'est-ce qu'un GPU?

Un GPU (pour graphics processing unit) est un processeur monopuce capable d'effectuer des calculs mathématiques rapidement pour produire des rendus d'images. Depuis quelques années, la puissance du GPU sert aussi à accélérer l'exécution de calculs intensifs dans plusieurs domaines de la recherche scientifique de pointe.

Qu'est-ce que CUDA?

CUDA est l'abréviation de compute unified device architecture. Il s'agit d'un environnement logiciel flexible et d'un modèle de programmation pour le traitement de calculs parallèles intensifs permettant l'accès aux instructions et à la mémoire d'un GPU.

Architecture du GPU

Un GPU se compose  :

• d'une mémoire globale
  • semblable à la mémoire du CPU
  • accessible par un CPU et un GPU
• des multiprocesseurs en continu (SMs pour streaming multiprocessors)
  • chaque SM est formé de plusieurs processeurs en continu (SPs pour streaming processors)
  • ils effectuent les calculs
  • chaque SM est doté d'une unité de contrôle, de registres, de pipelines d'exécution, etc. qui lui sont propres

Modèle de programmation

Voyons d'abord quelques termes importants :

• Hôte : désigne le CPU et sa mémoire (mémoire hôte).
• Carte graphique : désigne le GPU et sa mémoire (mémoire de la carte graphique).

Le modèle CUDA est un modèle hétérogène où à la fois le CPU et le GPU sont utilisés. Le code CUDA peut gérer les deux types de mémoires : la mémoire hôte CPU et la mémoire de la carte graphique. Le code exécute aussi les fonctions du GPU appelées kernels (noyaux). Ces fonctions sont exécutées en parallèle par plusieurs fils GPU. Le processus comporte cinq étapes :

1. Déclaration et allocation de la mémoire hôte CPU et de la mémoire de la carte graphique.
2. Initialisation de la mémoire hôte.
3. Transfert des données de la mémoire hôte à la mémoire de la carte graphique.
4. Exécution des fonctions GPU (kernels).
5. Retour des données à la mémoire hôte.

Modèle d'exécution

Le code CUDA simple exécuté dans un GPU s'appelle kernel (noyau). Il faut se demander :

• comment faire pour exécuter un kernel sur un groupe de multiprocesseurs en continu?
• comment faire pour que ce kernel soit exécuté de façon parallèle intensive?

Voici la recette en réponse à ces questions :

• chaque cœur GPU (processeur en continu) exécute un fil (thread) séquentiel, ce qui est le plus petit ensemble discret d'instructions géré par l'ordonnanceur du système d'exploitation
• tous les cœurs GPU exécutent le kernel de manière simultanée selon le modèle SIMT (single instruction, multiple threads)

La procédure suivante est recommandée ː

1. Copier les données en entrée de la mémoire CPU à la mémoire GPU.
2. Charger puis lancer le programme GPU (le kernel).
3. Copier les résultats de la mémoire GPU à la mémoire CPU.

Blocs de fils

Les fils sont groupés en blocs qui forment une grille.

Pour obtenir un parallélisme intensif, on doit utiliser le plus le fils possible; puisqu'un kernel CUDA comprend un très grand nombre de fils, il faut bien les organiser. Avec CUDA, les fils sont groupés en blocs de fils, eux-mêmes formant une grille. Diviser les fils fait en sorte que :

• les fils groupés coopèrent via la mémoire partagée,
• les fils d'un bloc ne coopèrent pas avec les fils des autres blocs.

Selon ce modèle, les fils dans un bloc travaillent sur le même groupe d'instructions (mais peut-être avec des jeux de données différents) et s'échangent les données via la mémoire partagée. Les fils dans les autres blocs font de même (voir la figure).

Intercommunication via la mémoire partagée des fils dans un bloc.

Chaque fil utilise des identifiants (IDs) pour décider quelles données utiliser :

• IDs des blocs : 1D ou 2D (blockIdx.x, blockIdx.y)
• IDs des fils : 1D, 2D, ou 3D (threadIdx.x, threadIdx.y, threadIdx.z)

Ce modèle simplifie l'adressage de la mémoire lors du traitement de données multidimensionnelles.

Ordonnancement des fils

Un processeur en continu (SM) exécute habituellement un bloc de fils à la fois. Le code est exécuté en groupes de 32 fils (appelés warps). Un ordonnanceur physique est libre d'assigner des blocs à tout SM en tout temps. De plus, quand un SM reçoit le bloc qui lui est assigné, ceci ne signifie pas que ce bloc en particulier sera exécuté sans arrêt. En fait, l'ordonnanceur peut retarder/suspendre l'exécution de tels blocs selon certaines conditions, par exemple si les données ne sont plus disponibles (en effet, la lecture de données à partir de la mémoire globale du GPU exige beaucoup de temps). Lorsque ceci se produit, l'ordonnanceur exécute un autre bloc de fils qui est prêt à être exécuté. Il s'agit en quelque sorte d'ordonnancement zero-overhead favorisant un flux d'exécution plus régulier afin que les SMs ne demeurent pas inactifs.

Types de mémoire GPU

Plusieurs types de mémoire sont disponibles aux opérations CUDA :

• mémoire globale
  • non sur la puce (off-chip), efficace pour opérations I/O, mais relativement lente
• mémoire partagée
  • sur la puce (on-chip), permet une bonne collaboration des fils, très rapide
• registres et mémoire locale
  • espace de travail des fils, très rapide
• mémoire constante

Quelques opérations de base

Allocation de la mémoire

• cudaMalloc((void**)&array, size)
  • Allocation d'objet dans la mémoire de la carte graphique. Exige l'adresse d'un pointeur vers les données allouées et la taille.
• cudaFree(array)
  • Désallocation de l'objet dans la mémoire. Exige uniquement le pointeur vers les données.

Transfert de données

• cudaMemcpy(array_dest, array_orig, size, direction)
  • Copie les données de la carte graphique vers l'hôte ou de l'hôte vers la carte graphique. Exige les pointeurs vers les données, la taille et le type de direction (cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, etc.)
• cudaMemcpyAsync
  • Identique à cudaMemcpy, mais transfère les données de manière asynchrone, ce qui signifie que l'exécution des autres processus n'est pas bloquée.

Exemple d’un programme CUDA C simple

Dans cet exemple, nous ajoutons deux nombres au GPU. Il s'agit d'un exemple très simple et il ne faut pas s'attendre à observer une grande accélération.

___global__ void add (int *a, int *b, int *c){
  *c = *a + *b;
}

int main(void){
  int a, b, c;
  int *dev_a, *dev_b, *dev_c;
  int size = sizeof(int);

  //  allocate device copies of a,b, c
  cudaMalloc ( (void**) &dev_a, size);
  cudaMalloc ( (void**) &dev_b, size);
  cudaMalloc ( (void**) &dev_c, size);

  a=2; b=7;
  //  copy inputs to device
  cudaMemcpy (dev_a, &a, size, cudaMemcpyHostToDevice);
  cudaMemcpy (dev_b, &b, size, cudaMemcpyHostToDevice);

  // launch add() kernel on GPU, passing parameters
  add <<< 1, 1 >>> (dev_a, dev_b, dev_c);

  // copy device result back to host
  cudaMemcpy (&c, dev_c, size, cudaMemcpyDeviceToHost);

  cudaFree ( dev_a ); cudaFree ( dev_b ); cudaFree ( dev_c );
}

Il manque certainement quelque chose; ce code n’a pas une allure parallèle… Comme solution, modifions le contenu du kernel entre les signes <<< et >>>.

add <<< N, 1 >>> (dev_a, dev_b, dev_c);

Ici, nous avons remplacé 1 par N pour que N blocs CUDA différents soient exécutés en même temps. Pour paralléliser cependant, il faut aussi faire des modifications au kernel :

__global__   void add (int *a, int *b, int *c){
  c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];

où blockIdx.x est le numéro unique identifiant un bloc CUDA. De cette manière, chaque bloc CUDA ajoute une valeur de a[ ] à b[ ].

Parallélisation basée sur les blocs.

Modifions à nouveau le contenu entre les crochets triples.

add <<< 1, '''N''' >>> (dev_a, dev_b, dev_c);

La tâche est maintenant distribuée sur des fils parallèles plutôt que sur des blocs. Quel est l'avantage des fils parallèles? Contrairement aux blocs, les fils peuvent communiquer ensemble; autrement dit, nous parallélisons sur plusieurs fils dans le bloc quand la communication est intense. Les portions de code qui peuvent être exécutées indépendamment, soit avec peu ou pas de communication, sont distribuées sur des blocs parallèles.

Avantages de la mémoire partagée

Jusqu'ici, tous les transferts en mémoire dans le kernel ont été via la mémoire régulière (globale) du GPU, ce qui est relativement lent. Il y a souvent tellement de communication entre fils que la performance est significativement diminuée. Pour contrer ce problème, nous pouvons utiliser la mémoire partagée qui peut accélérer les transferts en mémoire entre les fils. Le secret par contre est que seuls les fils du même bloc peuvent communiquer. Pour illustrer l'utilisation de cette mémoire, voyons l'exemple du produit scalaire où deux vecteurs sont multipliés élément par élément et additionnés par la suite, ainsi :

__global__   void dot(int *a, int *b, int *c){
  int temp = a[threadIdx.x]*b[threadIdx.x]; 
}

Après que chaque fil a exécuté sa portion, il faut tout additionner; chaque fils doit partager ses données. Toutefois, le problème est que chaque copie de la variable temporaire du fil est privée. La solution est d'utiliser la mémoire partagée avec les modifications suivantes au kernel :

#define N 512
__global__   void dot(int *a, int *b, int *c){
  __shared__ int temp[N];
  temp[threadIdx.x] = a[threadIdx.x]*b[threadIdx.x];
  __syncthreads();
  if(threadIdx.x==0){
    int sum; for(int i=0;i<N;i++) sum+= temp[i];
    *c=sum;
  }
}

Facteurs de performance de base

Transferts en mémoire

• PCI-e est extrêmement lent (4-6Go/s) en comparaison à la mémoire hôte et la mémoire de la carte graphique.
• Minimisez les copies de mémoire dans les deux directions.
• Gardez les données sur la carte graphique le plus longtemps possible.
• Il n'est parfois pas efficace d'utiliser l'hôte (le CPU) pour des tâches non optimales; il pourrait être plus rapide de les exécuter avec le GPU que de copier vers le CPU, exécuter et retourner le résultat.
• Utilisez les temps mémoire pour analyser les temps d'exécution.

Bande passante

• Tenez toujours compte des limites de la bande passante quand vous modifiez votre code.
• Connaissez la bande passante de pointe théorique des divers liens de données.
• Faites le décompte des octets écrits/lus et comparez avec la pointe théorique.
• Utilisez les divers types de mémoire selon le cas : globale, partagée, constante.

Stratégies usuelles de programmation

• La mémoire constante réside aussi dans DRAM; l'accès est beaucoup plus lent que pour la mémoire partagée.
  • MAIS, elle est cachée !!!
  • accès hautement efficace en lecture seule, transmission.
• Répartissez bien les données selon le mode d'accès :
  • lecture seule : mémoire constante (très rapide si dans la cache)
  • lecture/écriture dans le bloc : mémoire partagée (très rapide)
  • lecture/écriture dans le fil : registres (très rapide)
  • lecture/écriture en entrée/résultats : mémoire globale (très lent)