* The competing threads must then wait until the owner has written a 0 to the mutex beforfe they can attempt to modify the locked memory.\r
*\r
* Interet\r
- * On utilise un lock lorsque l'opération qui doit être synchroniser ne possède pas d'opérateur atomic (comme atomicADD, ...),\r
- * ou lorsqui'il s'agit de plusieurs opérations à synchroniser (joue alors le role de section critique)\r
+ * On utilise un lock lorsque l'op�ration qui doit �tre synchroniser ne poss�de pas d'op�rateur atomic (comme atomicADD, ...),\r
+ * ou lorsqui'il s'agit de plusieurs op�rations � synchroniser (joue alors le role de section critique)\r
*\r
* Note :\r
*\r
- * Lock ne laisse aucune trace coté host, il s'instancie only coté device: Code moins invasif\r
- * LockMixte laisse une trace coté host. Code plus invasif\r
+ * Lock ne laisse aucune trace cot� host, il s'instancie only cot� device: Code moins invasif\r
+ * LockMixte laisse une trace cot� host. Code plus invasif\r
*\r
* Use (Device side only)\r
*\r
*\r
* // Global variable of .cu\r
- * __device__ int mutex=0; // Attention à l'initialisation\r
+ * __device__ int mutex=0; // Attention � l'initialisation\r
*\r
* // variable local inside a kernel (same .cu as variable mutex)\r
* Lock lock=Lock(&mutex);\r
// Solution:\r
// atomicCAS = atomic Compare And Swap\r
// Prototype : c atomicCAS(ptr,a,b)\r
- // Action : compare ptr avec a, si egale affecte b à ptr, renvoie ptr\r
+ // Action : compare ptr avec a, si egale affecte b � ptr, renvoie ptr\r
\r
// Tant que ptrDev_mutex!=0 le thread cuda boucle sur while\r
- // Des qu'il vaut 0, il met le mutex à 1 et lock se termine\r
+ // Des qu'il vaut 0, il met le mutex � 1 et lock se termine\r
while (atomicCAS(ptrDevMutexGM, 0, 1) != 0);\r
}\r
\r
// Solution 1:\r
//\r
// *ptrDev_mutex=0;\r
- // Aucun thread en competition ici. L'affectation n'a pas besoin d'être atomique.\r
+ // Aucun thread en competition ici. L'affectation n'a pas besoin d'�tre atomique.\r
// Solution satisfaisante.\r
//\r
// Solution 2 (prefered for symetric approach)\r
int* ptrDevMutexGM;\r
};\r
\r
-#endif \r
+#endif\r
\r
/*----------------------------------------------------------------------*\\r
|* End *|\r
--- /dev/null
+#include <iostream>
+
+#include "Indice2D.h"
+#include "cudaTools.h"
+#include "Device.h"
+
+using std::cout;
+using std::endl;
+
+static __global__ void add(float* ptrDevV1, float* ptrDevV2, int n, float* ptrDevResult);
+static __device__ float work(float v1, float v2);
+
+__global__ void add(float* ptrDevV1, float* ptrDevV2, int n, float* ptrDevResult)
+ {
+ const int NB_THREAD = Indice2D::nbThread();
+ const int TID = Indice2D::tid();
+
+ int s = TID;
+
+ while (s < n)
+ {
+ ptrDevResult[s] = work(ptrDevV1[s], ptrDevV2[s]);
+ s += NB_THREAD;
+ }
+ }
+
+__device__ float work(float v1, float v2)
+ {
+ return v1 + v2;
+ }
+
+bool addVectors()
+ {
+ // Inputs (passé en paramètre de la fonction dans un cas général).
+ float v1[] = { 1, 2, 3 };
+ float v2[] = { 10, 20, 30 };
+
+ // Outputs (renvoyer de la fonction dans un cas général).
+ float vRes[3];
+
+ // Allocation coté GPU.
+ float* ptrDevV1, *ptrDevV2, *ptrDevVResult = 0;
+ const size_t vecSize = 3 * sizeof(float);
+ HANDLE_ERROR(cudaMalloc(&ptrDevV1, vecSize));
+ HANDLE_ERROR(cudaMalloc(&ptrDevV2, vecSize));
+ HANDLE_ERROR(cudaMalloc(&ptrDevVResult, vecSize));
+
+ HANDLE_ERROR(cudaMemset(ptrDevV1, 0, vecSize));
+ HANDLE_ERROR(cudaMemset(ptrDevV2, 0, vecSize));
+ HANDLE_ERROR(cudaMemset(ptrDevVResult, 0, vecSize));
+
+ HANDLE_ERROR(cudaMemcpy(ptrDevV1, v1, vecSize, cudaMemcpyHostToDevice));
+ HANDLE_ERROR(cudaMemcpy(ptrDevV2, v2, vecSize, cudaMemcpyHostToDevice));
+
+ const dim3 dg(2, 2, 1);
+ const dim3 db(2, 2, 1);
+ Device::assertDim(dg, db);
+
+ add<<<dg, db>>>(ptrDevV1, ptrDevV2, 3, ptrDevVResult);
+
+ // Barrière implicite de synchronisation ('cudaMemCpy').
+ HANDLE_ERROR(cudaMemcpy(vRes, ptrDevVResult, vecSize, cudaMemcpyDeviceToHost));
+
+ return vRes[0] == 11 && vRes[1] == 22 && vRes[2] == 33;
+ }
+#include "ProduitScalaire.h"
+
#include <iostream>
#include <cmath>
+#include <stdio.h>
using namespace std;
#include "Indice1D.h"
#include "cudaTools.h"
#include "Device.h"
+#include "Lock.h"
#define M_V 200
#define M_W 200
#define VI 1.4422495703074083
-#define WI 0.7390850782394409
+#define WI 0.739085133215160672293109200837
/*
* Renvoie la valeur du ième élément du vecteur v.
const double xCarre = x * x;
x = x - (xCarre * x - 3) / (3 * xCarre);
}
- return (x / VI) * sqrt(double(i));
+
+ /* Debug afin d'ajuster VI
+ if (Indice1D::tid() == 0)
+ printf("x: %.30f, VI: %.30f, x / VI: %.30f\n", x, VI, x / VI);
+ */
+
+ return (x / VI) * sqrt(double(i)); // x / VI doit être égal à 1.
}
/*
double x = abs(cos(double(i)));
for (long j = 1; j <= M_W; j++)
x = x - (cos(x) - x) / (-sin(x) - 1);
+
+ /* Debug afin d'ajuster WI
+ if (Indice1D::tid() == 0)
+ printf("x: %.30f, WI: %.30f, x / WI: %.30f\n", x, WI, x / WI); */
+
return (x / WI) * sqrt(double(i));
}
/*
+ * 1) Chaque thread calcule un résultat intermediaire qu'il va ensuite placer en shared memory.
* n: La taille des deux vecteurs.
*/
-__device__ void reductionIntraThread(int n, double* tabResultSM)
+__device__
+void reductionIntraThread(int n, double* tabSM)
{
const int NB_THREAD = Indice1D::nbThread();
const int TID = Indice1D::tid();
threadResult += v(s) * w(s);
s += NB_THREAD;
}
- tabResultSM[TID_LOCAL] = threadResult;
+
+ tabSM[TID_LOCAL] = threadResult;
}
/*
- * Combine les résultats de 'tabResultSM' dans 'tabResulSM[0]'
+ * Combine les résultats de 'tabSM' dans 'tabSM[0]'
*/
-__device__ void combine(double* tabResultSM, int middle)
+__device__
+void combine(double* tabSM, int middle)
{
- const int TID_LOCAL = Indice1D::tidLocal();
- const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
-
- int s = TID_LOCAL;
- while (s < middle)
- {
- tabResultSM[s] = tabResultSM[s] + tabResultSM[s + middle];
- s += NB_THREAD_LOCAL;
- }
+ const int TID_LOCAL = Indice1D::tidLocal();
+ const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
+
+ int s = TID_LOCAL;
+ while (s < middle)
+ {
+ tabSM[s] = tabSM[s] + tabSM[s + middle];
+ s += NB_THREAD_LOCAL;
+ }
}
-__device__ void reductionIntraBlock(double* tabResultSM)
+/*
+ * 2) La shared memory est réduite, le résultat est placé dans 'tabSM[0]'.
+ */
+__device__
+void reductionIntraBlock(double* tabSM)
{
const int TAB_SIZE = blockDim.x;
int middle = TAB_SIZE / 2;
while (middle > 0)
{
- combine(tabResultSM, middle);
+ combine(tabSM, middle);
middle /= 2;
__syncthreads();
}
}
-__device__ void reductionInterBlock(double* tabResultSM, float* ptrResult)
+__device__
+int mutexReductionInterBlock = 0;
+
+/*
+ * 3) Le 'tabSM[0]' de chaque bloc est reduit dans 'ptrResult' qui se trouve en global memory.
+ */
+__device__
+void reductionInterBlock(double* tabSM, double* ptrResult)
{
const int TID_LOCAL = Indice1D::tidLocal();
if (TID_LOCAL == 0)
{
- atomicAdd(ptrResult, float(tabResultSM[0]));
+ Lock lock(&mutexReductionInterBlock);
+ lock.lock();
+ (*ptrResult) += tabSM[0];
+ lock.unlock();
+
+ // Si on travail en float (pas besoin de mutex) :
+ // atomicAdd(ptrResult, float(tabSM[0]));
}
}
* ptrResult: Le resultat du produit scalaire.
*/
__global__
-void produitScalaire(int n, float* ptrResult)
+void produitScalaire(int n, double* ptrResult)
{
- extern __shared__ double tabResultSM[]; // Shared memory.
+ extern __shared__ double tabSM[]; // Dynamic shared memory.;
// 1) Réduction intra-thread.
- reductionIntraThread(n, tabResultSM);
+ reductionIntraThread(n, tabSM);
__syncthreads();
// 2) Réduction intra-block.
- reductionIntraBlock(tabResultSM);
+ reductionIntraBlock(tabSM);
// 3) Réduction inter-block.
- reductionInterBlock(tabResultSM, ptrResult);
+ reductionInterBlock(tabSM, ptrResult);
}
double resultatTheorique(long n)
-{
+ {
n -= 1;
- return (n / 2.0) * (n+1);
-}
+ return (n / 2.0) * (n + 1);
+ }
bool produitScalaire()
{
- const int N = 100000000; // Taille des deux vecteurs.
+ const int N = 10000000; // Taille des deux vecteurs : 10 * 10^6.
// Allocation coté GPU en global memory (GM).
- float* ptrDevResult = 0;
- HANDLE_ERROR(cudaMalloc(&ptrDevResult, sizeof(float)));
- HANDLE_ERROR(cudaMemset(ptrDevResult, 0, sizeof(float)));
+ double* ptrDevResult = 0;
+ HANDLE_ERROR(cudaMalloc(&ptrDevResult, sizeof(double)));
+ HANDLE_ERROR(cudaMemset(ptrDevResult, 0, sizeof(double)));
// Paramètre de l'appel de la fonction sur le device.
const dim3 dg(256, 1, 1);
produitScalaire<<<dg, db, SMSize>>>(N, ptrDevResult);
- float res;
+ cudaDeviceSynchronize(); // Utilisé pour flusher les prints sur le stdout à partir du device.
+
+ double res;
// Barrière implicite de synchronisation ('cudaMemCpy').
- HANDLE_ERROR(cudaMemcpy(&res, ptrDevResult, sizeof(float), cudaMemcpyDeviceToHost));
+ HANDLE_ERROR(cudaMemcpy(&res, ptrDevResult, sizeof(double), cudaMemcpyDeviceToHost));
- double resTheo = resultatTheorique(N);
+ const double resTheo = resultatTheorique(N);
- cout.precision(10);
+ cout.precision(30);
cout << "Résultat : " << res << endl;
cout << "Résultat théorique : " << resTheo << endl;
cout << "Différence absolue : " << resTheo - res << endl;
projects / GPU.git / commitdiff