817b37ba604d42b7e465849f578a5c55aee9b3e6
[GPU.git] / Convolution.cu
1 #include <iostream>
2 #include <assert.h>
3 #include <stdio.h>
4 using namespace std;
5
6 #include "Device.h"
7
8 #include "Convolution.h"
9
10 const float Convolution::kernel[9][9] =
11        {{0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828},
12        {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987},
13        {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705},
14        {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366},
15        {0.6063, 1.2765, 0.4061, -6.4829, -12.7462, -6.4829, 0.4061, 1.2765, 0.6063},
16        {0.5366, 1.1794, 1.0033, -2.8306, -6.4829, -2.8306, 1.0033, 1.1794, 0.5366},
17        {0.3705, 0.8646, 1.3475, 1.0033, 0.4061, 1.0033, 1.3475, 0.8646, 0.3705},
18        {0.1987, 0.4746, 0.8646, 1.1794, 1.2765, 1.1794, 0.8646, 0.4746, 0.1987},
19        {0.0828, 0.1987, 0.3705, 0.5366, 0.6063, 0.5366, 0.3705, 0.1987, 0.0828}};
20
21 Convolution::Convolution(int w, int h) :
22       w(w), h(h),
23       dg(8, 8, 1),
24       db(32, 32, 1),
25       title("Convolution")
26     {
27     Device::assertDim(dg, db);
28
29     // Allocation de la mémoire sur le GPU dédié à l'image source et bind avec la textue.
30     HANDLE_ERROR(cudaMalloc(&this->ptrDevImageSource, this->w * this->h * sizeof(uchar4)));
31     bindSouceAsTexture(this->ptrDevImageSource, this->w, this->h);
32
33     // Copie du kernel en constant memory.
34     ConstantMemoryLink cmKernelLink = constantMemoryKernelLink();
35     float* ptrDevKernel = (float*)cmKernelLink.ptrDevTab;
36     size_t sizeALL = cmKernelLink.sizeAll;
37     HANDLE_ERROR(cudaMemcpy(ptrDevKernel, kernel, sizeALL, cudaMemcpyHostToDevice));
38     }
39
40 Convolution::~Convolution()
41     {
42     unbindSouceAsTexture();
43     HANDLE_ERROR(cudaFree(this->ptrDevImageSource));
44     }
45
46 void Convolution::runGPU(uchar4* ptrDevPixels)
47     {
48     // Copie l'image donnée dans l'image source. La convolution sera effectuée de 'ptrImageSource' vers 'ptrDevPixels'.
49     HANDLE_ERROR(cudaMemcpy(this->ptrDevImageSource, ptrDevPixels, this->w * this->h * sizeof(uchar4), cudaMemcpyDeviceToDevice));
50
51     toGrayscale<<<dg,db>>>(this->ptrDevImageSource, this->w, this->h);
52
53     cudaDeviceSynchronize(); // Attend que toute l'image source ait été passée en niveau de gris.
54     convolution<<<dg,db>>>(ptrDevPixels, this->w, this->h);
55
56     //HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
57     }
58
59 void Convolution::animationStep()
60     {
61     this->t += 0.1; // TODO.
62     }
63
64 int Convolution::getW()
65     {
66     return this->w;
67     }
68
69 int Convolution::getH()
70     {
71     return this->h;
72     }
73
74 float Convolution::getT()
75     {
76     return this->t;
77     }
78
79 string Convolution::getTitle()
80     {
81     return this->title;
82     }
83