[CUDA][HIP] Fix ShouldDeleteSpecialMember for inherited constructors

[yaxunl]

Summary

ShouldDeleteSpecialMember is called upon inherited constructors.
It calls inferCUDATargetForImplicitSpecialMember.

Normally the special member enum passed to ShouldDeleteSpecialMember
matches the constructor. However this is not true when inherited
constructor is passed, where DefaultConstructor is passed to treat
the inherited constructor as DefaultConstructor. However
inferCUDATargetForImplicitSpecialMember expects the special
member enum argument to match the constructor, which results
in assertion when this expection is not satisfied.

This patch checks whether the constructor is inherited. If true it will
get the real special member enum for the constructor and pass it
to inferCUDATargetForImplicitSpecialMember.

[tra]

@jlebar Justin, can you take a look?

[yaxunl]

ping

[jlebar]

Sorry for missing tra's ping earlier, I get a lot of HIP email traffic that's 99% inactionable by me, so I didn't notice my username in tra's earlier email.

@tra, @timshen, and I debugged this IRL for a few hours this afternoon. The result of this is that we don't think the fix in this patch is correct.

Here's what we think is happening.

When clang sees using A::A inside of B, it has to check whether this constructor is legal in B. An example of where this constructor would *not* be legal is something like:

struct NoDefaultConstructor { NoDefaultConstructor() = delete; };
struct A { A(const int& x) {} }
struct B {
  using A::A;
  NoDefaultConstructor c;
};

The reason this using A::A is not legal here is because the using statement is equivalent to writing

B(const int& x) : A(x) {}

but this constructor is not legal, because NoDefaultConstructor is not default-constructible, and a constructor for B must explicitly initialize all non-default-initializable members.

Here is the code that checks whether the using statement is legal:

https://github.com/llvm-project/llvm-project-20170507/blob/51b65eeeab0d24268783d6246fd949d9a16e10e8/clang/lib/Sema/SemaDeclCXX.cpp#L11018

This code is kind of a lie! DerivedCtor is the constructor B(const int& x) : A(x) {} that we've created in response to the using declaration. Notice that it's not a default constructor! In fact, it's not even a special member (i.e. it's not a default constructor, copy constructor, move constructor, etc.). But notice that we pass CXXDefaultConstructor, and we call the function ShouldDeleteSpecialMember!

The reason we don't tell the truth here seems to be out of convenience. To determine whether we should delete the new constructor on B, it seems like we are trying to ask: Would a default constructor on B be legal, ignoring the fact that A has to be explicitly initialized? That is, the new constructor we're creating is just like a default constructor on B, except that first it initializes A. So we're trying to reuse the default constructor logic.

But eventually our tricks and dishonesty catch up to us, here in CUDA code. This patch fixes one instance where we do the wrong thing and hit an assertion, but who knows if the code is right in general; simply adding on another layer of hack does not seem like the right approach to us.

cc @rsmith

[yaxunl]

Revised by Justin's comments.

Handle the situation where an inherited constructor is faked to be a default constructor but it is really not.

[yaxunl]

fix a typo.

[jlebar]

lib/Sema/SemaDeclCXX.cpp

7231

LGTM, but perhaps we should use a new variable instead of modifying CSM in case someone adds code beneath this branch?

[yaxunl]

lib/Sema/SemaDeclCXX.cpp

7231

will do when committing.

lib/Sema/SemaDeclCXX.cpp

if (SMI.visit(SMI.IsAssignment ? SMI.VisitDirectBases
     return true;
 
   if (SMI.shouldDeleteForAllConstMembers())
     return true;
 
   if (getLangOpts().CUDA) {
     // We should delete the special member in CUDA mode if target inference
     // failed.
+    // For inherited constructors (non-null ICI), CSM may be passed so that MD
+    // is treated as certain special member, which may not reflect what special
+    // member MD really is. However inferCUDATargetForImplicitSpecialMember
+    // expects CSM to match MD, therefore recalculate CSM.
+    assert(ICI || CSM == getSpecialMember(MD));
+    if (ICI)
+      CSM = ICI->getSpecialMember(MD);
+
     return inferCUDATargetForImplicitSpecialMember(RD, CSM, MD, SMI.ConstArg,
                                                    Diagnose);
   }
 
   return false;
 }
 
 /// Perform lookup for a special member of the specified kind, and determine

test/SemaCUDA/implicit-member-target-inherited.cu

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: infer inherited default ctor to be host.
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
+
+// The inherited default constructor is inferred to be host, so we'll encounter
+// an error when calling it from a __device__ function, but not from a __host__
+// function.
+struct B1_with_implicit_default_ctor : A1_with_host_ctor {
+  using A1_with_host_ctor::A1_with_host_ctor;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+// expected-note@-6 2{{constructor from base class 'A1_with_host_ctor' inherited here}}
+
+void hostfoo() {
+  B1_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo() {
+  B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: infer inherited default ctor to be device.
+
+struct A2_with_device_ctor {
+  __device__ A2_with_device_ctor() {}
+};
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
+
+struct B2_with_implicit_default_ctor : A2_with_device_ctor {
+  using A2_with_device_ctor::A2_with_device_ctor;
+};
+
+// expected-note@-4 {{call to __device__ function from __host__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+// expected-note@-6 2{{constructor from base class 'A2_with_device_ctor' inherited here}}
+
+void hostfoo2() {
+  B2_with_implicit_default_ctor b;  // expected-error {{no matching constructor}}
+}
+
+__device__ void devicefoo2() {
+  B2_with_implicit_default_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: infer inherited copy ctor
+
+struct A3_with_device_ctors {
+  __host__ A3_with_device_ctors() {}
+  __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
+};
+
+struct B3_with_implicit_ctors : A3_with_device_ctors {
+  using A3_with_device_ctors::A3_with_device_ctors;
+};
+// expected-note@-3 2{{call to __device__ function from __host__ function}}
+// expected-note@-4 {{default constructor}}
+
+
+void hostfoo3() {
+  B3_with_implicit_ctors b;  // this is OK because the inferred inherited default ctor
+                             // here is __host__
+  B3_with_implicit_ctors b2 = b; // expected-error {{no matching constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 4: infer inherited default ctor from a field, not a base
+
+struct A4_with_host_ctor {
+  A4_with_host_ctor() {}
+};
+
+struct B4_with_inherited_host_ctor : A4_with_host_ctor{
+  using A4_with_host_ctor::A4_with_host_ctor;
+};
+
+struct C4_with_implicit_default_ctor {
+  B4_with_inherited_host_ctor field;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo4() {
+  C4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+  C4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: inherited copy ctor with non-const param
+
+struct A5_copy_ctor_constness {
+  __host__ A5_copy_ctor_constness() {}
+  __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
+};
+
+struct B5_copy_ctor_constness : A5_copy_ctor_constness {
+  using A5_copy_ctor_constness::A5_copy_ctor_constness;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
+// expected-note@-5 {{candidate constructor (the implicit default constructor) not viable}}
+
+void hostfoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg;
+}
+
+__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: explicitly defaulted ctor
+
+struct A6_with_device_ctor {
+  __device__ A6_with_device_ctor() {}
+};
+
+struct B6_with_defaulted_ctor : A6_with_device_ctor {
+  using A6_with_device_ctor::A6_with_device_ctor;
+  __host__ B6_with_defaulted_ctor() = default;
+};
+
+// expected-note@-3 {{explicitly defaulted function was implicitly deleted here}}
+// expected-note@-6 {{default constructor of 'B6_with_defaulted_ctor' is implicitly deleted because base class 'A6_with_device_ctor' has no default constructor}}
+
+void hostfoo6() {
+  B6_with_defaulted_ctor b; // expected-error {{call to implicitly-deleted default constructor}}
+}
+
+__device__ void devicefoo6() {
+  B6_with_defaulted_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 7: inherited copy assignment operator
+
+struct A7_with_copy_assign {
+  A7_with_copy_assign() {}
+  __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
+};
+
+struct B7_with_copy_assign : A7_with_copy_assign {
+  using A7_with_copy_assign::A7_with_copy_assign;
+};
+
+// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo7() {
+  B7_with_copy_assign b1, b2;
+  b1 = b2; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 8: inherited move assignment operator
+
+// definitions for std::move
+namespace std {
+inline namespace foo {
+template <class T> struct remove_reference { typedef T type; };
+template <class T> struct remove_reference<T&> { typedef T type; };
+template <class T> struct remove_reference<T&&> { typedef T type; };
+
+template <class T> typename remove_reference<T>::type&& move(T&& t);
+}
+}
+
+struct A8_with_move_assign {
+  A8_with_move_assign() {}
+  __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
+  __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
+};
+
+struct B8_with_move_assign : A8_with_move_assign {
+  using A8_with_move_assign::A8_with_move_assign;
+};
+
+// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo8() {
+  B8_with_move_assign b1, b2;
+  b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
+}

test/SemaCUDA/inherited-ctor.cu

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+
+// Inherit a valid non-default ctor.
+namespace NonDefaultCtorValid {
+  struct A {
+    A(const int &x) {}
+  };
+
+  struct B : A {
+    using A::A;
+  };
+
+  struct C {
+    struct B b;
+    C() : b(0) {}
+  };
+
+  void test() {
+    B b(0);
+    C c;
+  }
+}
+
+// Inherit an invalid non-default ctor.
+// The inherited ctor is invalid because it is unable to initialize s.
+namespace NonDefaultCtorInvalid {
+  struct S {
+    S() = delete;
+  };
+  struct A {
+    A(const int &x) {}
+  };
+
+  struct B : A {
+    using A::A;
+    S s;
+  };
+
+  struct C {
+    struct B b;
+    C() : b(0) {} // expected-error{{constructor inherited by 'B' from base class 'A' is implicitly deleted}}
+                  // expected-note@-6{{constructor inherited by 'B' is implicitly deleted because field 's' has a deleted corresponding constructor}}
+                  // expected-note@-15{{'S' has been explicitly marked deleted here}}
+  };
+}
+
+// Inherit a valid default ctor.
+namespace DefaultCtorValid {
+  struct A {
+    A() {}
+  };
+
+  struct B : A {
+    using A::A;
+  };
+
+  struct C {
+    struct B b;
+    C() {}
+  };
+
+  void test() {
+    B b;
+    C c;
+  }
+}
+
+// Inherit an invalid default ctor.
+// The inherited ctor is invalid because it is unable to initialize s.
+namespace DefaultCtorInvalid {
+  struct S {
+    S() = delete;
+  };
+  struct A {
+    A() {}
+  };
+
+  struct B : A {
+    using A::A;
+    S s;
+  };
+
+  struct C {
+    struct B b;
+    C() {} // expected-error{{call to implicitly-deleted default constructor of 'struct B'}}
+           // expected-note@-6{{default constructor of 'B' is implicitly deleted because field 's' has a deleted default constructor}}
+           // expected-note@-15{{'S' has been explicitly marked deleted here}}
+  };
+}