CI2 : Le flou...
Appliquer un filtre à une image et gestion de la mémoire.
Filtrage (∼60mn, – moyen)
Le but de cet exercice est d'appliquer un filtre à une image qui va la flouter. Pour cela, nous allons décomposer notre image RGB en trois composantes, séparées dans trois tableaux (R, G, et B). Puis nous appliquerons le filtre sur chacune des composantes.
Faites un git pull pour être sur d'être à jour.
Dans notre représentation actuelle des images, nous avons accès à un tableau de uchar4 qui est une structure de quatre chars représentatant les composantes rouges (R = x), vertes (G = y), bleues (B = z) et alpha (transparence, ignorée dans notre cas). Cela signifie que notre image est représentée dans la mémoire de manière continue comme une suite de R, G, et B : RGBRGBRGBRGB... Pour appliquer un filtre, nous devons l'appliquer de la même manière sur les trois composantes rouges, vertes et bleues. Or, cela n'est pas très efficace car nous devons accéder à des zones mémoires non continues. Nous allons donc écrire un premier kernel qui va séparer notre tableau représentant l'image et trois tableaux représentant de manière continue les couleurs rouges, vertes et bleues de chaque pixel.
Rendez-vous dans le dossier problem_sets/problem_set2. Dans ce TP, nous modifierons le fichier student_func.cu. Écrivez la fonction separateChannels effectuant la séparation des trois couleurs dans trois tableaux différents.
Rendez-vous dans le dossier problem_sets/problem_set2. Dans ce TP, nous modifierons le fichier student_func.cu. Écrivez la fonction separateChannels effectuant la séparation des trois couleurs dans trois tableaux différents.
Faites bien attention de ne pas sortir de l'image !
Calculer d'abord la position x et y dans l'image en fonction du block et du thread, puis déduisez-en la position actuelle en 1D. Faire un dessin peut aider.
__global__ void separateChannels(const uchar4* const inputImageRGBA,
int numRows,
int numCols,
unsigned char* const redChannel,
unsigned char* const greenChannel,
unsigned char* const blueChannel)
{
int pos_x = blockDim.x * blockIdx.x + threadIdx.x;
int pos_y = blockDim.y * blockIdx.y + threadIdx.y;
if (pos_x >= numCols || pos_y >= numRows) return;
int position = pos_y * numCols + pos_x;
redChannel[position] = inputImageRGBA[position].x;
greenChannel[position] = inputImageRGBA[position].y;
blueChannel[position] = inputImageRGBA[position].z;
}
Maintenant que nous avons fait cette séparation, nous allons appliquer un filtre sur chaque canal de couleur pour flouter l'image. L'idée est simple : pour chaque pixel, nous allons modifier sa valeur par la somme pondérée de ses voisins. Cette pondération est donnée par le filtre qui est centré sur le pixel considéré. Prenons par exemple le filtre :
Nous allons appliquer ce filtre à chaque élément de notre image. Par exemple, pour le pixel de valeur 6 dans l'exemple ci-dessous, nous pouvons calculer sa nouvelle valeur :
Dans cet exercice, nous considérons que les filtres sont carrés. Leur taille peut donc être décrite avec un seul entier filterWidth. Écrivez le kernel gaussian_blur qui appliquera le filtre à un pixel donnée.
0.0 0.2 0.0
0.2 0.2 0.2
0.0 0.2 0.0
1 2 5 2 0 3
-------
3 |2 5 1| 6 0 0.0*2 + 0.2*5 + 0.0*1 +
| |
4 |3 6 2| 1 4 -> 0.2*3 + 0.2*6 + 0.2*2 + -> 3.2
| |
0 |4 0 3| 4 2 0.0*4 + 0.2*0 + 0.0*3
-------
9 6 5 0 3 9
(1) (2) (3)
Faites bien attention de ne pas sortir de l'image quand vous appliquez le filtre. Si jamais c'est le cas, on fera en sorte que toutes les coordonnées négatives soient considérées comme étant égale à 0 (donc on prendra la valeur sur le bord), et les coordonnées plus grandes que le nombre de colonnes ou lignes comme étant la dernière coordonnées sur le bord. Par exemple :
1 2 5 2 0 3
3 2 5 1 6 0 0.0*0 + 0.2*0 + 0.0*4 +
4 3 6 2 1 4 -> 0.2*9 + 0.2*9 + 0.2*6 + -> 6.6
-------
| 0 4| 0 3 4 2 0.0*9 + 0.2*9 + 0.0*6
| |
| 9 6| 5 0 3 9
(1) (2) (3)
__global__ void gaussian_blur(const unsigned char* const inputChannel,
unsigned char* const outputChannel,
int numRows, int numCols,
const float* const filter, const int filterWidth)
{
int pos_x = blockDim.x * blockIdx.x + threadIdx.x;
int pos_y = blockDim.y * blockIdx.y + threadIdx.y;
if (pos_x >= numCols || pos_y >= numRows) return;
float res = 0.0f;
for (int i = 0; i < filterWidth; i++){
int new_pos_x = pos_x + i - filterWidth / 2;
new_pos_x = min(max(new_pos_x, 0), numCols - 1);
for (int j=0; j < filterWidth; j++){
int new_pos_y = pos_y + j - filterWidth / 2;
new_pos_y = min(max(new_pos_y, 0), numRows - 1);
res += filter[j * filterWidth + i] * inputChannel[new_pos_y * numCols + new_pos_x];
}
}
outputChannel[pos_y * numCols + pos_x] = res;
}
Notre filtre sera défini au niveau du CPU. Pour le transférer au GPU, il faut d'abord allouer la mémoire dans le GPU puis y copier le filtre. Modifiez la fonction allocateMemoryAndCopyToGPU pour y allouer la mémoire nécessaire, puis copiez-y le filtre.
Entourez votre allocation et copie de la fonction checkCudaErrors(...) afin de voir les erreurs éventuelle (oui, le GPU ne nous dit pas qu'il y a une erreur sauf si on lui demande).
cudaMalloc prend un pointeur de pointeur en entrée.
void allocateMemoryAndCopyToGPU(const size_t numRowsImage, const size_t numColsImage,
const float* const h_filter, const size_t filterWidth)
{
checkCudaErrors(cudaMalloc(&d_red, sizeof(unsigned char) * numRowsImage * numColsImage));
checkCudaErrors(cudaMalloc(&d_green, sizeof(unsigned char) * numRowsImage * numColsImage));
checkCudaErrors(cudaMalloc(&d_blue, sizeof(unsigned char) * numRowsImage * numColsImage));
checkCudaErrors(cudaMalloc(&d_filter, sizeof(float) * filterWidth * filterWidth));
checkCudaErrors(cudaMemcpy(d_filter, h_filter, sizeof(float) * filterWidth * filterWidth, cudaMemcpyHostToDevice));
}
Toute mémoire allouée devant être libérée, modifiez la fonction cleanup() afin de libérer la mémoire du filtre sur le GPU.
void cleanup() {
checkCudaErrors(cudaFree(d_red));
checkCudaErrors(cudaFree(d_green));
checkCudaErrors(cudaFree(d_blue));
checkCudaErrors(cudaFree(d_filter));
}
Nous avons maintenant tous les blocs pour exécuter le code. Il ne nous reste plus qu'à définir la taille des blocs et la taille de la grille (similaire au TP précédent). Ensuite, nous voulons appeler les kernels dans l'ordre (séparation des canaux et application du filtre trois fois, une fois sur chaque canal). Modifier la fonction your_gaussian_blur pour appeler les deux kernels avec une taille de grille et de bloc convenable.
La fonction appelle un dernier kernel pour recombiner les trois canaux dans une seule image. Vous n'avez pas à implémenter ce kernel.
void your_gaussian_blur(const uchar4 * const h_inputImageRGBA, uchar4 * const d_inputImageRGBA,
uchar4* const d_outputImageRGBA, const size_t numRows, const size_t numCols,
unsigned char *d_redBlurred,
unsigned char *d_greenBlurred,
unsigned char *d_blueBlurred,
const int filterWidth)
{
int block_size = 32;
const dim3 blockSize(block_size, block_size);
const dim3 gridSize(numCols / block_size + 1, numRows / block_size + 1);
separateChannels<<<gridSize, blockSize>>>(d_inputImageRGBA,
numRows,
numCols,
d_red, d_green, d_blue);
cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
gaussian_blur<<<gridSize, blockSize>>>(d_red,
d_redBlurred,
numRows, numCols,
d_filter, filterWidth);
gaussian_blur<<<gridSize, blockSize>>>(d_green,
d_greenBlurred,
numRows, numCols,
d_filter, filterWidth);
gaussian_blur<<<gridSize, blockSize>>>(d_blue,
d_blueBlurred,
numRows, numCols,
d_filter, filterWidth);
cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
recombineChannels<<<gridSize, blockSize>>>(d_redBlurred,
d_greenBlurred,
d_blueBlurred,
d_outputImageRGBA,
numRows,
numCols);
cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}
Testez votre code en générant le Makefile (cmake .) puis en l'exécutant (make). Lancer le programme avec la commande ./HW2 cinque_terre_small.jpg.
