Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-01-25 10:48:14 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[CUDA][HIP] Allow function-scope static const variable

Browse Source 
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
This commit is contained in:
Yaxun Liu 2018-07-28 03:05:25 +00:00
parent 39e5137f43
commit a4005e13f7
5 changed files with 40 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

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

4
clang/lib/CodeGen/CodeGenModule.cpp
View File

@ -3176,6 +3176,10 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
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;
}

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

@ -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<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();
}();
}
}

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

@ -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()

13
clang/test/SemaCUDA/device-var-init.cu
View File

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