Author: hliao Date: Wed May 29 10:23:27 2019 New Revision: 361994 URL: http://llvm.org/viewvc/llvm-project?rev=361994&view=rev Log: [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 Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp cfe/trunk/test/CodeGenCUDA/device-var-init.cu Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=361994&r1=361993&r2=361994&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original) +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Wed May 29 10:23:27 2019 @@ -3869,7 +3869,8 @@ void CodeGenModule::EmitGlobalVarDefinit // / 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 Modified: cfe/trunk/test/CodeGenCUDA/device-var-init.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/device-var-init.cu?rev=361994&r1=361993&r2=361994&view=diff ============================================================================== --- cfe/trunk/test/CodeGenCUDA/device-var-init.cu (original) +++ cfe/trunk/test/CodeGenCUDA/device-var-init.cu Wed May 29 10:23:27 2019 @@ -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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
