[PATCH] D49931: [CUDA][HIP] Allow function-scope static const variable
Yaxun Liu via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 27 12:39:53 PDT 2018
yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
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.
https://reviews.llvm.org/D49931
Files:
lib/CodeGen/CodeGenModule.cpp
lib/Sema/SemaDecl.cpp
test/CodeGenCUDA/device-var-init.cu
Index: test/CodeGenCUDA/device-var-init.cu
===================================================================
--- test/CodeGenCUDA/device-var-init.cu
+++ test/CodeGenCUDA/device-var-init.cu
@@ -112,6 +112,9 @@
// 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 @@
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()
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -11914,10 +11914,15 @@
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>() &&
+ // 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 &&
+ !(VD->hasAttr<CUDASharedAttr>() ||
+ VD->getType().isConstQualified()) &&
CUDADiagIfDeviceCode(VD->getLocation(),
diag::err_device_static_local_var)
<< CurrentCUDATarget())
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -3176,6 +3176,10 @@
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;
}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D49931.157738.patch
Type: text/x-patch
Size: 2805 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20180727/757bff34/attachment-0001.bin>
More information about the cfe-commits
mailing list