From: Artem Belevich Date: Mon, 9 May 2016 22:09:56 +0000 (+0000) Subject: [CUDA] Restrict init of local __shared__ variables to empty constructors only. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=e0c569d59703233005807ac4041210238ed18dfc;p=clang [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@268982 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index c39b980ef2..b6ab75f67c 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -371,8 +371,15 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D, llvm::GlobalVariable *var = cast(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(); // 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::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D, return; getOpenMPRuntime().emitUserDefinedReduction(CGF, D); } - diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp index a2c28de5f5..020b474c18 100644 --- a/lib/Sema/SemaDecl.cpp +++ b/lib/Sema/SemaDecl.cpp @@ -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() || VD->hasAttr() || VD->hasAttr())) { + assert((!VD->isStaticLocal() || VD->hasAttr())); bool AllowedInit = false; if (const CXXConstructExpr *CE = dyn_cast(Init)) AllowedInit = diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index 4670f46db9..449529bb24 100644 --- a/test/CodeGenCUDA/address-spaces.cu +++ b/test/CodeGenCUDA/address-spaces.cu @@ -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; diff --git a/test/CodeGenCUDA/device-var-init.cu b/test/CodeGenCUDA/device-var-init.cu index 23c9fe1376..8cf9278c8f 100644 --- a/test/CodeGenCUDA/device-var-init.cu +++ b/test/CodeGenCUDA/device-var-init.cu @@ -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@-1 {{initialization is not supported for __shared__ variables.}} + static __shared__ UC s_uc; + // 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; @@ -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.