cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

haahh
Adept II

OpenCL compiler bug

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

0 Likes
2 Replies
dipak
Big Boss

Thanks for reporting the issue.

Yes, I can reproduce the error using CodeXL also. I'll report the issue to the compiler team.

Regards,

0 Likes

Thanks!

0 Likes