Version actuelle datée du 23 juin 2021 à 08:09

Concepts

Le framework CUDA permet d'utiliser le GPU pour faire des calculs généraux, habituellement réalisés par le CPU. CUDA est une extension en C/C++ qui propose une API pour gérer le GPU et elle distingue deux entités:

l'entité host ou hôte pour le CPU
l'entité device ou périphérique pour le GPU

De manière générale on essayera de faire le traitement sérialisé sur le CPU et de décharger les calculs parallélisés sur le GPU.

La difficulté réside dans le faite que CPU et GPU utilisent deux espaces mémoire séparés... qu'il va falloir gérer par le biais de pointeurs et fonctions spécifiques !

Les programmes utilisant le GPU se déroule en 5 étapes:

instanciation des instructions et variables dans la mémoire centrale (RAM)
copie des instruction et variables dans la mémoire du GPU
exécution des instruction et lecture des variables par le GPU
modification des variables dans la mémoire du GPU
rapatriement des variables modifiées par le GPU dans la mémoire centrale

Les bases

L'API CUDA permet d'exécuter une fonction sur le GPU grâce à la notation chevrons <<<...>>> :

fonction<<<NB_BLOCK, NB_THREAD>>>();

fonction : correspond au nom de la fonction
NB_BLOCK : correspond au nombre de blocs utilisés
NB_THREAD : correspond au nombre de threads par bloc

La fonction appelée doit être précédée du spécificateur __global__ pour préciser au compilateur que sont exécution se fait sur le GPU:

__global__ void fonction(){
   // choses à faire
}

Les fichiers contenant du code CUDA doivent impérativement porter l'extension .cu et être compilés avec le compilateur CUDA : nvcc

# nvcc hello.cu -o hello

Programmation CPU vs. GPU

Ci-dessous l'exemple classique du Hello World utilisant le CPU:

#include <stdio.h>

int main(void) {
	printf("Hello World du CPU\n");
	return EXIT_SUCCESS;
}

Si nous voulons utiliser le GPU il faut sortir l'instruction à exécuter (ici le printf) dans une fonction __global__ puis lancer son exécution grâce à la notation chevrons:

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

__global__ void cuda_hello(){
	printf("Hello World du GPU\n");
}

int main(void) {
	printf("Hello World du CPU\n");
	cuda_hello<<<1,1>>>();
	return EXIT_SUCCESS;
}

Synchronisation CPU / GPU

Le programme précédent affichera le résultat ci-dessous :

Hello World du CPU

Le GPU n'a simplement rien fait. Il n'était pas informé de quoique ce soit car les instructions n'ont pas été copiées sur la mémoire du GPU. Cela se fait grâce à la fonction cudaDeviceSynchronize qu'il faut ajouter juste après l'instruction chevron:

...
	cuda_hello<<<1,1>>>();
	cudaDeviceSynchronize();
...

Et nous avons bien le résultat attendu:

Hello World du CPU
Hello World du GPU

Manipulation mémoire

Gestion manuellement

Le programme ci-dessous permet d'afficher le nom passé en paramètre :

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

__global__ void say_my_name(char * name){
	printf("GPU says : %s\n", name);
}

int main(int argc, char * argv[]) {
	if(argc < 2){
		printf("I don't know your name !");
		return EXIT_FAILURE;
	}
	printf("CPU says : %s\n", argv[1]);
	say_my_name<<<1,1>>>(argv[1]);
	cudaDeviceSynchronize();
	return EXIT_SUCCESS;
}

Le programme précédent affichera :

# name.bin tala
CPU says : tala

Il manque, encore une fois, l’exécution du code côté GPU... Ce qui est normal car la variable passée en paramètre argv[1] est dans la mémoire centrale donc inaccessible côté GPU !

Pour pouvoir l'utiliser, il faut copier le contenu de cette variable dans la mémoire du GPU.

Cela se fait en trois étapes:

on fait l'allocation mémoire avec cudaMalloc
on procède à la copie avec cudaMemCpy
on procède à la libération mémoire avec cudaFree

allocation

cudaMalloc prend en paramètre un pointeur côté CPU et la taille de l'allocation :

cudaMalloc(void ** ptr, size_t size);

Le pointeur ptr permet de faire le lien entre les variables dans la mémoire centrale CPU et celle dans la mémoire graphique.

copie

cudaMemCpy prend en paramètre :

un pointeur source
un pointeur destination
la taille de la copie
le sens de la copie

cudaMemCpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind direction);

Notez que le paramètre direction de type cudaMemcpyKind peut prendre les valeurs suivantes:

cudaMemcpyHostToDevice (CPU → GPU)
cudaMemcpyDeviceToHost (GPU → CPU)

libération

cudaFree prend en paramètre le pointeur côté GPU à libérer

cudaFree(void *ptr);

Ce qui donne les modifications suivantes :

...
	printf("CPU says : %s\n", argv[1]);
	char * name;
	int size = sizeof(char) * strlen(argv[1]) + 1;
	cudaMalloc(&name, size);
	cudaMemcpy(name, argv[1], size, cudaMemcpyHostToDevice);
	say_my_name<<<1,1>>>(name);
	cudaDeviceSynchronize();
	cudaFree(name);
...

Pour avoir le résultat attendu:

# ./name.bin tala
CPU says : tala
GPU says : tala

Gestion unifiée

L'architecture Maxwell a introduit le support de la mémoire unifiée. Les variables allouées grâce à la fonction cudaMallocManaged voit leurs espaces mémoire transiter entre CPU et GPU au besoin. Cella permet d'écrire les programmes plus simplement, sans avoir à appeler explicitement la fonction cudaMemcpy.

Dans l'exemple suivant on créer une chaîne de caractère que l'on modifie dans le kernel GPU avant de l'afficher avec le CPU:

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

__global__ void say_hello(char *str){
	printf("CPU said : %s\n", str);
	memcpy(str, "world", 5);
}

int main(int argc, char * argv[]) {
	char *str;
	cudaMallocManaged(&str, 5 * sizeof(char));
	say_hello<<<1,1>>>(str);
	cudaDeviceSynchronize();
	printf("GPU said : %s\n", str);
	return EXIT_SUCCESS;
}

La sortie permet de vérifier que l'espace mémoire est géré correctement:

CPU said : hello
GPU said : world

Affichage des erreurs

L'exemple précédent met en évidence le manque de visibilité sur ce qui se passe côté GPU. Pour cela on peut utiliser la fonction gpuAssert pour afficher les erreurs lors de l’exécution du code sur le GPU:

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) {
	if (code != cudaSuccess) {
		fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
		if (abort) exit(code);
	}
}

Le code précédent est à mettre au dessus de la fonction main et avant les fonctions chevrons. Si on reprend l'exemple précédent cela donnerait :

...
__global__ void say_my_name(char * name){
	printf("GPU says : %s\n", name);
}

int main(int argc, char * argv[]) {
	if(argc < 2){
		printf("I don't know your name !");
		return EXIT_FAILURE;
	}
	printf("CPU says : %s\n", argv[1]);
	say_my_name<<<1,1>>>(argv[1]);
	gpuErrchk(cudaDeviceSynchronize());
	return EXIT_SUCCESS;
}

Comme l'allocation n'est pas faite avec cudaMalloc, cela affiche l'erreur suivante:

# ./name.bin tala
CPU says : tala
GPUassert: an illegal memory access was encountered name.cu 30

Il est aussi possible d'utiliser la fonction cudaPeekAtLastError pour afficher la dernière erreur :

gpuErrchk(cudaPeekAtLastError());

Parallélisation

Le gros avantage du GPU c'est sa capacité à paralléliser l'exécution du code et pour cela on doit s'attarder sur la syntaxe chevron. On peut mixer le nombre de blocs ainsi que le nombre de thread par blocs jusqu'à un total de 1024, comme expliqué ici :

on peut faire <<<1024, 1>>> ou <<<128,8>>> (128 * 8 = 1024)
on ne peut pas faire <<<256, 10>>> car 256 * 10 = 2560 > 1024

Il faut comprendre que le code exécuté par les différentes instances est le même, il faudra donc modifier son comportement grâce au numéro de groupe ou de thread. Ces numéros sont accessible grâce aux variables :

threadIdx : id du thread
blockIdx : id du bloc
blockDim : dimension de chaque bloc
gridDim : dimension de la grille (nombre de blocs)

On peut schématiser la parallélisation effectuée par le GPU comme la matrice suivante:

Dans l'exemple suivant on affiche les différents indexes :

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>

__global__ void whoami() {
	printf("gridDim[%d], blockDim [%d], blockIdx [%d], threadIdx[%d]\n", gridDim.x, blockDim.x, blockIdx.x, threadIdx.x);
}

int main(int argc, char *argv[]) {
	whoami<<<2,3>>>();
	cudaDeviceSynchronize();
	return EXIT_SUCCESS;
}

Cette exemple donne la sortie suivante:

gridDim[2], blockDim [3], blockIdx [0], threadIdx[0]
gridDim[2], blockDim [3], blockIdx [0], threadIdx[1]
gridDim[2], blockDim [3], blockIdx [0], threadIdx[2]
gridDim[2], blockDim [3], blockIdx [1], threadIdx[0]
gridDim[2], blockDim [3], blockIdx [1], threadIdx[1]
gridDim[2], blockDim [3], blockIdx [1], threadIdx[2]

On peut remarquer que:

gridDim.x correspond au nombre de blocs (premier paramètre de la fonction chevron)
blockDim.x correspond au nombre de threads par blocs (deuxième paramètre de la fonction chevron)
blockIdx.x correspond à l'index du bloc (premier paramètre de la fonction chevron)
threadIdx.x correspond à l'index du thread (compris entre 0 et blockDim.x - 1)

On peut utiliser l'astuce suivante pour calculer un index unique par thread :

__global__ void whoami() {
	int index = blockDim.x * blockIdx.x + threadIdx.x;
	int total = gridDim.x * blockDim.x;
	printf("I am %d over %d\n", index, total);
}

Ce qui donne le résultat suivant:

I am 0 over 6
I am 1 over 6
I am 2 over 6
I am 3 over 6
I am 4 over 6
I am 5 over 6

Différences entre versions de « Cuda hello world »