From 4182eb3a07b7143afb8ebebfe77e8ef8e8abc266 Mon Sep 17 00:00:00 2001 From: gburri Date: Sat, 10 Jan 2015 22:49:13 +0100 Subject: [PATCH] =?utf8?q?D=C3=A9but=20du=20TP=20convolution.=20Pour=20l'i?= =?utf8?q?nstant=20uniquement=20lecture=20d'une=20vid=C3=A9o.?= MIME-Version: 1.0 Content-Type: text/plain; charset=utf8 Content-Transfer-Encoding: 8bit --- .../core/cudaImageTools/bitmap/cpp/Image.cpp | 2 +- .../core/cudaImageTools/bitmap/header/Image.h | 9 ++- .../02_Mandelbrot_Julia/moo/host/Fractal.cu | 19 ++++++ .../02_Mandelbrot_Julia/moo/host/Fractal.h | 1 + .../moo/device/RayTracingDevice.h | 2 - .../06_Convolution/ImageConvolutionCuda.cpp | 22 ++++++ .../06_Convolution/ImageConvolutionCuda.h | 24 +++++++ .../moo/device/ConvolutionDevice.cu | 40 +++++++++++ .../moo/device/ConvolutionDevice.h | 7 ++ .../06_Convolution/moo/host/Convolution.cu | 67 +++++++++++++++++++ .../06_Convolution/moo/host/Convolution.h | 39 +++++++++++ .../provider/ConvolutionProvider.cpp | 19 ++++++ .../provider/ConvolutionProvider.h | 14 ++++ .../Student_Cuda_Image/src/cpp/core/Viewer.h | 49 ++++++++++++++ .../src/cpp/core/mainGL.cpp | 46 ++++--------- 15 files changed, 318 insertions(+), 42 deletions(-) create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.cpp create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/Viewer.h diff --git a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/cpp/Image.cpp b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/cpp/Image.cpp index fa0eedf..11efd83 100755 --- a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/cpp/Image.cpp +++ b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/cpp/Image.cpp @@ -37,7 +37,7 @@ Image::Image(Animable_I* ptrAnimable,ColorRGB_01* ptrColorTitreRGB) : ImageMOOs_A(ptrAnimable->getW(), ptrAnimable->getH()) { this->ptrAnimable = ptrAnimable; - this->ptrColorTitreRGB=ptrColorTitreRGB; + this->ptrColorTitreRGB = ptrColorTitreRGB; this->valueNames = ptrAnimable->getNames(); this->values = new float[this->valueNames.size()]; diff --git a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Image.h b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Image.h index d68ec99..b599f8c 100755 --- a/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Image.h +++ b/WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Image.h @@ -41,16 +41,15 @@ class Image: public ImageMOOs_A public: - void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override - void animationStep(bool& isNeedUpdateView); // override - void paintPrimitives(Graphic2Ds& graphic2D); // override + virtual void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override + virtual void animationStep(bool& isNeedUpdateView); // override + virtual void paintPrimitives(Graphic2Ds& graphic2D); // override /*--------------------------------------*\ |* Attributs *| \*-------------------------------------*/ - private: - + protected: // Input Animable_I* ptrAnimable; ColorRGB_01* ptrColorTitreRGB; diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu index 67c8b74..9228000 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu @@ -79,6 +79,25 @@ FractalMandelbrot::FractalMandelbrot(int w, int h, int dn, bool multiGPU) : } } +FractalMandelbrot::~FractalMandelbrot() + { + if (this->multiGPU) + { + const int nbDevice = Device::getDeviceCount(); + + for (int i = 0; i < nbDevice - 1; ++i) + { + HANDLE_ERROR(cudaSetDevice(i + 1)); + HANDLE_ERROR(cudaFree(this->ptrDevPixelsMultGPU[i])); + } + + HANDLE_ERROR(cudaSetDevice(0)); + delete[] this->ptrDevPixelsMultGPU; + } + + cout << "OKOKOKOK" << endl; + } + void FractalMandelbrot::animationStep() { this->n = this->variateurAnimationN.varierAndGet(); diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h index 41813e2..e9299a0 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h @@ -37,6 +37,7 @@ class FractalMandelbrot : public Fractal { public: FractalMandelbrot(int w, int h, int dn, bool multiGPU = false); + ~FractalMandelbrot(); void animationStep(); std::vector getNames(); diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h index 5e4487d..1500c43 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h @@ -1,8 +1,6 @@ #ifndef RAY_TRACING_DEVICE_H #define RAY_TRACING_DEVICE_H -#include "DomaineMath.h" - __global__ void rayTracing(uchar4* ptrDevPixels, int w, int h, float t); diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp new file mode 100644 index 0000000..ffc30a1 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp @@ -0,0 +1,22 @@ +#include "ImageConvolutionCuda.h" + +ImageConvolutionCuda::ImageConvolutionCuda(Animable_I* ptrAnimable, CaptureVideo* captureur, ColorRGB_01* ptrColorTitreRGB) : + Image(ptrAnimable, ptrColorTitreRGB), + captureur(captureur) + { + + //Mat matImage = captureur.capturer(); // capture une image seulement ( à utiliser en boucle!) + //uchar4* image = CaptureVideo::castToUChar4(&matImage); // format cuda + } + +ImageConvolutionCuda::~ImageConvolutionCuda() + { + delete this->captureur; + } + +void ImageConvolutionCuda::fillImageGL(uchar4* ptrDevImageGL, int w, int h) + { + Mat matImage = this->captureur->capturer(); // capture une image seulement ( à utiliser en boucle!) + uchar4* image = CaptureVideo::castToUChar4(&matImage); // format cuda + HANDLE_ERROR(cudaMemcpy(ptrDevImageGL, image, sizeof(uchar4) * this->ptrAnimable->getW() * this->ptrAnimable->getH(), cudaMemcpyHostToDevice)); + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h new file mode 100644 index 0000000..fb35821 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h @@ -0,0 +1,24 @@ + +#ifndef IMAGE_CONVOLUTION_CUDA_H +#define IMAGE_CONVOLUTION_CUDA_H + +#include "Image.h" +#include "CaptureVideo.h" + +class ImageConvolutionCuda : public Image + { + public: + ImageConvolutionCuda(Animable_I* ptrAnimable, CaptureVideo* captureur ,ColorRGB_01* ptrColorTitreRGB = new ColorRGB_01(1, 0, 0)); + + virtual ~ImageConvolutionCuda(); + + public: + void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override + //void animationStep(bool& isNeedUpdateView); // override + //void paintPrimitives(Graphic2Ds& graphic2D); // override + + private: + CaptureVideo* captureur; + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu new file mode 100644 index 0000000..640d8a9 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu @@ -0,0 +1,40 @@ +#include +#include +using namespace std; + +#include "Indice2D.h" +#include "IndiceTools.h" +#include "cudaTools.h" +#include "Device.h" + +#include "ConvolutionDevice.h" + +__global__ +void convolution(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.w = 255; // Par défaut, l'image est opaque. + + double x, y; + int pixelI, pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) + + // (i,j) domaine écran. + // (x,y) domaine math. + // domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y). + + // newtonMath.colorXY(&color, x, y); + + ptrDevPixels[s] = color; + + s += NB_THREAD; + } + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h new file mode 100644 index 0000000..7044644 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h @@ -0,0 +1,7 @@ +#ifndef CONVOLUTION_DEVICE_H +#define CONVOLUTION_DEVICE_H + +__global__ +void convolution(uchar4* ptrDevPixels, int w, int h, float t); + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu new file mode 100644 index 0000000..ecc4464 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu @@ -0,0 +1,67 @@ +#include +#include +#include +using namespace std; + +#include "Device.h" + +#include "Convolution.h" +#include "ConvolutionDevice.h" + + +const float Convolution::kernel[9][9] = + {{0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828}, + {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987}, + {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705}, + {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366}, + {0.6063, 1.2765, 0.4061, -6.4829, -12.7462, -6.4829, 0.4061, 1.2765, 0.6063}, + {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366}, + {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705}, + {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987}, + {0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828}}; + +Convolution::Convolution(int w, int h) : + w(w), h(h), + dg(8, 8, 1), + db(32, 32, 1), + title("Convolution") + { + Device::assertDim(dg, db); + } + +Convolution::~Convolution() + { + } + +void Convolution::runGPU(uchar4* ptrDevPixels) + { + convolution<<>>(ptrDevPixels, this->w, this->h, this->t); + + // HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG). + } + +void Convolution::animationStep() + { + this->t += 0.1; // TODO. + } + +int Convolution::getW() + { + return this->w; + } + +int Convolution::getH() + { + return this->h; + } + +float Convolution::getT() + { + return this->t; + } + +string Convolution::getTitle() + { + return this->title; + } + diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h new file mode 100644 index 0000000..e67adc2 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h @@ -0,0 +1,39 @@ +#ifndef CONVOLUTION_H +#define CONVOLUTION_H + +#include "cudaTools.h" +#include "Animable_I.h" +#include "MathTools.h" +#include "AleaTools.h" + +class Convolution : public Animable_I + { + public: + Convolution(int w, int h); + ~Convolution(); + + void runGPU(uchar4* ptrDevPixels) /*override*/; + void animationStep() /*override*/; + + int getW() /*override*/; + int getH() /*override*/; + + float getT() /*override*/; + + std::string getTitle(void) /*override*/; + + private: + float t; + + const int w; + const int h; + + const dim3 dg; + const dim3 db; + + const std::string title; + + static const float kernel[9][9]; + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.cpp new file mode 100644 index 0000000..3ebd4d7 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.cpp @@ -0,0 +1,19 @@ +#include "ConvolutionProvider.h" + +#include +using namespace std; + +#include "CaptureVideo.h" + +Convolution* ConvolutionProvider::create(int w, int h) + { + return new Convolution(w, h); + } + +Image* ConvolutionProvider::createGL(const string& videoPath) + { + CaptureVideo* captureur = new CaptureVideo(videoPath, ""); + + ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0); + return new ImageConvolutionCuda(create(captureur->getW(), captureur->getH()), captureur, ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.h new file mode 100644 index 0000000..5ac81cc --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.h @@ -0,0 +1,14 @@ +#ifndef CONVOLUTION_PROVIDER_H +#define CONVOLUTION_PROVIDER_H + +#include "Convolution.h" +#include "ImageConvolutionCuda.h" + +class ConvolutionProvider + { + public: + static Convolution* create(int w, int h); + static Image* createGL(const string& videoPath); + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/Viewer.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/Viewer.h new file mode 100644 index 0000000..d77be2a --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/Viewer.h @@ -0,0 +1,49 @@ +#ifndef VIEWER_H +#define VIEWER_H + +/** + * Crée une instance automatiquement de 'TOutput' en appelant 'TProvider::createGL()'. + */ +template +class AutoViewer + { + private: + TOutput* ptrOutput; + GLUTImageViewers viewer; + + public: + AutoViewer(bool isAnimation, bool isSelection, int pxFrame, int pyFrame): + ptrOutput(TProvider::createGL()), + viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame) + { + } + + ~AutoViewer() + { + delete this->ptrOutput; + } + }; + + +template +class Viewer + { + private: + TOutput* ptrOutput; + GLUTImageViewers viewer; + + public: + Viewer(TOutput* output, bool isAnimation, bool isSelection, int pxFrame, int pyFrame): + ptrOutput(output), + viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame) + { + } + + ~Viewer() + { + delete this->ptrOutput; + } + }; + +#endif + diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp index b15ca7b..10a9393 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp @@ -14,44 +14,22 @@ using namespace std; #include "NewtonProvider.h" #include "HeatTransfertProvider.h" #include "RayTracingProvider.h" +#include "ConvolutionProvider.h" -template -class Viewer - { - private: - TOutput* ptrOutput; - GLUTImageViewers viewer; - - public: - Viewer(bool isAnimation, bool isSelection, int pxFrame, int pyFrame): - ptrOutput(TProvider::createGL()), - viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame) - { - } - - Viewer(TOutput* output, bool isAnimation, bool isSelection, int pxFrame, int pyFrame): - ptrOutput(output), - viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame) - { - } - - ~Viewer() - { - delete this->ptrOutput; - } - }; +#include "Viewer.h" int mainGL(void) { - // Viewer rippling0(true, true, 10, 10); - // Viewer rippling0(true, true, 10, 10); - Viewer fractalMandelbrot(MandelbrotProvider::createGL(true), true, true, 20, 20); - // Viewer fractalJulia(true, true, 30, 30); - // Viewer newtown(true, true, 20, 20); - // Viewer heatTransfert(true, false, 20, 20); - // Viewer rayTracing(true, true, 20, 20); - - GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte. + // AutoViewer rippling0(true, true, 10, 10); + // AutoViewer rippling0(true, true, 10, 10); + // Viewer fractalMandelbrot(MandelbrotProvider::createGL(true), true, true, 20, 20); + // AutoViewer fractalJulia(true, true, 30, 30); + // AutoViewer newtown(true, true, 20, 20); + // AutoViewer heatTransfert(true, false, 20, 20); + // AutoViewer rayTracing(true, true, 20, 20); + Viewer convolution(ConvolutionProvider::createGL("/media/Data/Video/nasaFHD_short.avi"), true, true, 20, 20); + + GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenêtre est ouverte. return EXIT_SUCCESS; } -- 2.43.0