HIP-CPU
HIP-CPU copied to clipboard
Types of functions in hip_api.h not consistent with HIP runtime and GCC
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?