Work with Halo boundary conditions / areas in CUDA

Question

I am working on image processing with CUDA, and I doubt the processing of pixels.

What is often done with the border pixels of an image when applying a convolution filter m x m?

At the core of the convolution 3 x 3, ignoring the pixel border of an 1image is easier to handle, especially when the code is improved using shared memory. In fact, in this case it is not necessary to check whether a given pixel has all the available neighborhoods (i.e., the Pixel is (0, 0)not left at the coordinate , upper left, upper neighbors). However, removing the pixel border of the 1original image may lead to partial results.

On the contrary, I would like to process all the pixels in the image, also when using shared memory improvements, that is, for example, loading pixels 16 x 16, but calculating the internal one 14 x 14. Also in this case, ignoring the boundary pixels generates a clearer code.

What is usually done in this case?

Does anyone usually use my approach ignoring border pixels?

Of course, I know that the answer depends on the type of problem, i.e. adding two images per pixel does not have this problem.

Thanks in advance.

Answer 1

. :

• (, )
• / ,
• (, [-1] = [1], [-2] = [2])
• (, [-1] = [-1], [-2] = [-2])

Answer 2

tl; dr: , , - , . , , , , , , .

( )

, , . , , ( "box" ). , ceil (n/2) ( ) . , "" ( , RGB - !). , , - . , . , , , , .

, . , , , "" . , / . , . , w=(1/(n*n)). , K 3 . , :

K*w + K*w + K*w  = K*3*w

, . : , . , ( : ..).

, , - . , "". , , ? - , , . , , . , , , , , , . : , , .

, , ? , . , , ... , ​​ , .

. , . / , . , . , . "", . "area" , , , . "area" ( ). ( : !) . , , , , , , , . , / , ab-initio/ , , "" .

Answer 3

, , :

• ;
• .

( ) , , 3 x 3. , .

( ) , , 16 x 16 , 14 x 14 .

, , : CUDA.

, , . CUDA .

, .

#include <stdio.h>

#include "TimingGPU.cuh"
#include "Utilities.cuh"

texture<float, 1, cudaReadModeElementType> signal_texture;

#define BLOCKSIZE 32

/*************************************************/
/* KERNEL FUNCTION FOR MEDIAN FILTER CALCULATION */
/*************************************************/
__global__ void median_filter_periodic_boundary(float * __restrict__ d_vec, const unsigned int N){

    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

    if (tid < N) {

        float signal_center = tex1D(signal_texture, tid - 0);
        float signal_before = tex1D(signal_texture, tid - 1);
        float signal_after  = tex1D(signal_texture, tid + 1);

        printf("%i %f %f %f\n", tid, signal_before, signal_center, signal_after);

        d_vec[tid] = (signal_center + signal_before + signal_after) / 3.f;

    }
}


/********/
/* MAIN */
/********/
int main() {

    const int N = 10;

    // --- Input host array declaration and initialization
    float *h_arr = (float *)malloc(N * sizeof(float));
    for (int i = 0; i < N; i++) h_arr[i] = (float)i;

    // --- Output host and device array vectors
    float *h_vec = (float *)malloc(N * sizeof(float));
    float *d_vec;   gpuErrchk(cudaMalloc(&d_vec, N * sizeof(float)));

    // --- CUDA array declaration and texture memory binding; CUDA array initialization
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    //Alternatively
    //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    cudaArray *d_arr;   gpuErrchk(cudaMallocArray(&d_arr, &channelDesc, N, 1));
    gpuErrchk(cudaMemcpyToArray(d_arr, 0, 0, h_arr, N * sizeof(float), cudaMemcpyHostToDevice));

    cudaBindTextureToArray(signal_texture, d_arr); 
    signal_texture.normalized = false; 
    signal_texture.addressMode[0] = cudaAddressModeWrap;

    // --- Kernel execution
    median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_vec, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_vec, d_vec, N * sizeof(float), cudaMemcpyDeviceToHost));

    for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);

    printf("Test finished\n");

    return 0;
}