使用Union内部寄存器的奇怪行为OpenCL

Question

我在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是一个选择，条件被定义时，*总分总是等于1，但如果我未定义条件，*总分得到正确的值，即0 ，1，..，或7，当然取决于哪个ni选择。

我还注意到，如果我停止使用union并只使用int8，它就解决了这个问题。

由于提前，

平台：Ubuntu的12.04 - 3.2.0-24 - 仿制PAE - 驱动程序的OpenCL 1.2 AMD-APP（923.1）

Answer 1

AMD与工会定位错误，可能这是原因？请参阅my report。我建议您尝试使用英特尔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;