Bonjour,
Je m’essaye à CUDA et au calcul sur GPU. J’avoue avoir un peu de mal encore avec ces histoires de blocs et de grille. Je me suis donc mis à faire l’exercice suivant : pour chaque élément d’une matrice, calculer la moyenne des voisins (en prenant en compte l’élément) et sortir donc la matrice des moyennes, c’est une sorte de flou je crois.
Pour le moment voici mon code :
#include <cuda.h>
#include <math.h>
#include <stdio.h>
#define MATRIX_ROWS (4)
#define MATRIX_COLS (4)
#define CU_THREADS_PER_BLOCK_X (16)
#define CU_THREADS_PER_BLOCK_Y (16)
__global__ void kernel_blur_matrix(float *input, float *output, int rows, int cols)
{
int tid = (threadIdx.x + threadIdx.y * blockDim.x) + (blockDim.x * blockDim.y) * blockIdx.x;
int size = rows * cols;
if (tid < size)
{
// Compute mean
int count = 1;
float mean = input[tid];
if ((tid - rows - 1) > 0 && (tid - rows - 1) < size)
{
mean += input[tid - rows - 1];
count++;
}
if ((tid - rows) > 0 && (tid - rows) < size)
{
mean += input[tid - rows];
count++;
}
if ((tid - rows + 1) > 0 && (tid - rows + 1) < size)
{
mean += input[tid - rows + 1];
count++;
}
if ((tid - 1) > 0 && (tid - 1) < size)
{
mean += input[tid - 1];
count++;
}
if ((tid + 1) > 0 && (tid + 1) < size)
{
mean += input[tid + 1];
count++;
}
if ((tid + rows - 1) > 0 && (tid + rows - 1) < size)
{
mean += input[tid + rows - 1];
count++;
}
if ((tid + rows) > 0 && (tid + rows) < size)
{
mean += input[tid + rows];
count++;
}
if ((tid + rows + 1) > 0 && (tid + rows + 1) < size)
{
mean += input[tid + rows + 1];
count++;
}
mean /= count;
output[tid] = mean;
}
}
int main()
{
// Host initializations
unsigned int size = MATRIX_COLS * MATRIX_ROWS * sizeof(float);
float *input = (float*)malloc(size);
float *output = (float*)malloc(size);
// a faire directement sur GPU
for (int i = 0; i < MATRIX_COLS * MATRIX_ROWS; ++i)
{
input[i] = i;
}
// Device initializations
float *d_input, *d_output;
cudaMalloc(&d_input, size);
cudaMalloc(&d_output, size);
cudaMemcpy(d_input, input, size, cudaMemcpyHostToDevice);
dim3 grid_dim(ceil((float)(MATRIX_COLS * MATRIX_ROWS) / (CU_THREADS_PER_BLOCK_X * CU_THREADS_PER_BLOCK_Y)), 1, 1);
dim3 block_dim(CU_THREADS_PER_BLOCK_X, CU_THREADS_PER_BLOCK_Y, 1);
kernel_blur_matrix<<<grid_dim, block_dim>>>(d_input, d_output, MATRIX_ROWS, MATRIX_COLS);
cudaDeviceSynchronize();
cudaMemcpy(output, d_output, size, cudaMemcpyDeviceToHost);
for (int i = 0; i < MATRIX_ROWS; ++i)
{
for (int j = 0; j < MATRIX_COLS; ++j)
printf("%8.2f ", input[j + i * MATRIX_ROWS]);
printf("\n");
}
printf("\n\n");
for (int i = 0; i < MATRIX_ROWS; ++i)
{
for (int j = 0; j < MATRIX_COLS; ++j)
printf("%8.2f ", output[j + i * MATRIX_ROWS]);
printf("\n");
}
cudaFree(d_input);
cudaFree(d_output);
free(input);
free(output);
return 0;
}
Première chose, ça ne fonctionne pas. Par exemple, sur cette entrée :
0.00 1.00 2.00 3.00
4.00 5.00 6.00 7.00
8.00 9.00 10.00 11.00
12.00 13.00 14.00 15.00
Le programme me sort:
2.60 3.60 4.00 5.00
5.29 5.62 6.00 7.00
8.00 9.00 10.00 10.38
10.71 11.00 12.00 12.40
La dernière ligne devrait être :
10.50 11.00 12.00 12.50
Les résultats sont proches mais incorrects. Est-ce du à une la gestion des flottants ? Ça m’étonnerai mais bon.
Deuxième chose, j’ai l’intuition que je peux remplacer ma suite de if par deux boucles for imbriqués, mais après avoir essayé pendant deux heures, j’ai abandonné. Est-ce possible ?
Merci pour votre aide. 
