c++ - cuda 多重图像侵 eclipse 不起作用

Question

我正在尝试使用 cuda 实现多个黑色 (0) 和白色 (255) 图像侵 eclipse ，我使用方形 (5x5) 结构元素。我实现的内核采用 < strong>unsigned char 数组缓冲区，其中存储 nImg 图像 200X200 像素。为了允许同时侵 eclipse 多个图像，我制作了一个具有 3D 结构的网格:

• 每个 block 的尺寸为 strel (5x5)
• 网格有 height = image_height/blockDim.y , width = image_width/blockDim.x , z = nImg

我已经尝试通过扩展 sample 来实现它.

问题是，如果我将一个线程 block 考虑的像素存储到 block 的线程之间共享的共享缓冲区； 为了允许快速内存访问，该算法无法正常工作。我尝试更改对我来说出错的 bindex，但我找不到解决方案。

有什么建议吗？

这是我的代码:

//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;
      }

}

这是我的内核设置:

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

我的内核调用:

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

我的图像输入和输出:

输入:输出(150 个增益元素):

没有共享内存的代码(低速):

__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;

}

更新的算法

我尝试遵循 samples做卷积。我改变了输入图像，现在有 512x512 大小，我写了那个算法:

#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;
      }

我的内核设置现在是:

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

输入:

输出:

似乎围裙的像素没有复制到输出缓冲区

Answer 1

您可能需要阅读以下链接以获得有关如何实现图像卷积 CUDA 核函数的更详细的描述和更好的示例代码。

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

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

基本上使用大小为 (5 x 5) 的卷积滤波器并不意味着将线程 block 的大小设置为 (5 x 5)。

通常，对于不可分离的卷积，您可以使用大小为 (16 x 16) 的线程 block 来计算输出图像上的 (16 x 16) 像素 block 。为实现这一点，您需要协同使用 (16 x 16) 个线程，从输入图像读取一 block ((2+16+2) x (2+16+2)) 像素到共享内存。