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

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. Clang already deals with device-side constructors (see D15305). This patch enforces similar rules for destructors. Differential Revision: http://reviews.llvm.org/D20140 llvm-svn: 270108
2016-05-19 20:13:53 +00:00 · 2016-05-19 20:13:53 +00:00 · 3650bbeebc
parent 85b6f63f42
commit 3650bbeebc
5 changed files with 181 additions and 5 deletions
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@ -9036,6 +9036,7 @@ public:
  /// \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
  //@{
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@ -372,7 +372,7 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
    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()))
@ -384,6 +384,54 @@ bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
  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
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@ -10442,6 +10442,12 @@ Sema::FinalizeDeclaration(Decl *ThisDecl) {
        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
--- a/clang/test/CodeGenCUDA/device-var-init.cu
+++ b/clang/test/CodeGenCUDA/device-var-init.cu
@ -24,6 +24,16 @@ struct EC {
  __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...) {}
@ -35,6 +45,12 @@ struct UC {
  __device__ UC();
 };

+// undefined destructor -- not allowed
+struct UD {
+  int ud;
+  __device__ ~UD();
+};
+
 // empty constructor w/ initializer list -- not allowed
 struct ECI {
  int eci;
@ -47,12 +63,23 @@ struct 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 {
@ -107,6 +134,20 @@ __shared__ EC s_ec;
 __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;
@ -180,6 +221,17 @@ 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

@ -190,16 +242,26 @@ __device__ void df() {
  // 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;
@ -226,6 +288,12 @@ __device__ void df() {
  // 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;
@ -234,9 +302,17 @@ __device__ void df() {
  // 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.
--- a/clang/test/SemaCUDA/device-var-init.cu
+++ b/clang/test/SemaCUDA/device-var-init.cu
@ -58,6 +58,13 @@ __shared__ UC s_uc;
 __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;
@ -72,6 +79,13 @@ __shared__ NEC s_nec;
 __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;
@ -79,6 +93,13 @@ __shared__ NCV s_ncv;
 __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;
@ -152,13 +173,37 @@ __shared__ T_FA_NEC s_t_fa_nec;
 __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;