[PATCH] D12739: [CUDA] Allow trivial constructors as initializer for __shared__ variables.

[Artem Belevich via cfe-commits]

tra created this revision.
tra added reviewers: jingyue, eliben, wengxt, jholewinski.
tra added a subscriber: cfe-commits.

r245786 (D12241) disabled all initializers for __shared__ variables. It also 
prevents placement of records in __shared__ space as they have constructor as 
their initializer.





http://reviews.llvm.org/D12739

Files:
  lib/CodeGen/CodeGenModule.cpp
  test/CodeGenCUDA/address-spaces.cu

Index: test/CodeGenCUDA/address-spaces.cu
===================================================================
--- test/CodeGenCUDA/address-spaces.cu
+++ test/CodeGenCUDA/address-spaces.cu
@@ -25,6 +25,8 @@
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
+// CHECK: @c = addrspace(3) global %struct.c undef
+// CHECK  @d = addrspace(3) global %struct.d undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
@@ -117,3 +119,15 @@
   return t.getData();
 // CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* 
addrspacecast (%struct.StructWithCtor addrspace(3)* 
@_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
 }
+
+// Make sure we allow __shared__ structures with default or empty constructors.
+struct c {
+  int i;
+};
+__shared__ struct c c;
+
+struct d {
+  int i;
+  d() {}
+};
+__shared__ struct d d;
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -2151,8 +2151,10 @@
   if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
       && D->hasAttr<CUDASharedAttr>()) {
     if (InitExpr) {
-      Error(D->getLocation(),
-            "__shared__ variable cannot have an initialization.");
+      const auto *C = dyn_cast<CXXConstructExpr>(InitExpr);
+      if (C == nullptr || !C->getConstructor()->hasTrivialBody())
+        Error(D->getLocation(),
+              "__shared__ variable cannot have an initialization.");
     }
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   } else if (!InitExpr) {



Index: test/CodeGenCUDA/address-spaces.cu
===================================================================
--- test/CodeGenCUDA/address-spaces.cu
+++ test/CodeGenCUDA/address-spaces.cu
@@ -25,6 +25,8 @@
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
+// CHECK: @c = addrspace(3) global %struct.c undef
+// CHECK  @d = addrspace(3) global %struct.d undef
 
 __device__ void foo() {
   // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
@@ -117,3 +119,15 @@
   return t.getData();
 // CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
 }
+
+// Make sure we allow __shared__ structures with default or empty constructors.
+struct c {
+  int i;
+};
+__shared__ struct c c;
+
+struct d {
+  int i;
+  d() {}
+};
+__shared__ struct d d;
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -2151,8 +2151,10 @@
   if (getLangOpts().CPlusPlus && getLangOpts().CUDAIsDevice
       && D->hasAttr<CUDASharedAttr>()) {
     if (InitExpr) {
-      Error(D->getLocation(),
-            "__shared__ variable cannot have an initialization.");
+      const auto *C = dyn_cast<CXXConstructExpr>(InitExpr);
+      if (C == nullptr || !C->getConstructor()->hasTrivialBody())
+        Error(D->getLocation(),
+              "__shared__ variable cannot have an initialization.");
     }
     Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   } else if (!InitExpr) {

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits