diff options
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 3 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/device-var-init.cu | 10 |
2 files changed, 12 insertions, 1 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 8c9e240a680..8cfb4e60e0d 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -3869,7 +3869,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D, // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())." if (GV && LangOpts.CUDA) { if (LangOpts.CUDAIsDevice) { - if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) + if (Linkage != llvm::GlobalValue::InternalLinkage && + (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>())) GV->setExternallyInitialized(true); } else { // Host-side shadows of external declarations of device-side 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 |
