AMD Compiler Crash

crash.py

KERNEL = r"""

#define PYOPENCL_DECLARE_COMPLEX_TYPE_INT(REAL_TP, REAL_3LTR, TPROOT, TP) \
  \
  REAL_TP TPROOT##_real(TP a) { return a.real; } \
  REAL_TP TPROOT##_imag(TP a) { return a.imag; } \
  \
  TP TPROOT##_new(REAL_TP real, REAL_TP imag) \
  { \
    TP result; \
    result.real = real; \
    result.imag = imag; \
    return result; \
  } \
  \
  TP TPROOT##_fromreal(REAL_TP real) \
  { \
    TP result; \
    result.real = real; \
    result.imag = 0; \
    return result; \
  } \
  \
  \
  TP TPROOT##_add(TP a, TP b) \
  { \
    return TPROOT##_new(a.real + b.real, a.imag + b.imag); \
    ; \
  } \
  \
  TP TPROOT##_mulr(TP a, REAL_TP b) \
  { \
    return TPROOT##_new(a.real*b, a.imag*b); \
  } \
  \

#define PYOPENCL_DECLARE_COMPLEX_TYPE(BASE, BASE_3LTR) \
  typedef union \
  { \
    struct { BASE x, y; }; \
    struct { BASE real, imag; }; \
  } c##BASE##_t; \
  \
  PYOPENCL_DECLARE_COMPLEX_TYPE_INT(BASE, BASE_3LTR, c##BASE, c##BASE##_t)

PYOPENCL_DECLARE_COMPLEX_TYPE(float, FLT);
#define cfloat_cast(a) cfloat_new((a).real, (a).imag)

__kernel void axpb(__global cfloat_t *z__base, long z__offset, cfloat_t a, __global float *x__base, long x__offset, float b, long n)
{
  int lid = get_local_id(0);
  int gsize = get_global_size(0);
  int work_group_start = get_local_size(0)*get_group_id(0);
  long i;

  __global cfloat_t *z = (__global cfloat_t *) ((__global char *) z__base + z__offset);
__global float *x = (__global float *) ((__global char *) x__base + x__offset);;
  //CL//
  for (i = work_group_start + lid; i < n; i += gsize)
  {
    z[i] = cfloat_add(cfloat_cast(cfloat_mulr(a, x[i])), cfloat_cast(cfloat_fromreal(b)));
  }

  ;
}
"""

import pyopencl as cl
ctx = cl.create_some_context()
cl.Program(ctx, KERNEL).build()