ROCm / HIP-CPU

An implementation of HIP that works on CPUs, across OSes.
MIT License
112 stars 19 forks source link

Types of functions in hip_api.h not consistent with HIP runtime and GCC #50

Open Naraenda opened 1 year ago

Naraenda commented 1 year ago

For example, in HIP-CPU __ffsll is defined with std::uint64_t:

std::uint32_t __ffsll(std::uint64_t x) noexcept
{
    return hip::detail::bit_scan_forward(x);
}

But in the HIP runtime (CLR) and in GCC it is defined with unsigned long long int:

__device__ static inline unsigned int __ffsll(unsigned long long int input) {
    return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
}

When calling HIP-CPU's implementation of __ffsll with unsigned long long int, it will throw a compiler error:

/workspaces/amd/libraries/rocRAND/library/include/rocrand/rocrand_sobol64.h:178:26: error: call to '__ffsll' is ambiguous
        unsigned int z = __ffsll(~x);
                         ^~~~~~~
/workspaces/amd/libraries/rocRAND/build/deps/hip-cpu/include/hip/hip_api.h:95:15: note: candidate function
std::uint32_t __ffsll(std::int64_t x) noexcept
              ^
/workspaces/amd/libraries/rocRAND/build/deps/hip-cpu/include/hip/hip_api.h:101:15: note: candidate function
std::uint32_t __ffsll(std::uint64_t x) noexcept

In the above code block, x is defined as unsigned long long int x.

This makes HIP-CPU not 100% compatible with otherwise valid HIP (GPU) code.


Should I make a PR for this or does this require some discussion first?

AlexVlx commented 1 year ago

Thank you for flagging this. I'd say that this is a historical misjudgement on my part, and that what you are suggesting is correct (matching the HIP public interface).