6.2 Le fichier calcul_cuda.cu

Précédent
Le fichier calcul_cuda.h		Parent
L'exemple : la multiplication de matrices carrées		SommaireSuivant
Le fichier main.c

Dans ce fichier, on commence par inclure le fichier calcul_cuda.h : 

1
 
#include "calcul_cuda.h"

Ensuite, on définit une structure qui contient tout les variables globales dont on a besoin pour le calcul : 

1
2
3
4
5
 
struct GlobalVarMatrix{
	float *matriceLeftd;
	float *matriceRightd;
	float *matriceResultd;
};

Ensuite on crée la variable globale proprement dite, et on l'initialise : 

1
 
GlobalVarMatrix globalVarMatrix = {NULL, NULL, NULL};

Maintenant, on peut créer la fonction qui fera le calcul : 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
 
__global__ void matrixMulKernel(float * matriceResultd, float * matriceLeftd, float * matriceRightd, int width){
	// identifiant de thread à deux dimensions, comme la matrice
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	// Pvaleur sert au stockage de la valeur calculée par le thread
	float pResult = 0;
	for(int i = 0; i < width; ++i){
		float mdElement = matriceLeftd[ty*width + i];
		float ndElement = matriceRightd[i*width + tx];
		pResult        += mdElement * ndElement;
	}
	// écrit la valeur calculée dans la matrice de résultat
	// chaque thread ne peut écrire qu'une valeur !
	matriceResultd[ty*width + tx] = pResult;
}

Heu... Il ne manque pas des boucles là ? On fait bien un produit matriciel avec 3 boucles.

Oui, il y a bien trois boucles dans un produit matriciel, mais qui vous dit qu'il en manque deux ?

Ben, il n'y a qu'un seul for

Voici une des subtilité de la programmation CUDA, dans notre programme il y aura bien trois boucles, mais c'est la carte graphique qui fera les deux boucles principales, en hardware et pas en software. C'est donc le matériel qui fait deux boucles à notre place. Donc le compte est bon.

En fait ce sont les variables threadIdx.x et threadIdx.y qui nous permettent de faire deux boucles avec le matériel.

Et rappeler-vous qu'une fonction __global__ ne retourne rien.

Nous pouvons donc ajouter la fonction qui appelle notre fonction CUDA : 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
 
///fonction qui appelle la fonction en cuda de multiplication de matrice
extern "C" void matrixMulOnDevice(float * matriceResult, float * matriceLeft, float * matriceRight, int width){
	//calcul de la taille des matrices
	int size = width*width*sizeof(float);
	//allocation des matrices et leur remplissage
	cudaMalloc(&globalVarMatrix.matriceLeftd, size);
	cudaMemcpy(globalVarMatrix.matriceLeftd, matriceLeft, size, cudaMemcpyHostToDevice) ;
	cudaMalloc(&globalVarMatrix.matriceRightd, size);
	cudaMemcpy(globalVarMatrix.matriceRightd, matriceRight, size, cudaMemcpyHostToDevice);
	//allocation de la matrice de résultat
	cudaMalloc(&globalVarMatrix.matriceResultd, size);
	//multiplication d'une seule matrice
	dim3 dimGrid(1, 1);
	//matrice carrée
	dim3 dimBlock(width, width);
	//produit matriciel proprement dit
	matrixMulKernel<<<dimGrid, dimBlock>>>(globalVarMatrix.matriceLeftd, globalVarMatrix.matriceRightd, globalVarMatrix.matriceResultd, width);
	//récupération du résultat du calcul
	cudaMemcpy(matriceResult, globalVarMatrix.matriceResultd, size, cudaMemcpyDeviceToHost);
	//destruction des matrices, désormais inutilisées
	cudaFree(globalVarMatrix.matriceLeftd);
	cudaFree(globalVarMatrix.matriceRightd);
	cudaFree(globalVarMatrix.matriceResultd);
}

N'oubliez pas que cudaMalloc prend un void** donc il faut un & devant le pointeur à allouer.

Voici le fichier calcul_cuda.cu en entier : 

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
 
#include "calcul_cuda.h"

///on définit une variable globale pour avoir les matrices pour la carte graphique 
struct GlobalVarMatrix{
	float *matriceLeftd;
	float *matriceRightd;
	float *matriceResultd;
};

//on n'oublie pas de l'initialisée à NULL
GlobalVarMatrix globalVarMatrix = {NULL, NULL, NULL};


///fonction qui va calculer la multiplication de la matrice dans la carte graphique
__global__ void matrixMulKernel(float * matriceResultd, float * matriceLeftd, float * matriceRightd, int width){
	// identifiant de thread à deux dimensions, comme la matrice
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	// Pvaleur sert au stockage de la valeur calculée par le thread
	float pResult = 0;
	for (int i = 0; i < width; ++i){
		float mdElement = matriceLeftd[ty*width + i];
		float ndElement = matriceRightd[i*width + tx];
		pResult        += mdElement * ndElement;
	}
	// écrit la valeur calculée dans la matrice de résultat
	// chaque thread ne peut écrire qu'une valeur !
	matriceResultd[ty*width + tx] = pResult;
}


///fonction qui appelle la fonction en cuda de multiplication de matrice
extern "C" void matrixMulOnDevice(float * matriceResult, float * matriceLeft, float * matriceRight, int width){
	//calcul de la taille des matrices
	int size = width*width*sizeof(float);
	//allocation des matrices et leur remplissage
	cudaMalloc(&globalVarMatrix.matriceLeftd, size);
	cudaMemcpy(globalVarMatrix.matriceLeftd, matriceLeft, size, cudaMemcpyHostToDevice) ;
	cudaMalloc(&globalVarMatrix.matriceRightd, size);
	cudaMemcpy(globalVarMatrix.matriceRightd, matriceRight, size, cudaMemcpyHostToDevice);
	//allocation de la matrice de résultat
	cudaMalloc(&globalVarMatrix.matriceResultd, size);
	//multiplication d'une seule matrice
	dim3 dimGrid(1, 1);
	//matrice carrée
	dim3 dimBlock(width, width);
	//produit matriciel proprement dit
	matrixMulKernel<<<dimGrid, dimBlock>>>(globalVarMatrix.matriceLeftd, globalVarMatrix.matriceRightd, globalVarMatrix.matriceResultd, width);
	//récupération du résultat du calcul
	cudaMemcpy(matriceResult, globalVarMatrix.matriceResultd, size, cudaMemcpyDeviceToHost);
	//destruction des matrices, désormais inutilisées
	cudaFree(globalVarMatrix.matriceLeftd);
	cudaFree(globalVarMatrix.matriceRightd);
	cudaFree(globalVarMatrix.matriceResultd);
}

Vous pouvez le récupérer ici.

Précédent
Le fichier calcul_cuda.h		Parent
L'exemple : la multiplication de matrices carrées		SommaireSuivant
Le fichier main.c