[cours]    [mise en pratique]

Table des matières

2. Programmer avec CUDA

2.1. Introduction

Jusqu'en 2006, les GPU étaient difficiles à programmer car il fallait utiliser les API (Application Programming Interface) des fonctions graphiques (OpenGL et Direct3D) ce qui limitait les performances et le type d'application parallélisable.

A partir de 2007, NVidia met à disposition des programmeurs : CUDA (Compute Unified Device Architecture)

"CUDA is a parallel computing platform and programming model invented by NVIDIA. It enables dramatic increases in computing performance by harnessing* the power of the graphics processing unit (GPU)." (NVidia)

* to harness = exploiter

La programmation en CUDA repose sur deux modèles :

Le problème majeur que l'on rencontre en CUDA est que les modèles logique et physique sont fortement couplés et notamment le modèle physique influe de manière importante sur l'organisation des threads

Trouver la solution la plus efficace pour résoudre un problème donné n'est pas simple comme nous le verrons. Il faut prendre en compte :

2.2. Spécificités

Le GPU est en fait un véritable ordinateur miniature car il possède de nombreux coeurs de calcul ainsi que de la mémoire globale, de la mémoire cache (L1, L2), de la mémoire partagée.

On appelle host l'ordinateur qui accueille le GPU.

Le GPU est quant à lui qualifié de device.

Lorsque l'on programme en CUDA on tente de différencier les données en mémoire centrale et celles en mémoire du GPU. On utilise en général un préfixe pour les noms de variables :

Pour que le GPU puisse réaliser les calculs parallèles, il faut créer une copie des données de la mémoire centrale de l'host vers la mémoire globale du GPU (device) (Excepté à partir de CUDA 6 qui profite d'un support pour la mémoire unifiée mais en architecture 64 bits). On dispose pour cela d'un certain nombre de fonctions d'allocation mémoire et de copie, dont les plus simples sont :

On notera que les fonctions CUDA retournent toutes un code d'erreur.

Seuls les kernels sont de type sous-programme (void) et ne retournent aucune valeur.

2.3. Premier programme CUDA

2.3.1. Extraction le kernel

Le kernel est le code qui sera exécuté en parallèle sur le GPU par chaque thread.

Nous allons commencer par écrire un premier programme très simple afin de présenter les concepts de base de CUDA. On désire réaliser la somme de deux vecteurs de float et placer le résultat de cette somme dans un troisième vecteur. Le code en langage C ressemble à cela :

  1. // compute z[i] = x[i] + y[i] for i in [0..size-1]
  2. void sum(float *x, float *y, float *z, int size) {
  3.     for (int i = 0; i < size; ++i) {
  4.         z[i] = x[i] + y[i]; // paralell part
  5.     }
  6. }
  7.  

Ce traitement est hautement parallèle puisque chaque somme z[i] = x[i] + y[i] est indépendante des autres, on peut donc les exécuter en parallèle.

Cette partie de code parallèle appelée kernel en CUDA sera exécutée par plusieurs threads et sera codée comme suit :

  1. __global__
  2. void kernel_sum(float *x, float *y, float *z, int size) {
  3.    
  4.     /* global thread index (depends on grid organization) */
  5.     int gtid =  ...;
  6.    
  7.     /* check if we are in the arrays x, y, z */
  8.     if (gtid < size) {
  9.         z[gtid] = x[gtid] + y[gtid]; // paralell part
  10.     }
  11. }
  12.  

Plusieurs préfixes de fonctions ont été ajoutés pour permettre au compilateur CUDA (nvcc) de savoir quelle fonction doit être codée et executée sur host ou device et donc comment elle doit être compilée :

2.3.2. Organisation des threads

Si on traite des vecteurs de taille size alors il faut pouvoir indiquer au kernel qu'on utilisera size threads.

L'organisation des threads est réalisée sous la forme d'une grille de blocs de threads.

On dispose d'une structure de données de type dim3 qui permet de définir la taille de la grille et du bloc :

struct dim3 {
    int x, y, z;
}

On utilisera une variable de type dim3 pour définir la taille de la grille et une autre variable de type dim3 pour définir la taille du bloc.

2.3.2.a  Exemple

On désire utiliser 1023 threads, plusieurs possibilités nous sont offertes :

Pour les cas suivants il faudra faire en sorte que le dernier thread n'exécute pas de traitement:

Le problème majeur est que la formule qui permet d'obtenir le global thread index est différente en fonction de la taille 1D, 2D ou 3D de la grille et du bloc. (cf TD 1)

Une contrainte supplémentaire est liée aux caractéristiques de la carte que l'on appelle Compute Capability : certaines cartes acceptent un maximum de 512, 1024 ou 2048 threads par bloc.

Il faut donc faire en sorte que :

$block.x × block.y × block.z ≤ 512$

ou 1024 ou 2048 selon les cartes.

2.3.3. Le code

Les différentes étapes à réaliser sont les suivantes

Heterogeneous computing

  1. // file sum.cu
  2.  
  3. #include <cuda.h>
  4. #include <algorithm>
  5. #include <cstdlib>
  6. using namespace std;
  7.  
  8. // ====================================
  9. // kernel declaration
  10. // ====================================
  11. __global__
  12. void sum(float *x, float *y, float *z, int size) {
  13.     // compute Global Thread InDex
  14.     int gtid = threadIdx.x ;
  15.    
  16.     if (gtid < size) {
  17.         z[gtid] = x[gtid] + y[gtid]; // paralell part
  18.     }
  19. }
  20.  
  21. // ====================================
  22. // main function
  23. // ====================================
  24. int main() {
  25.     const int SIZE = 512;
  26.  
  27.     // allocate data on computer global memory
  28.     float *x_cpu = new float[ SIZE ];
  29.     float *y_cpu = new float[ SIZE ];
  30.     float *z_cpu = new float[ SIZE ];
  31.  
  32.     // fill x, y
  33.     std::fill(&x_cpu[0], &x_cpu[SIZE], 1);
  34.     std::fill(&y_cpu[0], &y_cpu[SIZE], 2);
  35.  
  36.     // allocate data on GPU global memory
  37.     float *x_gpu, *y_gpu, *z_gpu;
  38.  
  39.     cudaMalloc( (void**) &x_gpu, SIZE * sizeof(float) );
  40.     cudaMalloc( (void**) &y_gpu, SIZE * sizeof(float) );
  41.     cudaMalloc( (void**) &z_gpu, SIZE * sizeof(float) );
  42.  
  43.     // copy x and y resp. to dev_x, dev_y
  44.     cudaMemcpy(x_gpu, x_cpu, SIZE * sizeof(float), cudaMemcpyHostToDevice);
  45.     cudaMemcpy(y_gpu, y_cpu, SIZE * sizeof(float), cudaMemcpyHostToDevice);
  46.  
  47.     // define number of threads
  48.     dim3 grid(1, 1, 1);
  49.     dim3 block(SIZE, 1, 1);
  50.     // call kernel
  51.     sum<<< grid, block >>>(x_gpu, y_gpu, z_gpu, SIZE);
  52.  
  53.     // copy result back to computer's global memory
  54.     cudaMemcpy(z_cpu, z_gpu, SIZE * sizeof(float), cudaMemcpyDeviceToHost);
  55.  
  56.     // free memory
  57.     cudaFree( x_gpu );
  58.     cudaFree( y_gpu );
  59.     cudaFree( z_gpu );
  60.  
  61.     delete [] x_cpu;
  62.     delete [] y_cpu;
  63.     delete [] z_cpu;
  64.    
  65.     exit(EXIT_FAILURE);
  66. }
  67.  

Dans le code précédent, l'appel au kernel est réalisé en passant le nombre de blocs dans une grille (1) ainsi que le nombre de threads par bloc, ici égal à la taille des tableaux (soit 512 éléments).

Les différentes étapes à réaliser

sont rébarbatives. Il est intéressant de pouvoir les encapsuler dans une classe spécifique qui réalise ces traitements de manière semi-automatique.

Téléchargement code source

Attention, avec certains portables du département il faut installer gcc/g++ 4.9 avec le CUDA Toolkit fourni et compiler avec l'option -D_FORCE_INLINES pour nvcc afin de pouvoir résoudre le problème avec memcpy.

2.4. Le compilateur NVCC

CUDA fournit au programmeur un ensemble d'outils dont le compilateur nvcc (NVidia C Compiler). Il agit comme un compilateur C++ et est capable d'interpréter et compiler les parties de code relatives au GPU.

document du référence : nvcc-2.0.pdf ou saisir dans le terminal nvcc -h

nvcc --compile --compiler-options -O2 -o sum.o sum.cu
nvcc --link --compiler-options -O2 -o sum.exe sum.o

Il n'y a pas avec CUDA de compatibilité ascendante des exécutables car les jeux d'instructions changent en fonction des architectures. Un programme compilé explicitement pour Fermi risque de ne pas fonctionner sous Kepler.

On peut indiquer que l'on compile le code pour une architecture virtuelle ou une architecture réelle :

 Architecture   virtuelle   réelle   caractéristiques 
    --gpu-architecture arch   --gpu-code code,...    
 Fermi   compute_20   sm_20   basique 
 Kepler   compute_30 compute_32   sm_30 sm_32   Unified Memory 
    compute_35   sm_35   Dynamic parallelism support 
 Maxwell   compute_50 compute_52 compute_53   sm_50, sm_52, and sm_53    
nvcc compilation architecture virtuelle ou réelle

On peut utiliser -gencode pour spécifier à la fois l'architecture virtuelle et l'architecture réelle.

D'après ce que j'ai compris (mais pas très clair du côté NVidia) :

Néanmoins on peut spécifier une architecture virtuelle pour gpu-code !?

Par exemple, lors de la compilation des samples du CUDA SDK (CUDA 6.5), on obtient :

/usr/local/cuda-6.5/bin/nvcc -ccbin g++ -I../../common/inc  -m64 -ftz=true
 -gencode arch=compute_11,code=sm_11
 -gencode arch=compute_13,code=sm_13
 -gencode arch=compute_20,code=sm_20
 -gencode arch=compute_30,code=sm_30
 -gencode arch=compute_35,code=sm_35
 -gencode arch=compute_37,code=sm_37
 -gencode arch=compute_50,code=sm_50
 -gencode arch=compute_50,code=compute_50
 -o bodysystemcuda.o -c bodysystemcuda.cu

Les parties qui s'exécutent sur le CPU seront compilées par un compilateur C/C++ et génèreront du code assembleur x86. Pour les kernels (__global__) et les fonctions de type __device__ appelées par le kernel, on utilisera la partie du compilateur qui produit du code assembleur PTX (Parallel Thread Execution).

2.5. Code exécuté sur le CPU et X Server

Attention, dans le cas où on utilise le GPU à la fois pour l'affichage (serveur X) et pour les calculs CUDA (cas d'une carte GTX par exemple), il se peut que l'exécution d'un kernel qui dure trop longtemps (>10s) provoque :

  • un blocage ou un ralentissement de l'affichage pendant un certain temps
  • et/ou un brusque arrêt du kernel reportant une erreur de type launch failure

2.6. Code GPU exécuté sur le CPU - Ocelot

Si on ne dispose pas de carte graphique on pouvait aupravant utiliser un simulateur appelé Ocelot. Celui ci n'est plus maintenu.

Ocelot is a modular dynamic compilation framework for heterogeneous system, providing various backend targets for CUDA programs and analysis modules for the PTX virtual instruction set. Ocelot currently allows CUDA programs to be executed on NVIDIA GPUs, AMD GPUs, and x86-CPUs at full speed without recompilation.

Télécharger sur le site d'Ocelot, le package suivant et l'installer : ocelot_2.1.1272_i386.deb.

Pour compiler : 

export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda/lib
nvcc ex1.cu -o ex1.cu.cpp -cuda
g++ -o ex1.cu.o -c ex1.cu.cpp
g++ -o ex1.cu.exe -Wl,--start-group ex1.cu.o -L/usr/local/cuda/lib/ -lcudart -Wl,--end-group -glut -locelot

Sous Ubuntu 14.04 64 bits:

g++ -o  mandel_c.cu.exe  -Wl,--start-group mandel_c.cu.o -L/usr/local/cuda-6.5/lib64 -lcudart -Wl,--end-group -L/usr/l/checkout/gpuocelot/ocelot/build_local/lib/ -locelot

Cependant Ocelot ne permet pas de faire autant de choses qu'avec un SDK CUDA.