1 #include "Histogramme.h"
10 #include "cudaTools.h"
16 * 0) Initialisation de la mémoire partagée.
19 void initSM(uint* tabSM)
21 const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
22 const int TID_LOCAL = Indice1D::tidLocal();
33 * 1) Chaque thread calcule un résultat intermediaire qu'il va ensuite placer en shared memory.
36 void reductionIntraThread(int n, uchar* ptrInput, uint* tabSM)
38 const int NB_THREAD = Indice1D::nbThread();
39 const int TID = Indice1D::tid();
40 const int TID_LOCAL = Indice1D::tidLocal();
45 atomicAdd(&tabSM[ptrInput[s]], 1);
51 * 2) Copie les résultats de block vers la global memory.
54 void reductionInterBlock(uint* tabSM, uint* ptrResult)
56 const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
57 const int TID_LOCAL = Indice1D::tidLocal();
59 // Un maximum de 256 threads par block s'occupe de la copie SM -> GM.
63 atomicAdd(&ptrResult[s], tabSM[s]);
67 // Un seul thread par block s'occupe de la copie SM -> GM.
68 /* Cette version ne présente pas d'amélioration en terme de temps.
71 for (int i = 0; i < 256; i++)
72 atomicAdd(&ptrResult[i], tabSM[i]);
77 * La taille de la shared memory (en terme de # de sizeof(float)) doit
78 * être égal à la taille des blocs.
79 * n: le nombre d'échantillon (taille du tableau d'entrée)
80 * ptrInput: les échantillons.
81 * ptrResult: Le resultat à peupler.
84 void histogramme(int n, uchar* ptrInput, uint* ptrResult)
86 extern __shared__ uint tabSM[]; // Dynamic shared memory.
88 // 0) Initialisation de la shared memory de manière parallèle.
93 // 1) Réduction intra-thread.
94 reductionIntraThread(n, ptrInput, tabSM);
98 // 2) Réduction inter-block.
99 reductionInterBlock(tabSM, ptrResult);
102 uchar* createTabInput(int size)
104 uchar* tab = new uchar[size];
106 // Génération de la suite.
107 for (int i = 1; i <= size; ++i)
110 // Permutations aléatoires.
112 for (int i = 0; i < size; ++i)
114 int i1 = rand() % size; // Il y a un biais mais ça n'a pas d'importance ici.
115 int i2 = rand() % size;
117 swap(tab[i1], tab[i2]);
125 cout << endl << "histogramme() ..." << endl;
130 // Taille du tableau en input, doit être un multiple de 256.
131 int N = 100000 * 256;
132 uchar* tabInput = createTabInput(N);
134 // Allocation coté GPU en global memory (GM) et copie du tableau.
136 HANDLE_ERROR(cudaMalloc(&ptrDevInput, N * sizeof(uchar)));
137 HANDLE_ERROR(cudaMemcpy(ptrDevInput, tabInput, N * sizeof(uchar), cudaMemcpyHostToDevice));
140 HANDLE_ERROR(cudaMalloc(&ptrDevResult, 256 * sizeof(uint)));
141 HANDLE_ERROR(cudaMemset(ptrDevResult, 0, 256 * sizeof(uint)));
143 // Paramètre de l'appel de la fonction sur le device.
144 const dim3 dg(256, 1, 1);
145 const dim3 db(256, 1, 1);
146 Device::assertDim(dg, db);
147 const size_t SMSize = db.x * sizeof(uint); // 256 uint;
149 histogramme<<<dg, db, SMSize>>>(N, ptrDevInput, ptrDevResult);
151 cudaDeviceSynchronize(); // Utilisé pour flusher les prints sur le stdout à partir du device (debug).
153 // Barrière implicite de synchronisation ('cudaMemCpy').
154 uint* ptrResult = new uint[256];
155 HANDLE_ERROR(cudaMemcpy(ptrResult, ptrDevResult, 256 * sizeof(uint), cudaMemcpyDeviceToHost));
157 uint expected = N / 256;
158 for (int i = 0; i < 256; ++i)
160 if (ptrResult[i] != expected)
162 cout << "Wrong result: " << ptrResult[i] << ", expected: " << expected << endl;
166 cout << "OK, all values equal the exptected value: " << expected << endl;
169 cout << "Time: " << chronos << endl;