如何读取结构数组（OpenCL 内核）答案

【问题标题】：How can I read an array of structure (OpenCL kernel)如何读取结构数组（OpenCL 内核）
【发布时间】：2018-01-21 23:27:19
【问题描述】：

要求：

假设我们有 1）五组颜色，每组有三种颜色（颜色在 CPU 中动态生成）和 2）1000 辆汽车的列表，每辆汽车在列表中由它的颜色表示（从组中挑选的颜色）。我们想将三个参数传递给 OpenCL 内核：1) 一组生成的颜色，2) 汽车的颜色数组 (1D)，以及 3) 一个整数数组 (1D)，用于根据颜色组测试汽车颜色 (做一个简单的计算）。

结构：

struct GeneratedColorGroup
{
   float4 Color1; //16 =2^4
   float4 Color2; //16 =2^4
   float4 Color3; //16 =2^4
   float4 Color4; //16 =2^4
}

struct ColorGroup
{
    GeneratedColorGroup Colors[8]; //512 = 2^9
}

内核代码：

__kernel void findCarColorRelation(
const __global ColorGroup *InColorGroups,
const __global float4* InCarColor,
const __global int* CarGroupIndicator
const int carsNumber)
{
    int globalID = get_global_id( 0 );
    if(globalID < carsNumber)
    {
        ColorGroup colorGroups;
        float4 carColor;
        colorGroups = InColorGroups[globalID];
        carColor = InCarColor[globalID];

        for(int groupIndex =0; groupIndex < 8; groupIndex++)
        {
            if(colorGroups[groupIndex].Color1 == carColor)
            {
                CarGroupIndicator[globalID] = groupIndex + 1 ;
                break;
            }

            if(colorGroups[groupIndex].Color2 == carColor)
            {
                CarGroupIndicator[globalID] = groupIndex * 2 + 2;
                break;
            }

            if(colorGroups[groupIndex].Color3 == carColor)
            {
                CarGroupIndicator[globalID] = groupIndex * 3 + 3;
                break;
            }
        }
    }

}

现在，我们有 1000 个项目，这意味着内核将被执行 1000 次。没关系。

问题：如您所见，我们有一个全局 ColorGroup 作为内核的输入，这个全局内存有五个“GeneratedColorGroup”类型的项。

我尝试访问这些项目，如上面的代码所示，但我得到了意外的结果。而且执行很慢。

我的代码有什么问题？非常感谢任何帮助。

【问题讨论】：

标签： opencl

【解决方案1】：

将结构从主机传递到设备时，请确保在主机和设备代码中都使用__attribute__ ((packed)) 声明结构类型。否则，主机和设备编译器可能会为结构创建不同的内存布局，即它们可以为填充使用不同的大小。

使用打包结构可能会导致性能下降，因为打包结构根本没有填充，因此结构内的数据可能无法正确对齐，未对齐的访问通常很慢。在这种情况下，您必须手动插入带有char[] 的填充，或者在结构字段（或结构本身）上使用__attribute__ ((aligned (N)))。

有关 packed 和 aligned 属性的详细信息，请参阅 OpenCL C 规范： https://www.khronos.org/registry/OpenCL/sdk/1.1/docs/man/xhtml/attributes-types.html

【讨论】：

感谢您的帮助。我会试试packed属性，让你知道结果。
我已经尝试过打包属性，并且我已经在结构中添加了一些新字段以进行自对齐，但不幸的是，这并没有解决问题。
那么问题可能不在于struct 的使用。您是否故意不检查 Color4 字段？
您的回答有助于我处理数据结构。

【解决方案2】：

我猜测问题是

... CarGroupIndicator[globalID] = groupIndex + 1 ;
... CarGroupIndicator[globalID] = groupIndex * 2 + 2;
... CarGroupIndicator[globalID] = groupIndex * 3 + 3;

... 这使得从结果CarGroupIndicator[globalID] 中无法分辨出完全匹配的内容。例如。第 5 组颜色 1 的匹配结果为 6，但第 2 组颜色 2 和第 1 组颜色 3 的结果也是 6。你想要的是这样的：

... CarGroupIndicator[globalID] = groupIndex;
... CarGroupIndicator[globalID] = groupIndex + 8;
... CarGroupIndicator[globalID] = groupIndex + 16;

..那么0-7就是color1，8-15 color2，16-24 color3。

【讨论】：