CUDA中二维三角形平滑的正确网格和块尺寸是多少

What are the right Grid and Block Dimensions for 2D triangular smooth in CUDA?

本文关键字：网格 多少 平滑 三角形 二维 CUDA      更新时间：2023-10-16

我有一个顺序平滑算法

void triangularSmooth(unsigned char *grayImage, unsigned char *smoothImage, const int width, const int height, const float *filter, NSTimer &timer, dim3 grid_size, dim3 block_size) {
for ( int y = 0; y < height; y++ ) {
    for ( int x = 0; x < width; x++ ) {
        unsigned int filterItem = 0;
        float filterSum = 0.0f;
        float smoothPix = 0.0f;
        for ( int fy = y - 2; fy < y + 3; fy++ ) {
            for ( int fx = x - 2; fx < x + 3; fx++ ) {
                if ( ((fy < 0) || (fy >= height)) || ((fx < 0) || (fx >= width)) ) {
                    filterItem++;
                    continue;
                }
                smoothPix += grayImage[(fy * width) + fx] * filter[filterItem];
                filterSum += filter[filterItem];
                filterItem++;
            }
        }
        smoothPix /= filterSum;
        smoothImage[(y * width) + x] = static_cast< unsigned char >(smoothPix);
    }
}
}

我正在CUDA中实现，并希望使用共享变量来保持grayImage中的像素。然而在此之前，我正在尝试按原样运行它

__global__ void smooth(unsigned char *grayImage, unsigned char *smoothImage, const int width, const int height, const float *filter)
{
        int x = blockIdx.x*blockDim.x + threadIdx.x;
        int y = blockIdx.y*blockDim.y + threadIdx.y;
        unsigned int filterItem = 0;
        float filterSum = 0.0f;
        float smoothPix = 0.0f;
        for ( int fy = y - 2; fy < y + 3; fy++ ) {
            for ( int fx = x - 2; fx < x + 3; fx++ ) {
                if ( ((fy < 0) || (fy >= height)) || ((fx < 0) || (fx >= width)) ) {
                    filterItem++;
                    continue;
                }
                smoothPix += grayImage[(fy * width) + fx] * filter[filterItem];
                filterSum += filter[filterItem];
                filterItem++;
            }
        }
        smoothPix /= filterSum;
        smoothImage[(y * width) + x] = static_cast< unsigned char >(smoothPix);
}

并呼叫：

const float filter[] = {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 2.0f, 2.0f, 2.0f, 1.0f, 1.0f, 2.0f, 3.0f, 2.0f, 1.0f, 1.0f, 2.0f, 2.0f, 2.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f};
dim3 gridSize((width*height)/1024,(width*height)/1024,1);
dim3 blockSize(256,256,1);
smooth <<< gridSize, blockSize >>> (grayImage, smoothImage, width, height, filter);
cudaDeviceSynchronize();

问题是，得到的平滑图像看起来像是像素都在错误的另一个（混淆）。这是来自网格和块的尺寸吗？我已经尝试了很多其他可能的维度。什么是正确的方式？

我使用的是GTX480，版本-2.x，线程块网格的最大维度-3，线程块的最大x、y或z维度65535，每个块的最大线程数1024

首先，维度完全无效。在这种情况下，以下内容应该有效；

dim3 blockSize(16, 16, 1);
dim3 gridSize((width + blockSize.x - 1)/ blockSize.x, (height + blockSize.y - 1) / blockSize.y, 1);
smooth <<< grid_size, block_size >>> (grayImage, smoothImage, width, height);

修正后，使用cuda memcheck产生了类似于的结果；

========= Invalid __global__ read of size 4
=========     at 0x00000120 in cudaFilter
=========     by thread (4,1,0) in block (1,0,0)
=========     Address 0x05100190 is out of bounds

这表明内核代码中的值越界（很可能是数组索引）。通过检查各种变量，可以确定筛选器[]为空。

最后，如果要将filter[]传递到内核，则应该使用之类的东西将其从CPU复制到GPU

cudaMemcpy(filterGpu, filter, 25 * sizeof(float), cudaMemcpyHostToDevice);

或者，如果在其他地方不需要过滤器（就像这里的情况一样），可以在内核中声明它。

看看这个与图像过滤相关的答案，我建议您为图像创建如下的块和网格：

dim3 blockSize(16,16,1);
dim3 gridSize((width + blockSize.x - 1)/blockSize.x,(height + blockSize.y - 1)/blockSize.y,1);

您正在犯的另一个非常常见的错误是，传递给内核的筛选器数组是在主机上分配的。在设备上创建一个相同大小的数组，并将系数从主机复制到设备。将该设备数组传递给内核。

此外，强烈建议在主机端计算滤波器系数的总和，并将其作为参数传递给内核，而不是在每个线程中一次又一次地计算总和。

边界条件可能导致超出范围的存储器访问。在内核中显式处理边界条件。或者，简单的方法是对输入图像使用CUDA纹理，以便自动处理边界条件。

相关文章：