maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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

llvm-svn: 268982
This commit is contained in:
Artem Belevich 2016-05-09 22:09:56 +00:00
parent a407f543c0
commit 4d430badeb
4 changed files with 23 additions and 39 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
clang/lib/CodeGen/CGDecl.cpp
View File

@ -371,8 +371,15 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D,
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::EmitOMPDeclareReduction(const OMPDeclareReductionDecl *D,
return;
getOpenMPRuntime().emitUserDefinedReduction(CGF, D);
}

9
clang/lib/Sema/SemaDecl.cpp
View File

@ -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 =

31
clang/test/CodeGenCUDA/address-spaces.cu
View File

@ -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;

12
clang/test/CodeGenCUDA/device-var-init.cu
View File

@ -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.