WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu

   1 #include <iostream>
   2 #include <assert.h>
   3 #include <stdio.h>
   4 using namespace std;
   5
   6 #include "Device.h"
   7
   8 #include "Convolution.h"
   9
  10 const float Convolution::kernel[9][9] =
  11        {{0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828},
  12        {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987},
  13        {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705},
  14        {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366},
  15        {0.6063, 1.2765, 0.4061, -6.4829, -12.7462, -6.4829, 0.4061, 1.2765, 0.6063},
  16        {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366},
  17        {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705},
  18        {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987},
  19        {0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828}};
  20
  21 Convolution::Convolution(int w, int h) :
  22       w(w), h(h),
  23       dg(8, 8, 1),
  24       db(32, 32, 1),
  25       title("Convolution")
  26     {
  27     Device::assertDim(dg, db);
  28
  29     // Allocation de la mémoire sur le GPU dédié à l'image source et bind avec la textue.
  30     HANDLE_ERROR(cudaMalloc(&this->ptrDevImageSource, this->w * this->h * sizeof(uchar4)));
  31     bindSouceAsTexture(this->ptrDevImageSource, this->w, this->h);
  32
  33     // Copie du kernel en constant memory.
  34     ConstantMemoryLink cmKernelLink = constantMemoryKernelLink();
  35     float* ptrDevKernel = (float*)cmKernelLink.ptrDevTab;
  36     size_t sizeALL = cmKernelLink.sizeAll;
  37     HANDLE_ERROR(cudaMemcpy(ptrDevKernel, kernel, sizeALL, cudaMemcpyHostToDevice));
  38     }
  39
  40 Convolution::~Convolution()
  41     {
  42     unbindSouceAsTexture();
  43     HANDLE_ERROR(cudaFree(this->ptrDevImageSource));
  44     }
  45
  46 void Convolution::runGPU(uchar4* ptrDevPixels)
  47     {
  48     // Copie l'image donnée dans l'image source. La convolution sera effectuée de 'ptrImageSource' vers 'ptrDevPixels'.
  49     HANDLE_ERROR(cudaMemcpy(this->ptrDevImageSource, ptrDevPixels, this->w * this->h * sizeof(uchar4), cudaMemcpyDeviceToDevice));
  50
  51     toGrayscale<<<dg,db>>>(this->ptrDevImageSource, this->w, this->h);
  52
  53     cudaDeviceSynchronize(); // Attend que toute l'image source ait été passée en niveau de gris.
  54     convolution<<<dg,db>>>(ptrDevPixels, this->w, this->h);
  55
  56     //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
  57     }
  58
  59 void Convolution::animationStep()
  60     {
  61     this->t += 0.1; // TODO.
  62     }
  63
  64 int Convolution::getW()
  65     {
  66     return this->w;
  67     }
  68
  69 int Convolution::getH()
  70     {
  71     return this->h;
  72     }
  73
  74 float Convolution::getT()
  75     {
  76     return this->t;
  77     }
  78
  79 string Convolution::getTitle()
  80     {
  81     return this->title;
  82     }
  83