HIP-CPU icon indicating copy to clipboard operation
HIP-CPU copied to clipboard

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

Open Naraenda opened this issue 9 months ago • 1 comments

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?

Naraenda avatar Sep 14 '23 16:09 Naraenda