X-Git-Url: http://git.euphorik.ch/?p=GPU.git;a=blobdiff_plain;f=WCudaMSE%2FStudent_Cuda_Image%2Fsrc%2Fcpp%2Fcore%2F06_Convolution%2Fmoo%2Fhost%2FConvolution.cu;h=817b37ba604d42b7e465849f578a5c55aee9b3e6;hp=ecc44643334e65a5baf8e8f24d690c1c9300059c;hb=52f4366920a116060ee62b84a2ed8b42f3c2d382;hpb=4182eb3a07b7143afb8ebebfe77e8ef8e8abc266

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
--- 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()