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

[Yaxun Liu via Phabricator via llvm-commits]

yaxunl added a comment.

In D67509#1679524 <https://reviews.llvm.org/D67509#1679524>, @tra wrote:

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


I will add a lit test to make sure we have the desired behavior.


Repository:
  rL LLVM

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

https://reviews.llvm.org/D67509