* Ajout d'un exemple CUDA non-openGL (AddVector.cu)

* Ajustement du dt pour rippling cuda (+warmup)

diff --git a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Animable_I.h b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Animable_I.h
index 3568641..eba8685 100755 (executable)
--- a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Animable_I.h
+++ b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Animable_I.h
@@ -29,10 +29,9 @@ class Animable_I
        /**\r
         * getParaAnimation\r
         */\r
-       virtual float getT()=0;\r
-\r
-       virtual string getTitle()=0;\r
+       virtual float getT() = 0;\r
 \r
+       virtual string getTitle() = 0;\r
     };\r
 \r
 #endif\r

diff --git a/WCudaMSE/RELEASE/doc/Student_Cuda_64.doc.tar.gz b/WCudaMSE/RELEASE/doc/Student_Cuda_64.doc.tar.gz
index a45f077..95bf645 100644 (file)
Binary files a/WCudaMSE/RELEASE/doc/Student_Cuda_64.doc.tar.gz and b/WCudaMSE/RELEASE/doc/Student_Cuda_64.doc.tar.gz differ

diff --git a/WCudaMSE/Student_Cuda/.project b/WCudaMSE/Student_Cuda/.project
index b87b289..8f211bf 100755 (executable)
--- a/WCudaMSE/Student_Cuda/.project
+++ b/WCudaMSE/Student_Cuda/.project
@@ -7,7 +7,7 @@
        <buildSpec>\r
                <buildCommand>\r
                        <name>org.eclipse.cdt.managedbuilder.core.genmakebuilder</name>\r
-                       <triggers>full,incremental,</triggers>\r
+                       <triggers>clean,full,incremental,</triggers>\r
                        <arguments>\r
                                <dictionary>\r
                                        <key>?name?</key>\r

diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/01_Hello/02_hello_add.cu b/WCudaMSE/Student_Cuda/src/cpp/core/01_Hello/02_hello_add.cu
index 70f3908..e4b3275 100755 (executable)
--- a/WCudaMSE/Student_Cuda/src/cpp/core/01_Hello/02_hello_add.cu
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/01_Hello/02_hello_add.cu
@@ -63,7 +63,7 @@ __host__ int addScalarGPU(int a, int b)
     int* ptrDev_c;     // on device (GPU)\r
 \r
     // Specifier nb thread : ici 1 thread au total !\r
-    dim3 dg = dim3(1,1,1);\r
+    dim3 dg = dim3(1, 1, 1);\r
     dim3 db = dim3(1, 1, 1);\r
 \r
     // Debug\r

diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/02_AddVector/AddVector.cu b/WCudaMSE/Student_Cuda/src/cpp/core/02_AddVector/AddVector.cu
new file mode 100644 (file)
index 0000000..766480e
--- /dev/null
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/02_AddVector/AddVector.cu
@@ -0,0 +1,65 @@
+#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;
+    }

diff --git a/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp b/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
index ad144a1..31af0c5 100755 (executable)
--- a/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
+++ b/WCudaMSE/Student_Cuda/src/cpp/core/mainCore.cpp
@@ -14,6 +14,7 @@ using std::endl;
  \*-------------------------------------*/
 
 extern bool useHello(void);
+extern bool addVectors();
 
 /*--------------------------------------*\
  |*            Public                  *|
@@ -39,6 +40,7 @@ int mainCore()
     {
     bool isOk = true;
     isOk &= useHello();
+    isOk &= addVectors();
 
     cout << "\nisOK = " << isOk << endl;
     cout << "\nEnd : mainCore" << endl;

diff --git a/WCudaMSE/Student_Cuda/src/cpp/main.cpp b/WCudaMSE/Student_Cuda/src/cpp/main.cpp
index 3fad450..ad4d735 100755 (executable)
--- a/WCudaMSE/Student_Cuda/src/cpp/main.cpp
+++ b/WCudaMSE/Student_Cuda/src/cpp/main.cpp
@@ -90,7 +90,7 @@ int start(void)
     {
     Device::printCurrent();
 
-    bool IS_TEST = true;
+    bool IS_TEST = false;
 
     if (IS_TEST)
        {

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/03_math/Rippling0Math.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/03_math/Rippling0Math.h
index 3879d1d..ca71c75 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/03_math/Rippling0Math.h
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/03_math/Rippling0Math.h
@@ -41,7 +41,7 @@ class Rippling0Math
        void color(int i, int j, float t, uchar4* ptrColor)\r
            {\r
            const double dxy10 = dxy(j, i) / 10.0;\r
-           const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - 100.0 * t / 7.0) / (dxy10 + 1);\r
+           const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0) / (dxy10 + 1);\r
            const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256);\r
 \r
            ptrColor->x = grayLevel;\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/Rippling0Provider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/Rippling0Provider.h
index 48b9ab3..d5472a2 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/Rippling0Provider.h
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/Rippling0Provider.h
@@ -18,7 +18,7 @@ class Rippling0Provider
 \r
        static Rippling0Image* createGL(void)\r
            {\r
-            float dt = 2 * PI / 1000; // animation para\r
+            float dt = 1;\r
 \r
             int dw = 16 * 60; // =32*30=960\r
             int dh = 16 * 60; // =32*30=960\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/math/RipplingMath.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/math/RipplingMath.h
index 655fa09..e90a060 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/math/RipplingMath.h
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/math/RipplingMath.h
@@ -40,7 +40,7 @@ class RipplingMath
        void color(int i, int j, float t, uchar4& color)\r
            {\r
            const double dxy10 = dxy(j, i) / 10.0;\r
-           const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0 / 10.0) / (dxy10 + 1);\r
+           const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0) / (dxy10 + 1);\r
            const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256);\r
 \r
            color.x = grayLevel;\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/ripplingDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/ripplingDevice.cu
index 3e3235f..3dc4557 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/ripplingDevice.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/device/ripplingDevice.cu
@@ -46,11 +46,10 @@ __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t)
 \r
     const int TID = Indice2D::tid();\r
     const int NB_THREAD = Indice2D::nbThread();\r
-\r
     const int WH = w * h;\r
 \r
     uchar4 color;\r
-    color.z = 255;\r
+    color.z = 255; // Par défaut, l'image est opaque.\r
 \r
     int pixelI;\r
     int pixelJ;\r
@@ -58,11 +57,9 @@ __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t)
     int s = TID;\r
     while (s < WH)\r
        {\r
-       IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ)\r
-\r
-       ripplingMath.color(pixelI, pixelJ, t, color);   // update color\r
+       IndiceTools::toIJ(s, w, &pixelI, &pixelJ);\r
+       ripplingMath.color(pixelI, pixelJ, t, color);\r
        ptrDevPixels[s] = color;\r
-\r
        s += NB_THREAD;\r
        }\r
     }\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/host/Rippling.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/host/Rippling.cu
index eb49b08..43068a2 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/host/Rippling.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/moo/host/Rippling.cu
@@ -38,17 +38,13 @@ extern __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t);
  \*-------------------------*/\r
 \r
 Rippling::Rippling(int w, int h, float dt)\r
-    : w(w), h(h), dt(dt), t(0)\r
+    : w(w), h(h), dt(dt), t(0),\r
+      dg(8, 8, 1),\r
+      db(16, 16, 1),\r
+      title("Rippling Cuda")\r
     {\r
     assert(w == h);\r
 \r
-    // Tools\r
-    this->dg = dim3(8, 8, 1); // disons a optimiser\r
-    this->db = dim3(16, 16, 1); // disons a optimiser\r
-\r
-    // Outputs\r
-    this->title = "Rippling Cuda";\r
-\r
     //print(dg, db);\r
     Device::assertDim(dg, db);\r
     }\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/provider/RipplingProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/provider/RipplingProvider.cpp
index 772073f..cade757 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/provider/RipplingProvider.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/01_Rippling/provider/RipplingProvider.cpp
@@ -31,10 +31,9 @@
 \r
 Rippling* RipplingProvider::createMOO()\r
     {\r
-    float dt = 1;\r
-\r
-    int dw = 16 * 60; // =32*30=960\r
-    int dh = 16 * 60; // =32*30=960\r
+    const float dt = 1;\r
+    const int dw = 16 * 60; // =32*30=960\r
+    const int dh = 16 * 60; // =32*30=960\r
 \r
     return new Rippling(dw, dh, dt);\r
     }\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
index 9fbf6e8..4b6a217 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
@@ -42,24 +42,22 @@ int mainGL(void);
 \r
 int mainGL(void)\r
     {\r
-    //Rippling0Image* ptrRippling0 = Rippling0Provider::createGL();\r
+    Rippling0Image* ptrRippling0 = Rippling0Provider::createGL();\r
     Image* ptrRippling = RipplingProvider::createGL();\r
     // TODO : Insert  autres Images ...\r
 \r
-    bool isAnimation = true;\r
-    bool isSelection = true;\r
+    const bool isAnimation = true;\r
+    const bool isSelection = true;\r
 \r
-    //GLUTImageViewers rippling0Viewer(ptrRippling0, isAnimation, isSelection, 0, 0);\r
+    GLUTImageViewers rippling0Viewer(ptrRippling0, isAnimation, isSelection, 0, 0);\r
     GLUTImageViewers ripplingViewer(ptrRippling, isAnimation, isSelection, 10, 10);\r
     // TODO : Insert here autres ImageViewers ...\r
 \r
     GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte\r
 \r
     // destruction\r
-       {\r
-       //delete ptrRippling0;\r
-       delete ptrRippling;\r
-       }\r
+    delete ptrRippling0;\r
+    delete ptrRippling;\r
 \r
     return EXIT_SUCCESS;\r
     }\r