Filtro medio de imagen Cuda

Contamos con la mejor solución que hallamos en todo internet. Nosotros deseamos que te sirva de ayuda y si deseas comentarnos cualquier detalle que nos pueda ayudar a perfeccionar nuestra información hazlo con total libertad.

Solución:

Este es un caso clásico de un problema de procesamiento de imágenes vergonzosamente paralelo que se puede asignar muy fácilmente al marco CUDA. El filtro de promedio se conoce como filtro de caja en los dominios de procesamiento de imágenes.

El enfoque más sencillo sería utilizar texturas CUDA para el proceso de filtrado, ya que las condiciones de contorno pueden manejarse muy fácilmente mediante texturas.

Suponiendo que tiene punteros de origen y destino asignados en el host. El procedimiento sería algo así.

1. Asigne memoria lo suficientemente grande para almacenar las imágenes de origen y destino en el dispositivo.
2. Copie la imagen de origen del host al dispositivo.
3. Vincula el puntero del dispositivo de la imagen de origen a la textura.
4. Especifique un tamaño de bloque apropiado y una cuadrícula lo suficientemente grande para cubrir cada píxel de la imagen.
5. Inicie el kernel de filtrado utilizando la cuadrícula y el tamaño de bloque especificados.
6. Copie los resultados al anfitrión.
7. Desenlazar la textura
8. Punteros de dispositivo gratuitos.

Implementación de muestra de filtro de caja

Núcleo

texture tex8u;

//Box Filter Kernel For Gray scale image with 8bit depth
__global__ void box_filter_kernel_8u_c1(unsigned char* output,const int width, const int height, const size_t pitch, const int fWidth, const int fHeight)

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

    const int filter_offset_x = fWidth/2;
    const int filter_offset_y = fHeight/2;

    float output_value = 0.0f;

    //Make sure the current thread is inside the image bounds
    if(xIndex(output_value);
    

Función de envoltura:

void box_filter_8u_c1(unsigned char* CPUinput, unsigned char* CPUoutput, const int width, const int height, const int widthStep, const int filterWidth, const int filterHeight)


    /*
     * 2D memory is allocated as strided linear memory on GPU.
     * The terminologies "Pitch", "WidthStep", and "Stride" are exactly the same thing.
     * It is the size of a row in bytes.
     * It is not necessary that width = widthStep.
     * Total bytes occupied by the image = widthStep x height.
     */

    //Declare GPU pointer
    unsigned char *GPU_input, *GPU_output;

    //Allocate 2D memory on GPU. Also known as Pitch Linear Memory
    size_t gpu_image_pitch = 0;
    cudaMallocPitch(&GPU_input,&gpu_image_pitch,width,height);
    cudaMallocPitch(&GPU_output,&gpu_image_pitch,width,height);

    //Copy data from host to device.
    cudaMemcpy2D(GPU_input,gpu_image_pitch,CPUinput,widthStep,width,height,cudaMemcpyHostToDevice);

    //Bind the image to the texture. Now the kernel will read the input image through the texture cache.
    //Use tex2D function to read the image
    cudaBindTexture2D(NULL,tex8u,GPU_input,width,height,gpu_image_pitch);

    /*
     * Set the behavior of tex2D for out-of-range image reads.
     * cudaAddressModeBorder = Read Zero
     * cudaAddressModeClamp  = Read the nearest border pixel
     * We can skip this step. The default mode is Clamp.
     */
    tex8u.addressMode[0] = tex8u.addressMode[1] = cudaAddressModeBorder;

    /*
     * Specify a block size. 256 threads per block are sufficient.
     * It can be increased, but keep in mind the limitations of the GPU.
     * Older GPUs allow maximum 512 threads per block.
     * Current GPUs allow maximum 1024 threads per block
     */

    dim3 block_size(16,16);

    /*
     * Specify the grid size for the GPU.
     * Make it generalized, so that the size of grid changes according to the input image size
     */

    dim3 grid_size;
    grid_size.x = (width + block_size.x - 1)/block_size.x;  /*< Greater than or equal to image width */
    grid_size.y = (height + block_size.y - 1)/block_size.y; /*< Greater than or equal to image height */

    //Launch the kernel
    box_filter_kernel_8u_c1<<>>(GPU_output,width,height,gpu_image_pitch,filterWidth,filterHeight);

    //Copy the results back to CPU
    cudaMemcpy2D(CPUoutput,widthStep,GPU_output,gpu_image_pitch,width,height,cudaMemcpyDeviceToHost);

    //Release the texture
    cudaUnbindTexture(tex8u);

    //Free GPU memory
    cudaFree(GPU_input);
    cudaFree(GPU_output);

La buena noticia es que no es necesario que implemente el filtro usted mismo. El kit de herramientas CUDA viene con una biblioteca gratuita de procesamiento de imágenes y señales llamada NVIDIA Performance Primitives, también conocida como NPP, fabricada por NVIDIA. NPP utiliza GPU habilitadas para CUDA para acelerar el procesamiento. El filtro de promediado ya está implementado en NPP. La versión actual de NPP (5.0) admite imágenes de 8 bits, 1 canal y 4 canales. Las funciones son:

• nppiFilterBox_8u_C1R para imagen de 1 canal.
• nppiFilterBox_8u_C4R para imagen de 4 canales.

Algunos pensamientos / pasos básicos:

1. Copie los datos de la imagen de la CPU a la GPU
2. Llame a un kernel para generar el promedio de cada línea (horizontal) y almacenarlo en la memoria compartida.
3. Llame a un kernel para generar el promedio de cada columna (vertical) y almacenarlo en la memoria global.
4. Copie los datos de nuevo a la memoria de la CPU.

Debería poder escalar esto con bastante facilidad con memoria 2D y llamadas de kernel multidimensionales.

Si el tamaño del filtro es normal y no enorme, el filtro promedio es un caso muy bueno para implementar con CUDA. Puede configurar esto usando bloques cuadrados y cada hilo del bloque es responsable del cálculo del valor de un píxel, sumando y promediando sus vecinos.

Si almacena la imagen en la memoria global, puede programarla fácilmente. Una posible optimización es que cargue bloques de la imagen en la memoria compartida del bloque. Usando elementos fantasmas (para que no exceda las dimensiones del bloque compartido al buscar píxeles vecinos) puede calcular el promedio de los píxeles dentro de un bloque.

Lo único que debe tener en cuenta es cómo se hará la “costura” al final, porque los bloques de memoria compartida se superpondrán (debido a los píxeles de “relleno” adicionales) y no desea calcular sus valores dos veces.

Comentarios y puntuaciones

Si te gustó nuestro trabajo, tienes el poder dejar un enunciado acerca de qué le añadirías a este enunciado.

---

---

Utiliza Nuestro Buscador