From: Yaxun Liu Date: Sat, 28 Jul 2018 03:05:25 +0000 (+0000) Subject: [CUDA][HIP] Allow function-scope static const variable X-Git-Tag: llvmorg-7.0.0-rc1~366 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=a4005e13f72ece696cce5d190fb73abec116d18a;p=platform%2Fupstream%2Fllvm.git [CUDA][HIP] Allow function-scope static const variable CUDA 8.0 E.3.9.4 says: Within the body of a __device__ or __global__ function, only __shared__ variables or variables without any device memory qualifiers may be declared with static storage class. It is unclear how a function-scope non-const static variable without device memory qualifier is implemented, therefore only static const variable without device memory qualifier is allowed, which can be emitted as a global variable in constant address space. Currently clang only allows function-scope static variable with __shared__ qualifier. This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space. Differential Revision: https://reviews.llvm.org/D49931 llvm-svn: 338188 --- diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index d0a2bec..b2ad273 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7129,7 +7129,8 @@ def err_shared_var_init : Error< "initialization is not supported for __shared__ variables.">; def err_device_static_local_var : Error< "within a %select{__device__|__global__|__host__|__host__ __device__}0 " - "function, only __shared__ variables may be marked 'static'">; + "function, only __shared__ variables or const variables without device " + "memory qualifier may be marked 'static'">; def err_cuda_vla : Error< "cannot use variable-length arrays in " "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ecdf78d..ae2c134 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3176,6 +3176,10 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) { return LangAS::cuda_constant; else if (D && D->hasAttr()) return LangAS::cuda_shared; + else if (D && D->hasAttr()) + return LangAS::cuda_device; + else if (D && D->getType().isConstQualified()) + return LangAS::cuda_constant; else return LangAS::cuda_device; } diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 5554282..6bee72a 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -11914,14 +11914,25 @@ void Sema::FinalizeDeclaration(Decl *ThisDecl) { 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 && !VD->hasAttr() && - CUDADiagIfDeviceCode(VD->getLocation(), - diag::err_device_static_local_var) - << CurrentCUDATarget()) - VD->setInvalidDecl(); + // CUDA 8.0 E.3.9.4: Within the body of a __device__ or __global__ + // function, only __shared__ variables or variables without any device + // memory qualifiers may be declared with static storage class. + // Note: It is unclear how a function-scope non-const static variable + // without device memory qualifier is implemented, therefore only static + // const variable without device memory qualifier is allowed. + [&]() { + if (!getLangOpts().CUDA) + return; + if (VD->hasAttr()) + return; + if (VD->getType().isConstQualified() && + !(VD->hasAttr() || VD->hasAttr())) + return; + if (CUDADiagIfDeviceCode(VD->getLocation(), + diag::err_device_static_local_var) + << CurrentCUDATarget()) + VD->setInvalidDecl(); + }(); } } diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu index 9f788b7..f96e42d 100644 --- a/clang/test/CodeGenCUDA/device-var-init.cu +++ b/clang/test/CodeGenCUDA/device-var-init.cu @@ -112,6 +112,9 @@ __constant__ EC_I_EC c_ec_i_ec; // CHECK: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef // CHECK: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef +// CHECK: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5] +// CHECK: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123 + // We should not emit global initializers for device-side variables. // CHECK-NOT: @__cxx_global_var_init @@ -234,6 +237,9 @@ __device__ void df() { static __shared__ ETC s_etc; // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*)) + static const int const_array[] = {1, 2, 3, 4, 5}; + static const int const_int = 123; + // anchor point separating constructors and destructors df(); // CHECK: call void @_Z2dfv() diff --git a/clang/test/SemaCUDA/device-var-init.cu b/clang/test/SemaCUDA/device-var-init.cu index 46cb90d..dd5d19a 100644 --- a/clang/test/SemaCUDA/device-var-init.cu +++ b/clang/test/SemaCUDA/device-var-init.cu @@ -207,17 +207,22 @@ __device__ void df_sema() { // expected-error@-1 {{initialization is not supported for __shared__ variables.}} static __device__ int ds; - // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}} + // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static __constant__ int dc; - // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}} + // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} static int v; - // expected-error@-1 {{within a __device__ function, only __shared__ variables may be marked 'static'}} + // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} + static const int cv = 1; + static const __device__ int cds = 1; + // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} + static const __constant__ int cdc = 1; + // expected-error@-1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} } __host__ __device__ void hd_sema() { static int x = 42; #ifdef __CUDA_ARCH__ - // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}} + // expected-error@-2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}} #endif }