--- /dev/null
+#include "Histogramme.h"
+
+#include <iostream>
+#include <cmath>
+#include <cstdlib>
+#include <stdio.h>
+using namespace std;
+
+#include "Indice1D.h"
+#include "cudaTools.h"
+#include "Device.h"
+#include "Lock.h"
+#include "Chronos.h"
+
+/**
+ * 0) Initialisation de la mémoire partagée.
+ */
+__device__
+void initSM(uint* tabSM)
+ {
+ const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
+ const int TID_LOCAL = Indice1D::tidLocal();
+
+ int s = TID_LOCAL;
+ while (s < 256)
+ {
+ tabSM[s] = 0;
+ s += NB_THREAD_LOCAL;
+ }
+ }
+
+/*
+ * 1) Chaque thread calcule un résultat intermediaire qu'il va ensuite placer en shared memory.
+ */
+__device__
+void reductionIntraThread(int n, uchar* ptrInput, uint* tabSM)
+ {
+ const int NB_THREAD = Indice1D::nbThread();
+ const int TID = Indice1D::tid();
+ const int TID_LOCAL = Indice1D::tidLocal();
+
+ int s = TID;
+ while (s < n)
+ {
+ atomicAdd(&tabSM[ptrInput[s]], 1);
+ s += NB_THREAD;
+ }
+ }
+
+/*
+ * 2) Copie les résultats de block vers la global memory.
+ */
+__device__
+void reductionInterBlock(uint* tabSM, uint* ptrResult)
+ {
+ const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
+ const int TID_LOCAL = Indice1D::tidLocal();
+
+ // Un maximum de 256 threads par block s'occupe de la copie SM -> GM.
+ int s = TID_LOCAL;
+ while (s < 256)
+ {
+ atomicAdd(&ptrResult[s], tabSM[s]);
+ s += NB_THREAD_LOCAL;
+ }
+
+ // Un seul thread par block s'occupe de la copie SM -> GM.
+ /* Cette version ne présente pas d'amélioration en terme de temps.
+ if (TID_LOCAL == 0)
+ {
+ for (int i = 0; i < 256; i++)
+ atomicAdd(&ptrResult[i], tabSM[i]);
+ }*/
+ }
+
+/**
+ * La taille de la shared memory (en terme de # de sizeof(float)) doit
+ * être égal à la taille des blocs.
+ * n: le nombre d'échantillon (taille du tableau d'entrée)
+ * ptrInput: les échantillons.
+ * ptrResult: Le resultat à peupler.
+ */
+__global__
+void histogramme(int n, uchar* ptrInput, uint* ptrResult)
+ {
+ extern __shared__ uint tabSM[]; // Dynamic shared memory.
+
+ // 0) Initialisation de la shared memory de manière parallèle.
+ initSM(tabSM);
+
+ __syncthreads();
+
+ // 1) Réduction intra-thread.
+ reductionIntraThread(n, ptrInput, tabSM);
+
+ __syncthreads();
+
+ // 2) Réduction inter-block.
+ reductionInterBlock(tabSM, ptrResult);
+ }
+
+uchar* createTabInput(int size)
+ {
+ uchar* tab = new uchar[size];
+
+ // Génération de la suite.
+ for (int i = 1; i <= size; ++i)
+ tab[i-1] = i % 256;
+
+ // Permutations aléatoires.
+ srand(time(NULL));
+ for (int i = 0; i < size; ++i)
+ {
+ int i1 = rand() % size; // Il y a un biais mais ça n'a pas d'importance ici.
+ int i2 = rand() % size;
+ if (i1 != i2)
+ swap(tab[i1], tab[i2]);
+ }
+
+ return tab;
+ }
+
+bool histogramme()
+ {
+ cout << "hisogramme() ..." << endl;
+
+ Chronos chronos;
+ chronos.start();
+
+ // Taille du tableau en input, doit être un multiple de 256.
+ int N = 100000 * 256;
+ uchar* tabInput = createTabInput(N);
+
+ // Allocation coté GPU en global memory (GM) et copie du tableau.
+ uchar* ptrDevInput;
+ HANDLE_ERROR(cudaMalloc(&ptrDevInput, N * sizeof(uchar)));
+ HANDLE_ERROR(cudaMemcpy(ptrDevInput, tabInput, N * sizeof(uchar), cudaMemcpyHostToDevice));
+
+ uint* ptrDevResult;
+ HANDLE_ERROR(cudaMalloc(&ptrDevResult, 256 * sizeof(uint)));
+ HANDLE_ERROR(cudaMemset(ptrDevResult, 0, 256 * sizeof(uint)));
+
+ // Paramètre de l'appel de la fonction sur le device.
+ const dim3 dg(256, 1, 1);
+ const dim3 db(256, 1, 1);
+ Device::assertDim(dg, db);
+ const size_t SMSize = db.x * sizeof(uint); // 256 uint;
+
+ histogramme<<<dg, db, SMSize>>>(N, ptrDevInput, ptrDevResult);
+
+ cudaDeviceSynchronize(); // Utilisé pour flusher les prints sur le stdout à partir du device (debug).
+
+ // Barrière implicite de synchronisation ('cudaMemCpy').
+ uint* ptrResult = new uint[256];
+ HANDLE_ERROR(cudaMemcpy(ptrResult, ptrDevResult, 256 * sizeof(uint), cudaMemcpyDeviceToHost));
+
+ uint expected = N / 256;
+ for (int i = 0; i < 256; ++i)
+ {
+ if (ptrResult[i] != expected)
+ {
+ cout << "Wrong result: " << ptrResult[i] << ", expected: " << expected << endl;
+ return false;
+ }
+ }
+ cout << "OK, all values equal the exptected value: " << expected << endl;
+
+ chronos.stop();
+ cout << "Time: " << chronos << endl;
+
+ return true;
+ }
projects / GPU.git / commitdiff