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

[GPU.git] / WCudaMSE / Student_Cuda_Image / src / cpp / core / 06_Convolution / moo / device / ConvolutionDevice.cu

diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu
index 640d8a9..0cfc48f 100644 (file)
--- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu
+++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.cu
@@ -9,17 +9,72 @@ using namespace std;
 
 #include "ConvolutionDevice.h"
 
+// Le kernel de la convolution en constant memory.
+__constant__ float TAB_KERNEL[KERNEL_SIZE][KERNEL_SIZE];
+
+// L'image source accédée comme une texture.
+texture<uchar4, 2, cudaReadModeElementType> textureImageSource;
+
+ConstantMemoryLink constantMemoryKernelLink()
+    {
+    float* ptrDevTabData;
+    size_t sizeAll = KERNEL_SIZE * KERNEL_SIZE * sizeof(float);
+    HANDLE_ERROR(cudaGetSymbolAddress((void**)&ptrDevTabData, TAB_KERNEL));
+    ConstantMemoryLink cmLink =
+        {
+          (void**)ptrDevTabData, KERNEL_SIZE * KERNEL_SIZE, sizeAll
+        };
+    return cmLink;
+    }
+
+void bindSouceAsTexture(uchar4* source, int w, int h)
+    {
+    // Propriétés de la texture (image en entrée).
+    textureImageSource.addressMode[0] = cudaAddressModeClamp;
+    textureImageSource.addressMode[1] = cudaAddressModeClamp;
+    textureImageSource.filterMode = cudaFilterModePoint;
+    textureImageSource.normalized = false;
+
+    const size_t pitch = w * sizeof(uchar4); // Taille d'une ligne.
+    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar4>();
+    HANDLE_ERROR(cudaBindTexture2D(NULL, textureImageSource, source, channelDesc, w, h, pitch));
+    }
+
+void unbindSouceAsTexture()
+    {
+    HANDLE_ERROR(cudaUnbindTexture(textureImageSource));
+    }
+
 __global__
-void convolution(uchar4* ptrDevPixels, int w, int h, float t)
+void toGrayscale(uchar4* ptrDevPixels, int w, int h)
     {
     const int TID = Indice2D::tid();
     const int NB_THREAD = Indice2D::nbThread();
     const int WH = w * h;
 
-    uchar4 color;
-    color.w = 255; // Par défaut, l'image est opaque.
+    int pixelI, pixelJ;
+
+    int s = TID;
+    while (s < WH)
+        {
+        IndiceTools::toIJ(s, w, &pixelI, &pixelJ);
+
+        const uchar average = (uchar)(((int)ptrDevPixels[s].x + (int)ptrDevPixels[s].y + (int)ptrDevPixels[s].z) / 3);
+        ptrDevPixels[s].x = average;
+        ptrDevPixels[s].y = average;
+        ptrDevPixels[s].z = average;
+
+        s += NB_THREAD;
+        }
+    }
+
+__global__
+void convolution(uchar4* ptrDevOutput, int w, int h)
+    {
+    const int TID = Indice2D::tid();
+    const int NB_THREAD = Indice2D::nbThread();
+    const int WH = w * h;
 
-    double x, y;
     int pixelI, pixelJ;
 
     int s = TID;
@@ -27,13 +82,22 @@ void convolution(uchar4* ptrDevPixels, int w, int h, float t)
         {
         IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ)
 
-        // (i,j) domaine écran.
-        // (x,y) domaine math.
-        // domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y).
+        // La somme des produits du kernel avec l'image source.
+        float sum = 0;
+        for (int i = -4; i <= 4; i++)
+            for (int j = -4; j <= 4; j++)
+                {
+                    uchar4 valueSource = tex2D(textureImageSource, pixelJ + j, pixelI + i);
+                    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.
+                }
+
+        sum /= 100; // Comme défini dans la donnée.
 
-        // newtonMath.colorXY(&color, x, y);
+        const float finalValue = (uchar)(sum * 50);
 
-        ptrDevPixels[s] = color;
+        ptrDevOutput[s].x = finalValue;
+        ptrDevOutput[s].y = finalValue;
+        ptrDevOutput[s].z = finalValue;
 
         s += NB_THREAD;
         }