r268982 - [CUDA] Restrict init of local __shared__ variables to empty constructors only.
Artem Belevich via cfe-commits
cfe-commits at lists.llvm.org
Mon May 9 15:09:56 PDT 2016
Author: tra
Date: Mon May 9 17:09:56 2016
New Revision: 268982
URL: http://llvm.org/viewvc/llvm-project?rev=268982&view=rev
Log:
[CUDA] Restrict init of local __shared__ variables to empty constructors only.
Allow only empty constructors for local __shared__ variables in a way
identical to restrictions imposed on dynamic initializers for global
variables on device.
Differential Revision: http://reviews.llvm.org/D20039
Modified:
cfe/trunk/lib/CodeGen/CGDecl.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/CodeGenCUDA/address-spaces.cu
cfe/trunk/test/CodeGenCUDA/device-var-init.cu
Modified: cfe/trunk/lib/CodeGen/CGDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDecl.cpp?rev=268982&r1=268981&r2=268982&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDecl.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDecl.cpp Mon May 9 17:09:56 2016
@@ -371,8 +371,15 @@ void CodeGenFunction::EmitStaticVarDecl(
llvm::GlobalVariable *var =
cast<llvm::GlobalVariable>(addr->stripPointerCasts());
+
+ // CUDA's local and local static __shared__ variables should not
+ // have any non-empty initializers. This is ensured by Sema.
+ // Whatever initializer such variable may have when it gets here is
+ // a no-op and should not be emitted.
+ bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ D.hasAttr<CUDASharedAttr>();
// If this value has an initializer, emit it.
- if (D.getInit())
+ if (D.getInit() && !isCudaSharedVar)
var = AddInitializerToStaticVarDecl(D, var);
var->setAlignment(alignment.getQuantity());
@@ -1874,4 +1881,3 @@ void CodeGenModule::EmitOMPDeclareReduct
return;
getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
}
-
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=268982&r1=268981&r2=268982&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Mon May 9 17:09:56 2016
@@ -10414,14 +10414,15 @@ Sema::FinalizeDeclaration(Decl *ThisDecl
// Perform check for initializers of device-side global variables.
// CUDA allows empty constructors as initializers (see E.2.3.1, CUDA
- // 7.5). CUDA also allows constant initializers for __constant__ and
- // __device__ variables.
+ // 7.5). We must also apply the same checks to all __shared__
+ // variables whether they are local or not. CUDA also allows
+ // constant initializers for __constant__ and __device__ variables.
if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice) {
const Expr *Init = VD->getInit();
- const bool IsGlobal = VD->hasGlobalStorage() && !VD->isStaticLocal();
- if (Init && IsGlobal &&
+ if (Init && VD->hasGlobalStorage() &&
(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
VD->hasAttr<CUDASharedAttr>())) {
+ assert((!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>()));
bool AllowedInit = false;
if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
AllowedInit =
Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=268982&r1=268981&r2=268982&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Mon May 9 17:09:56 2016
@@ -25,8 +25,6 @@ struct MyStruct {
// 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*)
@@ -94,32 +92,3 @@ __device__ float *func5() {
}
// CHECK: define float* @_Z5func5v()
// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
-
-struct StructWithCtor {
- __device__ StructWithCtor(): data(1) {}
- __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {}
- __device__ int getData() { return data; }
- int data;
-};
-
-__device__ int construct_shared_struct() {
-// CHECK-LABEL: define i32 @_Z23construct_shared_structv()
- __shared__ StructWithCtor s;
-// CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
- __shared__ StructWithCtor t(s);
-// CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
- 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;
Modified: cfe/trunk/test/CodeGenCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-var-init.cu?rev=268982&r1=268981&r2=268982&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Mon May 9 17:09:56 2016
@@ -63,6 +63,8 @@ struct NCF {
// static in-class field initializer. NVCC does not allow it, but
// clang generates static initializer for this, so we'll accept it.
+// We still can't use it on __shared__ vars as they don't allow *any*
+// initializers.
struct NCFS {
int ncfs = 3;
};
@@ -367,8 +369,13 @@ __device__ void df() {
T_B_NEC t_b_nec;
T_F_NEC t_f_nec;
T_FA_NEC t_fa_nec;
- static __shared__ UC s_uc;
+ static __shared__ EC s_ec;
+ static __shared__ ETC s_etc;
#if ERROR_CASE
+ static __shared__ NCFS s_ncfs;
+ // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
+ static __shared__ UC s_uc;
+ // expected-error at -1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
// expected-error at -1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}}
static __constant__ int dc;
@@ -394,7 +401,8 @@ __device__ void df() {
// CHECK: call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
// CHECK: call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
// CHECK: call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
-// CHECK: call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
+// CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+// CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
// CHECK: ret void
// We should not emit global init function.
More information about the cfe-commits
mailing list