CUDA侵蚀算法

CUDA erosion algorithm

本文关键字：算法 CUDA      更新时间：2023-10-16

我是CUDA的新手，我正在尝试开发具有结构元素3x3的简单(朴素)侵蚀算法。到目前为止，我已经开发了一个代码(它是基于nVidia演示):

#define bx (blockIdx.x)
#define by (blockIdx.y)
#define bdx (blockDim.x)
#define bdy (blockDim.y)
#define tx (threadIdx.x)
#define ty (threadIdx.y)
#define max( a, b ) ( ((a) > (b)) ? (a) : (b) )
#define min( a, b ) ( ((a) < (b)) ? (a) : (b) )
#define TILE_H 16
#define TILE_W 16
#define D 3    //structural element diameter
#define R 1    //structural element radius
#define BLOCK_W (TILE_W+D-1)
#define BLOCK_H (TILE_H+D-1)
__global__ void erosion(int *picture, unsigned int width, unsigned int height)
{    
    __shared__ int pixels[BLOCK_W*BLOCK_H];
    int x = bx*TILE_W + tx - R;
    int y = by*TILE_H + ty - R;
    x = max(0, x);
    x = min(x, (int)width-1);
    y = max(y,0);
    y = min(y, (int)height-1);
    unsigned int idx = y*width + x;
    unsigned int bidx = ty*bdy+tx;
    pixels[bidx] = picture[idx];
    __syncthreads();
    //compute pixels inside apron
    if (tx>=R && tx<BLOCK_W-R && ty>=R && ty < BLOCK_H-R)
    {
    //erode
    if (pixels[bidx] == 1)
        picture[idx] = pixels[ty*bdy+(tx+1)] & pixels[ty*bdy+(tx-1)] & pixels[(ty+1)*bdy+tx] & pixels[(ty-1)*bdy+tx];
    }
}

和main()函数:

int main()
{
    //...    
    int *pixels;
    int img_width=M; int img_height=N;
    cudaMemcpy(dev_pixels, pixels, M*N*sizeof(int), cudaMemcpyHostToDevice);
    dim3 blocks(img_width/BLOCK_W, img_height/BLOCK_H);
    erosion<<<blocks, D*D>>>(dev_pixels, img_width, img_height);
    cudaMemcpy(output, dev_pixels, M*N*sizeof(int), cudaMemcpyDeviceToHost);
}

我的问题是:似乎，erosion()从未达到if语句，我想在围裙内计算像素。你知道为什么会这样吗?我已经排除了img_widht/BLOCK_W划分(它可以返回0值，但目前我固定了img_widht=54和img_height=36)。

您正在启动一个内核，其网格由2D块数组组成，每个块数组都有一个1D线程数组:

dim3 blocks(img_width/BLOCK_W, img_height/BLOCK_H); // creates 2D blocks variable
erosion<<<blocks, D*D>>>(dev_pixels, img_width, img_height);
           ^       ^
           |       |
           |       1D array of threads
           2D array of blocks

因为你的threadblock是一个线程的1D数组，threadIdx.y总是零(对于每个block中的每个线程)。因此ty总是零，并且这个if-test总是失败:

if (tx>=R && tx<BLOCK_W-R && ty>=R && ty < BLOCK_H-R)

因为ty(==0)不大于或等于R(==1)

您可以通过定义适当的dim3数量在每个块中启动一个2D线程数组:

dim3 threads(D,D);

并将其传递到内核配置中:

erosion<<<blocks, threads>>>(dev_pixels, img_width, img_height);

我不能说这对你剩下的代码是否合理，但是通过这个修改，我可以说，你的if语句的内部(体)将被到达。

相关文章：