[clang-cuda] clang segfaults when trying to convert a float to a 128 bit integer on device

[miscco]

consider the following minimal reproducer:

__global__ void kernel() {
    float meow = 1.0f;
#if 1 
    [[maybe_unused]] auto val = static_cast<__int128_t>(meow);
#else
    [[maybe_unused]] auto val = static_cast<__int128_t>(static_cast<int>(meow));
#endif
}

int main() {
    float meow = 1.0f;
    auto val = __int128_t(meow);
    kernel<<<1, 1>>>();
    return 0;
}

This will fail with

fatal error: error in backend: unsupported library call operation
clang++: error: clang frontend command failed with exit code 70 (use -v to see invocation)

It would be awesome if the library call could be implemented for the 128 bit integers and all builtin types

[Artem-B]

It would be awesome if the library call could be implemented for the 128 bit integers and all builtin types

The thing is that what needs fixing is that we're missing the concept of the standard GPU-side libraries.

These days we can do GPU-side linking in principle, thanks to the new clang driver changes @jhuber6 has recently implemented, but we're still missing the runtime support libraries themselves, and we didn't even start figuring out how to make it all work out of the box for regular CUDA users.

So, for the time being, there's no easy fix for your issue, unless @jhuber6 has some magic trick to make it all work with the GPU-side libc.

[jhuber6]

Theoretically I have a build of compiler-rt for the GPU. The issue here (I think) is that we currently reject all runtime libcalls. We could probably just enable those in the runtime libcall interface, but the issue there is that we would then start extracting them from static libraries unconditionally because that's the genius logic we use for doing LTO with libcalls (since we don't know ahead of time if the backend will use them we extract all of them).

[Artem-B]

Yup. Before the new driver, disabling the libcalls was all we could do, but since the new driver landed, we've got more options, but didn't get to re-visit the issue yet. So the details of what/when/where still needs to be done are still up in the air.

Patches would be welcome.

[Artem-B]

That said, we could probably handle some of the simpler common cases like integer conversions without involving the runtime support, if involved types are legal for the current back-end. I'll take a look at what can be done about that.

[Artem-B]

The code for FP to int conversion is a bit hairy, but it may be just simple enough to teach LLVM back-end about.

llvm-project/compiler-rt/lib/builtins/fp_fixint_impl.inc

Line 16 in cd6311b

static __inline fixint_t __fixint(fp_t a) {