[HIP] Link with clang_rt.builtins

[yaxunl]

Summary

HIP supports _Float16 by default in host programs, which
may cause calls of conversion functions for _Float16
emitted e.g. __truncdfhf2. These functions are not
available in libgcc but in libclang_rt.builtins. Therefore
--hip-link needs to link with libclang_rt.builtins by
default.

[MaskRay]

clang/lib/Driver/ToolChains/MSVC.cpp

485 ↗

(On Diff #434556)

Note that the path is different with LLVM_ENABLE_PER_TARGET_RUNTIME_DIR=on. See D107799

[tra]

Looks OK syntax-wise. Library path should probably be fixed, though it appears to be a somewhat separate issue and could be done in its own patch, if the required change is not trivial.

clang/lib/Driver/ToolChains/Linux.cpp

684–690 ↗

(On Diff #434556)

Nit: Collapse all of these into a single append()

clang/test/Driver/hip-runtime-libs-msvc.hip

14

What are we matching here? A more verbose pattern or some comments would be helpful.

16

CHECK-SAME?

[yaxunl]

clang/lib/Driver/ToolChains/MSVC.cpp

485 ↗

(On Diff #434556)

will use getCompilerRT to get the compiler-rt builtin lib name so that it is always correct

clang/test/Driver/hip-runtime-libs-msvc.hip

14

this is for checking lib path for compiler lib.

As we will use getCompilerRT to get the complete path of compiler-rt builtin lib, this part will be changed. We will use a clang option to print the compiler-rt lib path and use it as a reference to check compiler-rt lib used by HIP.

16

will do

[yaxunl]

use getCompilerRT to get compiler-rt lib path

[MaskRay]

These functions are not available in libgcc but in libclang_rt.builtins. Therefore --hip-link needs to link with libclang_rt.builtins by default.

I think this is problematic.

The current link sequence is ... "-lamdhip64" "/tmp/Debug/lib/clang/15.0.0/lib/linux/libclang_rt.builtins-x86_64.a" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "-lc" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "/usr/lib/gcc/x86_64-linux-gnu/12/crtendS.o" "/lib/x86_64-linux-gnu/crtn.o"

Linking in libclang_rt.builtins is incompatible with libgcc. There may be subtle symbol resolving issues. I don't think silently changing the behavior is desired.
Perhaps --rocm-path should require --rtlib=compiler-rt. See the --rtlib=libgcc requires --unwindlib=libgcc diagnostic.

Your user will need to do a bit more work if CLANG_DEFAULT_RTLIB is not compiler-rt.

clang/test/Driver/hip-runtime-libs-linux.hip

6

Note: -target is legacy spelling. --target= is the recommended name.

[yaxunl]

In D127142#3568905, @MaskRay wrote:

These functions are not available in libgcc but in libclang_rt.builtins. Therefore --hip-link needs to link with libclang_rt.builtins by default.

I think this is problematic.

The current link sequence is ... "-lamdhip64" "/tmp/Debug/lib/clang/15.0.0/lib/linux/libclang_rt.builtins-x86_64.a" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "-lc" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "/usr/lib/gcc/x86_64-linux-gnu/12/crtendS.o" "/lib/x86_64-linux-gnu/crtn.o"

Linking in libclang_rt.builtins is incompatible with libgcc. There may be subtle symbol resolving issues. I don't think silently changing the behavior is desired.
Perhaps --rocm-path should require --rtlib=compiler-rt. See the --rtlib=libgcc requires --unwindlib=libgcc diagnostic.

Your user will need to do a bit more work if CLANG_DEFAULT_RTLIB is not compiler-rt.

If I use --rtlib=compiler-rt, does that also requires --unwindlib=unwindlib ?

[MaskRay]

In D127142#3570290, @yaxunl wrote:

In D127142#3568905, @MaskRay wrote:

These functions are not available in libgcc but in libclang_rt.builtins. Therefore --hip-link needs to link with libclang_rt.builtins by default.

I think this is problematic.

The current link sequence is ... "-lamdhip64" "/tmp/Debug/lib/clang/15.0.0/lib/linux/libclang_rt.builtins-x86_64.a" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "-lc" "-lgcc" "--as-needed" "-lgcc_s" "--no-as-needed" "/usr/lib/gcc/x86_64-linux-gnu/12/crtendS.o" "/lib/x86_64-linux-gnu/crtn.o"

Linking in libclang_rt.builtins is incompatible with libgcc. There may be subtle symbol resolving issues. I don't think silently changing the behavior is desired.
Perhaps --rocm-path should require --rtlib=compiler-rt. See the --rtlib=libgcc requires --unwindlib=libgcc diagnostic.

Your user will need to do a bit more work if CLANG_DEFAULT_RTLIB is not compiler-rt.

If I use --rtlib=compiler-rt, does that also requires --unwindlib=unwindlib ?

No. --unwindlib=libunwind requires --rtlib=compiler-rt. --rtlib=compiler-rt is compatible with both --unwindlib=libgcc and --unwindlib=libunwind.

clang/lib/Driver/ToolChains/Linux.cpp

684–690 ↗

(On Diff #434556)

Note: I think multiple push_back isn't that bad...

[tra]

clang/lib/Driver/ToolChains/Linux.cpp

684–690 ↗

(On Diff #434556)

That's why I've marked it as a "nit".

For arguments, push_back works best for single arguments. append(), when more than one is needed.
Here, I think an ideal arrangement would look like this:

CmdArgs.push_back(Args.MakeArgString(StringRef("-L") + RocmInstallation.getLibPath();
CmdArgs.append({"-rpath", Args.MakeArgString(RocmInstallation.getLibPath())});
CmdArgs.push_back("-lamdhip64");
CmdArgs.push_back(Args.MakeArgString(getCompilerRT(Args, "builtins")));

Single .append({}) works OK for me, too.

Passing a mix of single options and options with arguments via append() combined with passing multi-word arguments with push_back() looks rather odd in its inconsistency to me. It makes no practical difference either way, so it's merely a minor style/readability nit, as far as I'm concerned. I'm fine deferring to author's preferences.

[yaxunl]

use compiler-rt as runtime lib by default for --hip-link

[yaxunl]

In D127142#3571260, @MaskRay wrote:

In D127142#3570290, @yaxunl wrote:

If I use --rtlib=compiler-rt, does that also requires --unwindlib=unwindlib ?

No. --unwindlib=libunwind requires --rtlib=compiler-rt. --rtlib=compiler-rt is compatible with both --unwindlib=libgcc and --unwindlib=libunwind.

If only use -rtlib=compiler-rt without changing unwind lib, I will get missing symbol:

[2022-06-16T20:05:28.644Z] ld.lld: error: undefined symbol: _Unwind_Resume

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(generate_skipahead_file())

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(main)

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(std::function<void (float&, float, bool&)> reduce::ReduceOpFn2<float>(miopenReduceTensorOp_t))

[2022-06-16T20:05:28.644Z] >>> referenced 1246 more times

[2022-06-16T20:05:28.644Z]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Resume [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Backtrace [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_GetIP [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] clang-15: error: linker command failed with exit code 1 (use -v to see invocation)

[yaxunl]

In D127142#3590874, @yaxunl wrote:

In D127142#3571260, @MaskRay wrote:

In D127142#3570290, @yaxunl wrote:

If I use --rtlib=compiler-rt, does that also requires --unwindlib=unwindlib ?

No. --unwindlib=libunwind requires --rtlib=compiler-rt. --rtlib=compiler-rt is compatible with both --unwindlib=libgcc and --unwindlib=libunwind.

If only use -rtlib=compiler-rt without changing unwind lib, I will get missing symbol:

[2022-06-16T20:05:28.644Z] ld.lld: error: undefined symbol: _Unwind_Resume

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(generate_skipahead_file())

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(main)

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(std::function<void (float&, float, bool&)> reduce::ReduceOpFn2<float>(miopenReduceTensorOp_t))

[2022-06-16T20:05:28.644Z] >>> referenced 1246 more times

[2022-06-16T20:05:28.644Z]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Resume [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Backtrace [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_GetIP [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] clang-15: error: linker command failed with exit code 1 (use -v to see invocation)

If I use --rtlib=compiler-rt, do I have to explicitly add --unwindlib=libgcc or --unwindlib=libunwind? I suspect the linker is not linking the unwind library when -rtlib=compiler-rt is specified.

[MaskRay]

In D127142#3592207, @yaxunl wrote:

In D127142#3590874, @yaxunl wrote:

In D127142#3571260, @MaskRay wrote:

In D127142#3570290, @yaxunl wrote:

If I use --rtlib=compiler-rt, does that also requires --unwindlib=unwindlib ?

No. --unwindlib=libunwind requires --rtlib=compiler-rt. --rtlib=compiler-rt is compatible with both --unwindlib=libgcc and --unwindlib=libunwind.

If only use -rtlib=compiler-rt without changing unwind lib, I will get missing symbol:

[2022-06-16T20:05:28.644Z] ld.lld: error: undefined symbol: _Unwind_Resume

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(generate_skipahead_file())

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(main)

[2022-06-16T20:05:28.644Z] >>> referenced by main.cpp

[2022-06-16T20:05:28.644Z] >>> CMakeFiles/MIOpenDriver.dir/main.cpp.o:(std::function<void (float&, float, bool&)> reduce::ReduceOpFn2<float>(miopenReduceTensorOp_t))

[2022-06-16T20:05:28.644Z] >>> referenced 1246 more times

[2022-06-16T20:05:28.644Z]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Resume [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_Backtrace [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] ld.lld: error: ../lib/libMIOpen.so.1.0.50300: undefined reference to _Unwind_GetIP [--no-allow-shlib-undefined]

[2022-06-16T20:05:28.644Z] clang-15: error: linker command failed with exit code 1 (use -v to see invocation)

If I use --rtlib=compiler-rt, do I have to explicitly add --unwindlib=libgcc or --unwindlib=libunwind? I suspect the linker is not linking the unwind library when -rtlib=compiler-rt is specified.

Either is fine.

If --rtlib=libgcc, --unwindlib must be libgcc (platform defaults to libgcc on Linux).

--rtlib and --unwindlib are orthogonal. To use libunwind, specify --unwindlib.

[yaxunl]

add -unwindlib=libgcc by default for --hip-link since -rtlib=compiler-rt needs it

[MaskRay]

Magically deciding a default value for --unwindlib or --rtlib is not nice. You may emit a warning if the selected default happens to be incompatible with HIP.

[yaxunl]

In D127142#3600809, @MaskRay wrote:

Magically deciding a default value for --unwindlib or --rtlib is not nice. You may emit a warning if the selected default happens to be incompatible with HIP.

We build clang not just for compiling HIP programs. For C/C++ programs, we want to use the default runtime and unwind library.

We only want to change the default runtime and unwind library for HIP.

[MaskRay]

In D127142#3601810, @yaxunl wrote:

In D127142#3600809, @MaskRay wrote:

Magically deciding a default value for --unwindlib or --rtlib is not nice. You may emit a warning if the selected default happens to be incompatible with HIP.

We build clang not just for compiling HIP programs. For C/C++ programs, we want to use the default runtime and unwind library.

One choice is that ensure the user specify -DCMAKE_DEFAULT_RTLIB=compiler-rt.
(Since compiler-rt can be used with libgcc_s, -DCMAKE_DEFAULT_UNWINDLIB=libunwind is not strictly necessary.)

We only want to change the default runtime and unwind library for HIP.

The correct place is clang/lib/Driver/ToolChain.cpp ToolChain::GetUnwindLibType (see "platform").
HIP should have its derived ToolChain to override the member function.

clang/lib/Driver/ToolChain.cpp

void ToolChain::addProfileRTLibs(const llvm::opt::ArgList &Args,
   CmdArgs.push_back(getCompilerRTArgString(Args, "profile"));
 }
 
 ToolChain::RuntimeLibType ToolChain::GetRuntimeLibType(
     const ArgList &Args) const {
   if (runtimeLibType)
     return *runtimeLibType;
 
+  // HIP needs compier-rt for _Float16 conversion functions.
   const Arg* A = Args.getLastArg(options::OPT_rtlib_EQ);
-  StringRef LibName = A ? A->getValue() : CLANG_DEFAULT_RTLIB;
+  StringRef LibName =
+      A ? A->getValue()
+        : (Args.hasArg(options::OPT_hip_link) ? "compiler-rt"
+                                              : CLANG_DEFAULT_RTLIB);
 
   // Only use "platform" in tests to override CLANG_DEFAULT_RTLIB!
   if (LibName == "compiler-rt")
     runtimeLibType = ToolChain::RLT_CompilerRT;
   else if (LibName == "libgcc")
     runtimeLibType = ToolChain::RLT_Libgcc;
   else if (LibName == "platform")
     runtimeLibType = GetDefaultRuntimeLibType();
   else {
     if (A)
       getDriver().Diag(diag::err_drv_invalid_rtlib_name)
           << A->getAsString(Args);
 
     runtimeLibType = GetDefaultRuntimeLibType();
   }
 
   return *runtimeLibType;
 }
 
 ToolChain::UnwindLibType ToolChain::GetUnwindLibType(
     const ArgList &Args) const {
   if (unwindLibType)
     return *unwindLibType;
 
+  // HIP needs compiler-rt for _Float16 conversion functions, which in turn
+  // needs unwind library.
   const Arg *A = Args.getLastArg(options::OPT_unwindlib_EQ);
-  StringRef LibName = A ? A->getValue() : CLANG_DEFAULT_UNWINDLIB;
+  StringRef LibName = A ? A->getValue()
+                        : ((Args.hasArg(options::OPT_hip_link) &&
+                            !getTriple().isKnownWindowsMSVCEnvironment())
+                               ? "libgcc"
+                               : CLANG_DEFAULT_UNWINDLIB);
 
   if (LibName == "none")
     unwindLibType = ToolChain::UNW_None;
   else if (LibName == "platform" || LibName == "") {
     ToolChain::RuntimeLibType RtLibType = GetRuntimeLibType(Args);
     if (RtLibType == ToolChain::RLT_CompilerRT) {
       if (getTriple().isAndroid() || getTriple().isOSAIX())
         unwindLibType = ToolChain::UNW_CompilerRT;

clang/test/Driver/hip-runtime-libs-linux.hip

 // UNSUPPORTED: system-windows
 
 // RUN:  touch %t.o
 
+// Get compiler-rt library full path.
+// RUN: %clang -target x86_64-linux-gnu -rtlib=compiler-rt \
+// RUN:   -print-libgcc-file-name >%t.1
+
 // Test HIP runtime lib args specified by --rocm-path.
 // RUN: %clang -### --hip-link -target x86_64-linux-gnu \
-// RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
-// RUN:   | FileCheck -check-prefixes=ROCM-PATH %s
+// RUN:   --rocm-path=%S/Inputs/rocm %t.o >%t.2 2>&1
+// RUN: cat %t.1 %t.2 | FileCheck -check-prefixes=ROCM-PATH,COMMON %s
 
 // Test HIP runtime lib args specified by environment variable ROCM_PATH.
 // RUN: env ROCM_PATH=%S/Inputs/rocm %clang -### --hip-link \
-// RUN:   -target x86_64-linux-gnu %t.o 2>&1 \
-// RUN:   | FileCheck -check-prefixes=ROCM-PATH %s
+// RUN:   -target x86_64-linux-gnu %t.o >%t.2 2>&1
+// RUN: cat %t.1 %t.2 | FileCheck -check-prefixes=ROCM-PATH,COMMON %s
 
 // Test detecting latest /opt/rocm-{release} directory.
 // RUN: rm -rf %T/opt
 // RUN: mkdir -p %T/opt
 // RUN: cp -r %S/Inputs/rocm %T/opt/rocm-3.9.0-1234
 // RUN: cp -r %S/Inputs/rocm %T/opt/rocm-3.10.0
 // RUN: %clang -### --hip-link -target x86_64-linux-gnu \
-// RUN:   --sysroot=%T %t.o 2>&1 \
-// RUN:   | FileCheck -check-prefixes=ROCM-REL %s
+// RUN:   --sysroot=%T %t.o >%t.2 2>&1
+// RUN: cat %t.1 %t.2 | FileCheck -check-prefixes=ROCM-REL,COMMON %s
 
 // Test HIP runtime lib is not linked without --hip-link.
 // RUN: %clang -### -target x86_64-linux-gnu \
 // RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
 // RUN:   | FileCheck -check-prefixes=NOHIPRT %s
 
-// Test HIP runtime lib is not linked with -nostdlib.
-// RUN: %clang -### --hip-link -nostdlib -target x86_64-linux-gnu \
+// Test HIP runtime lib is not linked with -rtlib=gcc.
+// RUN: %clang -### --hip-link -rtlib=gcc -target x86_64-linux-gnu \
 // RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
-// RUN:   | FileCheck -check-prefixes=NOHIPRT %s
+// RUN:   | FileCheck -check-prefixes=LIBGCC %s
 
 // Test HIP runtime lib is not linked with -no-hip-rt.
 // RUN: %clang -### --hip-link -no-hip-rt -target x86_64-linux-gnu \
 // RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
 // RUN:   | FileCheck -check-prefixes=NOHIPRT %s
 
+// COMMON: [[COMPILER_RT:/.*/libclang_rt\.builtins.*]]
 // ROCM-PATH: "-L[[HIPRT:.*/Inputs/rocm/lib]]" "-rpath" "[[HIPRT]]" "-lamdhip64"
 // ROCM-REL: "-L[[HIPRT:.*/opt/rocm-3.10.0/lib]]" "-rpath" "[[HIPRT]]" "-lamdhip64"
+// COMMON-SAME: "[[COMPILER_RT]]"
+// COMMON-SAME: "-lgcc_{{s|eh}}"
+
 // NOHIPRT-NOT: "-L{{.*/Inputs/rocm/lib}}"
 // NOHIPRT-NOT: "-rpath" "{{.*/Inputs/rocm/lib}}"
 // NOHIPRT-NOT: "-lamdhip64"
+// LIBGCC-NOT: "{{/[^"]*clang_rt\.builtins[^"]*}}"

clang/test/Driver/hip-runtime-libs-msvc.hip

 // REQUIRES: system-windows
 
 // RUN:  touch %t.o
 
+// Get compiler-rt library full path.
+// RUN: %clang -target x86_64-pc-windows-msvc -rtlib=compiler-rt \
+// RUN:   -print-libgcc-file-name >%t.1
+
 // Test HIP runtime lib args specified by --rocm-path.
 // RUN: %clang -### --hip-link -target x86_64-pc-windows-msvc \
-// RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
-// RUN:   | FileCheck %s
+// RUN:   --rocm-path=%S/Inputs/rocm %t.o >%t.2 2>&1
+// RUN: cat %t.1 %t.2 | FileCheck %s
 
+// CHECK: [[COMPILER_RT:.*clang_rt\.builtins.*]]
 // CHECK: "-libpath:{{.*Inputs.*rocm.*lib}}" "amdhip64.lib"
+// CHECK-SAME: "[[COMPILER_RT]]"
+// CHECK-NOT: "-lgcc_{{s|eh}}"