AnsweredAssumed Answered

OpenCL compiler bug

Question asked by haahh on Sep 19, 2015
Latest reply on Sep 22, 2015 by haahh

Hi,

 

The following kernel gives segmentation fault (see transform algorithm with constant iterator · Issue #507 · boostorg/compute) on AMD (works on NVIDIA).

 

#define boost_pair_type(t1, t2) _pair_ ## t1 ## _ ## t2 ## _t
#define boost_pair_get(x, n) (n == 0 ? x.first ## x.second)
#define boost_make_pair(t1, x, t2, y) (boost_pair_type(t1, t2)) { x, y }
#define boost_tuple_get(x, n) (x.v ## n)
typedef struct {
   int v0;
   int v1;
} boost_tuple_int_int_t;

inline int fun(boost_tuple_int_int_t x){ return 10; }


__kernel void copy(uint count, __global int* _buf0, __global int* _buf1)
{
uint index = get_local_id(0) + (512 * get_group_id(0));
for(uint i = 0; i < 4; i++){
   if(index < count){
       _buf0[index]=fun((boost_tuple_int_int_t){ _buf1[index], 7});
       index += 128;
  }
}

}

 

However, when I replace

 

_buf0[index]=fun((boost_tuple_int_int_t){ _buf1[index], 7});

 

with

 

boost_tuple_int_int_t value = {_buf1[index], 7};
_buf0[index]=fun(value);

 

it works perfectly. 

 

It may be connect to Another bug in OpenCL compiler which was marked resolved.

 

Device: AMD Radeon HD7770

Env: Linux Mint 17.1 (based on ubuntu 14.04), AMD APP SDK 3.0, 15.20 drivers

Outcomes