cuda multiple image erosion not work

Question

i'm trying to implement multiple black(0) and white(255) image erosion with cuda,i use a square (5x5)structure element.The kernel that i had implemented take an unsigned char array buffer in which are stored nImg images 200X200 px . To allow erosion of multiple image simultaneosly i make a grid with 3D structure:

• each block has the dimension of the strel (5x5)
• the grid has height = image_height/blockDim.y , width = image_width/blockDim.x , z = nImg

i've try to implement it extending that sample.

the problem is that if i store the pixels that a block of threads consider into a shared buffer shared between the threads of the block; to allow fast memory access, the algorithm doesn't work properly.I try to change the bindex that for me make mistake,but i cannot found a solution.

any suggestion?

here's my code:

//strel size
#define STREL_W 5
#define STREL_H 5

// distance from the cente of strel to strel width or height
#define R (STREL_H/2)

//size of the 2D region that each block consider i.e all the  neighborns that each thread in a block consider
#define BLOCK_W (STREL_W+(2*R))
#define BLOCK_H (STREL_H+(2*R))

__global__ void erode_multiple_img_SM(unsigned char * buffer_in,
                                       unsigned char * buffer_out,
                                       int w,
                                       int h ){

    //array stored in shared memory,that contain all pixel neighborns that each thread in a block consider 
    __shared__ unsigned char fast_acc_arr[BLOCK_W*BLOCK_H];

    // map thread in a 3D structure 
    int col = blockIdx.x * STREL_W + threadIdx.x -R ;
    int row = blockIdx.y * STREL_H + threadIdx.y -R ;
    int plane = blockIdx.z * blockDim.z + threadIdx.z;

    // check if a foreground px of strel is not contain in a region of the image with size of strel (if only one px is not contain the image is eroded)
    bool is_contain = true;


    // clamp to edge of image
    col = max(0,col);
    col = min(col,w-1);

    row = max(0,row);
    row = min(row,h-1);

    //map each thread in one dim coord to map 3D structure(grid) with image buffer(1D) 
    unsigned int index =  (plane * h * w) + (row * w) + col;

    unsigned int bindex =  threadIdx.y * blockDim.y + threadIdx.x;


    //each thread copy its pixel of the block to shared memory (shared with thread of a block)
    fast_acc_arr[bindex] = buffer_in[index];

    __syncthreads();



    //the strel must be contain in image, thread.x and thread.y are the coords of the center of the mask that correspond to strel in image, and it must be contain in image
    if((threadIdx.x >= R) && (threadIdx.x < BLOCK_W-R) && (threadIdx.y >= R) && (threadIdx.y <BLOCK_H-R)){

        for(int dy=-R; dy<=R; dy++){
            if(is_contain == false)
            break;

            for (int dx = -R ; dx <= R; dx++) {

                //if only one element in mask is different from the value of strel el --> the strel is not contain in the mask --> the center of the mask is eroded (and it's no necessary to consider the other el of the mask this is the motivation of the break)
                if (fast_acc_arr[bindex + (dy * blockDim.x) + dx ] != 255 ){

                    buffer_out[index ] = 0;
                    is_contain = false;
                    break;
                 }
             }
         }
        // if the strel is contain into the image the the center is not eroded
        if(is_contain == true)
            buffer_out[index] = 255;
      }

}

that are my kernel settings:

 dim3 block(5,5,1);
 dim3 grid(200/(block.x),200/(block.y),nImg);

my kernel call:

 erode_multiple_img_SM<<<grid,block>>>(dimage_src,dimage_dst,200,200);

my image input and output:

input: output(150 buff element):

code without shared memory(low speed):

__global__ void erode_multiple_img(unsigned char * buffer_in,
                            unsigned char * buffer_out,
                            int w,int h ){



    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int plane = blockIdx.z * blockDim.z +threadIdx.z;

    bool is_contain = true;
    col = max(0,col);
    col = min(col,w-1);

    row = max(0,row);
    row = min(row,h-1);



    for(int dy=-STREL_H/2; dy<=STREL_H/2; dy++){
        if(is_contain == false)
            break;
        for (int dx = -STREL_W/2 ; dx <= STREL_W/2; dx++) {
            if (buffer_in[(plane * h * w) +( row + dy) * w + (col + dx) ] !=255 ){
                buffer_out[(plane * h * w) + row * w + col ] = 0;
                is_contain = false;
                break;
            }
        }
    }
    if(is_contain == true)
        buffer_out[(plane * h * w) + row * w +col ] = 255;

}

UPDATED ALGORITHM

i try to follow that samples to do convolution.I change the input image, now has 512x512 size and i wrote that algorithm:

#define STREL_SIZE 5


#define TILE_W 16
#define TILE_H 16

#define R (STREL_H/2)

#define BLOCK_W (TILE_W+(2*R))
#define BLOCK_H (TILE_H+(2*R))

__global__ void erode_multiple_img_SM_v2(unsigned char * buffer_in,
                            unsigned char * buffer_out,
                            int w,int h ){

    // Data cache: threadIdx.x , threadIdx.y
    __shared__ unsigned char data[TILE_W +STREL_SIZE ][TILE_W +STREL_SIZE ];

    // global mem address of this thread
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    int plane = blockIdx.z * blockDim.z +threadIdx.z;

    int gLoc = (plane*h/w)+ row*w +col;

    bool is_contain = true;

    // load cache (32x32 shared memory, 16x16 threads blocks)
    // each threads loads four values from global memory into shared mem
    int x, y;   // image based coordinate



    if((col<w)&&(row<h)) {
        data[threadIdx.x][threadIdx.y]=buffer_in[gLoc];

    if (threadIdx.y > (h-STREL_SIZE))
        data[threadIdx.x][threadIdx.y + STREL_SIZE]=buffer_in[gLoc + STREL_SIZE];

    if (threadIdx.x >(w-STREL_SIZE))
        data[threadIdx.x + STREL_SIZE][threadIdx.y]=buffer_in[gLoc+STREL_SIZE];

    if ((threadIdx.x >(w-STREL_SIZE)) && (threadIdx.y > (h-STREL_SIZE)))
        data[threadIdx.x+STREL_SIZE][threadIdx.y+STREL_SIZE] =   buffer_in[gLoc+2*STREL_SIZE];
    //wait for all threads to finish read
     __syncthreads();


     //buffer_out[gLoc] = data[threadIdx.x][threadIdx.y];

      unsigned char min_value = 255;
      for(x=0;x<STREL_SIZE;x++){
          for(y=0;y<STREL_SIZE;y++){
              min_value = min( (data[threadIdx.x+x][threadIdx.y+y]) , min_value);
              }
          }
      buffer_out[gLoc]= min_value;
      }

}

my kernel settings now are:

dim3 block(16,16);
dim3 grid(512/(block.x),512/(block.y),nImg);

input:

output:

seems that the pixels of the apron are not copyied in the ouput buffer

Answer 1

You may want to read the following links for more detailed description and better example code on how to implement an image convolution CUDA kernel function.

http://igm.univ-mlv.fr/~biri/Enseignement/MII2/Donnees/convolutionSeparable.pdf

https://www.evl.uic.edu/sjames/cs525/final.html

Basically using a convolution filter of the size (5 x 5) does not mean setting the size of the thread block to be (5 x 5).

Typically, for a non-separable convolution, you could use a thread block of the size (16 x 16), to calculate a block of (16 x 16) pixels on the output image. To achieve this you need to read a block of ((2+16+2) x (2+16+2)) pixels from the input image to the shared memory, using the (16 x 16) threads collaboratively.