Ajout du TP HeatTransfert.
authorgburri <gregory.burri@master.hes-so.ch>
Sun, 7 Dec 2014 15:29:08 +0000 (16:29 +0100)
committergburri <gregory.burri@master.hes-so.ch>
Sun, 7 Dec 2014 15:29:08 +0000 (16:29 +0100)
WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/IndiceTools_CPU.h
WCudaMSE/BilatTools_Cuda_Image/src/core/cudaImageTools/bitmap/header/Animable_I.h
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.cu [new file with mode: 0755]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/device/HeatTransfertDevice.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.cu [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatImage.h [new file with mode: 0644]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.cu [new file with mode: 0755]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/moo/host/HeatTransfert.h [new file with mode: 0755]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.cpp [new file with mode: 0755]
WCudaMSE/Student_Cuda_Image/src/cpp/core/05_HeatTransfert/provider/HeatTransfertProvider.h [new file with mode: 0755]
WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp

index 89e2f3e..a97f64b 100755 (executable)
 namespace cpu\r
     {\r
     class IndiceTools\r
-       {\r
+        {\r
 \r
-           /*--------------------------------------*\\r
+            /*--------------------------------------*\\r
            |*          Constructor                 *|\r
-            \*--------------------------------------*/\r
+             \*--------------------------------------*/\r
 \r
-       public:\r
+        public:\r
 \r
-           /*--------------------------------------*\\r
+            /*--------------------------------------*\\r
            |*          Methodes                    *|\r
-            \*--------------------------------------*/\r
+             \*--------------------------------------*/\r
 \r
-       public:\r
+        public:\r
 \r
-           /**\r
-            * s[0,W*H[ --> i[0,H[ j[0,W[\r
-            */\r
-           static void toIJ( int s, int W, int* ptri, int* ptrj);\r
+            /**\r
+             * s[0,W*H[ --> i[0,H[ j[0,W[\r
+             */\r
+            static void toIJ(int s, int W, int* ptri, int* ptrj);\r
 \r
-           /**\r
-            * i[0,H[ j[0,W[ --> s[0,W*H[\r
-            */\r
-           static  int toS(int W, int i, int j);\r
+            /**\r
+             * i[0,H[ j[0,W[ --> s[0,W*H[\r
+             */\r
+            static int toS(int W, int i, int j);\r
 \r
-\r
-           /*-------------------------------------*\\r
+            /*-------------------------------------*\\r
             |*         Attributs                   *|\r
-            \*-------------------------------------*/\r
+             \*-------------------------------------*/\r
 \r
-       private:\r
+        private:\r
 \r
-       };\r
+        };\r
     }\r
 \r
-#endif \r
+#endif\r
 \r
 /*----------------------------------------------------------------------*\\r
  |*                    End                                             *|\r
index eba8685..b421793 100755 (executable)
@@ -17,8 +17,7 @@ using std::string;
 class Animable_I\r
     {\r
     public:\r
-\r
-       //virtual ~Animable_I(void)=0;\r
+       virtual ~Animable_I(void) {};\r
 \r
        virtual void runGPU(uchar4* ptrDevPixels)=0;\r
        virtual void animationStep(void)=0;\r
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 (executable)
index 0000000..03225a5
--- /dev/null
@@ -0,0 +1,126 @@
+#include "HeatTransfertDevice.h"\r
+\r
+#include <iostream>\r
+using namespace std;\r
+\r
+#include "Indice2D.h"\r
+#include "IndiceTools.h"\r
+#include "cudaTools.h"\r
+#include "Device.h"\r
+\r
+#include "ColorTools.h"\r
+\r
+#include "HeatImage.h"\r
+\r
+__global__\r
+void display(HeatImage image, uchar4* ptrDevPixels, CalibreurF calibreur)\r
+    {\r
+    const int TID = Indice2D::tid();\r
+    const int NB_THREAD = Indice2D::nbThread();\r
+    const int WH = image.getWCuda() * image.getHCuda();\r
+\r
+    uchar4 color;\r
+\r
+    int pixelI;\r
+    int pixelJ;\r
+\r
+    int s = TID;\r
+    while (s < WH)\r
+       {\r
+       IndiceTools::toIJ(s, image.getWCuda(), &pixelI, &pixelJ);\r
+\r
+       float heatValue = image.getCuda(pixelJ, pixelI);\r
+       calibreur.calibrer(heatValue);\r
+       ColorTools::HSB_TO_RVB(heatValue, &color);\r
+\r
+       ptrDevPixels[s] = color;\r
+       s += NB_THREAD;\r
+       }\r
+    }\r
+\r
+__global__\r
+void copyHeaters(HeatImage heaters, HeatImage destination)\r
+    {\r
+    const int TID = Indice2D::tid();\r
+    const int NB_THREAD = Indice2D::nbThread();\r
+    const int WH = heaters.getWCuda() * heaters.getHCuda();\r
+\r
+    int pixelI;\r
+    int pixelJ;\r
+\r
+    int s = TID;\r
+    while (s < WH)\r
+        {\r
+        IndiceTools::toIJ(s, heaters.getWCuda(), &pixelI, &pixelJ);\r
+\r
+        float heatValue = heaters.getCuda(pixelJ, pixelI);\r
+        if (heatValue > 0.0)\r
+            destination.setCuda(pixelJ, pixelI, heatValue);\r
+\r
+        s += NB_THREAD;\r
+        }\r
+    }\r
+\r
+__global__\r
+void diffuseMethode1(HeatImage from, HeatImage to)\r
+    {\r
+    const int TID = Indice2D::tid();\r
+    const int NB_THREAD = Indice2D::nbThread();\r
+\r
+    const int W = from.getWCuda() - 2;\r
+    const int H = from.getHCuda() - 2;\r
+    const int WH = W * H;\r
+\r
+    const float k = 0.1;\r
+\r
+    int pixelI;\r
+    int pixelJ;\r
+\r
+    int s = TID;\r
+    while (s < WH)\r
+        {\r
+        IndiceTools::toIJ(s, W, &pixelI, &pixelJ);\r
+        const int x = pixelJ + 1;\r
+        const int y = pixelI + 1;\r
+        const float t_old = from.getCuda(x, y);\r
+        const float t_up = from.getCuda(x, y - 1);\r
+        const float t_down = from.getCuda(x, y + 1);\r
+        const float t_left = from.getCuda(x - 1, y);\r
+        const float t_right = from.getCuda(x + 1, y);\r
+\r
+        to.setCuda(x, y, t_old + k * (t_up + t_down + t_left + t_right - 4.0 * t_old));\r
+\r
+        s += NB_THREAD;\r
+        }\r
+    }\r
+\r
+__global__\r
+void diffuseMethode2(HeatImage from, HeatImage to)\r
+    {\r
+    const int TID = Indice2D::tid();\r
+    const int NB_THREAD = Indice2D::nbThread();\r
+\r
+    const int W = from.getWCuda() - 2;\r
+    const int H = from.getHCuda() - 2;\r
+    const int WH = W * H;\r
+\r
+    int pixelI;\r
+    int pixelJ;\r
+\r
+    int s = TID;\r
+    while (s < WH)\r
+        {\r
+        IndiceTools::toIJ(s, W, &pixelI, &pixelJ);\r
+        const int x = pixelJ + 1;\r
+        const int y = pixelI + 1;\r
+        const float t_old = from.getCuda(x, y);\r
+        const float t_up = from.getCuda(x, y - 1);\r
+        const float t_down = from.getCuda(x, y + 1);\r
+        const float t_left = from.getCuda(x - 1, y);\r
+        const float t_right = from.getCuda(x + 1, y);\r
+\r
+        to.setCuda(x, y, (t_up + t_down + t_left + t_right + 4.0 * t_old) / 8.0);\r
+\r
+        s += NB_THREAD;\r
+        }\r
+    }\r
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 (file)
index 0000000..d64c3cd
--- /dev/null
@@ -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 (file)
index 0000000..e383fb6
--- /dev/null
@@ -0,0 +1,12 @@
+#include "HeatImage.h"
+
+#include <cstring>
+
+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 (file)
index 0000000..00224c8
--- /dev/null
@@ -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 (executable)
index 0000000..7aa7cf6
--- /dev/null
@@ -0,0 +1,119 @@
+#include "HeatTransfert.h"\r
+\r
+#include <iostream>\r
+#include <cstring>\r
+#include <assert.h>\r
+#include <stdio.h>\r
+using namespace std;\r
+\r
+#include "Device.h"\r
+\r
+#include "HeatTransfertDevice.h"\r
+\r
+HeatTransfert::HeatTransfert() :\r
+      calibreur(IntervalF(0, 1), IntervalF(0.7, 0)),\r
+      devImageHeaters(800, 800),\r
+      devImageA(800, 800),\r
+      devImageB(800, 800),\r
+      t(0),\r
+      dg(8, 8, 1),\r
+      db(32, 32, 1),\r
+      title("Heat transfert")\r
+    {\r
+    const size_t IMAGE_SIZE = this->devImageA.getW() * this->devImageA.getH() * sizeof(float);\r
+    const int IMAGE_W = this->devImageA.getW();\r
+    const int IMAGE_H = this->devImageA.getH();\r
+\r
+    // Définition des heaters.\r
+    HeatImage imageHeaters(IMAGE_W, IMAGE_H, new float[IMAGE_W * IMAGE_H]);\r
+    imageHeaters.clear();\r
+    setHeaters(imageHeaters);\r
+\r
+    // Définition des températures initiales (toutes à zero).\r
+    HeatImage imageInit(IMAGE_W, IMAGE_H, new float[IMAGE_W * IMAGE_H]);\r
+    imageInit.clear();\r
+\r
+    // Copie des heaters en GM.\r
+    HANDLE_ERROR(cudaMalloc(&this->devImageHeaters.getRaw(), IMAGE_SIZE));\r
+    HANDLE_ERROR(cudaMemcpy(this->devImageHeaters.getRaw(), imageHeaters.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice));\r
+\r
+    // Allocation de image A et image B.\r
+    HANDLE_ERROR(cudaMalloc(&this->devImageA.getRaw(), IMAGE_SIZE));\r
+    HANDLE_ERROR(cudaMemcpy(this->devImageA.getRaw(), imageInit.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice));\r
+    HANDLE_ERROR(cudaMalloc(&this->devImageB.getRaw(), IMAGE_SIZE));\r
+    HANDLE_ERROR(cudaMemcpy(this->devImageB.getRaw(), imageInit.getRaw(), IMAGE_SIZE, cudaMemcpyHostToDevice));\r
+    }\r
+\r
+HeatTransfert::~HeatTransfert()\r
+    {\r
+    cudaFree(this->devImageHeaters.getRaw());\r
+    cudaFree(this->devImageA.getRaw());\r
+    cudaFree(this->devImageB.getRaw());\r
+    }\r
+\r
+void HeatTransfert::runGPU(uchar4* ptrDevPixels)\r
+    {\r
+    for (int i = 0; i < NB_ITERATION_AVEUGLE; i++)\r
+        {\r
+        copyHeaters<<<dg, db>>>(this->devImageHeaters, this->devImageA); // Copie les heaters dans l'image A.\r
+        diffuseMethode1<<<dg, db>>>(this->devImageA, this->devImageB); // Diffuse l'image A dans l'image B.\r
+        this->devImageA.swapWith(this->devImageB); // Swap l'image A et l'image B.\r
+        }\r
+\r
+    display<<<dg, db>>>(this->devImageA, ptrDevPixels, this->calibreur); // Affiche l'image A (pour avoir les heaters intacts   ).\r
+\r
+    //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).\r
+    }\r
+\r
+void HeatTransfert::animationStep()\r
+    {\r
+    this->t += NB_ITERATION_AVEUGLE;\r
+    }\r
+\r
+int HeatTransfert::getW()\r
+    {\r
+    return this->devImageA.getW();\r
+    }\r
+\r
+int HeatTransfert::getH()\r
+    {\r
+    return this->devImageA.getH();\r
+    }\r
+\r
+float HeatTransfert::getT()\r
+    {\r
+    return float(this->t);\r
+    }\r
+\r
+string HeatTransfert::getTitle()\r
+    {\r
+    return this->title;\r
+    }\r
+\r
+void HeatTransfert::setHeaters(HeatImage& image)\r
+    {\r
+    for (int x = 179; x <= 195; x++)\r
+        {\r
+        for (int y = 179; y <= 195; y++)\r
+            image.set(x, y, 0.2);\r
+        for (int y = 605; y <= 621; y++)\r
+            image.set(x, y, 0.2);\r
+        }\r
+\r
+    for (int x = 605; x <= 621; x++)\r
+        {\r
+        for (int y = 179; y <= 195; y++)\r
+            image.set(x, y, 0.2);\r
+        for (int y = 605; y <= 621; y++)\r
+            image.set(x, y, 0.2);\r
+        }\r
+\r
+    image.set(295, 400, 0.2);\r
+    image.set(400, 295, 0.2);\r
+    image.set(505, 400, 0.2);\r
+    image.set(400, 505, 0.2);\r
+\r
+    for (int x = 300; x <= 500; x++)\r
+        for (int y = 300; y <= 500; y++)\r
+            image.set(x, y, 1);\r
+    }\r
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 (executable)
index 0000000..795613f
--- /dev/null
@@ -0,0 +1,46 @@
+#ifndef HEAT_TRANSFERT_H_\r
+#define HEAT_TRANSFERT_H_\r
+\r
+#include "cudaTools.h"\r
+#include "Animable_I.h"\r
+#include "MathTools.h"\r
+#include "CalibreurF.h"\r
+\r
+#include "HeatImage.h"\r
+\r
+class HeatTransfert : public Animable_I\r
+    {\r
+        static const int NB_ITERATION_AVEUGLE = 1000;\r
+\r
+    public:\r
+        HeatTransfert();\r
+        virtual ~HeatTransfert();\r
+\r
+        void runGPU(uchar4* ptrDevPixels) /*override*/;\r
+        void animationStep() /*override*/;\r
+\r
+        int getW() /*override*/;\r
+        int getH() /*override*/;\r
+\r
+        float getT() /*override*/;\r
+\r
+        string getTitle(void) /*override*/;\r
+\r
+    private:\r
+        static void setHeaters(HeatImage& image);\r
+\r
+        CalibreurF calibreur;\r
+\r
+        HeatImage devImageHeaters;\r
+        HeatImage devImageA;\r
+        HeatImage devImageB;\r
+\r
+        int t;\r
+\r
+        const dim3 dg;\r
+        const dim3 db;\r
+\r
+        const string title;\r
+    };\r
+\r
+#endif\r
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 (executable)
index 0000000..3d95d87
--- /dev/null
@@ -0,0 +1,12 @@
+#include "HeatTransfertProvider.h"\r
+\r
+HeatTransfert* HeatTransfertProvider::create()\r
+    {\r
+    return new HeatTransfert();\r
+    }\r
+\r
+Image* HeatTransfertProvider::createGL()\r
+    {\r
+    ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0);\r
+    return new Image(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel\r
+    }\r
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 (executable)
index 0000000..e48afa3
--- /dev/null
@@ -0,0 +1,14 @@
+#ifndef HEAT_TRANSFERT_PROVIDER_H_\r
+#define HEAT_TRANSFERT_PROVIDER_H_\r
+\r
+#include "HeatTransfert.h"\r
+#include "Image.h"\r
+\r
+class HeatTransfertProvider\r
+    {\r
+    public:\r
+       static HeatTransfert* create();\r
+       static Image* createGL();\r
+    };\r
+\r
+#endif\r
index d2b6adf..6eed6e3 100755 (executable)
@@ -12,7 +12,7 @@ using namespace std;
 #include "RipplingProvider.h"\r
 #include "FractalProvider.h"\r
 #include "NewtonProvider.h"\r
-\r
+#include "HeatTransfertProvider.h"\r
 \r
 template <class TOutput, class TProvider>\r
 class Viewer\r
@@ -40,10 +40,8 @@ int mainGL(void)
     //Viewer<Image, RipplingProvider> rippling0(true, true, 10, 10);\r
     //Viewer<ImageFonctionel, MandelbrotProvider> fractalMandelbrot(true, true, 20, 20);\r
     //Viewer<ImageFonctionel, JuliaProvider> fractalJulia(true, true, 30, 30);\r
-\r
-\r
-    Viewer<ImageFonctionel, NewtonProvider> newtown(true, true, 20, 20);\r
-\r
+    //Viewer<ImageFonctionel, NewtonProvider> newtown(true, true, 20, 20);\r
+    Viewer<Image, HeatTransfertProvider> heatTransfert(true, false, 20, 20);\r
 \r
     GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte\r
 \r