CUDA 框过滤器索引错误

Question

我为图像的框过滤编写了一个简单的 CUDA 内核。

texture<unsigned char,2> tex8u;

#define FILTER_SIZE 7
#define FILTER_OFFSET (FILTER_SIZE/2)

__global__ void box_filter_8u_c1(unsigned char* out, int width, int height, int pitch)
{
   unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

   if(x>=width || y>=height)    return;

   float val = 0.0f;

   for(int i = -FILTER_OFFSET; i<= FILTER_OFFSET; i++)
     for(int j= -FILTER_OFFSET; j<= FILTER_OFFSET; j++)
        val += tex2D(tex8u,x + i, y + j);

   out[y * pitch + x] = static_cast<unsigned char>(val/(FILTER_SIZE  * FILTER_SIZE));   

}

上面代码的问题是图片的上下边框过滤不正确。它们分别包含来自底部和右侧边界的值。错误边框的宽度等于FILTER_OFFSET。

但是当我将 x 和 y 索引更改为 int 而不是 unsigned int 时，输出是完美的。

问题：为什么会这样？

P.S：x 和 y 方向的纹理寻址模式设置为cudaAddressModeClamp。

Answer 1

根本原因与 CUDA 无关，是基本的 C 类型转换规则导致了您看到的结果。 C99 标准说明了如何执行转换：

6.3.1.8 常用算术转换

1. 如果两个操作数的类型相同，则无需进一步转换。
2. 否则，如果两个操作数都具有有符号整数类型或都具有无符号整数类型，则具有较小整数类型的操作数 转换等级被转换为具有更大的操作数的类型 排名。
3. 否则，如果具有无符号整数类型的操作数的秩大于或等于另一个操作数类型的秩，则 带符号整数类型的操作数转换为 无符号整数类型的操作数。
4. 否则，如果带符号整数类型的操作数可以表示无符号操作数类型的所有值 整数类型，然后转换无符号整数类型的操作数 为带符号整数类型的操作数的类型。
5. 否则，两个操作数都将转换为与带符号整数类型的操作数的类型相对应的无符号整数类型。

第三点意味着有符号整数（在这种情况下是i 和j）首先转换为无符号整数，然后添加到无符号整数（x 和y)。将负符号整数转换为无符号整数的结果是特定于实现的，但在这里，简单的二进制补码表示会将一个小的负整数转换为一个非常大的无符号整数。纹理的读取模式将这个超出范围的坐标限制在纹理中允许的最大值，并且您的内核最终会从纹理的错误一侧读取。

如果您使用有符号整数，则不会发生转换，整个问题就会消失。这个故事的寓意可能是“了解你的编程语言”。