loopy icon indicating copy to clipboard operation
loopy copied to clipboard

conj() generates broken code for non-complex types

Open matthiasdiener opened this issue 3 years ago • 1 comments

Working code for complex conj():

#define lid(N) ((int) get_local_id(N))
#define gid(N) ((int) get_group_id(N))
#if __OPENCL_C_VERSION__ < 120
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif
#define PYOPENCL_DEFINE_CDOUBLE

#include <pyopencl-complex.h>

__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) _pt_kernel(__global cdouble_t const *__restrict__ _pt_in, __global cdouble_t *__restrict__ _pt_out)
{
  for (int _pt_out_dim0 = 0; _pt_out_dim0 <= 9; ++_pt_out_dim0)
    _pt_out[_pt_out_dim0] = cdouble_conj(_pt_in[_pt_out_dim0]);
}

For float conj():

#define lid(N) ((int) get_local_id(N))
#define gid(N) ((int) get_group_id(N))
#if __OPENCL_C_VERSION__ < 120
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif

__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) _pt_kernel(__global double const *__restrict__ _pt_in, __global double *__restrict__ _pt_out)
{
  for (int _pt_out_dim0 = 0; _pt_out_dim0 <= 9; ++_pt_out_dim0)
    _pt_out[_pt_out_dim0] = conj(_pt_in[_pt_out_dim0]);
}
pyopencl._cl.RuntimeError: clBuildProgram failed: BUILD_PROGRAM_FAILURE - clBuildProgram failed: BUILD_PROGRAM_FAILURE - clBuildProgram failed: BUILD_PROGRAM_FAILURE

Build on <pyopencl.Device 'pthread' on 'Portable Computing Language' at 0x7f933b720310>:

error: /Users/mdiener/.cache/pocl/kcache/tempfile-f5-be-04-9b-fb.cl:10:29: implicit declaration of function 'conj' is invalid in OpenCL

matthiasdiener avatar Jun 08 '21 04:06 matthiasdiener

Just add handling for the real-valued case here:

https://github.com/inducer/loopy/blob/718c503b68e0843d58b8e049a80e0072eaea88e6/loopy/target/pyopencl.py#L239-L266

inducer avatar Jun 08 '21 21:06 inducer