WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu

   1 #include <iostream>
   2 #include <stdio.h>
   3 using namespace std;
   4
   5 #include "Indice2D.h"
   6 #include "IndiceTools.h"
   7 #include "cudaTools.h"
   8 #include "Device.h"
   9
  10 #include "ConvolutionDevice.h"
  11
  12 // Le kernel de la convolution en constant memory.
  13 __constant__ float TAB_KERNEL[KERNEL_SIZE][KERNEL_SIZE];
  14
  15 // L'image source accédée comme une texture.
  16 texture<uchar4, 2, cudaReadModeElementType> textureImageSource;
  17
  18 ConstantMemoryLink constantMemoryKernelLink()
  19     {
  20     float* ptrDevTabData;
  21     size_t sizeAll = KERNEL_SIZE * KERNEL_SIZE * sizeof(float);
  22     HANDLE_ERROR(cudaGetSymbolAddress((void**)&ptrDevTabData, TAB_KERNEL));
  23     ConstantMemoryLink cmLink =
  24         {
  25           (void**)ptrDevTabData, KERNEL_SIZE * KERNEL_SIZE, sizeAll
  26         };
  27     return cmLink;
  28     }
  29
  30 void bindSouceAsTexture(uchar4* source, int w, int h)
  31     {
  32     // Propriétés de la texture (image en entrée).
  33     textureImageSource.addressMode[0] = cudaAddressModeClamp;
  34     textureImageSource.addressMode[1] = cudaAddressModeClamp;
  35     textureImageSource.filterMode = cudaFilterModePoint;
  36     textureImageSource.normalized = false;
  37
  38     const size_t pitch = w * sizeof(uchar4); // Taille d'une ligne.
  39     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar4>();
  40     HANDLE_ERROR(cudaBindTexture2D(NULL, textureImageSource, source, channelDesc, w, h, pitch));
  41     }
  42
  43 void unbindSouceAsTexture()
  44     {
  45     HANDLE_ERROR(cudaUnbindTexture(textureImageSource));
  46     }
  47
  48 __global__
  49 void toGrayscale(uchar4* ptrDevPixels, int w, int h)
  50     {
  51     const int TID = Indice2D::tid();
  52     const int NB_THREAD = Indice2D::nbThread();
  53     const int WH = w * h;
  54
  55     int pixelI, pixelJ;
  56
  57     int s = TID;
  58     while (s < WH)
  59         {
  60         IndiceTools::toIJ(s, w, &pixelI, &pixelJ);
  61
  62         const uchar average = (uchar)(((int)ptrDevPixels[s].x + (int)ptrDevPixels[s].y + (int)ptrDevPixels[s].z) / 3);
  63         ptrDevPixels[s].x = average;
  64         ptrDevPixels[s].y = average;
  65         ptrDevPixels[s].z = average;
  66
  67         s += NB_THREAD;
  68         }
  69     }
  70
  71 __global__
  72 void convolution(uchar4* ptrDevOutput, int w, int h)
  73     {
  74     const int TID = Indice2D::tid();
  75     const int NB_THREAD = Indice2D::nbThread();
  76     const int WH = w * h;
  77
  78     int pixelI, pixelJ;
  79
  80     int s = TID;
  81     while (s < WH)
  82         {
  83         IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ)
  84
  85         // La somme des produits du kernel avec l'image source.
  86         float sum = 0;
  87         for (int i = -4; i <= 4; i++)
  88             for (int j = -4; j <= 4; j++)
  89                 {
  90                     uchar4 valueSource = tex2D(textureImageSource, pixelJ + j, pixelI + i);
  91                     sum += (float)valueSource.x * TAB_KERNEL[i + 4][j + 4]; // Comme c'est une image en niveau de gris on ne prend qu'une composante.
  92                 }
  93
  94         sum /= 100; // Comme défini dans la donnée.
  95
  96         const float finalValue = (uchar)(sum * 50);
  97
  98         ptrDevOutput[s].x = finalValue;
  99         ptrDevOutput[s].y = finalValue;
 100         ptrDevOutput[s].z = finalValue;
 101
 102         s += NB_THREAD;
 103         }
 104     }