代码之家  ›  专栏  ›  技术社区  ›  A. Student

如何读取结构数组(OpenCL内核)

  •  0
  • A. Student  · 技术社区  · 7 年前

    要求:

    假设我们有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次。没关系。

    问题是: 如您所见,我们有一个全局颜色组作为内核的输入,这个全局内存有五个“GeneratedColorGroup”类型的项。

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

    我的代码有什么问题? 非常感谢您的帮助。

    2 回复  |  直到 7 年前
        1
  •  1
  •   Andrew Savonichev    7 年前

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

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

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

        2
  •  0
  •   mogu    7 年前

    我疯狂地猜测问题是

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

    ... 这让我们无法从结果中判断 CarGroupIndicator[globalID] 什么是完全匹配的。E、 g.组5颜色1的匹配结果为值6,但组2颜色2和组1颜色3的匹配结果也是值6。你想要的是这样的:

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

    .. 然后0-7为颜色1,8-15为颜色2,16-24为颜色3。