summaryrefslogtreecommitdiffstats
path: root/clang/test/CodeGenCUDA/device-var-init.cu
diff options
context:
space:
mode:
authorMichael Liao <michael.hliao@gmail.com>2019-05-29 17:23:27 +0000
committerMichael Liao <michael.hliao@gmail.com>2019-05-29 17:23:27 +0000
commit4b7a713accdce4e46a45f34bd8705e43b495ddec (patch)
tree24b70efb61c7c31df61306d94d90ad09883a6947 /clang/test/CodeGenCUDA/device-var-init.cu
parentc98b288b030cde85ae8d4e188677546efcbe0d29 (diff)
downloadbcm5719-llvm-4b7a713accdce4e46a45f34bd8705e43b495ddec.tar.gz
bcm5719-llvm-4b7a713accdce4e46a45f34bd8705e43b495ddec.zip
[CUDA][HIP] Skip setting `externally_initialized` for static device variables.
Summary: - By declaring device variables as `static`, we assume they won't be addressable from the host side. Thus, no `externally_initialized` is required. Reviewers: yaxunl Subscribers: cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D62603 llvm-svn: 361994
Diffstat (limited to 'clang/test/CodeGenCUDA/device-var-init.cu')
-rw-r--r--clang/test/CodeGenCUDA/device-var-init.cu10
1 files changed, 10 insertions, 0 deletions
diff --git a/clang/test/CodeGenCUDA/device-var-init.cu b/clang/test/CodeGenCUDA/device-var-init.cu
index af42e698cfe..fd236bb842d 100644
--- a/clang/test/CodeGenCUDA/device-var-init.cu
+++ b/clang/test/CodeGenCUDA/device-var-init.cu
@@ -33,6 +33,16 @@ __device__ int d_v_i = 1;
// DEVICE: @d_v_i = addrspace(1) externally_initialized global i32 1,
// HOST: @d_v_i = internal global i32 undef,
+// For `static` device variables, assume they won't be addressed from the host
+// side.
+static __device__ int d_s_v_i = 1;
+// DEVICE: @_ZL7d_s_v_i = internal addrspace(1) global i32 1,
+
+// Dummy function to keep static variables referenced.
+__device__ int foo() {
+ return d_s_v_i;
+}
+
// trivial constructor -- allowed
__device__ T d_t;
// DEVICE: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
OpenPOWER on IntegriCloud