[HIP] [AlwaysInliner] Disable AlwaysInliner to eliminate undefined symbols

[gandhi21299]

Summary

By default clang emits complete contructors as alias of base constructors if they are the same.
The backend is supposed to emit symbols for the alias, otherwise it causes undefined symbols.
@yaxunl observed that this issue is related to the llvm options -amdgpu-early-inline-all=true
and -amdgpu-function-calls=false. This issue is resolved by only inlining global values
with internal linkage. The getCalleeFunction() in AMDGPUResourceUsageAnalysis also had
to be extended to support aliases to functions. inline-calls.ll was corrected appropriately.

[arsenm]

While I think the early inliner is largely obsolete, it should still handle aliases correctly

[arsenm]

clang/lib/Driver/ToolChains/Clang.cpp

5094

This looks like an unrelated change?

[yaxunl]

We cannot disable early inline all since this will cause performance regressions. Instead, the early inline all pass should be fixed so it does not remove aliases.

[aeubanks]

should the amdgpu-early-inline-all flag be deleted?

[yaxunl]

In D109707#2998095, @aeubanks wrote:

should the amdgpu-early-inline-all flag be deleted?

No. It is still used by HIP. We will deprecate it in the feature, but it is not ready yet.

[gandhi21299]

clang/lib/Driver/ToolChains/Clang.cpp

5094

Ahh yes, I will get rid of it.

[gandhi21299]

• set GlobalOpt parameter to false by default to disallow alias elimination when the options EarlyInlineAll and EnableFunctionCalls are true and false, respectively.

clang/lib/Driver/ToolChains/Clang.cpp

5094

This was part of a revert that is required for this patch to function.

[gandhi21299]

• added the include header for HIP runtime

[gandhi21299]

• converted the HIP test into a CUDA test

[yaxunl]

I think you may try fixing the following line:

https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp#L97

[gandhi21299]

@yaxunl Under what criteria should an alias not be removed?

[yaxunl]

In D109707#3004108, @gandhi21299 wrote:

@yaxunl Under what criteria should an alias not be removed?

If the linkage of the alias is not internal, then it should not be removed.

[gandhi21299]

Internal linkage detection works great for our purposes but it causes a failure in llvm/test/CodeGen/AMDGPU/inline-calls.ll due to @func_alias unable to be casted into a Function. If we pass through that, the @kernel3 causes the error: scalar registers (98) exceeds limit (96) in function 'kernel3'.

[gandhi21299]

@yaxunl I think we have two ways to go from here:

1. If appropriate, reset the maximum number of scalar registers allowed in @kernel3 (inline-calls.ll) to fix the test.
2. Determine a stronger condition for inlining.

[arsenm]

In D109707#3004869, @gandhi21299 wrote:

Internal linkage detection works great for our purposes but it causes a failure in llvm/test/CodeGen/AMDGPU/inline-calls.ll due to @func_alias unable to be casted into a Function. If we pass through that, the @kernel3 causes the error: scalar registers (98) exceeds limit (96) in function 'kernel3'.

That almost sounds like using the wrong subtarget for the alias

llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp

694 ↗

(On Diff #372573)

This needs a backend test

[gandhi21299]

• Prevent removing alias if the GlobalAlias does not have internal linkage

[gandhi21299]

• replaced a cast with a dyn_cast since the return value from getCalleeFunction() is not always a Function
• RUN on line 2 was causing 2 more scalar registers to be used on tonga due to @func_alias not being inlined, hence I eliminated that test
• RUN on line 3 generated a call instruction to an aliased function which is not supported on r600 (according to @arsenm ), hence I eliminated that test as well

[gandhi21299]

• refreshing patch

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

68

I think this is not the right place for this. If we can determine the callee function, we should have directly set it in the instruction during call lowering

[gandhi21299]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

68

Which file would that be in?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

68

SIISelLowering and AMDGPUCallLowering

[gandhi21299]

• Declare an unhandled call lowering in SelectionDAG when a callee is encountered which cannot be casted into a Function
• I am still investigating the effects on GlobalISel side of things, there seems to be a problem when lowering a call to @func in @kernel as well.
• inline-calls.ll is expected to fail with this patch, we could turn it into a negative test depending on how the work goes.

[gandhi21299]

It does not look like function calls are supported yet in AMDGPUCallLowering, is that correct?

[yaxunl]

Pls update the description of the patch and also make sure it passes internal CI.

clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu

2

this test needs

// REQUIRES: amdgpu-registered-target, clang-driver

llvm/test/CodeGen/AMDGPU/inline-calls.ll

2–3

why these two lines are removed?

[gandhi21299]

In D109707#3016438, @gandhi21299 wrote:

• replaced a cast with a dyn_cast since the return value from getCalleeFunction() is not always a Function
• RUN on line 2 was causing 2 more scalar registers to be used on tonga due to @func_alias not being inlined, hence I eliminated that test
• RUN on line 3 generated a call instruction to an aliased function which is not supported on r600 (according to @arsenm ), hence I eliminated that test as well

@yaxunl

[gandhi21299]

• added the REQUIRES line as requested by Sam

[gandhi21299]

@yaxunl Should inline-calls.ll be converted into an expected failing test or removed? (to avoid cast failure in AMDGPUResourceAnalysis to break the test)

[gandhi21299]

• declare failure when lowering an accessor of a callee which is not a function, in GlobalISel

[yaxunl]

llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp

98–99

If we do this for older GPU's (e.g. Tonga/redwood), IR's using aliases will fail on them. I don't think it is acceptable.

Is it possible to restrict this change to gfx9 and above? Or should we introduce some feature to indicate 'alias support' and use that to restrict this change to subtargets supporting this feature.

llvm/test/CodeGen/AMDGPU/inline-calls.ll

0–3

need to add check for gfx906 and gfx1030

[gandhi21299]

llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp

98–99

Restricting this change to gfx9 and above sounds simpler and more relevant with the problem as well.

[gandhi21299]

Sorry, that was a mistake.

[gandhi21299]

inline-calls.ll failed on gfx908 due to the change in SIISelLowering.cpp, line 3015. Without the change, there is a failure in AMDGPUResourceAnalysis.cpp, line 65 because Op.getGlobal() is not a Function.

[gandhi21299]

• Since callees may alias to a function pointer, it makes sense for getCalleeFunction(...) to return a Function which is a cast of the operand of a GlobalAlias.

[gandhi21299]

• eliminated changes in SIISelLowering

[gandhi21299]

ping

[gandhi21299]

refreshing patch

[gandhi21299]

added -nogpulib and -nogpuinc flags to amdgpu-alias-undef-symbols.cu

[gandhi21299]

ping

[gandhi21299]

Passed internal CI

[gandhi21299]

llvm/test/CodeGen/AMDGPU/inline-calls.ll

3

@tstellar Is there a way to restrict the AlwaysInliner to only run on amdgcn architecture?

[gandhi21299]

add a restrictions to what architecture AlwaysInliner should run on, updated the inline-calls.ll test.

[gandhi21299]

Passed ePSDB

[gandhi21299]

ping

[yaxunl]

LGTM. Thanks.

[thakis]

This breaks tests on Mac: http://45.33.8.238/mac/37119/step_7.txt

Please take a look and revert for now if it takes a while to fix.

[yaxunl]

@gandhi21299 you may need to add "-target x86_64-unknown-linux-gnu" to your codegen test to avoid issue with Darwin toolchain.

[gandhi21299]

added -target option in the test amdgpu-alias-undef-symbols.cu

[gandhi21299]

@thakis can you please check if this solution is sufficient? Thanks for bringing it up

[MaskRay]

clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu

4

non-driver tests prefer %clang_cc1.

%clang invokes the driver and has varying behaviors on different platforms. Include paths/resource dir may be quite different.

[gandhi21299]

clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu

4

Alias is not generated when I make the change to:

// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -aux-triple x86_64-unknown-linux-gnu \
// RUN:   -x hip -fcuda-is-device -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true \
// RUN:   -mllvm -amdgpu-function-calls=false -emit-llvm %s -o - | FileCheck %s

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

67

I thought aliases could include embedded bitcasts of the function type, so the function wouldn't directly appear here

[gandhi21299]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

67

Can you please elaborate on "include embedded bitcasts of the function type"? It's a consequence of the AlwaysInliner where the callee gets replaced by the alias to a function, ie. @func_alias gets replaced by @func in the inline-calls.ll test.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

67

Something like this where the alias changes the type from the original function:

@add1alias3 = alias float (float), bitcast (i32 (i32)* @add1 to float(float)*)

[gandhi21299]

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

67

I see, that will probably break the compiler since a bitcast expression is not a Function.

clang/lib/Driver/ToolChains/Clang.cpp

if (Args.hasArg(options::OPT_fdebug_pass_structure)) {
     CmdArgs.push_back("Structure");
   }
   if (Args.hasArg(options::OPT_fdebug_pass_arguments)) {
     CmdArgs.push_back("-mdebug-pass");
     CmdArgs.push_back("Arguments");
   }
 
   // Enable -mconstructor-aliases except on darwin, where we have to work around
-  // a linker bug (see <rdar://problem/7651567>), and CUDA/AMDGPU device code,
-  // where aliases aren't supported.
-  if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX() && !RawTriple.isAMDGPU())
+  // a linker bug (see <rdar://problem/7651567>), and CUDA device code, where
+  // aliases aren't supported.
+  if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX())
     CmdArgs.push_back("-mconstructor-aliases");
 
   // Darwin's kernel doesn't support guard variables; just die if we
   // try to use them.
   if (KernelOrKext && RawTriple.isOSDarwin())
     CmdArgs.push_back("-fforbid-guard-variables");
 
   if (Args.hasFlag(options::OPT_mms_bitfields, options::OPT_mno_ms_bitfields,

clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu

+// REQUIRES: amdgpu-registered-target, clang-driver
+
+// RUN: %clang -target x86_64-unknown-linux-gnu --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) {
+}

llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp

 /// This pass marks all internal functions as always_inline and creates
 /// duplicates of all other functions and marks the duplicates as always_inline.
 //
 //===----------------------------------------------------------------------===//
 
 #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"
 
 using namespace llvm;
 
 namespace {
 
recursivelyVisitUsers(GlobalValue &GV,
   }
 }
 
 static bool alwaysInlineImpl(Module &M, bool GlobalOpt) {
   std::vector<GlobalAlias*> AliasesToRemove;
 
   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);
     }
 
     // FIXME: If the aliasee isn't a function, it's some kind of constant expr
     // cast that won't be inlined through.
   }
 

llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp

 //===----------------------------------------------------------------------===//
 
 #include "AMDGPUResourceUsageAnalysis.h"
 #include "AMDGPU.h"
 #include "GCNSubtarget.h"
 #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;
 using namespace llvm::AMDGPU;
 
 #define DEBUG_TYPE "amdgpu-resource-usage"
 
 char llvm::AMDGPUResourceUsageAnalysis::ID = 0;

 INITIALIZE_PASS(AMDGPUResourceUsageAnalysis, DEBUG_TYPE,
                 "Function register usage analysis", true, true)
 
 static const Function *getCalleeFunction(const MachineOperand &Op) {
   if (Op.isImm()) {
     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());
 }
 
 static bool hasAnyNonFlatUseOfReg(const MachineRegisterInfo &MRI,
                                   const SIInstrInfo &TII, unsigned Reg) {
   for (const MachineOperand &UseOp : MRI.reg_operands(Reg)) {
     if (!UseOp.isImplicit() || !TII.isFLAT(*UseOp.getParent()))
       return true;
   }

llvm/test/CodeGen/AMDGPU/inline-calls.ll

-; 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
+; 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
 
 ; ALL-NOT: {{^}}func:
 define internal i32 @func(i32 %a) {
 entry:
   %tmp0 = add i32 %a, 1
   ret i32 %tmp0
 }
 
-; ALL: {{^}}kernel:
+; CHECK: {{^}}kernel:
 ; GCN-NOT: s_swappc_b64
 define amdgpu_kernel void @kernel(i32 addrspace(1)* %out) {
 entry:
   %tmp0 = call i32 @func(i32 1)
   store i32 %tmp0, i32 addrspace(1)* %out
   ret void
 }
 
-; CHECK-NOT: func_alias
-; ALL-NOT: func_alias
+; CHECK: func_alias
+; R600-NOT: func_alias
 @func_alias = alias i32 (i32), i32 (i32)* @func
 
-; ALL: {{^}}kernel3:
+; CHECK-NOT: {{^}}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)
   store i32 %tmp0, i32 addrspace(1)* %out
   ret void
 }