#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},
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()