Fin du TP convolution. (Constant memory + Texture).

[GPU.git] / WCudaMSE / Student_Cuda_Image / src / cpp / core / 06_Convolution / moo / host / Convolution.cu

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
index ecc4464..817b37b 100644 (file)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu
@@ -6,8 +6,6 @@ using namespace std;
 #include "Device.h"
 
 #include "Convolution.h"
-#include "ConvolutionDevice.h"
-
 
 const float Convolution::kernel[9][9] =
        {{0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828},
@@ -27,17 +25,35 @@ Convolution::Convolution(int w, int h) :
       title("Convolution")
     {
     Device::assertDim(dg, db);
+
+    // Allocation de la mémoire sur le GPU dédié à l'image source et bind avec la textue.
+    HANDLE_ERROR(cudaMalloc(&this->ptrDevImageSource, this->w * this->h * sizeof(uchar4)));
+    bindSouceAsTexture(this->ptrDevImageSource, this->w, this->h);
+
+    // Copie du kernel en constant memory.
+    ConstantMemoryLink cmKernelLink = constantMemoryKernelLink();
+    float* ptrDevKernel = (float*)cmKernelLink.ptrDevTab;
+    size_t sizeALL = cmKernelLink.sizeAll;
+    HANDLE_ERROR(cudaMemcpy(ptrDevKernel, kernel, sizeALL, cudaMemcpyHostToDevice));
     }
 
 Convolution::~Convolution()
     {
+    unbindSouceAsTexture();
+    HANDLE_ERROR(cudaFree(this->ptrDevImageSource));
     }
 
 void Convolution::runGPU(uchar4* ptrDevPixels)
     {
-    convolution<<<dg,db>>>(ptrDevPixels, this->w, this->h, this->t);
+    // Copie l'image donnée dans l'image source. La convolution sera effectuée de 'ptrImageSource' vers 'ptrDevPixels'.
+    HANDLE_ERROR(cudaMemcpy(this->ptrDevImageSource, ptrDevPixels, this->w * this->h * sizeof(uchar4), cudaMemcpyDeviceToDevice));
+
+    toGrayscale<<<dg,db>>>(this->ptrDevImageSource, this->w, this->h);
+
+    cudaDeviceSynchronize(); // Attend que toute l'image source ait été passée en niveau de gris.
+    convolution<<<dg,db>>>(ptrDevPixels, this->w, this->h);
 
-    // HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
+    //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
     }
 
 void Convolution::animationStep()