Ajout du support du multi-GPU pour Mandelbrot.

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu
index a98ed36..1b326db 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu
@@ -12,11 +12,11 @@
 using std::cout;\r
 using std::endl;\r
 \r
-__device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath)\r
+__device__ void fractal(uchar4* ptrDevPixels, int w, int hFrom, int hTo, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath)\r
     {\r
     const int TID = Indice2D::tid();\r
     const int NB_THREAD = Indice2D::nbThread();\r
-    const int WH = w * h;\r
+    const int WH = w * (hTo - hFrom);\r
 \r
     uchar4 color;\r
     color.z = 255; // Par défaut, l'image est opaque.\r
@@ -31,7 +31,7 @@ __device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& d
 \r
         // (i,j) domaine écran\r
         // (x,y) domaine math\r
-        domaineMath.toXY(pixelI, pixelJ, &x, &y); //  (i,j) -> (x,y)\r
+        domaineMath.toXY(pixelI + hFrom, pixelJ, &x, &y); //  (i,j) -> (x,y)\r
 \r
         fractalMath.colorXY(&color, x, y);\r
 \r
@@ -41,15 +41,15 @@ __device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& d
         }\r
     }\r
 \r
-__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n)\r
+__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int hFrom, int hTo, DomaineMath domaineMath, int n)\r
     {\r
     FractalMandelbrotMath fractalMath(n);\r
-    fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath);\r
+    fractal(ptrDevPixels, w, hFrom, hTo, domaineMath, n, fractalMath);\r
     }\r
 \r
 __global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i)\r
     {\r
     FractalJuliaMath fractalMath(n, c_r, c_i);\r
-    fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath);\r
+    fractal(ptrDevPixels, w, 0, h, domaineMath, n, fractalMath);\r
     }\r
 \r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.h
index 137aab2..44bf422 100644 (file)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.h
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.h
@@ -3,7 +3,7 @@
 
 #include "DomaineMath.h"
 
-__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n);
+__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int hFrom, int hTo, DomaineMath domaineMath, int n);
 __global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i);
 
 #endif

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 bfaff4b..67c8b74 100755 (executable)
--- 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
@@ -1,17 +1,19 @@
 #include "Fractal.h"\r
 \r
 #include <iostream>\r
+#include <sstream>\r
 #include <assert.h>\r
 using namespace std;\r
 \r
+#include <omp.h>\r
+\r
 #include "FractalDevice.h"\r
 #include "Device.h"\r
 \r
 Fractal::Fractal(int w, int h) :\r
         w(w), h(h),\r
         dg(8, 8, 1),\r
-        db(16, 16, 1),\r
-        title("Fractal Cuda")\r
+        db(16, 16, 1)\r
     {\r
     //print(dg, db);\r
     Device::assertDim(dg, db);\r
@@ -43,22 +45,38 @@ DomaineMath* Fractal::getDomaineMathInit()
     return this->ptrDomaineMathInit;\r
     }\r
 \r
-/**\r
- * Override\r
- */\r
-string Fractal::getTitle()\r
-    {\r
-    return this->title;\r
-    }\r
-\r
 /////\r
 \r
-FractalMandelbrot::FractalMandelbrot(int w, int h, int dn) :\r
+FractalMandelbrot::FractalMandelbrot(int w, int h, int dn, bool multiGPU) :\r
         Fractal(w, h),\r
         variateurAnimationN(IntervalI(10, 100), dn),\r
-        n(0)\r
+        n(0),\r
+        multiGPU(multiGPU)\r
     {\r
+    // Constuit le titre dynamiquement.\r
+    ostringstream titleStream;\r
+    titleStream << "Fractal Mandelbrot (multi-GPU " << (this->multiGPU ? "activated" : "not activated") << ")";\r
+    this->title = titleStream.str();\r
+\r
     this->ptrDomaineMathInit = new DomaineMath(-2, -1.3, 0.8, 1.3);\r
+\r
+    if (this->multiGPU)\r
+        {\r
+        const int nbDevice = Device::getDeviceCount();\r
+\r
+        this->hDevices = h / nbDevice;\r
+        this->hFirstDevice = h - ((nbDevice - 1) * this->hDevices);\r
+\r
+        // Allocation de la mémoire sur chaque GPU (sauf le premier pour lequel 'ptrDevPixels' est automatiquement alloué à l'appel de 'runGPU(..)').\r
+        this->ptrDevPixelsMultGPU = new uchar4*[nbDevice - 1];\r
+        for (int i = 0; i < nbDevice - 1; ++i)\r
+            {\r
+            HANDLE_ERROR(cudaSetDevice(i + 1));\r
+            HANDLE_ERROR(cudaMalloc(&this->ptrDevPixelsMultGPU[i], sizeof(uchar4) * w * this->hDevices));\r
+            }\r
+\r
+        HANDLE_ERROR(cudaSetDevice(0));\r
+        }\r
     }\r
 \r
 void FractalMandelbrot::animationStep()\r
@@ -78,9 +96,36 @@ void FractalMandelbrot::getValues(float* values)
     values[0] = float(this->n);\r
     }\r
 \r
+string FractalMandelbrot::getTitle()\r
+    {\r
+    return this->title;\r
+    }\r
+\r
 void FractalMandelbrot::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath)\r
     {\r
-    fractalMandelbrot<<<dg,db>>>(ptrDevPixels, this->w, this->h, domaineMath, static_cast<int>(this->n));\r
+\r
+    if (this->multiGPU)\r
+        {\r
+        HANDLE_ERROR(cudaSetDevice(0));\r
+        fractalMandelbrot<<<dg,db>>>(ptrDevPixels, this->w, 0, this->hFirstDevice, domaineMath, this->n);\r
+\r
+        const int nbDevice = Device::getDeviceCount();\r
+\r
+        // Rend chaque tranche par un GPU différent puis copie chaque tranche dans la mémoire du premier GPU.\r
+        #pragma omp parallel for\r
+        for (int i = 0; i < nbDevice - 1; ++i)\r
+            {\r
+            HANDLE_ERROR(cudaSetDevice(i + 1));\r
+            fractalMandelbrot<<<dg,db>>>(this->ptrDevPixelsMultGPU[i], this->w, i * this->hDevices + this->hFirstDevice, (i + 1) * this->hDevices + this->hFirstDevice, domaineMath, this->n);\r
+            HANDLE_ERROR(cudaMemcpy(ptrDevPixels + this->w * this->hFirstDevice + i * this->w * this->hDevices, this->ptrDevPixelsMultGPU[i], sizeof(uchar4) * this->w * this->hDevices, cudaMemcpyDeviceToDevice));\r
+            }\r
+\r
+        HANDLE_ERROR(cudaSetDevice(0));\r
+        }\r
+    else\r
+        {\r
+        fractalMandelbrot<<<dg,db>>>(ptrDevPixels, this->w, 0, this->h, domaineMath, this->n);\r
+        }\r
     }\r
 \r
 /////\r
@@ -116,6 +161,11 @@ void FractalJulia::getValues(float* values)
     values[1] = this->z_i;\r
     }\r
 \r
+string FractalJulia::getTitle()\r
+    {\r
+    return "Fractal Julia";\r
+    }\r
+\r
 void FractalJulia::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath)\r
     {\r
     fractalJulia<<<dg,db>>>(ptrDevPixels, this->w, this->h, domaineMath, this->n, this->z_r, this->z_i);\r

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 5ad7b83..41813e2 100755 (executable)
--- 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
@@ -21,8 +21,6 @@ class Fractal : public AnimableFonctionel_I
         int getH() /*override*/;\r
         DomaineMath* getDomaineMathInit() /*override*/;\r
 \r
-        string getTitle(void) /*override*/;\r
-\r
     protected:\r
         // Inputs\r
         const int w;\r
@@ -33,25 +31,32 @@ class Fractal : public AnimableFonctionel_I
         const dim3 db;\r
 \r
         DomaineMath* ptrDomaineMathInit;\r
-\r
-        // Outputs\r
-        const string title;\r
     };\r
 \r
 class FractalMandelbrot : public Fractal\r
     {\r
     public:\r
-        FractalMandelbrot(int w, int h, int dn);\r
+        FractalMandelbrot(int w, int h, int dn, bool multiGPU = false);\r
         void animationStep();\r
 \r
         std::vector<std::string> getNames();\r
         void getValues(float* values);\r
 \r
+        std::string getTitle();\r
+\r
     private:\r
-        void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/;\r
+        void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath);\r
 \r
         VariateurI variateurAnimationN;\r
         int n;\r
+\r
+        string title;\r
+\r
+        // Utilisé uniquement dans le cadre du multi-GPU.\r
+        bool multiGPU;\r
+        uchar4** ptrDevPixelsMultGPU; // La mémoire alloué pour les GPU autres que le premier\r
+        int hFirstDevice; // Hauteur de l'image à traiter par le premier GPU.\r
+        int hDevices; // Hauteur de l'image à traiter par les autres GPU.\r
     };\r
 \r
 class FractalJulia : public Fractal\r
@@ -63,8 +68,10 @@ class FractalJulia : public Fractal
         std::vector<std::string> getNames();\r
         void getValues(float* values);\r
 \r
+        std::string getTitle();\r
+\r
     private:\r
-        void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/;\r
+        void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath);\r
 \r
         const int n;\r
 \r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp
index ba94d73..fc82c54 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp
@@ -1,11 +1,11 @@
 #include "FractalProvider.h"\r
 \r
-Fractal* MandelbrotProvider::create()\r
+Fractal* MandelbrotProvider::create(bool multiGPU)\r
     {\r
     int dw = 16 * 50;\r
     int dh = 16 * 30;\r
 \r
-    return new FractalMandelbrot(dw, dh, 1);\r
+    return new FractalMandelbrot(dw, dh, 1, multiGPU);\r
     }\r
 \r
 Fractal* JuliaProvider::create()\r
@@ -16,10 +16,10 @@ Fractal* JuliaProvider::create()
     return new FractalJulia(dw, dh, 300, -0.745, -0.32, -0.09, 0.1);\r
     }\r
 \r
-ImageFonctionel* MandelbrotProvider::createGL()\r
+ImageFonctionel* MandelbrotProvider::createGL(bool multiGPU)\r
     {\r
     ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 100);\r
-    return new ImageFonctionel(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel\r
+    return new ImageFonctionel(create(multiGPU), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel\r
     }\r
 \r
 ImageFonctionel* JuliaProvider::createGL()\r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h
index eca7b3c..f29e4cb 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h
@@ -7,8 +7,8 @@
 class MandelbrotProvider\r
     {\r
     public:\r
-       static Fractal* create();\r
-       static ImageFonctionel* createGL();\r
+       static Fractal* create(bool multiGPU);\r
+       static ImageFonctionel* createGL(bool multiGPU);\r
     };\r
 \r
 \r

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
index fe29a86..b15ca7b 100755 (executable)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp
@@ -19,19 +19,25 @@ template <class TOutput, class TProvider>
 class Viewer\r
     {\r
     private:\r
-        TOutput* ptrProvider;\r
+        TOutput* ptrOutput;\r
         GLUTImageViewers viewer;\r
 \r
     public:\r
         Viewer(bool isAnimation, bool isSelection, int pxFrame, int pyFrame):\r
-            ptrProvider(TProvider::createGL()),\r
-            viewer(ptrProvider, isAnimation, isSelection, pxFrame, 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->ptrProvider;\r
+            delete this->ptrOutput;\r
             }\r
     };\r
 \r
@@ -39,13 +45,13 @@ int mainGL(void)
     {\r
     // Viewer<Rippling0Image, Rippling0Provider> rippling0(true, true, 10, 10);\r
     // Viewer<Image, RipplingProvider> rippling0(true, true, 10, 10);\r
-    Viewer<ImageFonctionel, MandelbrotProvider> fractalMandelbrot(true, true, 20, 20);\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
+    GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte.\r
 \r
     return EXIT_SUCCESS;\r
     }\r