From: gburri <gregory.burri@master.hes-so.ch>
Date: Mon, 12 Jan 2015 16:55:39 +0000 (+0100)
Subject: TP Saucisson.
X-Git-Url: http://git.euphorik.ch/?p=GPU.git;a=commitdiff_plain;h=fba9c1cf9ca6c4064f01cae73145fd3286f4a316

TP Saucisson.
---

diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.cu b/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.cu
new file mode 100644
index 0000000..c4175be
--- /dev/null
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.cu
@@ -0,0 +1,138 @@
+#include "Saucisson.h"
+
+#include <iostream>
+#include <cmath>
+#include <stdio.h>
+using namespace std;
+
+#include "Indice1D.h"
+#include "cudaTools.h"
+#include "Device.h"
+#include "Lock.h"
+
+/*
+ * 1) Chaque thread calcule un résultat intermediaire qu'il va ensuite placer en shared memory.
+ * n: Nombre d'échantillon.
+ */
+__device__
+void reductionIntraThread(int n, float deltaX, float* tabSM)
+    {
+    const int NB_THREAD = Indice1D::nbThread();
+    const int TID = Indice1D::tid();
+    const int TID_LOCAL = Indice1D::tidLocal();
+
+    float threadResult = 0.0;
+    int s = TID;
+    while (s < n)
+        {
+        const float i = s + 1;
+        const float xi = -1 + i * deltaX;
+        threadResult += sqrtf(1-xi*xi);
+
+        s += NB_THREAD;
+        }
+
+    tabSM[TID_LOCAL] = threadResult;
+    }
+
+/*
+ * Combine les résultats de 'tabSM' dans 'tabSM[0]'
+ */
+__device__
+void combine(float* tabSM, int middle)
+    {
+    const int TID_LOCAL = Indice1D::tidLocal();
+    const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
+
+    int s = TID_LOCAL;
+    while (s < middle)
+        {
+        tabSM[s] += tabSM[s + middle];
+        s += NB_THREAD_LOCAL;
+        }
+    }
+
+/*
+ * 2) La shared memory est réduite, le résultat est placé dans 'tabSM[0]'.
+ */
+__device__
+void reductionIntraBlock(float* tabSM)
+    {
+    const int TAB_SIZE = blockDim.x;
+    int middle = TAB_SIZE / 2;
+
+    while (middle > 0)
+        {
+        combine(tabSM, middle);
+        middle /= 2;
+        __syncthreads(); // Synchronisation des threads au niveau du bloc.
+        }
+    }
+
+/*
+ * 3) Le 'tabSM[0]' de chaque bloc est reduit dans 'ptrResult' qui se trouve en global memory.
+ */
+__device__
+void reductionInterBlock(float* tabSM, float* ptrResult)
+    {
+    const int TID_LOCAL = Indice1D::tidLocal();
+    if (TID_LOCAL == 0)
+        {
+        atomicAdd(ptrResult, float(tabSM[0]));
+        }
+    }
+
+/**
+ * La taille de la shared memory (en terme de # de sizeof(float)) doit
+ * être égal à la taille des blocs.
+ * n: le nombre d'échantillon
+ * ptrResult: Le resultat du calcul de pi.
+ */
+__global__
+void saucisson(int n, float deltaX, float* ptrResult)
+    {
+    extern __shared__ float tabSM[]; // Dynamic shared memory.
+
+    // 1) Réduction intra-thread.
+    reductionIntraThread(n, deltaX, tabSM);
+
+    __syncthreads();
+
+    // 2) Réduction intra-block.
+    reductionIntraBlock(tabSM);
+
+    // 3) Réduction inter-block.
+    reductionInterBlock(tabSM, ptrResult);
+    }
+
+bool saucisson()
+    {
+    // Nombre d'échantillon. Au-delà, la qualité du résultat n'est pas meilleure. Il faudrait employé des doubles à la place de floats.
+    const int N = 100000;
+
+    // Allocation coté GPU en global memory (GM).
+    float* ptrDevResult = 0;
+    HANDLE_ERROR(cudaMalloc(&ptrDevResult, sizeof(float)));
+    HANDLE_ERROR(cudaMemset(ptrDevResult, 0, sizeof(float)));
+
+    // 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(float); // 256 floats;
+
+    const float deltaX = 2.0f / N;
+    saucisson<<<dg, db, SMSize>>>(N, deltaX, ptrDevResult);
+
+    // cudaDeviceSynchronize(); // Utilisé pour flusher les prints sur le stdout à partir du device (debug).
+
+    float pi;
+    // Barrière implicite de synchronisation ('cudaMemCpy').
+    HANDLE_ERROR(cudaMemcpy(&pi, ptrDevResult, sizeof(float), cudaMemcpyDeviceToHost));
+    pi *= 2 * deltaX;
+
+    cout.precision(20);
+    cout << "Approximation de PI : " << pi << endl;
+
+    return true;
+    }
diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.h b/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.h
new file mode 100644
index 0000000..0c26ba2
--- /dev/null
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/01c_Saucisson/Saucisson.h
@@ -0,0 +1,6 @@
+#ifndef SAUCISSON_H
+#define SAUCISSON_H
+
+bool saucisson();
+
+#endif
diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp b/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
index d16b105..be769f8 100755
--- a/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
@@ -2,6 +2,7 @@
 #include <stdlib.h>
 using namespace std;
 
+#include "01c_Saucisson/Saucisson.h"
 #include "02_ProduitScalaire/ProduitScalaire.h"
 
 extern bool useHello();
@@ -11,8 +12,9 @@ int mainCore()
     {
     bool isOk = true;
     /*isOk &= useHello();
-    isOk &= addVectors();*/
-    isOk &= produitScalaire();
+    isOk &= addVectors();
+    isOk &= produitScalaire();*/
+    isOk &= saucisson();
 
     cout << "\nisOK = " << isOk << endl;
     cout << "\nEnd : mainCore" << endl;
diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
index 817b37b..9bcc12b 100644
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
@@ -51,6 +51,7 @@ void Convolution::runGPU(uchar4* ptrDevPixels)
     toGrayscale<<<dg,db>>>(this->ptrDevImageSource, this->w, this->h);
 
     cudaDeviceSynchronize(); // Attend que toute l'image source ait été passée en niveau de gris.
+
     convolution<<<dg,db>>>(ptrDevPixels, this->w, this->h);
 
     //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).