diff options
author | Michael Liao <michael.hliao@gmail.com> | 2019-05-29 17:23:27 +0000 |
---|---|---|
committer | Michael Liao <michael.hliao@gmail.com> | 2019-05-29 17:23:27 +0000 |
commit | 4b7a713accdce4e46a45f34bd8705e43b495ddec (patch) | |
tree | 24b70efb61c7c31df61306d94d90ad09883a6947 /clang/test/CodeGenCUDA/device-var-init.cu | |
parent | c98b288b030cde85ae8d4e188677546efcbe0d29 (diff) | |
download | bcm5719-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.cu | 10 |
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 |