def _post_kernel_reduction_task(self, nelems, reduction_operator):
        assert [INC, MIN, MAX] in reduction_operator

        def generate_code():
            def headers():
                if self.dtype == np.dtype('float64'):
                    return """
#if defined(cl_khr_fp64)
#if defined(cl_amd_fp64)
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#else
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#endif
#elif defined(cl_amd_fp64)
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif

"""
                else:
                    return ""

            op = {INC: 'INC', MIN: 'min', MAX: 'max'}[reduction_operator]

            return """
%(headers)s
#define INC(a,b) ((a)+(b))
__kernel
void global_%(type)s_%(dim)s_post_reduction (
  __global %(type)s* dat,
  __global %(type)s* tmp,
  __private int count
)
{
  __private %(type)s accumulator[%(dim)d];
  for (int j = 0; j < %(dim)d; ++j)
  {
    accumulator[j] = dat[j];
  }
  for (int i = 0; i < count; ++i)
  {
    for (int j = 0; j < %(dim)d; ++j)
    {
      accumulator[j] = %(op)s(accumulator[j], *(tmp + i * %(dim)d + j));
    }
  }
  for (int j = 0; j < %(dim)d; ++j)
  {
    dat[j] = accumulator[j];
  }
}
""" % {'headers': headers(), 'dim': self.cdim, 'type': self._cl_type, 'op': op}

        src, kernel = _reduction_task_cache.get(
            (self.dtype, self.cdim, reduction_operator), (None, None))
        if src is None:
            src = generate_code()
            prg = cl.Program(_ctx, src).build(options="-Werror")
            name = "global_%s_%s_post_reduction" % (self._cl_type, self.cdim)
            kernel = prg.__getattr__(name)
            _reduction_task_cache[
                (self.dtype, self.cdim, reduction_operator)] = (src, kernel)

        kernel.set_arg(0, self._array.data)
        kernel.set_arg(1, self._d_reduc_array.data)
        kernel.set_arg(2, np.int32(nelems))
        cl.enqueue_task(_queue, kernel).wait()
        self._array.get(queue=_queue, ary=self._data)
        self.state = DeviceDataMixin.BOTH
        del self._d_reduc_array
dataset/ETHPy150Open OP2/PyOP2/pyop2/opencl.py/Global._post_kernel_reduction_task/swapped operands of `in`