1

我有一个顺序平滑算法

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

4

2 回答 2

1

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

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

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

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

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

于 2013-02-07T08:23:00.900 回答
1

首先，尺寸完全无效。在这种情况下，以下应该起作用；

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[] 为空。

最后，如果要将 filter[] 传递给内核，则应使用类似的方法将其从 CPU 复制到 GPU

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

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

于 2013-02-07T09:14:13.460 回答