CUDA侵蚀算法

Question

我是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)。

Answer 1

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

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

因为您的线程块是一个一维线程数组，所以threadIdx.y总是为零(对于每个块中的每个线程)。因此，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-语句的内部(主体)将被访问。