From: gburri Date: Thu, 9 Oct 2014 08:31:32 +0000 (+0200) Subject: * Ajout d'un exemple CUDA non-openGL (AddVector.cu) X-Git-Url: http://git.euphorik.ch/?p=GPU.git;a=commitdiff_plain;h=3e601cb6c0cc2c5b3a9b30ebf3ad1102e53c0e0b * 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 --- 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 /** * getParaAnimation */ - virtual float getT()=0; - - virtual string getTitle()=0; + virtual float getT() = 0; + virtual string getTitle() = 0; }; #endif 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 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 --- a/WCudaMSE/Student_Cuda/.project +++ b/WCudaMSE/Student_Cuda/.project @@ -7,7 +7,7 @@ org.eclipse.cdt.managedbuilder.core.genmakebuilder - full,incremental, + clean,full,incremental, ?name? 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 --- 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) // Specifier nb thread : ici 1 thread au total ! - dim3 dg = dim3(1,1,1); + dim3 dg = dim3(1, 1, 1); dim3 db = dim3(1, 1, 1); // Debug 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 index 0000000..766480e --- /dev/null +++ b/WCudaMSE/Student_Cuda/src/cpp/core/02_AddVector/AddVector.cu @@ -0,0 +1,65 @@ +#include + +#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<<>>(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 --- 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 --- 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 --- 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) { 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 double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0) / (dxy10 + 1); const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256); ptrColor->x = grayLevel; 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 --- 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 static Rippling0Image* createGL(void) { - float dt = 2 * PI / 1000; // animation para + float dt = 1; int dw = 16 * 60; // =32*30=960 int dh = 16 * 60; // =32*30=960 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 --- 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) { 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 double grayLevelFloat = 128.0 + 127.0 * cos(dxy10 - t / 7.0) / (dxy10 + 1); const uchar grayLevel = (uchar)(long(grayLevelFloat) % 256); color.x = grayLevel; 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 --- 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) const int TID = Indice2D::tid(); const int NB_THREAD = Indice2D::nbThread(); - const int WH = w * h; uchar4 color; - color.z = 255; + color.z = 255; // Par défaut, l'image est opaque. int pixelI; int pixelJ; @@ -58,11 +57,9 @@ __global__ void rippling(uchar4* ptrDevPixels, int w, int h, float t) int s = TID; while (s < WH) { - IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) - - ripplingMath.color(pixelI, pixelJ, t, color); // update color + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); + ripplingMath.color(pixelI, pixelJ, t, 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 eb49b08..43068a2 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,17 +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) + : w(w), h(h), dt(dt), t(0), + dg(8, 8, 1), + db(16, 16, 1), + title("Rippling Cuda") { assert(w == h); - // Tools - this->dg = dim3(8, 8, 1); // disons a optimiser - this->db = dim3(16, 16, 1); // disons a optimiser - - // Outputs - this->title = "Rippling Cuda"; - //print(dg, db); Device::assertDim(dg, db); } 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 --- 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 @@ Rippling* RipplingProvider::createMOO() { - float dt = 1; - - int dw = 16 * 60; // =32*30=960 - int dh = 16 * 60; // =32*30=960 + const float dt = 1; + const int dw = 16 * 60; // =32*30=960 + const int dh = 16 * 60; // =32*30=960 return new Rippling(dw, dh, dt); } 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 --- 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); 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; + const bool isAnimation = true; + const 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 ... GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte // destruction - { - //delete ptrRippling0; - delete ptrRippling; - } + delete ptrRippling0; + delete ptrRippling; return EXIT_SUCCESS; }