Cuda hello world

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:

1. instanciation des instructions et variables dans la mémoire centrale (RAM)
2. copie des instruction et variables dans la mémoire du GPU
3. exécution des instruction et lecture des variables par le GPU
4. modification des variables dans la mémoire du GPU
5. 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<<<CORE, THREAD>>>();

• fonction : correspond au nom de la fonction
• CORE : correspond au nombre de cœur CUDA utilisés
• THREAD : correspond au nombre de threads par cœur CUDA

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

Exemple de programme

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

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 normale car la variable passée en paramètre est argv[1] qui est dans la mémoire centrale !

Pour cela il faudrait copier le contenu de cette variable dans la mémoire du GPU. Cela se fait en deux étapes:

1. on fait l'allocation mémoire avec cudaMalloc
2. on procède à la copie avec cudaMemCpy