[BUG]: Including <type_traits> header breaks unrelated code in CUDA >=12.2

[andipeer]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

libcu++

Describe the bug

Including the <cuda/std/type_traits> header breaks code that doesn't use any type traits at all. This behavior started with CUDA 12.2 (it compiles fine on CUDA 12.1 and below). The following error is given:

no type named ‘Type’ in ‘cuda::std::__4::__is_primary_template

I dived a bit into it an saw that in the respective code in is_primary_template.h, a special path for MSVC has been added. When using this path, the code compiles fine also on my Linux system.

How to Reproduce

1. Open godbold link below
2. See the error.
3. Comment out/remove the include on the first line.
4. Now, the code compiles.

Also switching to CUDA 11.8 will resolve the compile error. In contrast to my system, switching to CUDA 12.0 or 12.1 does not solve the problem on godbolt (maybe this changes only the nvcc version, but not the libcu++ version?).

Expected behavior

The code should compile fine independently of std headers that are included.

Reproduction link

https://godbolt.org/z/P8zzqx4Tn

Operating System

Ubuntu Linux 22.04

nvidia-smi output

No response

NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Wed_Nov_22_10:17:15_PST_2023
Cuda compilation tools, release 12.3, V12.3.107
Build cuda_12.3.r12.3/compiler.33567101_0

[miscco]

This seems like a nvcc compiler bug. Note that in the reproducer it suffices to name the alias as type instead of Type.

I am forwarding this internally to the nvcc team

[miscco]

I have filed an internal ticket with the following reproducer:

struct true_type {
  static constexpr bool value = true;
};
struct false_type {
  static constexpr bool value = false;
};

template <class _Tp>
using __test_for_primary_template = void;

template <template <class...> class _Templ, class... _Args,
          class = _Templ<_Args...> >
__host__ __device__ true_type __sfinae_test_impl(int);
template <template <class...> class, class...>
__host__ __device__ false_type __sfinae_test_impl(...);

template <template <class...> class _Templ, class... _Args>
using _IsValidExpansion = decltype(__sfinae_test_impl<_Templ, _Args...>(0));

template <class _Tp>
using __is_primary_template = _IsValidExpansion<__test_for_primary_template, _Tp>;

template <typename T>
struct S {
  using Type = T;
};

template <typename T>
__host__ __device__ auto func(S<T> s) -> S<T> {
  using ReturnType = decltype(func(s));
  using Type = typename ReturnType::Type;
  return {};
}

int main(int, char**) {
  func(S<int>());
  return 0;
}

[miscco]

Also as a side note, I prepared a patch to fix this potential issue by using the msvc branch but I still got the same error, so I am waiting for guidance from the compiler team on what to do here

[miscco]

@andipeer I believe the actual issue is this line:

 using ReturnType = decltype(func(s));

This requires the function to know its own return type while being compiled. I played around a bit and even fully removing __is_primary_template gave me an error at a different position. At some point it complained about

error: ‘_ForwardIterator2’ was not declared in this scope

Given that this is completely unrelated to any cccl code I am closing this issue

That said a possible solution to your problem would be to replace the offending line with

    using ReturnType = S<T>;

[jrhemstad]

@andipeer this looks very similar to another nvcc bug as I helped someone triage recently.

This was the repro I had created before: https://godbolt.org/z/s9dMfchG3

And here was the workaround I had found: https://godbolt.org/z/j6Mhj4dcP

[miscco]

@andipeer and future readers.

The bug has been fixed and will be released in an upcoming release.

As @jrhemstad mentioned a workaround is to defer the extraction of the type to an alias see this godbolt

[andipeer]

@miscco thanks for the really fast reaction on this, that's impressive! Applying a workaround would be bit of a hassle in our case, as the respective code resides inside a header of a 3rd party library. We will happily wait for the next CUDA release :)

[miscco]

@andipeer unfortunately it did not make it into the next cuda release due to scheduling, so it will be the one after the next