Author: Anshil Gandhi Date: 2021-10-15T16:16:18-06:00 New Revision: 1830ec94ac022ae0b6d6876fc2251e6b91e5931e
URL: https://github.com/llvm/llvm-project/commit/1830ec94ac022ae0b6d6876fc2251e6b91e5931e DIFF: https://github.com/llvm/llvm-project/commit/1830ec94ac022ae0b6d6876fc2251e6b91e5931e.diff LOG: Revert "[HIP] [AlwaysInliner] Disable AlwaysInliner to eliminate undefined symbols" This reverts commit 03375a3fb33b11e1249d9c88070b7f33cb97802a. Added: Modified: clang/lib/Driver/ToolChains/Clang.cpp llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp llvm/test/CodeGen/AMDGPU/inline-calls.ll Removed: clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu ################################################################################ diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 316c6026adf5c..83afbc3952d84 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5089,9 +5089,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, } // Enable -mconstructor-aliases except on darwin, where we have to work around - // a linker bug (see <rdar://problem/7651567>), and CUDA device code, where - // aliases aren't supported. - if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX()) + // a linker bug (see <rdar://problem/7651567>), and CUDA/AMDGPU device code, + // where aliases aren't supported. + if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX() && !RawTriple.isAMDGPU()) CmdArgs.push_back("-mconstructor-aliases"); // Darwin's kernel doesn't support guard variables; just die if we diff --git a/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu b/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu deleted file mode 100644 index f75088f8e1415..0000000000000 --- a/clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu +++ /dev/null @@ -1,17 +0,0 @@ -// REQUIRES: amdgpu-registered-target, clang-driver - -// RUN: %clang --offload-arch=gfx906 --cuda-device-only -nogpulib -nogpuinc -x hip -emit-llvm -S -o - %s \ -// RUN: -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false | \ -// RUN: FileCheck %s - -#include "Inputs/cuda.h" - -// CHECK: %struct.B = type { i8 } -struct B { - - // CHECK: @_ZN1BC1Ei = hidden unnamed_addr alias void (%struct.B*, i32), void (%struct.B*, i32)* @_ZN1BC2Ei - __device__ B(int x); -}; - -__device__ B::B(int x) { -} diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp index 2e24e9f929d2a..7ff24d1e9c62b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp @@ -15,7 +15,6 @@ #include "AMDGPU.h" #include "AMDGPUTargetMachine.h" #include "Utils/AMDGPUBaseInfo.h" -#include "llvm/CodeGen/CommandFlags.h" #include "llvm/IR/Module.h" #include "llvm/Pass.h" #include "llvm/Support/CommandLine.h" @@ -91,13 +90,9 @@ static bool alwaysInlineImpl(Module &M, bool GlobalOpt) { SmallPtrSet<Function *, 8> FuncsToAlwaysInline; SmallPtrSet<Function *, 8> FuncsToNoInline; - Triple TT(M.getTargetTriple()); for (GlobalAlias &A : M.aliases()) { if (Function* F = dyn_cast<Function>(A.getAliasee())) { - if (TT.getArch() == Triple::amdgcn && - A.getLinkage() != GlobalValue::InternalLinkage) - continue; A.replaceAllUsesWith(F); AliasesToRemove.push_back(&A); } diff --git a/llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp b/llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp index 3c5cb6e190850..e841e939ef34b 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp @@ -29,8 +29,6 @@ #include "SIMachineFunctionInfo.h" #include "llvm/Analysis/CallGraph.h" #include "llvm/CodeGen/TargetPassConfig.h" -#include "llvm/IR/GlobalAlias.h" -#include "llvm/IR/GlobalValue.h" #include "llvm/Target/TargetMachine.h" using namespace llvm; @@ -63,8 +61,7 @@ static const Function *getCalleeFunction(const MachineOperand &Op) { assert(Op.getImm() == 0); return nullptr; } - if (auto *GA = dyn_cast<GlobalAlias>(Op.getGlobal())) - return cast<Function>(GA->getOperand(0)); + return cast<Function>(Op.getGlobal()); } diff --git a/llvm/test/CodeGen/AMDGPU/inline-calls.ll b/llvm/test/CodeGen/AMDGPU/inline-calls.ll index 134cd301b9743..233485a202057 100644 --- a/llvm/test/CodeGen/AMDGPU/inline-calls.ll +++ b/llvm/test/CodeGen/AMDGPU/inline-calls.ll @@ -1,6 +1,6 @@ -; RUN: llc -mtriple amdgcn-unknown-linux-gnu -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s -; RUN: llc -mtriple amdgcn-unknown-linux-gnu -mcpu=tonga -verify-machineinstrs < %s | FileCheck %s -; RUN: llc -mtriple r600-unknown-linux-gnu -mcpu=redwood -verify-machineinstrs < %s | FileCheck %s --check-prefix=R600 +; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s +; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck %s +; RUN: llc -march=r600 -mcpu=redwood -verify-machineinstrs < %s | FileCheck %s ; ALL-NOT: {{^}}func: define internal i32 @func(i32 %a) { @@ -9,7 +9,7 @@ entry: ret i32 %tmp0 } -; CHECK: {{^}}kernel: +; ALL: {{^}}kernel: ; GCN-NOT: s_swappc_b64 define amdgpu_kernel void @kernel(i32 addrspace(1)* %out) { entry: @@ -18,13 +18,12 @@ entry: ret void } -; CHECK: func_alias -; R600-NOT: func_alias +; CHECK-NOT: func_alias +; ALL-NOT: func_alias @func_alias = alias i32 (i32), i32 (i32)* @func -; CHECK-NOT: {{^}}kernel3: +; ALL: {{^}}kernel3: ; GCN-NOT: s_swappc_b64 -; R600: {{^}}kernel3: define amdgpu_kernel void @kernel3(i32 addrspace(1)* %out) { entry: %tmp0 = call i32 @func_alias(i32 1) _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
