CUDA Box过滤器索引错误

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,输出是完美的.

问题:为什么会这样？

PS:纹理寻址模式设置cudaAddressModeClamp为x和y方向.

Answer 1

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

6.3.1.8 常用算术转换

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

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

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