From fd0031be0a39a5d902750affaff6322fcd5229b1 Mon Sep 17 00:00:00 2001 From: gburri Date: Sat, 10 Jan 2015 18:20:27 +0100 Subject: [PATCH] Ajout du support du multi-GPU pour Mandelbrot. --- .../moo/device/FractalDevice.cu | 12 +-- .../moo/device/FractalDevice.h | 2 +- .../02_Mandelbrot_Julia/moo/host/Fractal.cu | 76 +++++++++++++++---- .../02_Mandelbrot_Julia/moo/host/Fractal.h | 23 ++++-- .../provider/FractalProvider.cpp | 8 +- .../provider/FractalProvider.h | 4 +- .../src/cpp/core/mainGL.cpp | 18 +++-- 7 files changed, 103 insertions(+), 40 deletions(-) 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 --- 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; using std::endl; -__device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath) +__device__ void fractal(uchar4* ptrDevPixels, int w, int hFrom, int hTo, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath) { const int TID = Indice2D::tid(); const int NB_THREAD = Indice2D::nbThread(); - const int WH = w * h; + const int WH = w * (hTo - hFrom); uchar4 color; color.z = 255; // Par défaut, l'image est opaque. @@ -31,7 +31,7 @@ __device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& d // (i,j) domaine écran // (x,y) domaine math - domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y) + domaineMath.toXY(pixelI + hFrom, pixelJ, &x, &y); // (i,j) -> (x,y) fractalMath.colorXY(&color, x, y); @@ -41,15 +41,15 @@ __device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& d } } -__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) { FractalMandelbrotMath fractalMath(n); - fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath); + fractal(ptrDevPixels, w, hFrom, hTo, domaineMath, n, fractalMath); } __global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i) { FractalJuliaMath fractalMath(n, c_r, c_i); - fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath); + fractal(ptrDevPixels, w, 0, h, domaineMath, n, fractalMath); } 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 --- 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 --- 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" #include +#include #include using namespace std; +#include + #include "FractalDevice.h" #include "Device.h" Fractal::Fractal(int w, int h) : w(w), h(h), dg(8, 8, 1), - db(16, 16, 1), - title("Fractal Cuda") + db(16, 16, 1) { //print(dg, db); Device::assertDim(dg, db); @@ -43,22 +45,38 @@ DomaineMath* Fractal::getDomaineMathInit() return this->ptrDomaineMathInit; } -/** - * Override - */ -string Fractal::getTitle() - { - return this->title; - } - ///// -FractalMandelbrot::FractalMandelbrot(int w, int h, int dn) : +FractalMandelbrot::FractalMandelbrot(int w, int h, int dn, bool multiGPU) : Fractal(w, h), variateurAnimationN(IntervalI(10, 100), dn), - n(0) + n(0), + multiGPU(multiGPU) { + // Constuit le titre dynamiquement. + ostringstream titleStream; + titleStream << "Fractal Mandelbrot (multi-GPU " << (this->multiGPU ? "activated" : "not activated") << ")"; + this->title = titleStream.str(); + this->ptrDomaineMathInit = new DomaineMath(-2, -1.3, 0.8, 1.3); + + if (this->multiGPU) + { + const int nbDevice = Device::getDeviceCount(); + + this->hDevices = h / nbDevice; + this->hFirstDevice = h - ((nbDevice - 1) * this->hDevices); + + // Allocation de la mémoire sur chaque GPU (sauf le premier pour lequel 'ptrDevPixels' est automatiquement alloué à l'appel de 'runGPU(..)'). + this->ptrDevPixelsMultGPU = new uchar4*[nbDevice - 1]; + for (int i = 0; i < nbDevice - 1; ++i) + { + HANDLE_ERROR(cudaSetDevice(i + 1)); + HANDLE_ERROR(cudaMalloc(&this->ptrDevPixelsMultGPU[i], sizeof(uchar4) * w * this->hDevices)); + } + + HANDLE_ERROR(cudaSetDevice(0)); + } } void FractalMandelbrot::animationStep() @@ -78,9 +96,36 @@ void FractalMandelbrot::getValues(float* values) values[0] = float(this->n); } +string FractalMandelbrot::getTitle() + { + return this->title; + } + void FractalMandelbrot::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) { - fractalMandelbrot<<>>(ptrDevPixels, this->w, this->h, domaineMath, static_cast(this->n)); + + if (this->multiGPU) + { + HANDLE_ERROR(cudaSetDevice(0)); + fractalMandelbrot<<>>(ptrDevPixels, this->w, 0, this->hFirstDevice, domaineMath, this->n); + + const int nbDevice = Device::getDeviceCount(); + + // Rend chaque tranche par un GPU différent puis copie chaque tranche dans la mémoire du premier GPU. + #pragma omp parallel for + for (int i = 0; i < nbDevice - 1; ++i) + { + HANDLE_ERROR(cudaSetDevice(i + 1)); + fractalMandelbrot<<>>(this->ptrDevPixelsMultGPU[i], this->w, i * this->hDevices + this->hFirstDevice, (i + 1) * this->hDevices + this->hFirstDevice, domaineMath, this->n); + HANDLE_ERROR(cudaMemcpy(ptrDevPixels + this->w * this->hFirstDevice + i * this->w * this->hDevices, this->ptrDevPixelsMultGPU[i], sizeof(uchar4) * this->w * this->hDevices, cudaMemcpyDeviceToDevice)); + } + + HANDLE_ERROR(cudaSetDevice(0)); + } + else + { + fractalMandelbrot<<>>(ptrDevPixels, this->w, 0, this->h, domaineMath, this->n); + } } ///// @@ -116,6 +161,11 @@ void FractalJulia::getValues(float* values) values[1] = this->z_i; } +string FractalJulia::getTitle() + { + return "Fractal Julia"; + } + void FractalJulia::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) { fractalJulia<<>>(ptrDevPixels, this->w, this->h, domaineMath, this->n, this->z_r, this->z_i); 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 --- 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*/; DomaineMath* getDomaineMathInit() /*override*/; - string getTitle(void) /*override*/; - protected: // Inputs const int w; @@ -33,25 +31,32 @@ class Fractal : public AnimableFonctionel_I const dim3 db; DomaineMath* ptrDomaineMathInit; - - // Outputs - const string title; }; class FractalMandelbrot : public Fractal { public: - FractalMandelbrot(int w, int h, int dn); + FractalMandelbrot(int w, int h, int dn, bool multiGPU = false); void animationStep(); std::vector getNames(); void getValues(float* values); + std::string getTitle(); + private: - void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/; + void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath); VariateurI variateurAnimationN; int n; + + string title; + + // Utilisé uniquement dans le cadre du multi-GPU. + bool multiGPU; + uchar4** ptrDevPixelsMultGPU; // La mémoire alloué pour les GPU autres que le premier + int hFirstDevice; // Hauteur de l'image à traiter par le premier GPU. + int hDevices; // Hauteur de l'image à traiter par les autres GPU. }; class FractalJulia : public Fractal @@ -63,8 +68,10 @@ class FractalJulia : public Fractal std::vector getNames(); void getValues(float* values); + std::string getTitle(); + private: - void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/; + void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath); const int n; 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 --- 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" -Fractal* MandelbrotProvider::create() +Fractal* MandelbrotProvider::create(bool multiGPU) { int dw = 16 * 50; int dh = 16 * 30; - return new FractalMandelbrot(dw, dh, 1); + return new FractalMandelbrot(dw, dh, 1, multiGPU); } Fractal* JuliaProvider::create() @@ -16,10 +16,10 @@ Fractal* JuliaProvider::create() return new FractalJulia(dw, dh, 300, -0.745, -0.32, -0.09, 0.1); } -ImageFonctionel* MandelbrotProvider::createGL() +ImageFonctionel* MandelbrotProvider::createGL(bool multiGPU) { ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 100); - return new ImageFonctionel(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + return new ImageFonctionel(create(multiGPU), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel } ImageFonctionel* JuliaProvider::createGL() 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 --- 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 { public: - static Fractal* create(); - static ImageFonctionel* createGL(); + static Fractal* create(bool multiGPU); + static ImageFonctionel* createGL(bool multiGPU); }; 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 --- 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 Viewer { private: - TOutput* ptrProvider; + TOutput* ptrOutput; GLUTImageViewers viewer; public: Viewer(bool isAnimation, bool isSelection, int pxFrame, int pyFrame): - ptrProvider(TProvider::createGL()), - viewer(ptrProvider, isAnimation, isSelection, pxFrame, 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->ptrProvider; + delete this->ptrOutput; } }; @@ -39,13 +45,13 @@ int mainGL(void) { // Viewer rippling0(true, true, 10, 10); // Viewer rippling0(true, true, 10, 10); - Viewer fractalMandelbrot(true, true, 20, 20); + 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 + GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte. return EXIT_SUCCESS; } -- 2.45.2