Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
Naraenda opened this issue Sep 14, 2023 · 1 comment
Open
Labels
bug Something isn't working

Comments

@Naraenda
Copy link
Member

Naraenda commented Sep 14, 2023

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 Naraenda added the bug Something isn't working label Sep 14, 2023
@AlexVlx
Copy link
Collaborator

AlexVlx commented Sep 25, 2023

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).

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working
Projects
None yet
Development

No branches or pull requests

2 participants