r338188 - [CUDA][HIP] Allow function-scope static const variable
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 27 20:05:25 PDT 2018
Author: yaxunl
Date: Fri Jul 27 20:05:25 2018
New Revision: 338188
URL: http://llvm.org/viewvc/llvm-project?rev=338188&view=rev
Log:
[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
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/CodeGenCUDA/device-var-init.cu
cfe/trunk/test/SemaCUDA/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=338188&r1=338187&r2=338188&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Jul 27 20:05:25 2018
@@ -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">;
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=338188&r1=338187&r2=338188&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Jul 27 20:05:25 2018
@@ -3176,6 +3176,10 @@ LangAS CodeGenModule::GetGlobalVarAddres
return LangAS::cuda_constant;
else if (D && D->hasAttr<CUDASharedAttr>())
return LangAS::cuda_shared;
+ else if (D && D->hasAttr<CUDADeviceAttr>())
+ return LangAS::cuda_device;
+ else if (D && D->getType().isConstQualified())
+ return LangAS::cuda_constant;
else
return LangAS::cuda_device;
}
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=338188&r1=338187&r2=338188&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Fri Jul 27 20:05:25 2018
@@ -11914,14 +11914,25 @@ void Sema::FinalizeDeclaration(Decl *Thi
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<CUDASharedAttr>() &&
- 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<CUDASharedAttr>())
+ return;
+ if (VD->getType().isConstQualified() &&
+ !(VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+ return;
+ if (CUDADiagIfDeviceCode(VD->getLocation(),
+ diag::err_device_static_local_var)
+ << CurrentCUDATarget())
+ VD->setInvalidDecl();
+ }();
}
}
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=338188&r1=338187&r2=338188&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Fri Jul 27 20:05:25 2018
@@ -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()
Modified: cfe/trunk/test/SemaCUDA/device-var-init.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/device-var-init.cu?rev=338188&r1=338187&r2=338188&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/device-var-init.cu (original)
+++ cfe/trunk/test/SemaCUDA/device-var-init.cu Fri Jul 27 20:05:25 2018
@@ -207,17 +207,22 @@ __device__ void df_sema() {
// expected-error at -1 {{initialization is not supported for __shared__ variables.}}
static __device__ int ds;
- // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error at -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 at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error at -1 {{within a __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
static int v;
- // expected-error at -1 {{within a __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error at -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 at -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 at -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 at -2 {{within a __host__ __device__ function, only __shared__ variables may be marked 'static'}}
+ // expected-error at -2 {{within a __host__ __device__ function, only __shared__ variables or const variables without device memory qualifier may be marked 'static'}}
#endif
}
More information about the cfe-commits
mailing list