From: gburri Date: Sun, 7 Dec 2014 15:29:08 +0000 (+0100) Subject: Ajout du TP HeatTransfert. X-Git-Url: http://git.euphorik.ch/index.cgi?a=commitdiff_plain;h=cb39d6a91b65d2862018430d65e633d2a8fdc818;p=GPU.git Ajout du TP HeatTransfert. --- diff --git a/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/IndiceTools_CPU.h b/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/IndiceTools_CPU.h index 89e2f3e..a97f64b 100755 --- a/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/IndiceTools_CPU.h +++ b/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/IndiceTools_CPU.h @@ -12,41 +12,40 @@ namespace cpu { class IndiceTools - { + { - /*--------------------------------------*\ + /*--------------------------------------*\ |* Constructor *| - \*--------------------------------------*/ + \*--------------------------------------*/ - public: + public: - /*--------------------------------------*\ + /*--------------------------------------*\ |* Methodes *| - \*--------------------------------------*/ + \*--------------------------------------*/ - public: + public: - /** - * s[0,W*H[ --> i[0,H[ j[0,W[ - */ - static void toIJ( int s, int W, int* ptri, int* ptrj); + /** + * s[0,W*H[ --> i[0,H[ j[0,W[ + */ + static void toIJ(int s, int W, int* ptri, int* ptrj); - /** - * i[0,H[ j[0,W[ --> s[0,W*H[ - */ - static int toS(int W, int i, int j); + /** + * i[0,H[ j[0,W[ --> s[0,W*H[ + */ + static int toS(int W, int i, int j); - - /*-------------------------------------*\ + /*-------------------------------------*\ |* Attributs *| - \*-------------------------------------*/ + \*-------------------------------------*/ - private: + private: - }; + }; } -#endif +#endif /*----------------------------------------------------------------------*\ |* End *| 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 eba8685..b421793 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 @@ -17,8 +17,7 @@ using std::string; class Animable_I { public: - - //virtual ~Animable_I(void)=0; + virtual ~Animable_I(void) {}; virtual void runGPU(uchar4* ptrDevPixels)=0; virtual void animationStep(void)=0; diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.cu new file mode 100755 index 0000000..03225a5 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.cu @@ -0,0 +1,126 @@ +#include "HeatTransfertDevice.h" + +#include +using namespace std; + +#include "Indice2D.h" +#include "IndiceTools.h" +#include "cudaTools.h" +#include "Device.h" + +#include "ColorTools.h" + +#include "HeatImage.h" + +__global__ +void display(HeatImage image, uchar4* ptrDevPixels, CalibreurF calibreur) + { + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + const int WH = image.getWCuda() * image.getHCuda(); + + uchar4 color; + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, image.getWCuda(), &pixelI, &pixelJ); + + float heatValue = image.getCuda(pixelJ, pixelI); + calibreur.calibrer(heatValue); + ColorTools::HSB_TO_RVB(heatValue, &color); + + ptrDevPixels[s] = color; + s += NB_THREAD; + } + } + +__global__ +void copyHeaters(HeatImage heaters, HeatImage destination) + { + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + const int WH = heaters.getWCuda() * heaters.getHCuda(); + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, heaters.getWCuda(), &pixelI, &pixelJ); + + float heatValue = heaters.getCuda(pixelJ, pixelI); + if (heatValue > 0.0) + destination.setCuda(pixelJ, pixelI, heatValue); + + s += NB_THREAD; + } + } + +__global__ +void diffuseMethode1(HeatImage from, HeatImage to) + { + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + + const int W = from.getWCuda() - 2; + const int H = from.getHCuda() - 2; + const int WH = W * H; + + const float k = 0.1; + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, W, &pixelI, &pixelJ); + const int x = pixelJ + 1; + const int y = pixelI + 1; + const float t_old = from.getCuda(x, y); + const float t_up = from.getCuda(x, y - 1); + const float t_down = from.getCuda(x, y + 1); + const float t_left = from.getCuda(x - 1, y); + const float t_right = from.getCuda(x + 1, y); + + to.setCuda(x, y, t_old + k * (t_up + t_down + t_left + t_right - 4.0 * t_old)); + + s += NB_THREAD; + } + } + +__global__ +void diffuseMethode2(HeatImage from, HeatImage to) + { + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + + const int W = from.getWCuda() - 2; + const int H = from.getHCuda() - 2; + const int WH = W * H; + + int pixelI; + int pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, W, &pixelI, &pixelJ); + const int x = pixelJ + 1; + const int y = pixelI + 1; + const float t_old = from.getCuda(x, y); + const float t_up = from.getCuda(x, y - 1); + const float t_down = from.getCuda(x, y + 1); + const float t_left = from.getCuda(x - 1, y); + const float t_right = from.getCuda(x + 1, y); + + to.setCuda(x, y, (t_up + t_down + t_left + t_right + 4.0 * t_old) / 8.0); + + s += NB_THREAD; + } + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.h new file mode 100644 index 0000000..d64c3cd --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.h @@ -0,0 +1,20 @@ +#ifndef HEAT_TRANSFERT_DEVICE_H +#define HEAT_TRANSFERT_DEVICE_H + +#include "CalibreurF.h" + +#include "HeatImage.h" + +__global__ +void display(HeatImage image, uchar4* ptrDevPixels, CalibreurF calibreur); + +__global__ +void copyHeaters(HeatImage heaters, HeatImage destination); + +__global__ +void diffuseMethode1(HeatImage from, HeatImage to); + +__global__ +void diffuseMethode2(HeatImage from, HeatImage to); + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.cu new file mode 100644 index 0000000..e383fb6 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.cu @@ -0,0 +1,12 @@ +#include "HeatImage.h" + +#include + +HeatImage::HeatImage(int w, int h, float* image) : + w(w), h(h), image(image) + { + } + +HeatImage::~HeatImage() + { + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.h new file mode 100644 index 0000000..00224c8 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.h @@ -0,0 +1,73 @@ +#ifndef HEAT_IMAGE_H +#define HEAT_IMAGE_H + +/* + * Une classe représentant une image de float. + */ +class HeatImage + { + public: + HeatImage(int w, int h, float* image = 0); + ~HeatImage(); + + inline void clear() + { + memset(this->image, 0, this->w * this->h * sizeof(float)); + } + + __device__ + inline void setCuda(int x, int y, float value) + { + this->image[x + this->w * y] = value; + } + + inline void set(int x, int y, float value) + { + this->image[x + this->w * y] = value; + } + + __device__ + inline float getCuda(int x, int y) + { + return this->image[x + this->w * y]; + } + + __device__ + inline int getWCuda() + { + return this->w; + } + + inline int getW() + { + return this->w; + } + + __device__ + inline int getHCuda() + { + return this->h; + } + + inline int getH() + { + return this->h; + } + + inline void swapWith(HeatImage& other) + { + float* tmp = this->image; + this->image = other.image; + other.image = tmp; + } + + inline float*& getRaw() { return this->image; } + inline void setRaw(float* image) { this->image = image; } + + private: + int w, h; + float* image; + }; + + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.cu new file mode 100755 index 0000000..7aa7cf6 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.cu @@ -0,0 +1,119 @@ +#include "HeatTransfert.h" + +#include +#include +#include +#include +using namespace std; + +#include "Device.h" + +#include "HeatTransfertDevice.h" + +HeatTransfert::HeatTransfert() : + calibreur(IntervalF(0, 1), IntervalF(0.7, 0)), + devImageHeaters(800, 800), + devImageA(800, 800), + devImageB(800, 800), + t(0), + dg(8, 8, 1), + db(32, 32, 1), + title("Heat transfert") + { + const size_t IMAGE_SIZE = this->devImageA.getW() * this->devImageA.getH() * sizeof(float); + const int IMAGE_W = this->devImageA.getW(); + const int IMAGE_H = this->devImageA.getH(); + + // Définition des heaters. + HeatImage imageHeaters(IMAGE_W, IMAGE_H, new float[IMAGE_W * IMAGE_H]); + imageHeaters.clear(); + setHeaters(imageHeaters); + + // Définition des températures initiales (toutes à zero). + HeatImage imageInit(IMAGE_W, IMAGE_H, new float[IMAGE_W * IMAGE_H]); + imageInit.clear(); + + // Copie des heaters en GM. + HANDLE_ERROR(cudaMalloc(&this->devImageHeaters.getRaw(), IMAGE_SIZE)); + HANDLE_ERROR(cudaMemcpy(this->devImageHeaters.getRaw(), imageHeaters.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice)); + + // Allocation de image A et image B. + HANDLE_ERROR(cudaMalloc(&this->devImageA.getRaw(), IMAGE_SIZE)); + HANDLE_ERROR(cudaMemcpy(this->devImageA.getRaw(), imageInit.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice)); + HANDLE_ERROR(cudaMalloc(&this->devImageB.getRaw(), IMAGE_SIZE)); + HANDLE_ERROR(cudaMemcpy(this->devImageB.getRaw(), imageInit.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice)); + } + +HeatTransfert::~HeatTransfert() + { + cudaFree(this->devImageHeaters.getRaw()); + cudaFree(this->devImageA.getRaw()); + cudaFree(this->devImageB.getRaw()); + } + +void HeatTransfert::runGPU(uchar4* ptrDevPixels) + { + for (int i = 0; i < NB_ITERATION_AVEUGLE; i++) + { + copyHeaters<<>>(this->devImageHeaters, this->devImageA); // Copie les heaters dans l'image A. + diffuseMethode1<<>>(this->devImageA, this->devImageB); // Diffuse l'image A dans l'image B. + this->devImageA.swapWith(this->devImageB); // Swap l'image A et l'image B. + } + + display<<>>(this->devImageA, ptrDevPixels, this->calibreur); // Affiche l'image A (pour avoir les heaters intacts ). + + //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG). + } + +void HeatTransfert::animationStep() + { + this->t += NB_ITERATION_AVEUGLE; + } + +int HeatTransfert::getW() + { + return this->devImageA.getW(); + } + +int HeatTransfert::getH() + { + return this->devImageA.getH(); + } + +float HeatTransfert::getT() + { + return float(this->t); + } + +string HeatTransfert::getTitle() + { + return this->title; + } + +void HeatTransfert::setHeaters(HeatImage& image) + { + for (int x = 179; x <= 195; x++) + { + for (int y = 179; y <= 195; y++) + image.set(x, y, 0.2); + for (int y = 605; y <= 621; y++) + image.set(x, y, 0.2); + } + + for (int x = 605; x <= 621; x++) + { + for (int y = 179; y <= 195; y++) + image.set(x, y, 0.2); + for (int y = 605; y <= 621; y++) + image.set(x, y, 0.2); + } + + image.set(295, 400, 0.2); + image.set(400, 295, 0.2); + image.set(505, 400, 0.2); + image.set(400, 505, 0.2); + + for (int x = 300; x <= 500; x++) + for (int y = 300; y <= 500; y++) + image.set(x, y, 1); + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.h new file mode 100755 index 0000000..795613f --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.h @@ -0,0 +1,46 @@ +#ifndef HEAT_TRANSFERT_H_ +#define HEAT_TRANSFERT_H_ + +#include "cudaTools.h" +#include "Animable_I.h" +#include "MathTools.h" +#include "CalibreurF.h" + +#include "HeatImage.h" + +class HeatTransfert : public Animable_I + { + static const int NB_ITERATION_AVEUGLE = 1000; + + public: + HeatTransfert(); + virtual ~HeatTransfert(); + + void runGPU(uchar4* ptrDevPixels) /*override*/; + void animationStep() /*override*/; + + int getW() /*override*/; + int getH() /*override*/; + + float getT() /*override*/; + + string getTitle(void) /*override*/; + + private: + static void setHeaters(HeatImage& image); + + CalibreurF calibreur; + + HeatImage devImageHeaters; + HeatImage devImageA; + HeatImage devImageB; + + int t; + + const dim3 dg; + const dim3 db; + + const string title; + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.cpp new file mode 100755 index 0000000..3d95d87 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.cpp @@ -0,0 +1,12 @@ +#include "HeatTransfertProvider.h" + +HeatTransfert* HeatTransfertProvider::create() + { + return new HeatTransfert(); + } + +Image* HeatTransfertProvider::createGL() + { + ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0); + return new Image(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.h new file mode 100755 index 0000000..e48afa3 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.h @@ -0,0 +1,14 @@ +#ifndef HEAT_TRANSFERT_PROVIDER_H_ +#define HEAT_TRANSFERT_PROVIDER_H_ + +#include "HeatTransfert.h" +#include "Image.h" + +class HeatTransfertProvider + { + public: + static HeatTransfert* create(); + static Image* createGL(); + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp index d2b6adf..6eed6e3 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp @@ -12,7 +12,7 @@ using namespace std; #include "RipplingProvider.h" #include "FractalProvider.h" #include "NewtonProvider.h" - +#include "HeatTransfertProvider.h" template class Viewer @@ -40,10 +40,8 @@ int mainGL(void) //Viewer rippling0(true, true, 10, 10); //Viewer fractalMandelbrot(true, true, 20, 20); //Viewer fractalJulia(true, true, 30, 30); - - - Viewer newtown(true, true, 20, 20); - + //Viewer newtown(true, true, 20, 20); + Viewer heatTransfert(true, false, 20, 20); GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte