[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

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

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

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

10

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

12

CHECK-SAME?

[yaxunl]

clang/lib/Driver/ToolChains/MSVC.cpp

485

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

10

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.

12

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

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

[tra]

clang/lib/Driver/ToolChains/Linux.cpp

684–690

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/ToolChains/Linux.cpp

 
 void Linux::AddHIPIncludeArgs(const ArgList &DriverArgs,
                               ArgStringList &CC1Args) const {
   RocmInstallation.AddHIPIncludeArgs(DriverArgs, CC1Args);
 }
 
 void Linux::AddHIPRuntimeLibArgs(const ArgList &Args,
                                  ArgStringList &CmdArgs) const {
   CmdArgs.append(
       {Args.MakeArgString(StringRef("-L") + RocmInstallation.getLibPath()),
        "-rpath", Args.MakeArgString(RocmInstallation.getLibPath())});
 
   CmdArgs.push_back("-lamdhip64");
+  CmdArgs.push_back(
+      Args.MakeArgString("-lclang_rt.builtins-" + getTriple().getArchName()));
 }
 
 void Linux::AddIAMCUIncludeArgs(const ArgList &DriverArgs,
                                 ArgStringList &CC1Args) const {
   if (GCCInstallation.isValid()) {
     CC1Args.push_back("-isystem");
     CC1Args.push_back(DriverArgs.MakeArgString(
         GCCInstallation.getParentLibPath() + "/../" +

clang/lib/Driver/ToolChains/MSVC.cpp

void MSVCToolChain::AddHIPIncludeArgs(const ArgList &DriverArgs,
   RocmInstallation.AddHIPIncludeArgs(DriverArgs, CC1Args);
 }
 
 void MSVCToolChain::AddHIPRuntimeLibArgs(const ArgList &Args,
                                          ArgStringList &CmdArgs) const {
   CmdArgs.append({Args.MakeArgString(StringRef("-libpath:") +
                                      RocmInstallation.getLibPath()),
                   "amdhip64.lib"});
+  CmdArgs.push_back(Args.MakeArgString("clang_rt.builtins-" +
+                                       getTriple().getArchName() + ".lib"));
 }
 
 void MSVCToolChain::printVerboseInfo(raw_ostream &OS) const {
   CudaInstallation.print(OS);
   RocmInstallation.print(OS);
 }
 
 std::string

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

 // UNSUPPORTED: system-windows
 
 // RUN:  touch %t.o
 
 // 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:   | 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:   | 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:   | 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 \
 // RUN:   --rocm-path=%S/Inputs/rocm %t.o 2>&1 \
 // RUN:   | FileCheck -check-prefixes=NOHIPRT %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: "-L{{.*}}lib/x86_64-unknown-linux-gnu"
 // ROCM-PATH: "-L[[HIPRT:.*/Inputs/rocm/lib]]" "-rpath" "[[HIPRT]]" "-lamdhip64"
 // ROCM-REL: "-L[[HIPRT:.*/opt/rocm-3.10.0/lib]]" "-rpath" "[[HIPRT]]" "-lamdhip64"
+// COMMON: "-lclang_rt.builtins-x86_64"
 // NOHIPRT-NOT: "-L{{.*/Inputs/rocm/lib}}"
 // NOHIPRT-NOT: "-rpath" "{{.*/Inputs/rocm/lib}}"
 // NOHIPRT-NOT: "-lamdhip64"
+// NOHIPRT-NOT: "-lclang_rt.builtins-x86_64"

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

 // REQUIRES: system-windows
 
 // RUN:  touch %t.o
 
 // 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
 
+// CHECK: "-libpath:{{.*lib.*windows}}"
 // CHECK: "-libpath:{{.*Inputs.*rocm.*lib}}" "amdhip64.lib"
+// CHECK: "clang_rt.builtins-x86_64.lib"