From: gburri Date: Wed, 8 Oct 2014 22:20:25 +0000 (+0200) Subject: Rippling CUDA Warmup et Smart. X-Git-Url: https://git.euphorik.ch/?a=commitdiff_plain;h=b1138708dfa104f784170fcf2d50989156e5c5d0;p=GPU.git 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 --- 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); \*-------------------------------------*/ Rippling0Image::Rippling0Image(unsigned int w, unsigned int h, float dt) : - ImageMOOs_A(w, h) + ImageMOOs_A(w, h), dt(dt), t(0) { - assert(getW() == getH()); // image carrer - - // Input - this->dt = dt; - - // Tools - this->t = 0; } Rippling0Image::~Rippling0Image(void) @@ -71,7 +64,7 @@ void Rippling0Image::animationStep(bool& isNeedUpdateView) // Override */ void Rippling0Image::fillImageGL(uchar4* ptrDevImageGL, int w, int h) // Override { - launchKernelRippling0(ptrDevImageGL, w, h, t); + launchKernelRippling0(ptrDevImageGL, w, h, this->t); } /** 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 --- 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); void launchKernelRippling0(uchar4* ptrDevPixels, int w, int h, float t) { - dim3 dg = dim3(8, 8, 1); // disons, a optimiser - dim3 db = dim3(16, 16, 1); // disons, a optimiser + dim3 dg = dim3(4, 4, 1); // disons, a optimiser + dim3 db = dim3(8, 8, 1); // disons, a optimiser //Device::print(dg, db); - Device::checkDimError(dg,db); + Device::checkDimError(dg,db); rippling0<<>>(ptrDevPixels,w,h,t); Device::checkKernelError("rippling0"); @@ -57,7 +57,7 @@ void launchKernelRippling0(uchar4* ptrDevPixels, int w, int h, float t) __global__ void rippling0(uchar4* ptrDevPixels, int w, int h, float t) { - Rippling0Math rippling0Math = Rippling0Math(w, h); + Rippling0Math rippling0Math(w, h); const int TID = Indice2D::tid(); const int NB_THREAD = Indice2D::nbThread(); 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 --- 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_ #define RIPPLING_0_MATH_H_ -#include +#include #include "MathTools.h" /*----------------------------------------------------------------------*\ @@ -23,14 +23,8 @@ class Rippling0Math __device__ Rippling0Math(int w, int h) + : dim2(w / 2.0) { - // TODO - } - - __device__ - Rippling0Math(const Rippling0Math& source) - { - // TODO } /*--------------------------------------*\ @@ -46,33 +40,32 @@ class Rippling0Math __device__ void color(int i, int j, float t, uchar4* ptrColor) { - // Debug (a mettre en commentaire) - { - unsigned char levelGris = 128; //in [0.255] (debug image) - ptrColor->x = levelGris; - ptrColor->y = levelGris; - ptrColor->z = levelGris; - } - - // Vrai problem - { - // TODO - } - - //color.w = 255; // opaque + const double dxy10 = dxy(j, i) / 10.0; + const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - 100.0 * t / 7.0) / (dxy10 + 1); + const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256); + + ptrColor->x = grayLevel; + ptrColor->y = grayLevel; + ptrColor->z = grayLevel; } private: + __device__ + double dxy(int x, int y) + { + return sqrt(pow(x - this->dim2, 2.0) + pow(y - this->dim2, 2.0)); + } /*--------------------------------------*\ |* Attributs *| \*-------------------------------------*/ private: + const double dim2; }; -#endif +#endif /*----------------------------------------------------------------------*\ |* End *| 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 --- 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 static Rippling0Image* createGL(void) { - float dt = 2*PI/1000; // animation para + float dt = 2 * PI / 1000; // animation para - int dw = 16 * 60; // =32*30=960 - int dh = 16 * 60; // =32*30=960 + int dw = 16 * 60; // =32*30=960 + int dh = 16 * 60; // =32*30=960 - return new Rippling0Image(dw, dh, dt); + return new Rippling0Image(dw, dh, dt); } }; 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 --- 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 __device__ RipplingMath(int w, int h) + : dim2(w / 2.0) { - // TODO - } - - __device__ - RipplingMath(const RipplingMath& source) - { - // TODO } /*--------------------------------------*\ @@ -45,35 +39,35 @@ class RipplingMath __device__ void color(int i, int j, float t, uchar4& color) { - // Debug (a mettre en commentaire) - { - unsigned char levelGris = 128; //in [0.255] (debug image) - color.x = levelGris; - color.y = levelGris; - color.z = levelGris; - } - - // Vrai problem - { - // TODO - } - - //color.w = 255; // opaque + const double dxy10 = dxy(j, i) / 10.0; + const double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0 / 10.0) / (dxy10 + 1); + const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256); + + color.x = grayLevel; + color.y = grayLevel; + color.z = grayLevel; } private: + __device__ + double dxy(int x, int y) + { + return sqrt(pow(x - this->dim2, 2.0) + pow(y - this->dim2, 2.0)); + } + /*--------------------------------------*\ |* Attributs *| \*-------------------------------------*/ private: + const double dim2; // Tools }; -#endif +#endif /*----------------------------------------------------------------------*\ |* End *| 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 --- 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 #include "Indice2D.h" +#include "IndiceTools.h" #include "cudaTools.h" #include "Device.h" @@ -41,9 +42,29 @@ __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t); __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t) { - RipplingMath ripplingMath = RipplingMath(w, h); + RipplingMath ripplingMath(w, h); - // TODO pattern entrelacement + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + + const int WH = w * h; + + uchar4 color; + color.z = 255; + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) + + ripplingMath.color(pixelI, pixelJ, t, color); // update color + ptrDevPixels[s] = color; + + s += NB_THREAD; + } } /*----------------------------------------------------------------------*\ 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 --- 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); \*-------------------------*/ Rippling::Rippling(int w, int h, float dt) + : w(w), h(h), dt(dt), t(0) { assert(w == h); - // Inputs - this->w = w; - this->h = h; - this->dt = dt; - // Tools - //this->dg = // TODO - //this->db = // TODO - this->t = 0; + this->dg = dim3(8, 8, 1); // disons a optimiser + this->db = dim3(16, 16, 1); // disons a optimiser // Outputs this->title = "Rippling Cuda"; @@ -72,7 +67,7 @@ Rippling::~Rippling() */ void Rippling::animationStep() { - // TODO + this->t += dt; } /** @@ -80,7 +75,7 @@ void Rippling::animationStep() */ void Rippling::runGPU(uchar4* ptrDevPixels) { - // TODO lancer le kernel avec <<>> + rippling<<>>(ptrDevPixels, this->w, this->h, this->t); } /*--------------*\ 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 --- 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); int mainGL(void) { - Rippling0Image* ptrRippling0 = Rippling0Provider::createGL(); + //Rippling0Image* ptrRippling0 = Rippling0Provider::createGL(); Image* ptrRippling = RipplingProvider::createGL(); // TODO : Insert autres Images ... bool isAnimation = true; bool isSelection = true; - GLUTImageViewers rippling0Viewer(ptrRippling0, isAnimation, isSelection, 0, 0); + //GLUTImageViewers rippling0Viewer(ptrRippling0, isAnimation, isSelection, 0, 0); GLUTImageViewers ripplingViewer(ptrRippling, isAnimation, isSelection, 10, 10); // TODO : Insert here autres ImageViewers ... @@ -57,11 +57,8 @@ int mainGL(void) // destruction { - delete ptrRippling0; + //delete ptrRippling0; delete ptrRippling; - - ptrRippling0 = NULL; - ptrRippling = NULL; } return EXIT_SUCCESS; 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 --- 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); \*-------------------------------------*/ Vague0Image::Vague0Image(unsigned int w, unsigned int h, float dt) : - ImageMOOs_A(w, h) + ImageMOOs_A(w, h), dt(dt), t(0) { - assert(getW() == getH()); // image carrer - - // Input - this->dt = dt; - - // Tools - this->t = 0; + assert(getW() == getH()); // image carrée } Vague0Image::~Vague0Image(void) 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 --- 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 \*-------------------------------------*/ public: - Vague0Image(unsigned int w, unsigned int h, float dt = 2 * PI / 1000); virtual ~Vague0Image(void); @@ -28,7 +27,6 @@ class Vague0Image: public ImageMOOs_A \*-------------------------------------*/ public: - /*----------------*\ |* Override *| \*---------------*/ @@ -53,10 +51,8 @@ class Vague0Image: public ImageMOOs_A \*-------------------------------------*/ private: - double dt; double t; - }; #endif