首页 文章

CUDA中2D三角形光滑的网格和块尺寸是什么？

提问于
浏览 次
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

c++ cuda

2 回答

  • 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);

    或者，如果其他地方不需要过滤器（如此处所示），则可以在内核中声明它 .

    回复于
  • 1

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

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

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

    此外，强烈建议计算主机端的滤波器系数之和，并将其作为参数传递给内核，而不是在每个线程中反复计算和 .

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

    回复于

相关问题