Rippling CUDA Warmup et Smart.

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/01_imageAPI/Rippling0Image.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/01_imageAPI/Rippling0Image.cpp
index dc9e211..310fad7 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/01_imageAPI/Rippling0Image.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/01_imageAPI/Rippling0Image.cpp
@@ -37,15 +37,8 @@ extern void launchKernelRippling0(uchar4* ptrDevPixels, int w, int h, float t);
  \*-------------------------------------*/\r
 \r
 Rippling0Image::Rippling0Image(unsigned int w, unsigned int h, float dt) :\r
-       ImageMOOs_A(w, h)\r
+       ImageMOOs_A(w, h), dt(dt), t(0)\r
     {\r
-    assert(getW() == getH()); // image carrer\r
-\r
-           // Input\r
-    this->dt = dt;\r
-\r
-    // Tools\r
-    this->t = 0;\r
     }\r
 \r
 Rippling0Image::~Rippling0Image(void)\r
@@ -71,7 +64,7 @@ void Rippling0Image::animationStep(bool& isNeedUpdateView) // Override
  */\r
 void Rippling0Image::fillImageGL(uchar4* ptrDevImageGL, int w, int h) // Override\r
     {\r
-    launchKernelRippling0(ptrDevImageGL, w, h, t);\r
+    launchKernelRippling0(ptrDevImageGL, w, h, this->t);\r
     }\r
 \r
 /**\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/02_cuda/rippling0.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/02_cuda/rippling0.cu
index eb778bf..376d4c3 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/02_cuda/rippling0.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/02_cuda/rippling0.cu
@@ -41,11 +41,11 @@ static __global__ void rippling0(uchar4* ptrDevPixels,int w, int h,float t);
 \r
 void launchKernelRippling0(uchar4* ptrDevPixels, int w, int h, float t)\r
     {\r
-    dim3 dg = dim3(8, 8, 1); // disons, a optimiser\r
-    dim3 db = dim3(16, 16, 1); // disons, a optimiser\r
+    dim3 dg = dim3(4, 4, 1); // disons, a optimiser\r
+    dim3 db = dim3(8, 8, 1); // disons, a optimiser\r
 \r
     //Device::print(dg, db);\r
-     Device::checkDimError(dg,db);\r
+    Device::checkDimError(dg,db);\r
 \r
     rippling0<<<dg,db>>>(ptrDevPixels,w,h,t);\r
     Device::checkKernelError("rippling0");\r
@@ -57,7 +57,7 @@ void launchKernelRippling0(uchar4* ptrDevPixels, int w, int h, float t)
 \r
 __global__ void rippling0(uchar4* ptrDevPixels, int w, int h, float t)\r
     {\r
-    Rippling0Math rippling0Math = Rippling0Math(w, h);\r
+    Rippling0Math rippling0Math(w, h);\r
 \r
     const int TID = Indice2D::tid();\r
     const int NB_THREAD = Indice2D::nbThread();\r

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 67b4387..3879d1d 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
@@ -1,7 +1,7 @@
 #ifndef RIPPLING_0_MATH_H_\r
 #define RIPPLING_0_MATH_H_\r
 \r
-#include <math.h>\r
+#include <cmath>\r
 #include "MathTools.h"\r
 \r
 /*----------------------------------------------------------------------*\\r
@@ -23,14 +23,8 @@ class Rippling0Math
 \r
        __device__\r
        Rippling0Math(int w, int h)\r
+           : dim2(w / 2.0)\r
            {\r
-           // TODO\r
-           }\r
-\r
-       __device__\r
-       Rippling0Math(const Rippling0Math& source)\r
-           {\r
-           // TODO\r
            }\r
 \r
        /*--------------------------------------*\\r
@@ -46,33 +40,32 @@ class Rippling0Math
        __device__\r
        void color(int i, int j, float t, uchar4* ptrColor)\r
            {\r
-           // Debug (a mettre en commentaire)\r
-               {\r
-               unsigned char levelGris = 128; //in [0.255]  (debug image)\r
-               ptrColor->x = levelGris;\r
-               ptrColor->y = levelGris;\r
-               ptrColor->z = levelGris;\r
-               }\r
-\r
-           // Vrai problem\r
-               {\r
-               // TODO\r
-               }\r
-\r
-           //color.w = 255; // opaque\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 uchar grayLevel = (uchar)(long(grayLevelFloat) % 256);\r
+\r
+           ptrColor->x = grayLevel;\r
+           ptrColor->y = grayLevel;\r
+           ptrColor->z = grayLevel;\r
            }\r
 \r
     private:\r
+       __device__\r
+       double dxy(int x, int y)\r
+           {\r
+           return sqrt(pow(x - this->dim2, 2.0) + pow(y - this->dim2, 2.0));\r
+           }\r
 \r
        /*--------------------------------------*\\r
        |*              Attributs               *|\r
         \*-------------------------------------*/\r
 \r
     private:\r
+       const double dim2;\r
 \r
     };\r
 \r
-#endif \r
+#endif\r
 \r
 /*----------------------------------------------------------------------*\\r
  |*                    End                                             *|\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 bc730be..48b9ab3 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,12 +18,12 @@ class Rippling0Provider
 \r
        static Rippling0Image* createGL(void)\r
            {\r
-            float dt = 2*PI/1000; // animation para\r
+            float dt = 2 * PI / 1000; // animation para\r
 \r
-           int dw = 16 * 60; // =32*30=960\r
-           int dh = 16 * 60; // =32*30=960\r
+            int dw = 16 * 60; // =32*30=960\r
+            int dh = 16 * 60; // =32*30=960\r
 \r
-           return new Rippling0Image(dw, dh, dt);\r
+            return new Rippling0Image(dw, dh, dt);\r
            }\r
 \r
     };\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 74c4dfc..655fa09 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
@@ -22,14 +22,8 @@ class RipplingMath
 \r
        __device__\r
        RipplingMath(int w, int h)\r
+           : dim2(w / 2.0)\r
            {\r
-           // TODO\r
-           }\r
-\r
-       __device__\r
-       RipplingMath(const RipplingMath& source)\r
-           {\r
-           // TODO\r
            }\r
 \r
        /*--------------------------------------*\\r
@@ -45,35 +39,35 @@ class RipplingMath
        __device__\r
        void color(int i, int j, float t, uchar4& color)\r
            {\r
-           // Debug (a mettre en commentaire)\r
-               {\r
-               unsigned char levelGris = 128; //in [0.255]  (debug image)\r
-               color.x = levelGris;\r
-               color.y = levelGris;\r
-               color.z = levelGris;\r
-               }\r
-\r
-           // Vrai problem\r
-               {\r
-               // TODO\r
-               }\r
-\r
-           //color.w = 255; // opaque\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 uchar grayLevel = (uchar)(long(grayLevelFloat) % 256);\r
+\r
+           color.x = grayLevel;\r
+           color.y = grayLevel;\r
+           color.z = grayLevel;\r
            }\r
 \r
     private:\r
+       __device__\r
+       double dxy(int x, int y)\r
+           {\r
+           return sqrt(pow(x - this->dim2, 2.0) + pow(y - this->dim2, 2.0));\r
+           }\r
+\r
 \r
        /*--------------------------------------*\\r
        |*              Attributs               *|\r
         \*-------------------------------------*/\r
 \r
     private:\r
+       const double dim2;\r
 \r
        // Tools\r
 \r
     };\r
 \r
-#endif \r
+#endif\r
 \r
 /*----------------------------------------------------------------------*\\r
  |*                    End                                             *|\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 d599127..3e3235f 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
@@ -1,6 +1,7 @@
 #include <iostream>\r
 \r
 #include "Indice2D.h"\r
+#include "IndiceTools.h"\r
 #include "cudaTools.h"\r
 #include "Device.h"\r
 \r
@@ -41,9 +42,29 @@ __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t);
 \r
 __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t)\r
     {\r
-    RipplingMath ripplingMath = RipplingMath(w, h);\r
+    RipplingMath ripplingMath(w, h);\r
 \r
-    // TODO pattern entrelacement\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
+\r
+    int pixelI;\r
+    int pixelJ;\r
+\r
+    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
+       ptrDevPixels[s] = color;\r
+\r
+       s += NB_THREAD;\r
+       }\r
     }\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 6bd9e0c..eb49b08 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,18 +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
     {\r
     assert(w == h);\r
 \r
-    // Inputs\r
-    this->w = w;\r
-    this->h = h;\r
-    this->dt = dt;\r
-\r
     // Tools\r
-    //this->dg = // TODO\r
-    //this->db = // TODO\r
-    this->t = 0;\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
@@ -72,7 +67,7 @@ Rippling::~Rippling()
  */\r
 void Rippling::animationStep()\r
     {\r
-    // TODO\r
+    this->t += dt;\r
     }\r
 \r
 /**\r
@@ -80,7 +75,7 @@ void Rippling::animationStep()
  */\r
 void Rippling::runGPU(uchar4* ptrDevPixels)\r
     {\r
-    // TODO lancer le kernel avec <<<dg,db>>>\r
+    rippling<<<dg,db>>>(ptrDevPixels, this->w, this->h, this->t);\r
     }\r
 \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 44ee751..9fbf6e8 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
@@ -42,14 +42,14 @@ 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
 \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
@@ -57,11 +57,8 @@ int mainGL(void)
 \r
     // destruction\r
        {\r
-       delete ptrRippling0;\r
+       //delete ptrRippling0;\r
        delete ptrRippling;\r
-\r
-       ptrRippling0 = NULL;\r
-       ptrRippling = NULL;\r
        }\r
 \r
     return EXIT_SUCCESS;\r

diff --git a/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.cpp b/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.cpp
index c91d361..3b8a207 100755 (executable)
--- a/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.cpp
+++ b/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.cpp
@@ -36,15 +36,9 @@ extern void launchKernelVague0(uchar4* ptrDevPixels, int w, int h, float t);
  \*-------------------------------------*/\r
 \r
 Vague0Image::Vague0Image(unsigned int w, unsigned int h, float dt) :\r
-       ImageMOOs_A(w, h)\r
+       ImageMOOs_A(w, h), dt(dt), t(0)\r
     {\r
-    assert(getW() == getH()); // image carrer\r
-\r
-           // Input\r
-    this->dt = dt;\r
-\r
-    // Tools\r
-    this->t = 0;\r
+    assert(getW() == getH()); // image carrée\r
     }\r
 \r
 Vague0Image::~Vague0Image(void)\r

diff --git a/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.h b/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.h
index dba0119..4970f5f 100755 (executable)
--- a/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.h
+++ b/WCudaMSE/Tuto_Image_Cuda/src/cpp/core/00_Vague_warmup/01_imageAPI/Vague0Image.h
@@ -19,7 +19,6 @@ class Vague0Image: public ImageMOOs_A
         \*-------------------------------------*/\r
 \r
     public:\r
-\r
        Vague0Image(unsigned int w, unsigned int h, float dt = 2 * PI / 1000);\r
        virtual ~Vague0Image(void);\r
 \r
@@ -28,7 +27,6 @@ class Vague0Image: public ImageMOOs_A
         \*-------------------------------------*/\r
 \r
     public:\r
-\r
        /*----------------*\\r
        |*  Override      *|\r
        \*---------------*/\r
@@ -53,10 +51,8 @@ class Vague0Image: public ImageMOOs_A
         \*-------------------------------------*/\r
 \r
     private:\r
-\r
        double dt;\r
        double t;\r
-\r
     };\r
 \r
 #endif\r