error: no matching function for call to 'atomicMax'

[linehill]

A case discovered from HeCBench/lebesgue-hip. Reduced case:

#include <hip/hip_runtime.h>
__device__ double Foo = 0;
__global__ void k(double Val) { atomicMax(&Foo, Val); }

This compiles with ROCm's hipcc but chipStar's hipcc fails:

$ ../install/bin/hipcc repro-atomicmax.hip -c
repro-atomicmax.hip:3:33: error: no matching function for call to 'atomicMax'
__global__ void k(double Val) { atomicMax(&Foo, Val); }
                                ^~~~~~~~~
/mnt/md1/linehill/ws-chip-spv-2/install/include/hip/devicelib/atomics.hh:216:36: note: candidate function not viable: no known conversion from 'double *' to 'int *' for 1st argument
extern "C++" inline __device__ int atomicMax(int *address, int val) {
                                   ^
/mnt/md1/linehill/ws-chip-spv-2/install/include/hip/devicelib/atomics.hh:222:45: note: candidate function not viable: no known conversion from 'double *' to 'unsigned int *' for 1st argument
extern "C++" inline __device__ unsigned int atomicMax(unsigned int *address,
                                            ^
/mnt/md1/linehill/ws-chip-spv-2/install/include/hip/devicelib/atomics.hh:230:1: note: candidate function not viable: no known conversion from 'double *' to 'unsigned long long *' for 1st argument
atomicMax(unsigned long long *address, unsigned long long val) {
^
1 error generated when compiling for .

failed to execute:/mnt/md1/linehill/ws-chip-spv-2/install/bin/clang++ -x hip repro-atomicmax.hip -D__HIP_PLATFORM_SPIRV__= --offload=spirv64 -Xclang -no-opaque-pointers -nohipwrapperinc --hip-path=/mnt/md1/linehill/ws-chip-spv-2/install --target=x86_64-unknown-linux-gnu -include /mnt/md1/linehill/ws-chip-spv-2/install/include/hip/spirv_fixups.h -I//mnt/md1/linehill/ws-chip-spv-2/install/include   -c

[pvelesko]

Does this work for CUDA? Officially not supported. I see a solution here:

https://stackoverflow.com/questions/18950732/atomic-max-for-floats-in-opencl

[pjaaskel]

Seems to be the same as atomicAdd(), the long version is not in HIP officially, but it's used.

[pvelesko]

@pjaaskel Does the implementation in the link I sent good enough for us?

[linehill]

Does this work for CUDA? Officially not supported. I see a solution here: https://stackoverflow.com/questions/18950732/atomic-max-for-floats-in-opencl

LGTM but it could be simplified applying atomic_compare_exchange_strong directly on the floating-point values without need for bitcasting.