14 * 1) Chaque thread calcule un résultat intermediaire qu'il va ensuite placer en shared memory.
15 * n: Nombre d'échantillon.
18 void reductionIntraThread(int n, float deltaX, float* tabSM)
20 const int NB_THREAD = Indice1D::nbThread();
21 const int TID = Indice1D::tid();
22 const int TID_LOCAL = Indice1D::tidLocal();
24 float threadResult = 0.0;
28 const float i = s + 1;
29 const float xi = -1 + i * deltaX;
30 threadResult += sqrtf(1 - xi * xi);
35 tabSM[TID_LOCAL] = threadResult;
39 * Combine les résultats de 'tabSM' dans 'tabSM[0]'
42 void combine(float* tabSM, int middle)
44 const int TID_LOCAL = Indice1D::tidLocal();
45 const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
50 tabSM[s] += tabSM[s + middle];
56 * 2) La shared memory est réduite, le résultat est placé dans 'tabSM[0]'.
59 void reductionIntraBlock(float* tabSM)
61 const int TAB_SIZE = blockDim.x;
62 int middle = TAB_SIZE / 2;
66 combine(tabSM, middle);
68 __syncthreads(); // Synchronisation des threads au niveau du bloc.
73 * 3) Le 'tabSM[0]' de chaque bloc est reduit dans 'ptrResult' qui se trouve en global memory.
76 void reductionInterBlock(float* tabSM, float* ptrResult)
78 const int TID_LOCAL = Indice1D::tidLocal();
81 atomicAdd(ptrResult, tabSM[0]);
86 * La taille de la shared memory (en terme de # de sizeof(float)) doit
87 * être égal à la taille des blocs.
88 * n: le nombre d'échantillon
89 * ptrResult: Le resultat du calcul de pi.
92 void saucisson(int n, float deltaX, float* ptrResult)
94 extern __shared__ float tabSM[]; // Dynamic shared memory.
96 // 1) Réduction intra-thread.
97 reductionIntraThread(n, deltaX, tabSM);
101 // 2) Réduction intra-block.
102 reductionIntraBlock(tabSM);
104 // 3) Réduction inter-block.
105 reductionInterBlock(tabSM, ptrResult);
110 cout << endl << "saucisson() ..." << endl;
112 // Nombre d'échantillon. Au-delà, la qualité du résultat n'est pas meilleure. Il faudrait employé des doubles à la place de floats.
113 const int N = 100000;
115 // Allocation coté GPU en global memory (GM).
117 HANDLE_ERROR(cudaMalloc(&ptrDevResult, sizeof(float)));
118 HANDLE_ERROR(cudaMemset(ptrDevResult, 0, sizeof(float)));
120 // Paramètre de l'appel de la fonction sur le device.
121 const dim3 dg(256, 1, 1);
122 const dim3 db(256, 1, 1);
123 Device::assertDim(dg, db);
124 const size_t SMSize = db.x * sizeof(float); // 256 floats;
126 const float deltaX = 2.0f / N;
127 saucisson<<<dg, db, SMSize>>>(N, deltaX, ptrDevResult);
129 // cudaDeviceSynchronize(); // Utilisé pour flusher les prints sur le stdout à partir du device (debug).
132 // Barrière implicite de synchronisation ('cudaMemCpy').
133 HANDLE_ERROR(cudaMemcpy(&pi, ptrDevResult, sizeof(float), cudaMemcpyDeviceToHost));
137 cout << "Approximation de PI : " << pi << endl;