5

我为图像的框过滤编写了一个简单的 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。

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

问：为什么会这样？

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

4

1 回答 1

2

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

6.3.1.8 常用算术转换

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

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

If you use signed integers, no conversion occurs, and this whole problem disappears. The moral of this story is probably "know thy programming language".

于 2012-10-12T09:02:12.063 回答