Author: tra Date: Mon May 9 14:36:08 2016 New Revision: 268962 URL: http://llvm.org/viewvc/llvm-project?rev=268962&view=rev Log: [CUDA] Only __shared__ variables can be static local on device side.
According to CUDA programming guide (v7.5): > E.2.9.4: Within the body of a device or global function, only > shared variables may be declared with static storage class. Differential Revision: http://reviews.llvm.org/D20034 Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaDecl.cpp cfe/trunk/test/CodeGenCUDA/address-spaces.cu cfe/trunk/test/CodeGenCUDA/device-var-init.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=268962&r1=268961&r2=268962&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Mon May 9 14:36:08 2016 @@ -6554,7 +6554,9 @@ def err_dynamic_var_init : Error< "__device__, __constant__, and __shared__ variables.">; def err_shared_var_init : Error< "initialization is not supported for __shared__ variables.">; - +def err_device_static_local_var : Error< + "Within a __device__/__global__ function, " + "only __shared__ variables may be marked \"static\"">; def warn_non_pod_vararg_with_format_string : Warning< "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic " "%select{function|block|method|constructor}2; expected type from format " Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=268962&r1=268961&r2=268962&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Mon May 9 14:36:08 2016 @@ -10391,15 +10391,24 @@ Sema::FinalizeDeclaration(Decl *ThisDecl } } - // Static locals inherit dll attributes from their function. if (VD->isStaticLocal()) { if (FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(VD->getParentFunctionOrMethod())) { + // Static locals inherit dll attributes from their function. if (Attr *A = getDLLAttr(FD)) { auto *NewAttr = cast<InheritableAttr>(A->clone(getASTContext())); NewAttr->setInherited(true); VD->addAttr(NewAttr); } + // CUDA E.2.9.4: Within the body of a __device__ or __global__ + // function, only __shared__ variables may be declared with + // static storage class. + if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice && + (FD->hasAttr<CUDADeviceAttr>() || FD->hasAttr<CUDAGlobalAttr>()) && + !VD->hasAttr<CUDASharedAttr>()) { + Diag(VD->getLocation(), diag::err_device_static_local_var); + VD->setInvalidDecl(); + } } } Modified: cfe/trunk/test/CodeGenCUDA/address-spaces.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/address-spaces.cu?rev=268962&r1=268961&r2=268962&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/address-spaces.cu (original) +++ cfe/trunk/test/CodeGenCUDA/address-spaces.cu Mon May 9 14:36:08 2016 @@ -38,14 +38,6 @@ __device__ void foo() { // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*) k++; - static int li; - // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*) - li++; - - __constant__ int lj; - // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*) - lj++; - __shared__ int lk; // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*) lk++; 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=268962&r1=268961&r2=268962&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original) +++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Mon May 9 14:36:08 2016 @@ -368,6 +368,14 @@ __device__ void df() { T_F_NEC t_f_nec; T_FA_NEC t_fa_nec; static __shared__ UC s_uc; +#if ERROR_CASE + static __device__ int ds; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} + static __constant__ int dc; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} + static int v; + // expected-error@-1 {{Within a __device__/__global__ function, only __shared__ variables may be marked "static"}} +#endif } // CHECK: call void @_ZN2ECC1Ev(%struct.EC* %ec) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
