search

CHIP-SPV / chipStar

chipStar is a tool for compiling and running HIP/CUDA on SPIR-V via OpenCL or Level Zero APIs.
Other
185 stars 30 forks source link

atomicAdd for unsigned long lost #700

Closed pjaaskel closed 10 months ago

pjaaskel commented 10 months ago

hipCUB used to build when CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE` was defined, which brought in the atomicAdd() for longs (see 71b5c82fd0235c1) which is not defined in the HIP spec.

It's seems it's missing again:

/home/pjaaskel/src/hipCUB/build/deps/rocprim/include/rocprim/intrinsics/../intrinsics/atomic.hpp:51:16: error: no matching function for call to 'atomicAdd'
        return ::atomicAdd(address, value);
linehill commented 10 months ago

The only atomic function I see the CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE covers is atomicAdd(unsigned long*, unsigned long) which is the undocumented one. However, the function seems to be callable with the recent chipStar.

$ cat atomicAddULong.hip
#include <hip/hip_runtime.h>
__device__ unsigned long x = 0;
__global__ void k() { atomicAdd(&x, 1); }
$ ../install/bin/hipcc -DCHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE atomicAddULong.hip -c

$ echo $?
0
pjaaskel commented 10 months ago

Seems I forgot that configuring chipStar doesn't then make hipcc automatically pass that macro (perhaps it should). Passing it separately moves the build forward, bumping to #701.