Filtro medio Cuda Image

Question

Il filtro medio è un filtro windowed di classe lineare, che uniforma il segnale (immagine). Il filtro funziona come uno passa-basso. L'idea alla base del filtro è che ogni elemento del segnale (immagine) abbia una media sul suo vicinato.

---

se abbiamo una matrice m x n e vogliamo applicare il filtro media con dimensioni k su di esso, quindi per ogni punto nella matrice p:(i,j) il valore del punto sarebbe la media di tutti i punti in piazza

Questa cifra è per il kernel piazza di filtraggio con il formato 2, che la casella gialla è il pixel da media, e tutta la griglia è il quadrato di pixel vicini, il nuovo valore del pixel sarà la media di loro.

Il problema è che questo algoritmo è molto lento, specialmente su immagini di grandi dimensioni, quindi ho pensato di usare GPGPU.

La domanda ora è, Come può essere implementato in cuda, se è possibile?

Answer 1

Questo è un caso classico di problema di elaborazione delle immagini embarrassingly parallel che può essere facilmente mappato al framework CUDA. Il filtro di calcolo della media è noto come Box Filter nei domini di elaborazione delle immagini.

L'approccio più semplice sarebbe utilizzare le trame CUDA per il processo di filtraggio in quanto le condizioni al contorno possono essere gestite molto facilmente dalle trame.

Supponendo di disporre di puntatori di origine e destinazione allocati sull'host. La procedura sarebbe simile a questa.

1. Assegna memoria sufficiente per contenere le immagini di origine e di destinazione sul dispositivo.
2. Copia l'immagine di origine dall'host al dispositivo.
3. Associare il puntatore del dispositivo di immagine di origine alla trama.
4. Specificare una dimensione di blocco appropriata e una griglia sufficientemente grande da coprire ogni pixel dell'immagine.
5. Avvia il kernel di filtraggio utilizzando la griglia e le dimensioni del blocco specificate.
6. Copia i risultati di nuovo su host.
7. Annulla la trama
8. Puntatori di dispositivo gratuiti.

esempio di implementazione scatola filtro

Kernel

texture<unsigned char, cudaTextureType2D> 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<width && yIndex<height) 
    { 
     //Sum the window pixels 
     for(int i= -filter_offset_x; i<=filter_offset_x; i++) 
     { 
      for(int j=-filter_offset_y; j<=filter_offset_y; j++) 
      { 
       //No need to worry about Out-Of-Range access. tex2D automatically handles it. 
       output_value += tex2D(tex8u,xIndex + i,yIndex + j); 
      } 
     } 

     //Average the output value 
     output_value /= (fWidth * fHeight); 

     //Write the averaged value to the output. 
     //Transform 2D index to 1D index, because image is actually in linear memory 
     int index = yIndex * pitch + xIndex; 

     output[index] = static_cast<unsigned char>(output_value); 
    } 
} 

Wrapper Funzione:

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<unsigned char>(&GPU_input,&gpu_image_pitch,width,height); 
    cudaMallocPitch<unsigned char>(&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<<<grid_size,block_size>>>(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 buona notizia è che non c'è bisogno di attuare il filtrati. Il toolkit CUDA viene fornito con una libreria di elaborazione di segnali e immagini gratuita denominata NVIDIA Performance Primitives, ovvero NPP, creata da NVIDIA. NPP utilizza GPU abilitate CUDA per accelerare l'elaborazione. Il filtro di calcolo della media è già implementato in NPP. La versione corrente di NPP (5.0) supporta le immagini a 8 bit, 1 canale e 4 canali. Le funzioni sono: immagine del canale

• nppiFilterBox_8u_C1R per 1.
• nppiFilterBox_8u_C4R per l'immagine a 4 canali.

Answer 2

alcuni pensieri di base/fasi:

1. copiare i dati di immagine dalla CPU alla GPU
2. chiamata un kernel per costruire la media per ogni riga (orizzontale) e conservarla in memoria condivisa.
3. Chiama un kernel per creare la media per ogni colonna (verticale) e memorizzarla nella memoria globale.
4. Copia i dati nella memoria della CPU.

Si dovrebbe essere in grado di ridimensionare questo abbastanza facilmente con chiamate di memoria 2D e multidimensionali del kernel.

Answer 3

Se la dimensione del filtro è normale e non eccessiva, il filtro medio è un ottimo esempio per l'implementazione con CUDA. Puoi impostarlo usando blocchi quadrati e ogni thread del blocco è responsabile del calcolo del valore di un pixel, sommando e facendo la media dei suoi vicini.

Se si memorizza l'immagine nella memoria globale, può essere programmata facilmente, ma si verificano numerosi conflitti bancari. Una possibile ottimizzazione consiste nel caricare i blocchi dell'immagine nella memoria condivisa del blocco. Utilizzando gli elementi fantasma (in modo da non superare le dimensioni del blocco condiviso quando si cercano i pixel vicini) è possibile calcolare la media dei pixel all'interno di un blocco.

L'unico pensiero che bisogna fare attenzione è come verrà eseguita la "cucitura" alla fine, perché i blocchi di memoria condivisa si sovrappongono (a causa dei pixel di "riempimento" extra) e non si desidera per calcolare i loro valori due volte.