[CUDA] Do not allow non-empty destructors for global device-side variables.

[tra]

Summary

According to Cuda Programming guide (v7.5, E2.3.1):

device, constant and shared variables defined in namespace
scope, that are of class type, cannot have a non-empty constructor or a
non-empty destructor.

A destructor for a class is considered empty at a point in the translation unit, if it is either a trivial destructor or it satisfies all of the following conditions:

• The destructor function has been defined.
• The destructor function body is an empty compound statement.
• Its class has no virtual functions and no virtual base classes.
• The destructors of all base classes of its class can be considered empty.
• For all the nonstatic data members of its class that are of class type (or array thereof), the destructor can be considered empty.

Clang already deals with device-side constructors (see D15305).
This patch enforces similar rules for destructors.

[jlebar]

lgtm, but I'd like Richard to sign off on this too.

lib/Sema/SemaDecl.cpp

10438 ↗

(On Diff #56829)

if there

cfe/trunk/include/clang/Sema/Sema.h

bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                                CXXSpecialMember CSM,
                                                CXXMethodDecl *MemberDecl,
                                                bool ConstRHS,
                                                bool Diagnose);
 
   /// \return true if \p CD can be considered empty according to CUDA
   /// (E.2.3.1 in CUDA 7.5 Programming guide).
   bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
+  bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
 
   /// \name Code completion
   //@{
   /// \brief Describes the context in which code completion occurs.
   enum ParserCompletionContext {
     /// \brief Code completion occurs at top-level or namespace context.
     PCC_Namespace,
     /// \brief Code completion occurs within a class, struct, or union.

cfe/trunk/lib/Sema/SemaCUDA.cpp

bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
   if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
     return false;
 
   // Its class has no virtual functions and no virtual base classes.
   if (CD->getParent()->isDynamicClass())
     return false;
 
   // The only form of initializer allowed is an empty constructor.
-  // This will recursively checks all base classes and member initializers
+  // This will recursively check all base classes and member initializers
   if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
         if (const CXXConstructExpr *CE =
                 dyn_cast<CXXConstructExpr>(CI->getInit()))
           return isEmptyCudaConstructor(Loc, CE->getConstructor());
         return false;
       }))
     return false;
 
   return true;
 }
 
+bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
+  // No destructor -> no problem.
+  if (!DD)
+    return true;
+
+  if (!DD->isDefined() && DD->isTemplateInstantiation())
+    InstantiateFunctionDefinition(Loc, DD->getFirstDecl());
+
+  // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
+  // empty at a point in the translation unit, if it is either a
+  // trivial constructor
+  if (DD->isTrivial())
+    return true;
+
+  // ... or it satisfies all of the following conditions:
+  // The destructor function has been defined.
+  // and the function body is an empty compound statement.
+  if (!DD->hasTrivialBody())
+    return false;
+
+  const CXXRecordDecl *ClassDecl = DD->getParent();
+
+  // Its class has no virtual functions and no virtual base classes.
+  if (ClassDecl->isDynamicClass())
+    return false;
+
+  // Only empty destructors are allowed. This will recursively check
+  // destructors for all base classes...
+  if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
+        if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
+          return isEmptyCudaDestructor(Loc, RD->getDestructor());
+        return true;
+      }))
+    return false;
+
+  // ... and member fields.
+  if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
+        if (CXXRecordDecl *RD = Field->getType()
+                                    ->getBaseElementTypeUnsafe()
+                                    ->getAsCXXRecordDecl())
+          return isEmptyCudaDestructor(Loc, RD->getDestructor());
+        return true;
+      }))
+    return false;
+
+  return true;
+}
+
 // With -fcuda-host-device-constexpr, an unattributed constexpr function is
 // treated as implicitly __host__ __device__, unless:
 //  * it is a variadic function (device-side variadic functions are not
 //    allowed), or
 //  * a __device__ function with this signature was already declared, in which
 //    case in which case we output an error, unless the __device__ decl is in a
 //    system header, in which case we leave the constexpr function unattributed.
 void Sema::maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *NewD,

cfe/trunk/lib/Sema/SemaDecl.cpp

if (Init && VD->hasGlobalStorage() &&
       // We'll allow constant initializers even if it's a non-empty
       // constructor according to CUDA rules. This deviates from NVCC,
       // but allows us to handle things like constexpr constructors.
       if (!AllowedInit &&
           (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
         AllowedInit = VD->getInit()->isConstantInitializer(
             Context, VD->getType()->isReferenceType());
 
+      // Also make sure that destructor, if there is one, is empty.
+      if (AllowedInit)
+        if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+          AllowedInit =
+              isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+
       if (!AllowedInit) {
         Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
                                     ? diag::err_shared_var_init
                                     : diag::err_dynamic_var_init)
             << Init->getSourceRange();
         VD->setInvalidDecl();
       }
     }

cfe/trunk/test/CodeGenCUDA/device-var-init.cu

 
 // empty constructor
 struct EC {
   int ec;
   __device__ EC() {}     // -- allowed
   __device__ EC(int) {}  // -- not allowed
 };
 
+// empty destructor
+struct ED {
+  __device__ ~ED() {}     // -- allowed
+};
+
+struct ECD {
+  __device__ ECD() {}     // -- allowed
+  __device__ ~ECD() {}    // -- allowed
+};
+
 // empty templated constructor -- allowed with no arguments
 struct ETC {
   template <typename... T> __device__ ETC(T...) {}
 };
 
 // undefined constructor -- not allowed
 struct UC {
   int uc;
   __device__ UC();
 };
 
+// undefined destructor -- not allowed
+struct UD {
+  int ud;
+  __device__ ~UD();
+};
+
 // empty constructor w/ initializer list -- not allowed
 struct ECI {
   int eci;
   __device__ ECI() : eci(1) {}
 };
 
 // non-empty constructor -- not allowed
 struct NEC {
   int nec;
   __device__ NEC() { nec = 1; }
 };
 
+// non-empty destructor -- not allowed
+struct NED {
+  int ned;
+  __device__ ~NED() { ned = 1; }
+};
+
 // no-constructor,  virtual method -- not allowed
 struct NCV {
   int ncv;
   __device__ virtual void vm() {}
 };
 
+// virtual destructor -- not allowed.
+struct VD {
+  __device__ virtual ~VD() {}
+};
+
 // dynamic in-class field initializer -- not allowed
 __device__ int f();
 struct NCF {
   int ncf = f();
 };
 
 // static in-class field initializer.  NVCC does not allow it, but
 // clang generates static initializer for this, so we'll accept it.

 
 __device__ EC d_ec;
 // CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
 __shared__ EC s_ec;
 // CHECK: @s_ec = addrspace(3) global %struct.EC undef,
 __constant__ EC c_ec;
 // CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
 
+__device__ ED d_ed;
+// CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
+__shared__ ED s_ed;
+// CHECK: @s_ed = addrspace(3) global %struct.ED undef,
+__constant__ ED c_ed;
+// CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+
+__device__ ECD d_ecd;
+// CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
+__shared__ ECD s_ecd;
+// CHECK: @s_ecd = addrspace(3) global %struct.ECD undef,
+__constant__ ECD c_ecd;
+// CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+
 __device__ ETC d_etc;
 // CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
 __shared__ ETC s_etc;
 // CHECK: @s_etc = addrspace(3) global %struct.ETC undef,
 __constant__ ETC c_etc;
 // CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
 
 __device__ NCFS d_ncfs;

 struct T_B_NEC : NEC {};
 struct T_F_NEC {
   NEC nec;
 };
 struct T_FA_NEC {
   NEC nec[2];
 };
 
+
+// Inherited from or incapsulated class with non-empty desstructor --
+// not allowed
+struct T_B_NED : NED {};
+struct T_F_NED {
+  NED ned;
+};
+struct T_FA_NED {
+  NED ned[2];
+};
+
 // We should not emit global initializers for device-side variables.
 // CHECK-NOT: @__cxx_global_var_init
 
 // Make sure that initialization restrictions do not apply to local
 // variables.
 __device__ void df() {
   T t;
   // CHECK-NOT: call
   EC ec;
   // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
+  ED ed;
+  // CHECK-NOT: call
+  ECD ecd;
+  // CHECK:   call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
   ETC etc;
   // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
   UC uc;
   // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
+  UD ud;
+  // CHECK-NOT: call
   ECI eci;
   // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
   NEC nec;
   // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  NED ned;
+  // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
   NCV ncv;
   // CHECK-NOT: call
+  VD vd;
+  // CHECK:   call void @_ZN2VDC1Ev(%struct.VD* %vd)
   NCF ncf;
   // CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
   NCFS ncfs;
   // CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
   UTC utc;
   // CHECK:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
   NETC netc;
   // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
__device__ void df() {
   T_V_T t_v_t;
   // CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
   T_B_NEC t_b_nec;
   // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
   T_F_NEC t_f_nec;
   // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
   T_FA_NEC t_fa_nec;
   // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+  T_B_NED t_b_ned;
+  // CHECK-NOT: call
+  T_F_NED t_f_ned;
+  // CHECK-NOT: call
+  T_FA_NED t_fa_ned;
+  // CHECK-NOT: call
   static __shared__ EC s_ec;
   // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
   static __shared__ ETC s_etc;
   // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
 
   // anchor point separating constructors and destructors
   df(); // CHECK: call void @_Z2dfv()
 
-  // CHECK-NOT: call
+  // Verify that we only call non-empty destructors
+  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned) #6
+  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned) #6
+  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned) #6
+  // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
+  // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
+  // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
+  // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
+  // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
 
-  // CHECK: ret void
+  // CHECK-NEXT: ret void
 }
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I

cfe/trunk/test/SemaCUDA/device-var-init.cu

 
 __device__ UC d_uc;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ UC s_uc;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ UC c_uc;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
+__device__ UD d_ud;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ UD s_ud;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ UD c_ud;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
 __device__ ECI d_eci;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ ECI s_eci;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ ECI c_eci;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
 __device__ NEC d_nec;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ NEC s_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NEC c_nec;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
+__device__ NED d_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ NED s_ned;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ NED c_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
 __device__ NCV d_ncv;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ NCV s_ncv;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCV c_ncv;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
+__device__ VD d_vd;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ VD s_vd;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ VD c_vd;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
 __device__ NCF d_ncf;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ NCF s_ncf;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ NCF c_ncf;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
 __shared__ NCFS s_ncfs;

 
 __device__ T_FA_NEC d_t_fa_nec;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 __shared__ T_FA_NEC s_t_fa_nec;
 // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
 __constant__ T_FA_NEC c_t_fa_nec;
 // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
 
-// Make sure that initialization restrictions do not apply to local
-// variables.
+__device__ T_B_NED d_t_b_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_B_NED s_t_b_ned;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_B_NED c_t_b_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_F_NED d_t_f_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_F_NED s_t_f_ned;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_F_NED c_t_f_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+__device__ T_FA_NED d_t_fa_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+__shared__ T_FA_NED s_t_fa_ned;
+// expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+__constant__ T_FA_NED c_t_fa_ned;
+// expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, and __shared__ variables.}}
+
+// Verify that only __shared__ local variables may be static on device
+// side and that they are not allowed to be initialized.
 __device__ void df_sema() {
   static __shared__ NCFS s_ncfs;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
   static __shared__ UC s_uc;
   // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+  static __shared__ NED s_ned;
+  // expected-error@-1 {{initialization is not supported for __shared__ variables.}}
+
   static __device__ int ds;
   // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
   static __constant__ int dc;
   // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
   static int v;
   // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
 }