X-Git-Url: http://git.euphorik.ch/?p=GPU.git;a=blobdiff_plain;f=WCudaMSE%2FStudent_Cuda_Image%2Fsrc%2Fcpp%2Fcore%2F05_HeatTransfert%2Fmoo%2Fdevice%2FHeatTransfertDevice.cu;fp=WCudaMSE%2FStudent_Cuda_Image%2Fsrc%2Fcpp%2Fcore%2F05_HeatTransfert%2Fmoo%2Fdevice%2FHeatTransfertDevice.cu;h=03225a5251d2dcc89f2021feaebb97248b7d9404;hp=0000000000000000000000000000000000000000;hb=cb39d6a91b65d2862018430d65e633d2a8fdc818;hpb=2d95edd9a2d09421e5eae56755bdf3105e12edf7

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 <iostream>
+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;
+        }
+    }