矢量指针在 openCL 中的工作原理

Question

我正在编写一个将 RGB 图像转换为灰度图像的示例程序。因此，图像作为一维数组从我的主机复制到设备，在我的代码中称为 imgIn。由于 imgIn 是 RGB 图像，每个像素由 3 个无符号字符组件（R、G 和 B）组成。由于输出 (imgOut) 是灰度图像，因此它仅由一个通道（亮度）组成。代码如下：

__kernel void rgbToGray(__global const uchar* restrict imgIn, 
                        __global uchar* restrict imgOut) {
    //Get two indexes of the work item
    int x = get_global_id(0);
    int y = get_global_id(1);
    //rgb average is luminosity
    //uchar3 channels = *(((__global uchar3 *) imgIn) + (x+640*y));
    uchar3 channels = *((__global uchar3 *) (imgIn+3*(x+640*y)));
    channels = channels/(uchar3)(3);
    imgOut[x+640*y] = channels.s0 + channels.s1 + channels.s2;
}

我想明白为什么uchar3 通道的注释声明与未注释的不等同。当我将 uchar 指针移动到正确的像素，然后将其转换为 uchar3 指针时，channels 变量具有正确的值，并且我的输出图像是完美的。但是当我将指针指向一个 uchar3 指针然后将指针移动到正确的像素（据推测）时，我的图像有一个奇怪的图案，如下一行所示。

http://i.stack.imgur.com/pjHuR.jpg

Answer 1

一个uchar3（其实就是任意三分量向量类型）与对应类型的四分量向量具有相同的对齐方式和大小。所以 uchar3 实际上只是一个 uchar4 上面有语法糖，以防止您访问最后一个组件，它的大小仍然是 4 个字节。

所以你的第一行在这里

uchar3 channels = *(((__global uchar3 *) imgIn) + (x+640*y));

失败是因为当你用你的 uchar3 * 做指针运算时，你最终增加了 4 个字节次 (x+640*y)，而你只想增加 3 个字节，所以你最终每次跳过一个字节像素，这会给您屏幕截图中显示的失真结果。

然而你的第二行

uchar3 channels = *((__global uchar3 *) (imgIn+3*(x+640*y)));

工作正常，因为您正在手动计算正确的偏移量，然后将偏移量指针转换为 uchar3 *，这很好，并且可以为您提供正确的像素字节。但是我相信如果 imgIn + 3*(x+640*y) 没有与 4 字节边界对齐，它在技术上仍然是未定义的。如果我错了（这是很有可能的），有人可以纠正我，但除此之外，我建议您使用未使用的填充字节一次传递 4 个字节的像素，或者手动将 3 个字节解压缩到 uchar3 中而不用通过指针重新解释转换（编辑：或者，更确切地说，按照 prunge 的建议使用 vload3，忘了那个）。

我的建议是拒绝将三个分量向量作为内核输入和输出。您可以在内核中使用它们，但将它们视为文字 uchar[3] 类型只会让人感到困惑。

Answer 2

根据规范中关于 data types 的部分：

For 3-component vector data types, the size of the data type is 4 * sizeof(component). This means that a 3-component vector data type will be aligned to a 4 * sizeof(component) boundary. The vload3 and vstore3 built-in functions can be used to read and write, respectively, 3-component vector data types from an array of packed scalar data type.

如果需要读取 3 分量矢量值，请使用 vload3。这个 doco 明确指出它只会从内存中读取 3 个值：

vload3 and vload_half3 read x, y, z components from address (p + (offset * 3)) into a 3-component vector.

所以这样的事情应该可行：

uchar3 channels = vload3(x + 640 * y, imgIn);

矢量指针在 openCL 中的工作原理

How vector pointers work in openCL

pointers

vector

opencl