[HIP] support 128 bit int division

clang/include/clang/CodeGen/CodeGenAction.h

@@ -9,12 +9,16 @@
 #ifndef LLVM_CLANG_CODEGEN_CODEGENACTION_H
 #define LLVM_CLANG_CODEGEN_CODEGENACTION_H
 
+#include "clang/Basic/CodeGenOptions.h"
 #include "clang/Frontend/FrontendAction.h"
 #include <memory>
 
 namespace llvm {
   class LLVMContext;
   class Module;
+  namespace object {
+  class Archive;
+  }
 }
 
 namespace clang {
@@ -54,8 +58,15 @@ class CodeGenAction : public ASTFrontendAction {
   std::unique_ptr<llvm::Module> loadModule(llvm::MemoryBufferRef MBRef);
 
   /// Load bitcode modules to link into our module from the options.
+  /// \returns true if error happens.
   bool loadLinkModules(CompilerInstance &CI);
 
+  /// Add bitcode modules in an archive to LinkModules.
+  /// \returns true if error happens.
+  bool addArchiveToLinkModules(llvm::object::Archive *Archive,
+                               const CodeGenOptions::BitcodeFileToLink &F,
+                               CompilerInstance &CI);
+
 protected:
   /// Create a new code generation action.  If the optional \p _VMContext
   /// parameter is supplied, the action uses it without taking ownership,

clang/lib/CodeGen/CodeGenAction.cpp

@@ -41,6 +41,7 @@
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/LTO/LTOBackend.h"
 #include "llvm/Linker/Linker.h"
+#include "llvm/Object/Archive.h"
 #include "llvm/Pass.h"
 #include "llvm/Support/MemoryBuffer.h"
 #include "llvm/Support/SourceMgr.h"
@@ -940,33 +941,120 @@ CodeGenAction::~CodeGenAction() {
     delete VMContext;
 }
 
+bool CodeGenAction::addArchiveToLinkModules(
+    llvm::object::Archive *Archive, const CodeGenOptions::BitcodeFileToLink &F,
+    CompilerInstance &CI) {
+  Error Err = Error::success();
+
+  for (auto &Child : Archive->children(Err)) {
+    Expected<llvm::MemoryBufferRef> ChildBufOrErr = Child.getMemoryBufferRef();
+    if (!ChildBufOrErr) {
+      handleAllErrors(ChildBufOrErr.takeError(),
+                      [&](const llvm::ErrorInfoBase &EIB) {
+                        CI.getDiagnostics().Report(diag::err_cannot_open_file)
+                            << F.Filename << EIB.message();
+                      });
+      LinkModules.clear();
+      return true;
+    }
+    auto ChildBuffer = llvm::MemoryBuffer::getMemBufferCopy(
+        ChildBufOrErr->getBuffer(), ChildBufOrErr->getBufferIdentifier());
+
+    if (!ChildBuffer) {
+      handleAllErrors(ChildBufOrErr.takeError(),
+                      [&](const llvm::ErrorInfoBase &EIB) {
+                        CI.getDiagnostics().Report(diag::err_cannot_open_file)
+                            << F.Filename << EIB.message();
+                      });
+      LinkModules.clear();
+      return true;
+    }
+
+    Expected<std::unique_ptr<llvm::Module>> ChildModuleOrErr =
+        getOwningLazyBitcodeModule(std::move(ChildBuffer), *VMContext);
+    if (!ChildModuleOrErr) {
+      handleAllErrors(ChildModuleOrErr.takeError(),
+                      [&](const llvm::ErrorInfoBase &EIB) {
+                        CI.getDiagnostics().Report(diag::err_cannot_open_file)
+                            << F.Filename << EIB.message();
+                      });
+      LinkModules.clear();
+      return true;
+    }
+
+    LinkModules.push_back({std::move(ChildModuleOrErr.get()), F.PropagateAttrs,
+                           F.Internalize, F.LinkFlags});
+  }
+  if (Err) {
+    CI.getDiagnostics().Report(diag::err_cannot_open_file)
+        << F.Filename << toString(std::move(Err));
+    LinkModules.clear();
+    return true;
+  }
+  return false;
+}
+
 bool CodeGenAction::loadLinkModules(CompilerInstance &CI) {
   if (!LinkModules.empty())
     return false;
 
   for (const CodeGenOptions::BitcodeFileToLink &F :
        CI.getCodeGenOpts().LinkBitcodeFiles) {
-    auto BCBuf = CI.getFileManager().getBufferForFile(F.Filename);
-    if (!BCBuf) {
+
+    auto BCBufOrErr = CI.getFileManager().getBufferForFile(F.Filename);
+    if (!BCBufOrErr) {
       CI.getDiagnostics().Report(diag::err_cannot_open_file)
-          << F.Filename << BCBuf.getError().message();
+          << F.Filename << BCBufOrErr.getError().message();
       LinkModules.clear();
       return true;
     }
 
+    auto &BCBuf = *BCBufOrErr;
+
     Expected<std::unique_ptr<llvm::Module>> ModuleOrErr =
-        getOwningLazyBitcodeModule(std::move(*BCBuf), *VMContext);
-    if (!ModuleOrErr) {
-      handleAllErrors(ModuleOrErr.takeError(), [&](ErrorInfoBase &EIB) {
-        CI.getDiagnostics().Report(diag::err_cannot_open_file)
-            << F.Filename << EIB.message();
-      });
+        getOwningLazyBitcodeModule(std::move(BCBuf), *VMContext);
+
+    if (ModuleOrErr) {
+      LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
+                             F.Internalize, F.LinkFlags});
+      continue;
+    }
+
+    // If parsing as bitcode failed, clear the error and try to parse as an
+    // archive.
+    handleAllErrors(ModuleOrErr.takeError(),
+                    [&](const llvm::ErrorInfoBase &EIB) {});
+
+    Expected<std::unique_ptr<llvm::object::Binary>> BinOrErr =
+        llvm::object::createBinary(BCBuf->getMemBufferRef(), VMContext);
+
+    if (!BinOrErr) {
+      handleAllErrors(BinOrErr.takeError(),
+                      [&](const llvm::ErrorInfoBase &EIB) {
+                        CI.getDiagnostics().Report(diag::err_cannot_open_file)
+                            << F.Filename << EIB.message();
+                      });
+      LinkModules.clear();
+      return true;
+    }
+
+    std::unique_ptr<llvm::object::Binary> &Bin = *BinOrErr;
+
+    if (Bin->isArchive()) {
+      llvm::object::Archive *Archive =
+          llvm::cast<llvm::object::Archive>(Bin.get());
+      if (addArchiveToLinkModules(Archive, F, CI))
+        return true;
+    } else {
+      // It's not an archive, and we failed to parse it as bitcode, so report
+      // an error.
+      CI.getDiagnostics().Report(diag::err_cannot_open_file)
+          << F.Filename << "Unrecognized file format";
       LinkModules.clear();
       return true;
     }
-    LinkModules.push_back({std::move(ModuleOrErr.get()), F.PropagateAttrs,
-                           F.Internalize, F.LinkFlags});
   }
+
   return false;
 }
 

clang/lib/Driver/ToolChains/HIPAMD.cpp

@@ -21,6 +21,7 @@
 #include "llvm/Support/Alignment.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/Path.h"
+#include "llvm/Support/VirtualFileSystem.h"
 #include "llvm/TargetParser/TargetParser.h"
 
 using namespace clang::driver;
@@ -403,6 +404,10 @@ HIPAMDToolChain::getDeviceLibs(const llvm::opt::ArgList &DriverArgs) const {
         BCLibs.emplace_back(AsanRTL, /*ShouldInternalize=*/false);
     }
 
+    auto BuiltinCRT = getCompilerRT(DriverArgs, "builtins");
+    if (getVFS().exists(BuiltinCRT))
+      BCLibs.emplace_back(BuiltinCRT, /*ShouldInternalize=*/false);
+
     // Add the HIP specific bitcode library.
     BCLibs.push_back(RocmInstallation->getHIPPath());
 

clang/test/CodeGenCUDA/link-bitcode-archive.cu

@@ -0,0 +1,63 @@
+// Prepare archive of bitcode file.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm-bc \
+// RUN:    -fcuda-is-device \
+// RUN:    -disable-llvm-passes -DIS_LIB -o %t.bc -xhip %s
+
+// RUN: rm -f %t.a
+// RUN: llvm-ar rcs %t.a %t.bc
+
+// Link archive of bitcode file.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:    -mlink-builtin-bitcode %t.a  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - -xhip %s \
+// RUN:    | FileCheck %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:    -mlink-bitcode-file %t.a  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - -xhip %s \
+// RUN:    | FileCheck %s
+
+// Test empty file as arhive.
+
+// RUN: rm -f %t.a
+// RUN: touch %t.a
+
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:    -mlink-builtin-bitcode %t.a  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - -xhip %s 2>&1\
+// RUN:    | FileCheck %s -check-prefix=INVLID
+
+// Test invalid arhive.
+
+// RUN: echo "!<arch>\nfake" >%t.a
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:    -mlink-builtin-bitcode %t.a  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - -xhip %s 2>&1 \
+// RUN:    | FileCheck %s -check-prefix=INVLID
+
+// Test archive of invalid bitcode file.
+
+// RUN: echo "BC\xC0\xDE" >%t.bc
+// RUN: rm -f %t.a
+// RUN: llvm-ar rcs %t.a %t.bc
+// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:    -mlink-builtin-bitcode %t.a  -emit-llvm \
+// RUN:    -disable-llvm-passes -o - -xhip %s 2>&1 \
+// RUN:    | FileCheck %s -check-prefix=INVLID-BC
+
+#include "Inputs/cuda.h"
+
+#ifdef IS_LIB
+__device__ void libfun() {}
+#else
+__device__ void libfun();
+__global__ void kern() {
+ libfun();
+}
+#endif
+
+// CHECK: define {{.*}}void @_Z6libfunv()
+// INVLID: fatal error: cannot open file {{.*}}: The file was not recognized as a valid object file
+// INVLID-BC: fatal error: cannot open file {{.*}}: Invalid bitcode signature

clang/test/Driver/Inputs/device_compiler_rt_resource_dir/lib/amdgcn-amd-amdhsa/libclang_rt.builtins.a

@@ -0,0 +1 @@
+!<arch>

clang/test/Driver/Inputs/device_compiler_rt_resource_dir/lib64/amdgcn-amd-amdhsa/libclang_rt.builtins.a

@@ -0,0 +1 @@
+!<arch>

clang/test/Driver/hip-device-compiler-rt.hip

@@ -0,0 +1,15 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: amdgpu-registered-target
+
+// Check device compiler-rt is linked when available.
+
+// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx1010 \
+// RUN:   -nogpuinc -nostdinc --offload-device-only --rocm-path=%S/Inputs/rocm \
+// RUN:   -resource-dir=%S/Inputs/device_compiler_rt_resource_dir \
+// RUN:   %s 2>&1 | FileCheck %s
+
+// CHECK: "-mlink-bitcode-file" "{{[^"]+(/|\\\\)device_compiler_rt_resource_dir(/|\\\\)lib(64)?(/|\\\\)amdgcn-amd-amdhsa(/|\\\\).*}}libclang_rt.builtins.a"
+// CHECK-SAME: "-mlink-builtin-bitcode" "[[DEVICELIB_DIR:[^"]+(/|\\\\)rocm(/|\\\\)amdgcn(/|\\\\).*]]hip.bc"
+// CHECK-SAME: "-mlink-builtin-bitcode" "[[DEVICELIB_DIR]]ocml.bc"
+// CHECK-SAME: "-mlink-builtin-bitcode" "[[DEVICELIB_DIR]]ockl.bc"
+// CHECK-SAME: "-mlink-builtin-bitcode" "[[DEVICELIB_DIR]]oclc_isa_version_{{[0-9]+}}.bc"

compiler-rt/cmake/Modules/CompilerRTUtils.cmake

@@ -456,6 +456,8 @@ function(get_compiler_rt_target arch variable)
       endif()
     endif()
     set(target "${arch}${triple_suffix}")
+  elseif(${arch} STREQUAL "amdgcn")
+    set(target "amdgcn-amd-amdhsa")
   else()
     set(target "${arch}${triple_suffix}")
   endif()

compiler-rt/cmake/base-config-ix.cmake

@@ -194,6 +194,11 @@ macro(test_targets)
     endif()
   endif()
 
+  set(COMPILER_RT_ENABLE_TARGET_AMDGCN OFF CACHE BOOL "Option to enable AMDGCN in Compiler RT")
+  if (COMPILER_RT_ENABLE_TARGET_AMDGCN)
+    add_default_target_arch("amdgcn")
+  endif()
+
   # Generate the COMPILER_RT_SUPPORTED_ARCH list.
   if(ANDROID)
     # Examine compiler output to determine target architecture.

compiler-rt/cmake/builtin-config-ix.cmake

@@ -65,6 +65,7 @@ set(SPARCV9 sparcv9)
 set(WASM32 wasm32)
 set(WASM64 wasm64)
 set(VE ve)
+set(AMDGCN amdgcn)
 
 if(APPLE)
   set(ARM64 arm64 arm64e)
@@ -76,7 +77,7 @@ set(ALL_BUILTIN_SUPPORTED_ARCH
   ${X86} ${X86_64} ${ARM32} ${ARM64} ${AVR}
   ${HEXAGON} ${MIPS32} ${MIPS64} ${PPC32} ${PPC64}
   ${RISCV32} ${RISCV64} ${SPARC} ${SPARCV9}
-  ${WASM32} ${WASM64} ${VE} ${LOONGARCH64})
+  ${WASM32} ${WASM64} ${VE} ${LOONGARCH64} ${AMDGCN})
 
 include(CompilerRTUtils)
 include(CompilerRTDarwinUtils)

compiler-rt/lib/builtins/CMakeLists.txt

@@ -560,6 +560,13 @@ set(aarch64_SOURCES
   aarch64/fp_mode.c
 )
 
+set(amdgcn_SOURCES
+  divti3.c
+  udivmodti4.c
+  truncdfbf2.c
+  truncsfbf2.c
+)
+
 if(COMPILER_RT_HAS_ASM_SME AND (COMPILER_RT_HAS_AUXV OR COMPILER_RT_BAREMETAL_BUILD))
   list(APPEND aarch64_SOURCES aarch64/sme-abi.S aarch64/sme-abi-init.c)
   message(STATUS "AArch64 SME ABI routines enabled")
@@ -846,6 +853,15 @@ else ()
         list(APPEND BUILTIN_CFLAGS_${arch} -fomit-frame-pointer -DCOMPILER_RT_ARMHF_TARGET)
       endif()
 
+      if (${arch} STREQUAL "amdgcn")
+        list(APPEND BUILTIN_CFLAGS_${arch}
+             --target=amdgcn-amd-amdhsa
+             -emit-llvm
+             -nogpuinc
+             -nogpulib
+             -Xclang -mcode-object-version=none )
+      endif()
+
       # For RISCV32, we must force enable int128 for compiling long
       # double routines.
       if(COMPILER_RT_ENABLE_SOFTWARE_INT128 OR "${arch}" STREQUAL "riscv32")

compiler-rt/lib/builtins/int_lib.h

@@ -26,6 +26,8 @@
 #else
 #define COMPILER_RT_ABI __attribute__((__pcs__("aapcs")))
 #endif
+#elif __AMDGPU__
+#define COMPILER_RT_ABI __attribute__((amdgpu_lib_fun, weak))
 #else
 #define COMPILER_RT_ABI
 #endif