2 Replies Latest reply on Sep 22, 2015 6:00 AM by haahh

    OpenCL compiler bug

    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