#include "ImageConvolutionCuda.h"
+/**
+ * 'captureur' est supprimé par le destructeur de cette classe.
+ */
ImageConvolutionCuda::ImageConvolutionCuda(Animable_I* ptrAnimable, CaptureVideo* captureur, ColorRGB_01* ptrColorTitreRGB) :
Image(ptrAnimable, ptrColorTitreRGB),
captureur(captureur)
{
-
- //Mat matImage = captureur.capturer(); // capture une image seulement ( Ã utiliser en boucle!)
- //uchar4* image = CaptureVideo::castToUChar4(&matImage); // format cuda
}
ImageConvolutionCuda::~ImageConvolutionCuda()
void ImageConvolutionCuda::fillImageGL(uchar4* ptrDevImageGL, int w, int h)
{
- Mat matImage = this->captureur->capturer(); // capture une image seulement ( Ã utiliser en boucle!)
+ Mat matImage = this->captureur->capturer(); // Capture l'image suivante.
uchar4* image = CaptureVideo::castToUChar4(&matImage); // format cuda
HANDLE_ERROR(cudaMemcpy(ptrDevImageGL, image, sizeof(uchar4) * this->ptrAnimable->getW() * this->ptrAnimable->getH(), cudaMemcpyHostToDevice));
+
+ Image::fillImageGL(ptrDevImageGL, w, h);
}
public:
void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override
- //void animationStep(bool& isNeedUpdateView); // override
- //void paintPrimitives(Graphic2Ds& graphic2D); // override
private:
CaptureVideo* captureur;
#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;
}
#ifndef CONVOLUTION_DEVICE_H
#define CONVOLUTION_DEVICE_H
+#include "ConstantMemoryLink.h"
+
+const int KERNEL_SIZE = 9;
+
+/**
+ * Accès à la mémoire constante.
+ */
+ConstantMemoryLink constantMemoryKernelLink();
+
+/**
+ * Associe l'image source en tant que texture.
+ */
+void bindSouceAsTexture(uchar4* source, int w, int h);
+void unbindSouceAsTexture();
+
+/**
+ * Transforme l'image fournit en niveau de gris.
+ */
+__global__
+void toGrayscale(uchar4* ptrDevPixels, int w, int h);
+
+/**
+ * Convolution de l'image source 'textureImageSource' vers 'ptrDevOutput'.
+ */
__global__
-void convolution(uchar4* ptrDevPixels, int w, int h, float t);
+void convolution(uchar4* ptrDevOutput, int w, int h);
#endif
#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()
#include "MathTools.h"
#include "AleaTools.h"
+#include "ConvolutionDevice.h"
+
class Convolution : public Animable_I
{
public:
int getW() /*override*/;
int getH() /*override*/;
-
float getT() /*override*/;
std::string getTitle(void) /*override*/;
const std::string title;
- static const float kernel[9][9];
+ uchar4* ptrDevImageSource;
+
+ static const float kernel[KERNEL_SIZE][KERNEL_SIZE];
};
#endif
projects / GPU.git / commitdiff