#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;
{
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;
}