amd_hip_fp16.h: candidate function not viable: call to __device__ function from __host__ function

[jinz2014]

I am not sure if the issues were reported before.

/opt/rocm-6.2.0/lib/llvm/bin/../../../include/hip/amd_detail/amd_hip_fp16.h:280:24: note: candidate function not viable: call to device function from host function
280 | __half operator*(const __half& x, const __half& y)

/opt/rocm-6.2.0/lib/llvm/bin/../../../include/hip/amd_detail/amd_hip_fp16.h:287:24: note: candidate function not viable: call to device function from host function
287 | __half operator/(const __half& x, const __half& y)
| ^

[cjatin]

can you share sample as well.

[jinz2014]

https://github.com/zjin-lcf/HeCBench/tree/master/src/nlll-hip

[cjatin]

I think the issue is the attempt to use * operator for __half on host. Which on HIP is not supported. There is just __device__ variants available: https://github.com/ROCm/clr/blob/c47f9dda5886b8062cad1b9c9736671241ab6f9a/hipamd/include/hip/amd_detail/amd_hip_fp16.h#L198

One thing you might need to do (for now) here is pass on _Float16 for hosts and __half for device. Something like this: https://gist.github.com/cjatin/b92a5b51adc2e0e4a60d4ad7c4d2fdc2 (although there might be better design possible than what I wrote here)

I think we should add the host operators, will create a task for it internally to track and add these.