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

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

<p>i've try to implement it extending that <a href="http://www.nvidia.com/content/nvision2008/tech_presentations/Game_Developer_Track/NVISION08-Image_Processing_and_Video_with_CUDA.pdf" rel="nofollow noreferrer">sample</a>.</p>

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

<p>any suggestion? </p>

<p>here's my code:</p>

<pre><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 &gt;= R) &amp;&amp; (threadIdx.x &lt; BLOCK_W-R) &amp;&amp; (threadIdx.y &gt;= R) &amp;&amp; (threadIdx.y &lt;BLOCK_H-R)){

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

            for (int dx = -R ; dx &lt;= R; dx++) {

                //if only one element in mask is different from the value of strel el --&gt; the strel is not contain in the mask --&gt; 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;
      }

}
</code></pre>

<p>that are my kernel settings:</p>

<pre><code> dim3 block(5,5,1);
 dim3 grid(200/(block.x),200/(block.y),nImg);
</code></pre>

<p>my kernel call: </p>

<pre><code> erode_multiple_img_SM&lt;&lt;&lt;grid,block&gt;&gt;&gt;(dimage_src,dimage_dst,200,200);
</code></pre>

<p>my image input and output:</p>

<p>input: <a href="https://i.stack.imgur.com/8VBkG.jpg" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/8VBkG.jpg" alt="input"></a> output(150 buff element): <a href="https://i.stack.imgur.com/kPTCt.jpg" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/kPTCt.jpg" alt="output"></a></p>

<p>code without shared memory(low speed):</p>

<pre><code>__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&lt;=STREL_H/2; dy++){
        if(is_contain == false)
            break;
        for (int dx = -STREL_W/2 ; dx &lt;= 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;

}
</code></pre>

<p><strong>UPDATED ALGORITHM</strong></p>

<p>i try to follow that <a href="https://bluewaters.ncsa.illinois.edu/documents/10157/63094/NCSA02a-window-example-cpp-08-14-13.pdf" rel="nofollow noreferrer">samples</a> to do convolution.I change the input image, now has <strong>512x512 size</strong> and i wrote that algorithm:</p>

<pre><code>#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&lt;w)&amp;&amp;(row&lt;h)) {
        data[threadIdx.x][threadIdx.y]=buffer_in[gLoc];

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

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

    if ((threadIdx.x &gt;(w-STREL_SIZE)) &amp;&amp; (threadIdx.y &gt; (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&lt;STREL_SIZE;x++){
          for(y=0;y&lt;STREL_SIZE;y++){
              min_value = min( (data[threadIdx.x+x][threadIdx.y+y]) , min_value);
              }
          }
      buffer_out[gLoc]= min_value;
      }
</code></pre>

<p>}</p>

<p>my kernel settings <strong>now</strong> are:</p>

<pre><code>dim3 block(16,16);
dim3 grid(512/(block.x),512/(block.y),nImg);
</code></pre>

<p>input:
    <a href="https://i.stack.imgur.com/DAci4.jpg" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/DAci4.jpg" alt="input"></a></p>

<p>output:
<a href="https://i.stack.imgur.com/OVVvp.jpg" rel="nofollow noreferrer"><img src="https://i.stack.imgur.com/OVVvp.jpg" alt="output"></a></p>

<p><strong>seems that the pixels of the apron are not copyied in the ouput buffer</strong> </p>