|
Commençons par les paramètres :
1
|
int main(int argc, char **argv){
|
Les quelques variables que l'on utilisera :
1
2
3
|
float sum;
size_t nthreads, nblocks = 16, ntot = 1024;
nthreads = ntot/(2*nblocks);
|
On affiche la taille des vecteurs, le nombre de blocks et le nombre de threads par bloc :
1
2
3
|
printf("Ntot : %lu\n", ntot);
printf("nthreads : %lu\n", nthreads);
printf("nblocks : %lu\n", nblocks);
|
On définit les pointeurs que l'on utilisera, le préfixe h_ indique que le pointeur est sur l'hôte, et d_ signifie que le pointeur pointe sur le GPU (Device) :
1
|
float *d_sum, *d_bl, *h_bl, *d_in, *h_in;
|
On alloue les deux tableaux de valeurs sur l'hôte :
1
2
|
cudaMallocHost((void**)&h_in, ntot*sizeof(float));
cudaMallocHost((void**)&h_bl, nblocks*sizeof(float));
|
On utilise la fonction cudaMallocHost pour punaiser la page allouée à l'endroit où elle sera allouée au départ.
On alloue les pointeurs sur le GPU :
1
2
3
|
cudaMalloc((void**)&d_sum, sizeof(float));
cudaMalloc((void**)&d_bl, nblocks*sizeof(float));
cudaMalloc((void**)&d_in, ntot*sizeof(float));
|
On test si tout s'est bien passé jusque là :
1
|
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
|
On initialise le vecteur h_in, on calcule la somme de tout ces éléments, et on affiche h_in :
1
2
3
|
init_vec(h_in, ntot);
float sommeCPU = reductionCPU(h_in, ntot);
printf_vec(h_in, ntot);
|
On copie le vecteur h_in sur le GPU et on teste si tout c'est bien passé :
1
2
|
cudaMemcpy(d_in, h_in, ntot*sizeof(float), cudaMemcpyHostToDevice);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
|
On définit les dimensions de la grille et des blocks pour faire une réduction partielle, sur tout les blocks :
1
2
|
dim3 dimGrid(nblocks, 1, 1);
dim3 dimBlock(nthreads, 1, 1);
|
On définit les dimensions de la grille et du blocks pour faire une réduction, sur tout les valeurs de chaque bloc que l'on aura avec d_bl :
1
2
|
dim3 dimGridHalf(nblocks/2, 1, 1);
dim3 unique(1,1,1);
|
On lance la réduction partielle dans chaque bloc :
1
|
kernelReduction<<<dimGrid, dimBlock>>>(d_in, ntot, d_bl, nblocks);
|
On fait la réduction entre les blocks pour finir le calcul :
1
|
kernelReduction<<<unique, dimGridHalf>>>(d_bl, nblocks, d_sum, 1);
|
On test si il n'y as pas eu de problème, et on synchronise les threads avant de continuer :
1
2
3
|
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
cudaThreadSynchronize();
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
|
On récupère les vecteurs que nous avons envoyer sur le GPU :
1
2
3
4
5
|
cudaMemcpy(h_in, d_in, ntot*sizeof(float), cudaMemcpyDeviceToHost);
cudaMemcpy(h_bl, d_bl, nblocks*sizeof(float), cudaMemcpyDeviceToHost);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
cudaMemcpy(&sum, d_sum, sizeof(float), cudaMemcpyDeviceToHost);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
|
Nous sommes d'accord que seul sum est utile pour le résultat du calcul, mais nous devons copier les autres vecteurs pour faire du débogage, mais seul sum sera retourné dans une éventuelle version finale.
On affiche le résultat :
1
2
3
4
|
printf("\nAprès le calcul h_in :\n");
printf_vec(h_in, ntot);
printf("\nAprès le calcul h_dl :\n");
printf_vec(h_bl, nblocks);
|
On vérifie que le résultat est juste :
1
|
verifCalcul(sum, sommeCPU);
|
On désalloue tout le monde
1
2
3
4
5
6
7
|
cudaFree(d_sum);
cudaFree(d_bl);
cudaFreeHost(h_bl);
cudaFree(d_in);
cudaFreeHost(h_in);
return 0;
}
|
Voilà.
Et voici le main complet :
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
|
int main(int argc, char **argv){
float sum;
size_t nthreads, nblocks = 16, ntot = 1024;
nthreads = ntot/(2*nblocks);
printf("Ntot : %lu\n", ntot);
printf("nthreads : %lu\n", nthreads);
printf("nblocks : %lu\n", nblocks);
float *d_sum, *d_bl, *h_bl, *d_in, *h_in;
cudaMallocHost((void**)&h_in, ntot*sizeof(float));
cudaMallocHost((void**)&h_bl, nblocks*sizeof(float));
cudaMalloc((void**)&d_sum, sizeof(float));
cudaMalloc((void**)&d_bl, nblocks*sizeof(float));
cudaMalloc((void**)&d_in, ntot*sizeof(float));
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
init_vec(h_in, ntot);
float sommeCPU = reductionCPU(h_in, ntot);
printf_vec(h_in, ntot);
cudaMemcpy(d_in, h_in, ntot*sizeof(float), cudaMemcpyHostToDevice);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
dim3 dimGrid(nblocks, 1, 1);
dim3 dimBlock(nthreads, 1, 1);
dim3 dimGridHalf(nblocks/2, 1, 1);
dim3 unique(1,1,1);
kernelReduction<<<dimGrid, dimBlock>>>(d_in, ntot, d_bl, nblocks); kernelReduction<<<unique, dimGridHalf>>>(d_bl, nblocks, d_sum, 1); checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
cudaThreadSynchronize();
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
cudaMemcpy(h_in, d_in, ntot*sizeof(float), cudaMemcpyDeviceToHost);
cudaMemcpy(h_bl, d_bl, nblocks*sizeof(float), cudaMemcpyDeviceToHost);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
cudaMemcpy(&sum, d_sum, sizeof(float), cudaMemcpyDeviceToHost);
checkNoErr(__LINE__, d_sum, d_bl, d_in, h_in);
printf("\nAprès le calcul h_in :\n");
printf_vec(h_in, ntot);
printf("\nAprès le calcul h_dl :\n");
printf_vec(h_bl, nblocks);
verifCalcul(sum, sommeCPU);
cudaFree(d_sum);
cudaFree(d_bl);
cudaFreeHost(h_bl);
cudaFree(d_in);
cudaFreeHost(h_in);
return 0;
}
|
Et voilà, pour la fonction main.
|