12 #define VI 1.4422495703074083
13 #define WI 0.7390850782394409
16 * Renvoie la valeur du ième élément du vecteur v.
21 double x = 1.5 + abs(cos(double(i)));
22 for (long j = 1; j <= M_V; j++)
24 const double xCarre = x * x;
25 x = x - (xCarre * x - 3) / (3 * xCarre);
27 return (x / VI) * sqrt(double(i));
31 * Renvoie la valeur du ième élément du vecteur w.
36 double x = abs(cos(double(i)));
37 for (long j = 1; j <= M_W; j++)
38 x = x - (cos(x) - x) / (-sin(x) - 1);
39 return (x / WI) * sqrt(double(i));
43 * n: La taille des deux vecteurs.
45 __device__ void reductionIntraThread(int n, double* tabResultSM)
47 const int NB_THREAD = Indice1D::nbThread();
48 const int TID = Indice1D::tid();
49 const int TID_LOCAL = Indice1D::tidLocal();
51 double threadResult = 0.0;
55 threadResult += v(s) * w(s);
58 tabResultSM[TID_LOCAL] = threadResult;
62 * Combine les résultats de 'tabResultSM' dans 'tabResulSM[0]'
64 __device__ void combine(double* tabResultSM, int middle)
66 const int TID_LOCAL = Indice1D::tidLocal();
67 const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
72 tabResultSM[s] = tabResultSM[s] + tabResultSM[s + middle];
77 __device__ void reductionIntraBlock(double* tabResultSM)
79 const int TAB_SIZE = blockDim.x;
80 int middle = TAB_SIZE / 2;
84 combine(tabResultSM, middle);
90 __device__ void reductionInterBlock(double* tabResultSM, float* ptrResult)
92 const int TID_LOCAL = Indice1D::tidLocal();
95 atomicAdd(ptrResult, float(tabResultSM[0]));
100 * La taille de la shared memory (en terme de # de sizeof(double)) doit
101 * être égal à la taille des blocs.
102 * n: La taille des deux vecteurs.
103 * ptrResult: Le resultat du produit scalaire.
106 void produitScalaire(int n, float* ptrResult)
108 extern __shared__ double tabResultSM[]; // Shared memory.
110 // 1) Réduction intra-thread.
111 reductionIntraThread(n, tabResultSM);
115 // 2) Réduction intra-block.
116 reductionIntraBlock(tabResultSM);
118 // 3) Réduction inter-block.
119 reductionInterBlock(tabResultSM, ptrResult);
122 double resultatTheorique(long n)
125 return (n / 2.0) * (n+1);
128 bool produitScalaire()
130 const int N = 100000000; // Taille des deux vecteurs.
132 // Allocation coté GPU en global memory (GM).
133 float* ptrDevResult = 0;
134 HANDLE_ERROR(cudaMalloc(&ptrDevResult, sizeof(float)));
135 HANDLE_ERROR(cudaMemset(ptrDevResult, 0, sizeof(float)));
137 // Paramètre de l'appel de la fonction sur le device.
138 const dim3 dg(256, 1, 1);
139 const dim3 db(256, 1, 1);
140 Device::assertDim(dg, db);
141 const size_t SMSize = db.x * sizeof(double); // 256 doubles;
143 produitScalaire<<<dg, db, SMSize>>>(N, ptrDevResult);
146 // Barrière implicite de synchronisation ('cudaMemCpy').
147 HANDLE_ERROR(cudaMemcpy(&res, ptrDevResult, sizeof(float), cudaMemcpyDeviceToHost));
149 double resTheo = resultatTheorique(N);
152 cout << "Résultat : " << res << endl;
153 cout << "Résultat théorique : " << resTheo << endl;
154 cout << "Différence absolue : " << resTheo - res << endl;
155 cout << "Différence relatif : " << 100 * (resTheo - res) / (resTheo + res) << " %" << endl;