CUDA compilation fails if <complex> is included

[llvmbot]

Bugzilla Link	50383
Resolution	FIXED
Resolved on	Jun 01, 2021 17:33
Version	11.0
OS	Linux
Reporter	LLVM Bugzilla Contributor
CC	@fodinabor,@rgreenblatt,@Artem-B,@jwakely
Fixed by commit(s)	9a75c06

Extended Description

The file: "/usr/include/c++/11.1.0/complex" contains function named "polar", which does the following test:

__glibcxx_assert( __rho >= 0 );

which breaks CUDA compilation. Commenting the above line solves the issue. The error is as follows:

error: __host__ __device__ function '__failed_assertion' cannot overload __host__ function '__failed_assertion'
      __glibcxx_assert( __rho >= 0 );
      ^
/usr/bin/../lib64/gcc/x86_64-pc-linux-gnu/11.1.0/../../../../include/c++/11.1.0/x86_64-pc-linux-gnu/bits/c++config.h:754:5: note: expanded from macro '__glibcxx_assert'
    __glibcxx_assert_1(_Condition)        \
    ^
/usr/bin/../lib64/gcc/x86_64-pc-linux-gnu/11.1.0/../../../../include/c++/11.1.0/x86_64-pc-linux-gnu/bits/c++config.h:743:13: note: expanded from macro '__glibcxx_assert_1'
       void __failed_assertion();       \
            ^
/usr/bin/../lib64/gcc/x86_64-pc-linux-gnu/11.1.0/../../../../include/c++/11.1.0/bits/atomic_base.h:266:7: note: previous declaration is here
      __glibcxx_assert(__b != memory_order_consume);
      ^
/usr/bin/../lib64/gcc/x86_64-pc-linux-gnu/11.1.0/../../../../include/c++/11.1.0/x86_64-pc-linux-gnu/bits/c++config.h:754:5: note: expanded from macro '__glibcxx_assert'
    __glibcxx_assert_1(_Condition)        \
    ^
/usr/bin/../lib64/gcc/x86_64-pc-linux-gnu/11.1.0/../../../../include/c++/11.1.0/x86_64-pc-linux-gnu/bits/c++config.h:743:13: note: expanded from macro '__glibcxx_assert_1'
       void __failed_assertion();       \
            ^

[llvmbot]

assigned to @Artem-B

[Artem-B]

I'll need more details.
Fo far I'm unable to reproduce the issue.
AFAICT #include <complex> works: https://godbolt.org/z/fjheqs1Kh

If you can reproduce the issue on godbolt.org and post a link, that would help a lot.

[llvmbot]

I'll need more details.
Fo far I'm unable to reproduce the issue.
AFAICT #include <complex> works: https://godbolt.org/z/fjheqs1Kh

If you can reproduce the issue on godbolt.org and post a link, that would
help a lot.

Thanks for the quick response. Please check out the following link:

https://godbolt.org/z/WTzKcEKqh

TRUNK -> not working
11.0 -> working
11.1.0 on my PC -> not working

[fodinabor]

Well, the actual issue keeps persisting with 11.0.0: https://cuda.godbolt.org/z/fW7jxs4Eh

As far as I dug into it, this is due to libstdc++11 using this awful construct as constant evaluated assertion, where the inline declared __failed_assertion() function will be marked as __host__ __device__ due to the #pragma clang force_cuda_host_device begin in Clang's cuda_wrappers/complex. Not really sure what the Clang side could do differently here, though. Redefining __glibcxx_assert inside the complex header like this https://cuda.godbolt.org/z/e44qcaah5 might work.. but seems brittle..

#if _GLIBCXX_HAVE_BUILTIN_IS_CONSTANT_EVALUATED
# define __glibcxx_assert_1(_Condition)		\
    if (__builtin_is_constant_evaluated())	\
     {						\
       void __failed_assertion();	\
       if (!bool(_Condition))			\
	 __failed_assertion();	\
     }						\
    else
#else
# define __glibcxx_assert_1(_Condition)
#endif

# define __glibcxx_assert(_Condition)	  \
  do {					  \
    __glibcxx_assert_1(_Condition)	  \
    { __glibcxx_assert_2(_Condition); }   \
  } while (false)

[llvmbot]

What about using using push_macro and pop_macro pragmas in /cuda_wrappers/complex, since it is already using it for _GLIBCXX_USE_C99_COMPLEX macro (https://godbolt.org/z/hWWsd43hj).

#pragma push_macro("__glibcxx_assert")
#define __glibcxx_assert(_Condition)

#include_next <complex>

#pragma pop_macro("__glibcxx_assert")

[Artem-B]

AFAICT it's an issue with the libstdc++ 11 or newer. It does not affect v10 I have on my machine.

The root cause is that libstdc++ appears to redeclare __failed_assertion.
If you notice, the error reports the same location of the conflicting declarations: c++config.h:743:13

I'm not sure why c++config happens to be included again. Did they miss include guards? If we can make it to be included only once, then all we need is to provide device __failed_assertion.

If multiple inclusion is needed, then one way to work around that would be to redeclare the original __failed_assertion() with __attribute__((enable_if(false))) and then provide our own declaration with host device.

[fodinabor]

It's not a multiple inclusion problem, but instead they redeclare __failed_assertion, every time the __glibcxx_assert macro is used. Thus it is declared in every header where that macro is used.
E.g. atomic and complex. And yes, the commit leading to this change is new in version 11.

[fodinabor]

Linking the corresponding libstdc++ bug: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=100676

Seems it might be fixed on their side in 11.2 by using __builtin_unreachable instead of redeclaring void __failed_assertion();.

[jwakely]

I'll change the libstdc++ code, but AFAICT our code is valid, so this seems like a bug in the CUDA compiler:

/opt/compiler-explorer/gcc-snapshot/lib/gcc/x86_64-linux-gnu/12.0.0/../../../../include/c++/12.0.0/complex:703:7: error: host device function '__failed_assertion' cannot overload host function '__failed_assertion'

It's not an overload, it's a redeclaration of the same function. Shouldn't the declarations be merged?

[Artem-B]

It's not an overload, it's a redeclaration of the same function. Shouldn't the declarations be merged?

CUDA is a bit unusual in this regard. In our attempt to be able to use parts of the standard library on GPU, we have to resort to some rather gross hacks, with various degrees of success.

In a nutshell, CUDA allows function overloading based on host and device attributes. If a function has no attributes, host is implied.

In order to be able use on GPU we force host device on all the functions in it:
https://github.com/llvm/llvm-project/blob/main/clang/lib/Headers/cuda_wrappers/complex#L43

What happens in this case is that __failed_assertion() is seen multiple times.
When it's seen while we're including , it has host device attributes, otherwise, it's assumed to be host.

As far as CUDA compilation is concerned, those are two different overloads, with the conflicting attributes.

So, yes, libstdc++ does not do anything wrong. It just happens to make it a bit harder to use from CUDA sources on the GPU side -- something which the standard library was never intended for.

[jwakely]

Ah, so the first include of is what creates the host declaration. That's the part I was missing in understanding it, thanks.

The change should be on GCC trunk in the next 24 hours, and then the gcc-11 branch next week.

[Artem-B]

Considering that this change has already been released, we should probably still have a work around.

Do you foresee any issues if we were to do something like this in the clang's wrapper for the :

#pragma clang force_cuda_host_device begin
#pragma push_macro ("__failed_assertion")
#if libstdcxx_version_within_range 
#define __failed_assertion __cuda_failed_assertion
__host__ __device__ __cuda_failed_assertion() { crash_gpu(); }
#endif
#include_next <complex>
#pragma pop_macro ("__failed_assertion")
#pragma clang force_cuda_host_device end

This should get rid of the attribute conflict if the user is unlucky enough to be stuck with the version that has the issue.

[jwakely]

Considering that this change has already been released, we should probably
still have a work around.

Do you foresee any issues if we were to do something like this in the
clang's wrapper for the :

#pragma clang force_cuda_host_device begin
#pragma push_macro ("__failed_assertion")
#if libstdcxx_version_within_range 
#define __failed_assertion __cuda_failed_assertion
__host__ __device__ __cuda_failed_assertion() { crash_gpu(); }

Do you need this definition? How do you crash the GPU at constexpr time?

If you do need it, it should be inside namespace std.

[Artem-B]

#define __failed_assertion __cuda_failed_assertion
host device __cuda_failed_assertion() { crash_gpu(); }

Do you need this definition?

Maybe. What is __failed_assertion() expected to do?

How do you crash the GPU at constexpr time?

At compile-time -- the same way we'd error out in the regular C++ compilation.
If we do emit code that would get executed on a GPU, we'd need to do something sensible there to mimic the behavior on the host to the extent possible.
We may need two separate overloads for __cuda_failed_assertion() -- one for the host which would just forward to __failed_assert() and one for device which would do... something different.

If you do need it, it should be inside namespace std.

Ack.

[jwakely]

Its sole purpose is to be a non-constexpr function that can be called during constant evaluation to cause a compilation error. It can't be called at runtime, because compilation will fail if control flow ever results in calling it. So no definition is needed.