使用 CUDA 和 libpng 进行图像处理

Question

我一直在尝试使用 CUDA 和 libpng 库处理图像（即应用黑白滤镜）。但是，我不确定如何将图像数据传递给设备。

据我了解，libpng 库读取 png 并将信息存储在一个 png_bytep 结构中，使用此函数逐行调用 row_pointers。

    void read_png_file(char* file_name){
        char header[8];    // 8 is the maximum size that can be checked

        FILE *fp = fopen(file_name, "rb");
        if (!fp)
                abort_("[read_png_file] File %s could not be opened for reading" file_name);
        fread(header, 1, 8, fp);
        // if (png_sig_cmp(header, 0, 8))
        //         abort_("[read_png_file] File %s is not recognized as a PNG file", file_name);


        //Inicializa variables necesarias para libpng
        png_ptr =   png_create_read_struct (PNG_LIBPNG_VER_STRING, NULL, NULL, NULL);

        if (!png_ptr)
                abort_("[read_png_file] png_create_read_struct failed");

        info_ptr = png_create_info_struct(png_ptr);
        if (!info_ptr)
                abort_("[read_png_file] png_create_info_struct failed");

        if (setjmp(png_jmpbuf(png_ptr)))
                abort_("[read_png_file] Error during init_io");

        //Inicializa el input/output para el archivo PNG
        png_init_io(png_ptr, fp);
        png_set_sig_bytes(png_ptr, 8);

        //Lee la información anterior a los datos de los píxeles como tal
        png_read_info(png_ptr, info_ptr);

        //Almacena información del archivo PNG 
        width = png_get_image_width(png_ptr, info_ptr);
        height = png_get_image_height(png_ptr, info_ptr);
        color_type = png_get_color_type(png_ptr, info_ptr);
        bit_depth = png_get_bit_depth(png_ptr, info_ptr);

        number_of_passes = png_set_interlace_handling(png_ptr);
        png_read_update_info(png_ptr, info_ptr);


        // Lectura del archivo PNG
        if (setjmp(png_jmpbuf(png_ptr)))
                abort_("[read_png_file] Error during read_image");

        // Reserva el espacio necesario para almacenar los datos del archivo PNG por filas
        row_pointers = (png_bytep*) malloc(sizeof(png_bytep) * height);
        for (y=0; y<height; y++)
                row_pointers[y] = (png_byte*) malloc(png_get_rowbytes(png_ptr,info_ptr));

        // Y para la copia para el device
        d_row_pointers = (png_bytep*) malloc(sizeof(png_bytep) * height);
        for (y=0; y<height; y++)
                d_row_pointers[y] = (png_byte*) malloc(png_get_rowbytes(png_ptr,info_ptr));


        png_read_image(png_ptr, row_pointers);
        
        fclose(fp);
}

我一直在尝试传递存储在 row_pointers 中的信息，创建一个名为 d_row_pointers 的副本并使用 cudaMalloc 和 cudaMemcpy 函数，如下所示：

png_bytep * row_pointers;
png_bytep * d_row_pointers;

int main(int argc, char **argv)
{       
        
        // Verifica los parámetros para ejecutar el programa
        if (argc != 3)
                abort_("Uso: ./Nombre_del_Programa <file_in> <file_out>");

        read_png_file(argv[1]);

        // CUDA
        int size = sizeof(png_bytep);
        int int_size = sizeof(int);
        
        cudaMalloc((void **)&d_row_pointers, size);
        for (y=0; y<height; y++)
                 cudaMalloc((void **)&d_row_pointers[y],png_get_rowbytes(png_ptr,info_ptr));

        cudaMalloc((void **)&d_width, int_size);
        cudaMalloc((void **)&d_height, int_size);


        cudaMemcpy(d_row_pointers, &row_pointers, size, cudaMemcpyHostToDevice);
        for (y=0; y<height; y++)
                 cudaMemcpy(d_row_pointers[y], &row_pointers[y], png_get_rowbytes(png_ptr,info_ptr), cudaMemcpyHostToDevice);
        cudaMemcpy(d_width, &width, int_size, cudaMemcpyHostToDevice);
        cudaMemcpy(d_height, &height, int_size, cudaMemcpyHostToDevice);

        // Lanzar el kernel
        process_file<<<1,1>>>(d_row_pointers, d_width, d_height);

        // Copiar los resultados de vuelta al host
        cudaMemcpy(&row_pointers, d_row_pointers, size, cudaMemcpyDeviceToHost);

        // Limpieza
        for (y=0; y<height; y++)
                cudaFree(d_row_pointers[y]);
        cudaFree(d_row_pointers);
        cudaFree(d_width);
        cudaFree(d_height);

        // Escritura de la imagen con los resultados
        write_png_file(argv[2]);

        return 0;
}

这是应该应用过滤器的 process_file 函数

__global__ void process_file(png_bytep * d_row_pointers, int * d_width, int * d_height)
{
        // Se realizan los cambios deseados en la imagen
        
        //Verificar los datos recibidos
        printf("Width = %d , Height = %d ", *d_width, *d_height);

        int rgb_total = 0;  
        float rgb_average = 0.0;
        int x = 0;
        int y = 0;
        png_byte *row;
        png_byte *ptr;

        for (y=0; y<*d_height-1; y++) {

                for (x=0; x<*d_width; x++) {
                
                        rgb_total = 0;
                        rgb_average = 0;

                        row             = d_row_pointers[y];
                        ptr             = &(row[x*3]);
                        
                        printf("Pixel  %d - %d, Rgb values: %d - %d - %d \n", x, y, ptr[0], ptr[1], ptr[2]); 
                        rgb_total      += ptr[0] + ptr[1] + ptr[2];
                        
                        // Calculando el promedios RGB
                        rgb_average = rgb_total / 3;
                        // printf("Average: %d \n", (int)rgb_average);
                        
                        ptr[0]  = (int)rgb_average;
                        ptr[1]  = (int)rgb_average;
                        ptr[2]  = (int)rgb_average;
                        
                        // printf("Changed to  %d - %d - %d \n",ptr[0], ptr[1], ptr[2]); 
                        // printf("Pixel  %d - %d done\n",x,y);  
                }

        }
        // printf("Para la imagen de resolución: %d x %d - ", *width, *height);
        
}

但我无法让它工作，即使内核似乎正在接收结构，当我访问数据时它全为零，当它应该是 RGB 值时。我非常感谢有关将此数据传递给内核的正确方法的任何帮助。谢谢！

PD：可以找到完整代码here

Answer 1

我很确定是这一行引起了问题：

cudaMemcpy(d_row_pointers[y], &row_pointers[y], 
    png_get_rowbytes(png_ptr,info_ptr), cudaMemcpyHostToDevice);

您无法通过 cpu 代码访问 d_row_pointers[y] 处的数据。 cudaMalloc 需要一个指向 cpu 上的指针的指针，并将这样对待参数。

您需要小心处理 CUDA 中的多维数组。大多数人选择使用平面阵列，即使是图像。

这就是我的建议：将数据复制到平面（一维）缓冲区，然后将其复制到设备。这样就更难犯错了。

啊，还有一件事：您不需要为内核显式指定 cudaMalloc 和 cudaMemcpy 的宽度和高度参数。您可以按值传递它们，就好像它是一个普通函数一样。您只需要数组的那些函数。

使用 CUDA 和 libpng 进行图像处理

Image processing with CUDA and libpng

c

cuda

image-processing

libpng