Début du TP convolution. Pour l'instant uniquement lecture d'une vidéo.
authorgburri <gregory.burri@master.hes-so.ch>
Sat, 10 Jan 2015 21:49:13 +0000 (22:49 +0100)
committergburri <gregory.burri@master.hes-so.ch>
Sat, 10 Jan 2015 21:49:13 +0000 (22:49 +0100)
15 files changed:
WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/cpp/Image.cpp
WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Image.h
WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu
WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h
WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.cpp [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/provider/ConvolutionProvider.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/Viewer.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp

index fa0eedf..11efd83 100755 (executable)
@@ -37,7 +37,7 @@ Image::Image(Animable_I* ptrAnimable,ColorRGB_01* ptrColorTitreRGB) :
        ImageMOOs_A(ptrAnimable->getW(), ptrAnimable->getH())\r
     {\r
     this->ptrAnimable = ptrAnimable;\r
-    this->ptrColorTitreRGB=ptrColorTitreRGB;\r
+    this->ptrColorTitreRGB = ptrColorTitreRGB;\r
 \r
     this->valueNames = ptrAnimable->getNames();\r
     this->values = new float[this->valueNames.size()];\r
index d68ec99..b599f8c 100755 (executable)
@@ -41,16 +41,15 @@ class Image: public ImageMOOs_A
 \r
     public:\r
 \r
-       void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override\r
-       void animationStep(bool& isNeedUpdateView); // override\r
-       void paintPrimitives(Graphic2Ds& graphic2D); // override\r
+       virtual void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override\r
+       virtual void animationStep(bool& isNeedUpdateView); // override\r
+       virtual void paintPrimitives(Graphic2Ds& graphic2D); // override\r
 \r
        /*--------------------------------------*\\r
        |*              Attributs               *|\r
        \*-------------------------------------*/\r
 \r
-    private:\r
-\r
+    protected:\r
        // Input\r
        Animable_I* ptrAnimable;\r
        ColorRGB_01* ptrColorTitreRGB;\r
index 67c8b74..9228000 100755 (executable)
@@ -79,6 +79,25 @@ FractalMandelbrot::FractalMandelbrot(int w, int h, int dn, bool multiGPU) :
         }\r
     }\r
 \r
+FractalMandelbrot::~FractalMandelbrot()\r
+    {\r
+    if (this->multiGPU)\r
+        {\r
+        const int nbDevice = Device::getDeviceCount();\r
+\r
+        for (int i = 0; i < nbDevice - 1; ++i)\r
+            {\r
+            HANDLE_ERROR(cudaSetDevice(i + 1));\r
+            HANDLE_ERROR(cudaFree(this->ptrDevPixelsMultGPU[i]));\r
+            }\r
+\r
+        HANDLE_ERROR(cudaSetDevice(0));\r
+        delete[] this->ptrDevPixelsMultGPU;\r
+        }\r
+\r
+    cout << "OKOKOKOK" << endl;\r
+    }\r
+\r
 void FractalMandelbrot::animationStep()\r
     {\r
     this->n = this->variateurAnimationN.varierAndGet();\r
index 41813e2..e9299a0 100755 (executable)
@@ -37,6 +37,7 @@ class FractalMandelbrot : public Fractal
     {\r
     public:\r
         FractalMandelbrot(int w, int h, int dn, bool multiGPU = false);\r
+        ~FractalMandelbrot();\r
         void animationStep();\r
 \r
         std::vector<std::string> getNames();\r
index 5e4487d..1500c43 100644 (file)
@@ -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 (file)
index 0000000..ffc30a1
--- /dev/null
@@ -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 (file)
index 0000000..fb35821
--- /dev/null
@@ -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 (file)
index 0000000..640d8a9
--- /dev/null
@@ -0,0 +1,40 @@
+#include <iostream>
+#include <stdio.h>
+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 (file)
index 0000000..7044644
--- /dev/null
@@ -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 (file)
index 0000000..ecc4464
--- /dev/null
@@ -0,0 +1,67 @@
+#include <iostream>
+#include <assert.h>
+#include <stdio.h>
+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<<<dg,db>>>(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 (file)
index 0000000..e67adc2
--- /dev/null
@@ -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 (file)
index 0000000..3ebd4d7
--- /dev/null
@@ -0,0 +1,19 @@
+#include "ConvolutionProvider.h"
+
+#include <iostream>
+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 (file)
index 0000000..5ac81cc
--- /dev/null
@@ -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 (file)
index 0000000..d77be2a
--- /dev/null
@@ -0,0 +1,49 @@
+#ifndef VIEWER_H
+#define VIEWER_H
+
+/**
+ * Crée une instance automatiquement de 'TOutput' en appelant 'TProvider::createGL()'.
+ */
+template <class TOutput, class TProvider>
+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 TOutput>
+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
+
index b15ca7b..10a9393 100755 (executable)
@@ -14,44 +14,22 @@ using namespace std;
 #include "NewtonProvider.h"\r
 #include "HeatTransfertProvider.h"\r
 #include "RayTracingProvider.h"\r
+#include "ConvolutionProvider.h"\r
 \r
-template <class TOutput, class TProvider>\r
-class Viewer\r
-    {\r
-    private:\r
-        TOutput* ptrOutput;\r
-        GLUTImageViewers viewer;\r
-\r
-    public:\r
-        Viewer(bool isAnimation, bool isSelection, int pxFrame, int pyFrame):\r
-        ptrOutput(TProvider::createGL()),\r
-            viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame)\r
-            {\r
-            }\r
-\r
-        Viewer(TOutput* output, bool isAnimation, bool isSelection, int pxFrame, int pyFrame):\r
-            ptrOutput(output),\r
-            viewer(ptrOutput, isAnimation, isSelection, pxFrame, pyFrame)\r
-            {\r
-            }\r
-\r
-        ~Viewer()\r
-            {\r
-            delete this->ptrOutput;\r
-            }\r
-    };\r
+#include "Viewer.h"\r
 \r
 int mainGL(void)\r
     {\r
-    // Viewer<Rippling0Image, Rippling0Provider> rippling0(true, true, 10, 10);\r
-    // Viewer<Image, RipplingProvider> rippling0(true, true, 10, 10);\r
-    Viewer<ImageFonctionel, MandelbrotProvider> fractalMandelbrot(MandelbrotProvider::createGL(true), true, true, 20, 20);\r
-    // Viewer<ImageFonctionel, JuliaProvider> fractalJulia(true, true, 30, 30);\r
-    // Viewer<ImageFonctionel, NewtonProvider> newtown(true, true, 20, 20);\r
-    // Viewer<Image, HeatTransfertProvider> heatTransfert(true, false, 20, 20);\r
-    // Viewer<ImageFonctionel, RayTracingProvider> rayTracing(true, true, 20, 20);\r
-\r
-    GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte.\r
+    // AutoViewer<Rippling0Image, Rippling0Provider> rippling0(true, true, 10, 10);\r
+    // AutoViewer<Image, RipplingProvider> rippling0(true, true, 10, 10);\r
+    // Viewer<ImageFonctionel> fractalMandelbrot(MandelbrotProvider::createGL(true), true, true, 20, 20);\r
+    // AutoViewer<ImageFonctionel, JuliaProvider> fractalJulia(true, true, 30, 30);\r
+    // AutoViewer<ImageFonctionel, NewtonProvider> newtown(true, true, 20, 20);\r
+    // AutoViewer<Image, HeatTransfertProvider> heatTransfert(true, false, 20, 20);\r
+    // AutoViewer<ImageFonctionel, RayTracingProvider> rayTracing(true, true, 20, 20);\r
+    Viewer<Image> convolution(ConvolutionProvider::createGL("/media/Data/Video/nasaFHD_short.avi"), true, true, 20, 20);\r
+\r
+    GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenêtre est ouverte.\r
 \r
     return EXIT_SUCCESS;\r
     }\r