From 52f4366920a116060ee62b84a2ed8b42f3c2d382 Mon Sep 17 00:00:00 2001 From: gburri Date: Sun, 11 Jan 2015 19:06:04 +0100 Subject: [PATCH] Fin du TP convolution. (Constant memory + Texture). --- .../06_Convolution/ImageConvolutionCuda.cpp | 10 ++- .../06_Convolution/ImageConvolutionCuda.h | 2 - .../moo/device/ConvolutionDevice.cu | 82 +++++++++++++++++-- .../moo/device/ConvolutionDevice.h | 26 +++++- .../06_Convolution/moo/host/Convolution.cu | 24 +++++- .../06_Convolution/moo/host/Convolution.h | 7 +- 6 files changed, 129 insertions(+), 22 deletions(-) diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp index ffc30a1..aeeadce 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.cpp @@ -1,12 +1,12 @@ #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() @@ -16,7 +16,9 @@ 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); } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h index fb35821..ee9c8fe 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/ImageConvolutionCuda.h @@ -14,8 +14,6 @@ class ImageConvolutionCuda : public Image public: void fillImageGL(uchar4* ptrDevImageGL, int w, int h); // override - //void animationStep(bool& isNeedUpdateView); // override - //void paintPrimitives(Graphic2Ds& graphic2D); // override private: CaptureVideo* captureur; 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 --- 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 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(); + 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; } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h index 7044644..957015b 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/device/ConvolutionDevice.h @@ -1,7 +1,31 @@ #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 diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu index ecc4464..817b37b 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.cu @@ -6,8 +6,6 @@ using namespace std; #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}, @@ -27,17 +25,35 @@ Convolution::Convolution(int w, int h) : 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<<>>(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<<>>(this->ptrDevImageSource, this->w, this->h); + + cudaDeviceSynchronize(); // Attend que toute l'image source ait été passée en niveau de gris. + convolution<<>>(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() diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h index e67adc2..d952df8 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/06_Convolution/moo/host/Convolution.h @@ -6,6 +6,8 @@ #include "MathTools.h" #include "AleaTools.h" +#include "ConvolutionDevice.h" + class Convolution : public Animable_I { public: @@ -17,7 +19,6 @@ class Convolution : public Animable_I int getW() /*override*/; int getH() /*override*/; - float getT() /*override*/; std::string getTitle(void) /*override*/; @@ -33,7 +34,9 @@ class Convolution : public Animable_I const std::string title; - static const float kernel[9][9]; + uchar4* ptrDevImageSource; + + static const float kernel[KERNEL_SIZE][KERNEL_SIZE]; }; #endif -- 2.45.2