X-Git-Url: http://git.euphorik.ch/index.cgi?a=blobdiff_plain;f=WCudaMSE%2FStudent_Cuda_Image%2Fsrc%2Fcpp%2Fcore%2F00_Rippling_warmup%2F02_cuda%2Frippling0.cu;fp=WCudaMSE%2FStudent_Cuda_Image%2Fsrc%2Fcpp%2Fcore%2F00_Rippling_warmup%2F02_cuda%2Frippling0.cu;h=eb778bf0bff1a3787ec53ff9cf24baa423a7a05f;hb=8d08c12b29c2a14684f35c023ee39e694bb80d25;hp=0000000000000000000000000000000000000000;hpb=226de81f7e1f1fbf4ac79d0d089e8a05ec7159a0;p=GPU.git 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 new file mode 100755 index 0000000..eb778bf --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/00_Rippling_warmup/02_cuda/rippling0.cu @@ -0,0 +1,91 @@ +#include +#include + +#include "Indice2D.h" +#include "IndiceTools.h" +#include "cudaTools.h" +#include "Device.h" + +#include "Rippling0Math.h" + +using std::cout; +using std::endl; + +/*----------------------------------------------------------------------*\ + |* Declaration *| + \*---------------------------------------------------------------------*/ + +/*--------------------------------------*\ + |* Imported *| + \*-------------------------------------*/ + +/*--------------------------------------*\ + |* Public *| + \*-------------------------------------*/ + + + +/*--------------------------------------*\ + |* Private *| + \*-------------------------------------*/ + +static __global__ void rippling0(uchar4* ptrDevPixels,int w, int h,float t); + +/*----------------------------------------------------------------------*\ + |* Implementation *| + \*---------------------------------------------------------------------*/ + +/*--------------------------------------*\ + |* Public *| + \*-------------------------------------*/ + +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 + + //Device::print(dg, db); + Device::checkDimError(dg,db); + + rippling0<<>>(ptrDevPixels,w,h,t); + Device::checkKernelError("rippling0"); + } + +/*--------------------------------------*\ + |* Private *| + \*-------------------------------------*/ + +__global__ void rippling0(uchar4* ptrDevPixels, int w, int h, float t) + { + Rippling0Math rippling0Math = Rippling0Math(w, h); + + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + + const int WH=w*h; + + uchar4 color; + color.w = 255; // alpha + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) + + rippling0Math.color(pixelI, pixelJ, t, &color); // update color + ptrDevPixels[s] = color; + + s += NB_THREAD; + } + + } + + + +/*----------------------------------------------------------------------*\ + |* End *| + \*---------------------------------------------------------------------*/ +