[PATCH] D67509: [CUDA][HIP] Fix hostness of defaulted constructor

[Artem Belevich via Phabricator via llvm-commits]

tra added a comment.

In D67509#1678394 <https://reviews.llvm.org/D67509#1678394>, @yaxunl wrote:

> A reduced test case is
>
>   struct A {
>     A();
>   };
>  
>   template <class T>
>   struct B
>   {
>     T a;
>     constexpr B() = default;
>   };
>  
>   B<A> x;
>  
>
>
> `B<A>::B()` got implicit `__host__ __device__`  attrs due to constexpr before entering Sema::inferCUDATargetForImplicitSpecialMember.
>  In Sema::inferCUDATargetForImplicitSpecialMember, the inferred hostness of `B<A>::B()` is host since `A::A()` is host. This causes discrepancy between the inferred hostness and the existing hostness.


On one hand `inferCUDATargetForImplicitSpecialMember` is correct here. 
On the other hand, `constexpr` being implicitly `__host__ __device__` also works OK, with the error popping up only if we need to instantiate the B<A> on device side.

So, what we want is:

  __host__ void f() {B<A> x;} // This should compile
  __device__ void f() {B<A> x;} // This should produce an error.
  
  struct foo {
    __host__ foo() { B<A> x; } // should compile
    __device__ foo() { B<A> x; } // ???
  };

We could remove the implicit '__device__' attribute from the function. This should make `__device__ foo()` fail to compile regardless of whether `struct foo` is instantiated on device.
Or we can keep the defaulted constexpr function as `__host__ __device__` and catch the error only if/when `struct foo` is instantiated on device side.

NVCC (and clang as it is right now) appear to follow the latter -- there's no error if we don't generate code for the function.
https://godbolt.org/z/aVhvVn

For the sake of avoiding surprises, I think we should preserve this behavior and just relax the assertion here. We should be OK to infer stricter set of attributes, but not to relax them.


Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D67509/new/

https://reviews.llvm.org/D67509