diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -1089,6 +1089,28 @@ llvm::Function *CGNVCUDARuntime::finalizeModule() { if (CGM.getLangOpts().CUDAIsDevice) { transformManagedVars(); + + // Mark ODR-used device variables as compiler used to prevent it from being + // eliminated by optimization. This is necessary for device variables + // ODR-used by host functions. Sema correctly marks them as ODR-used no + // matter whether they are ODR-used by device or host functions. + // + // We do not need to do this if the variable has used attribute since it + // has already been added. + // + // Static device variables have been externalized at this point, therefore + // variables with LLVM private or internal linkage need not be added. + for (auto &&Info : DeviceVars) { + auto Kind = Info.Flags.getKind(); + if (!Info.Var->isDeclaration() && + !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) && + (Kind == DeviceVarFlags::Variable || + Kind == DeviceVarFlags::Surface || + Kind == DeviceVarFlags::Texture) && + Info.D->isUsed() && !Info.D->hasAttr()) { + CGM.addCompilerUsedGlobal(Info.Var); + } + } return nullptr; } return makeModuleCtorFunction(); diff --git a/clang/test/CodeGenCUDA/host-used-device-var.cu b/clang/test/CodeGenCUDA/host-used-device-var.cu new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/host-used-device-var.cu @@ -0,0 +1,47 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \ +// RUN: -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \ +// RUN: | FileCheck %s + +#include "Inputs/cuda.h" + +// Check device variables used by neither host nor device functioins are not kept. + +// CHECK-NOT: @v1 +__device__ int v1; + +// CHECK-NOT: @v2 +__constant__ int v2; + +// CHECK-NOT: @_ZL2v3 +static __device__ int v3; + +// Check device variables used by host functions are kept. + +// CHECK-DAG: @u1 +__device__ int u1; + +// CHECK-DAG: @u2 +__constant__ int u2; + +// Check host-used static device var is in llvm.compiler.used. +// CHECK-DAG: @_ZL2u3 +static __device__ int u3; + +// Check device-used static device var is emitted but is not in llvm.compiler.used. +// CHECK-DAG: @_ZL2u4 +static __device__ int u4; + +// Check device variables with used attribute are always kept. +// CHECK-DAG: @u5 +__device__ __attribute__((used)) int u5; + +int fun1() { + return u1 + u2 + u3; +} + +__global__ void kern1(int **x) { + *x = &u4; +} +// Check the exact list of variables to ensure @_ZL2u4 is not among them. +// CHECK: @llvm.compiler.used = {{[^@]*}} @_ZL2u3 {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5