1

我在 opencl 中使用变量 union(private) 得到一个奇怪的行为。编码：

v = { 0, 1, 2, ... } // Defined in host and load to Device

typedef union svec8{
    int  word[8];
    int8 word8;
} vec8;

__kernel
void red( global vec8 *v, global int *out ){
    uint sizeBin = 8;
    vec8 binning = {0}; // Every Thread has a 8-space bin, initialized with 0
    uint gID     = get_global_id(0);
    int temp;

    binning.word8 = v[ dID ].word8;

    #ifdef CONDITION
        temp = 0;
        for ( uint i = 0; i < sizeBin; i++ ){
            temp += binning.word[ i ];
        }
    #endif

    if ( dID == 0 ){
        *out = binning.word[n]; // n belongs to [0, 7]
    }
}

问题是，对于一个选择的每一个 n，当 CONDITION 被定义时，*out 总是等于 1，但是如果 i undef CONDITION，*out 得到正确的值，即 0,1,.. 或 7,当然取决于你选择哪个。

我还注意到，如果我停止使用 union 而只使用 int8，它可以解决问题。

提前致谢，

平台：Ubuntu 12.04 - 3.2.0-24-generic-pae - 驱动程序 OpenCL 1.2 AMD-APP (923.1)

4

1 回答 1

3

AMD 在联合对齐方面存在错误，这可能是原因吗？见我的报告。我建议您尝试使用英特尔 SDK 运行相同的代码，看看它是否有所作为。

编辑：在我自己的代码（许多不同的联合）中，我使用这个解决方案，直到它在上游得到修复：

#define AMD_UNION_ALIGN_BUG_WORKAROUND() __attribute__((aligned(32)))
//            this is the biggest alignment value in the union ^^^^^
// in your case sizeof(int)*8=32 (in bytes)

// define unions like this
union svec8 { /*...*/ } AMD_UNION_ALIGN_BUG_WORKAROUND() vec8;
于 2012-06-12T14:30:48.750 回答