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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
