Bonjour,

J'ai conscience que je ne suis pas forcément dans la bonne sous catégorie pour parler de mon problème mais je vous serais reconnaissant pour l'assistance apportée. Je débute en cuda.  

J'ai un code en c que je dois mettre en cuda. Le voici :

#include <unistd.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <stdarg.h>
#include <png.h>
#include <math.h>

#define CHANNEL 4
#define BLUR 20

int * sizes;
int * pixel_val;

int i,x,y;
int width, height;

png_bytep* rows;
png_bytep* rows_target;
png_structp img_struct;
png_infop img_info;
png_byte color, depth;

void box_blur_h(png_bytep* scl, png_bytep* tcl,int r){
  double iarr = 1.0/(r+r+1);
  pixel_val = (int *)malloc(sizeof(int)*CHANNEL);
  //int i=0;
  for (int a = 0; a < height; a++) {
    png_bytep row = scl[a];
    png_bytep row_target = tcl[a];
    png_bytep fv = &(row[0]);
    png_bytep lv = &(row[(width-1)*CHANNEL]);
    int ri=r;
    int ti = 0;
    int li = 0;
    for(i=0;i<CHANNEL;i++)
      pixel_val[i] = (r+1)*fv[i];

    for (int x = 0;x < r;x++) {
      for(i=0;i<CHANNEL;i++){
        png_bytep ptr = &(row[x*4]);
        pixel_val[i] += ptr[i];
      }
    }

    for(x=0;x<=r;x++){
      png_bytep ptr = &(row[(ri++)*4]);
      png_bytep ptr_target = &(row_target[(ti++)*4]);
      for(i=0;i<CHANNEL;i++){
        pixel_val[i] += ptr[i]-fv[i];
        ptr_target[i] = round(pixel_val[i]*iarr);
      }
    }

    for(x=r+1;x<=width-r;x++){
      png_bytep ptr = &(row[(ri++)*4]);
      png_bytep ptr2 = &(row[(li++)*4]);
      png_bytep ptr_target = &(row_target[(ti++)*4]);
      for(i = 0;i<CHANNEL;i++){
        pixel_val[i] += ptr[i]-ptr2[i];
        ptr_target[i] = round(pixel_val[i]*iarr);
      }
    }

    for(x=width-r;x<width;x++){
      png_bytep ptr = &(row[(li++)*4]);
      png_bytep ptr_target = &(row_target[(ti++)*4]);
      for(i=0;i<CHANNEL;i++){
        pixel_val[i] += lv[i] - ptr[i];
        ptr_target[i] = round(pixel_val[i]*iarr);
      }
    }
  }
  free(pixel_val);
}

Pour plus ce détails sur ce que fait le code, il s'agit d'appliquer un filtre gaussien sur une image : https://github.com/Lescurel/gaussian_blur/blob/master/blur.c

L'objectif est maintenant de le transformer en cuda tout en le parallélisant pour réduire le temps d'exécution. J'ai donc fais ce code en cuda :

#include <unistd.h>
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <stdarg.h>
#include <png.h>
#include <math.h>


#define CHANNEL 4
#define BLUR 20

int * sizes;
int *pixel_val;

int i,x,y;
int width, height;

png_bytep* rows;
png_bytep* rows_target;
png_structp img_struct;
png_infop img_info;
png_byte color, depth;



void checkCudaError(cudaError_t error) {
    if (error != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(error));
        exit(-1);
    }
    else
      printf("CUDA error: %s\n", cudaGetErrorString(error));
}

__global__ void box_blur_h_kernel(png_bytep* scl, png_bytep* tcl, int width, int height, int r, int* pixel_val) {
    int i;
    double iarr = 1.0 / (r + r + 1);

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    if (x < width && y < height) {
        png_bytep row = scl[y];
        png_bytep row_target = tcl[y];

        png_bytep fv = &(row[0]);
        png_bytep lv = &(row[(width - 1) * CHANNEL]);

        int ri = r;
        int ti = 0;
        int li = 0;

        for (i = 0; i < CHANNEL; i++)
            pixel_val[i] = (r + 1) * fv[i];

        for (int i = 0; i < r; i++) {
            for (int j = 0; j < CHANNEL; j++) {
              png_bytep ptr = &(row[i * 4]);
              pixel_val[j] += ptr[j];
            }
        }

        for (int i = 0; i <= r; i++) {
            png_bytep ptr = &(row[(ri++) * 4]);
            png_bytep ptr_target = &(row_target[(ti++) * 4]);
            for (int j = 0; j < CHANNEL; j++) {
                pixel_val[j] += ptr[j] - fv[j];
                ptr_target[j] = round(pixel_val[j] * iarr);
            }
        }

        for (int i = r + 1; i <= width - r; i++) {
            png_bytep ptr = &(row[(ri++) * 4]);
            png_bytep ptr2 = &(row[(li++) * 4]);
            png_bytep ptr_target = &(row_target[(ti++) * 4]);
            for (int j = 0; j < CHANNEL; j++) {
                pixel_val[j] += ptr[j] - ptr2[j];
                ptr_target[j] = round(pixel_val[j] * iarr);
            }
        }

        for (int i = width - r; i < width; i++) {
            png_bytep ptr = &(row[(li++) * 4]);
            png_bytep ptr_target = &(row_target[(ti++) * 4]);
            for (int j = 0; j < CHANNEL; j++) {
                pixel_val[j] += lv[j] - ptr[j];
                ptr_target[j] = round(pixel_val[j] * iarr);
            }
        }
    }
}

 int d_width;
 int d_height;
 int d_r;

void box_blur_h(png_bytep* scl, png_bytep* tcl, int r) {
    // Allouer de la mémoire pour pixel_val sur le GPU
    int* d_pixel_val;
    cudaMalloc((void**)&d_pixel_val, CHANNEL * sizeof(int));

    png_bytep* d_scl;
    png_bytep* d_tcl;
    cudaMalloc((void**)&d_scl, sizeof(scl) * sizeof(png_bytep));
    cudaMalloc((void**)&d_tcl, sizeof(tcl) * sizeof(png_bytep));

    // Define block and grid dimensions
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x,
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y);

    // Transférer la valeur de pixel_val de CPU à GPU
    cudaMemcpy(d_pixel_val, pixel_val, CHANNEL * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_scl, scl, sizeof(scl) * sizeof(png_bytep), cudaMemcpyHostToDevice);
    cudaMemcpy(d_tcl, tcl, sizeof(tcl) * sizeof(png_bytep), cudaMemcpyHostToDevice);

    // Transférer la largeur et la hauteur de l'image du CPU au GPU
    cudaMemcpyToSymbol(d_width, &width, sizeof(int));
    cudaMemcpyToSymbol(d_height, &height, sizeof(int));
    cudaMemcpyToSymbol(d_r, &r, sizeof(int));

    // Appeler le noyau CUDA avec pixel_val comme argument
    box_blur_h_kernel<<<numBlocks, threadsPerBlock>>>(d_scl, d_tcl, d_width, d_height, d_r, d_pixel_val);
        cudaDeviceSynchronize();


    cudaMemcpy(pixel_val, d_pixel_val, CHANNEL * sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(scl, d_scl, sizeof(scl) * sizeof(png_bytep), cudaMemcpyDeviceToHost);
    cudaMemcpy(tcl, d_tcl, sizeof(tcl) * sizeof(png_bytep), cudaMemcpyDeviceToHost);
    
    // Synchroniser le GPU et le CPU
    cudaDeviceSynchronize();

    

    // Libérer la mémoire allouée sur le GPU
    cudaFree(d_pixel_val);
    cudaFree(d_scl);
    cudaFree(d_tcl);
    /*cudaFree(d_width);
    cudaFree(d_height);
    cudaFree(d_r);*/
}

Le problème c'est que lorsque je remplace le bout de code C par ce que j'ai fait dans mon fichier cuda en faisant les ajustements nécessaires, mon imageblur en sortie ne contient que 3ko et rien ne s'affiche quand j'ouvre le fichier alors que quand ça marche avec le code C, j'ai bien l'image floutée et une taille de 222ko.

Auriez vous des pistes sur ce pourquoi j'ai ce problème ? Aussi, mon code est il bien parallélisé ?

Merci à vous,

-
Edité par Zekeee 13 mars 2024 à 16:05:47

J'ai remarqué que mon problème venait de l'allocation mémoire pour les types de données complexes (structure, tableau 3D...). *

Pour ceux qui rencontreront plus tard le même problème que moi concernant l'allocation mémoire, il est nécessaire de réussir à placer l'ensemble de vos données dans un tableau 1D. Les quelques recherches et réponses sur lesquelles je suis tombé m'ont fait comprendre qu'il n'est pas possible de travailler directement sur le GPU avec des types de données complexes. Dans mon cas, j'ai converti l'ensemble des données de ma structure png_bytep en un tableau de données de type primitif, ici "byte" qui est un "unsigned char". Cela a donné : png_bytep* rows; --> byte **rows; et j'ai stocké mes données dedans. J'espère que vous pourrez avancer avec ces explications. 

//Je suis passé de ça 
png_bytep* rows;

//à ça
typedef unsigned char byte;

byte rows;

//Il est tout de même possible d'écrire directement 
unsigned char rows;

-
Edité par Zekeee 16 mars 2024 à 13:02:37