An issue with function pointer - `__global__` function interaction

[linehill]

The following should compile (tested with ROCm 6.0.0) but is doesn't currently on chipStar:

#include <hip/hip_runtime.h>

__global__ void foo(float *x);

// chipStar: error: cannot initialize a variable of type 'void (*)(float *)' 
// with an lvalue of type 'void (float *)'
void bar1(float *x) {
  void (*kernel)(float *x) = foo;
  kernel<<<1, 1>>>(x);
}

void baz(void (*kernel)(float *), float *x) { kernel<<<1, 1>>>(x); }
// chipStar: 
//   error: no matching function for call to 'baz'
//   ...
//   note: candidate function not viable: no known conversion from
//   'void (float *)' to 'void (*)(float *)' for 1st argument
void bar2(float *x) { baz(foo, x); }

Until a fix is provided, the errors can be worked around:

#include <hip/hip_runtime.h>

__global__ void foo(float *x);

using WorkAround = __global__ void (*)(float *x);
void bar1_workaround(float *x) {
  WorkAround kernel = foo;
  // Or just: auto kernel = foo;
  kernel<<<1, 1>>>(x);
}

void baz_workaround(WorkAround kernel, float *x) { kernel<<<1, 1>>>(x); }
void bar2_workaround(float *x) { baz_workaround(foo, x); }

[pvelesko]

Is this an issue in LLVM or our declarations? @linehill

[pvelesko]

I think this was related to me getting a bunch of compilation issues with HIP rebase.

[linehill]

Is this an issue in LLVM or our declarations?

Very likely in the LLVM. I doubt there is any chipStar declaration causing the error in the description.

I think this was related to me getting a bunch of compilation issues with HIP rebase.

The first HIP sample in the description was derived from a compilation error seen in the HIP rebase. It should compile (it compiles for AMD platform) but it doesn't currently.