Merge pull request #299589 from mschwaig/compress-composable-kernel

rocmPackages.composable_kernel: compress output

pkgs/development/rocm-modules/6/llvm/base.nix

@@ -72,7 +72,9 @@ in stdenv.mkDerivation (finalAttrs: {
     "info" # Avoid `attribute 'info' missing` when using with wrapCC
   ];
 
-  patches = extraPatches;
+  patches = [
+    ./add-compression-to-clang-offload-bundler.patch
+  ] ++ extraPatches;
 
   src = fetchFromGitHub {
     owner = "ROCm";
@@ -133,7 +135,14 @@ in stdenv.mkDerivation (finalAttrs: {
     "-DLLVM_EXTERNAL_LIT=${lit}/bin/.lit-wrapped"
   ] ++ extraCMakeFlags;
 
-  postPatch = lib.optionalString finalAttrs.passthru.isLLVM ''
+  prePatch = ''
+    cd ../
+    chmod -R u+w .
+  '';
+
+  postPatch = ''
+    cd ${targetDir}
+  '' + lib.optionalString finalAttrs.passthru.isLLVM ''
     patchShebangs lib/OffloadArch/make_generated_offload_arch_h.sh
   '' + lib.optionalString (buildTests && finalAttrs.passthru.isLLVM) ''
     # FileSystem permissions tests fail with various special bits
@@ -141,7 +150,7 @@ in stdenv.mkDerivation (finalAttrs: {
     rm unittests/Support/Path.cpp
 
     substituteInPlace unittests/Support/CMakeLists.txt \
-      --replace "Path.cpp" ""
+      --replace-fail "Path.cpp" ""
   '' + extraPostPatch;
 
   doCheck = buildTests;

pkgs/development/rocm-modules/6/llvm/stage-1/clang-unwrapped.nix

@@ -2,13 +2,14 @@
 , callPackage
 , rocmUpdateScript
 , llvm
+, makeWrapper
 }:
 
 callPackage ../base.nix rec {
   inherit stdenv rocmUpdateScript;
   targetName = "clang-unwrapped";
   targetDir = "clang";
-  extraBuildInputs = [ llvm ];
+  extraBuildInputs = [ llvm makeWrapper ];
 
   extraCMakeFlags = [
     "-DCLANG_INCLUDE_DOCS=ON"
@@ -20,7 +21,7 @@ callPackage ../base.nix rec {
     ln -s ../cmake/Modules/FindLibEdit.cmake cmake/modules
 
     substituteInPlace CMakeLists.txt \
-      --replace "include(CheckIncludeFile)" "include(CheckIncludeFile)''\nfind_package(LibEdit)"
+      --replace-fail "include(CheckIncludeFile)" "include(CheckIncludeFile)''\nfind_package(LibEdit)"
 
     # `No such file or directory: '/build/source/clang/tools/scan-build/bin/scan-build'`
     rm test/Analysis/scan-build/*.test
@@ -41,6 +42,12 @@ callPackage ../base.nix rec {
 
   extraPostInstall = ''
     mv bin/clang-tblgen $out/bin
+    # add wrapper to compress embedded accelerator-specific code
+    # this makes the output of composable_kernel significantly smaller right now
+    # TODO: remove this once ROCm does it out of the box
+    mv $out/bin/clang-offload-bundler $out/bin/clang-offload-bundler-unwrapped
+    makeWrapper $out/bin/clang-offload-bundler-unwrapped $out/bin/clang-offload-bundler \
+      --add-flags '-compress'
   '';
 
   requiredSystemFeatures = [ "big-parallel" ];

pkgs/development/rocm-modules/6/llvm/stage-2/libc.nix

@@ -14,11 +14,11 @@ callPackage ../base.nix rec {
     # `Failed to match ... against ...` `Match value not within tolerance value of MPFR result:`
     # We need a better way, but I don't know enough sed magic and patching `CMakeLists.txt` isn't working...
     substituteInPlace ../libc/test/src/math/log10_test.cpp \
-      --replace "i < N" "i < 0" \
-      --replace "test(mpfr::RoundingMode::Nearest);" "" \
-      --replace "test(mpfr::RoundingMode::Downward);" "" \
-      --replace "test(mpfr::RoundingMode::Upward);" "" \
-      --replace "test(mpfr::RoundingMode::TowardZero);" ""
+      --replace-fail "i < N" "i < 0" \
+      --replace-fail "test(mpfr::RoundingMode::Nearest);" "" \
+      --replace-fail "test(mpfr::RoundingMode::Downward);" "" \
+      --replace-fail "test(mpfr::RoundingMode::Upward);" "" \
+      --replace-fail "test(mpfr::RoundingMode::TowardZero);" ""
   '';
 
   checkTargets = [ "check-${targetName}" ];

pkgs/development/rocm-modules/6/llvm/stage-3/0000-mlir-fix-debugtranslation.patch

@@ -1,36 +0,0 @@
-From f1d1e10ec7e1061bf0b90abbc1e298d9438a5e74 Mon Sep 17 00:00:00 2001
-From: Scott Linder <Scott.Linder@amd.com>
-Date: Mon, 11 Sep 2023 18:37:37 +0000
-Subject: [PATCH] [HeterogeneousDWARF] Update MLIR DI Metadata handling
-
-Pass a default DW_MSPACE_LLVM_none to satisfy new API
-
-Change-Id: I50df461f00b5510a715f55f61107122318102d22
----
- lib/Target/LLVMIR/DebugTranslation.cpp | 6 ++++--
- 1 file changed, 4 insertions(+), 2 deletions(-)
-
-diff --git a/lib/Target/LLVMIR/DebugTranslation.cpp b/lib/Target/LLVMIR/DebugTranslation.cpp
-index 2053f5bcef06aa6..635ee5d7e5fefdc 100644
---- a/lib/Target/LLVMIR/DebugTranslation.cpp
-+++ b/lib/Target/LLVMIR/DebugTranslation.cpp
-@@ -148,7 +148,8 @@ llvm::DIDerivedType *DebugTranslation::translateImpl(DIDerivedTypeAttr attr) {
-       /*File=*/nullptr, /*Line=*/0,
-       /*Scope=*/nullptr, translate(attr.getBaseType()), attr.getSizeInBits(),
-       attr.getAlignInBits(), attr.getOffsetInBits(),
--      /*DWARFAddressSpace=*/std::nullopt, /*Flags=*/llvm::DINode::FlagZero);
-+      /*DWARFAddressSpace=*/std::nullopt, llvm::dwarf::DW_MSPACE_LLVM_none,
-+      /*Flags=*/llvm::DINode::FlagZero);
- }
- 
- llvm::DIFile *DebugTranslation::translateImpl(DIFileAttr attr) {
-@@ -185,7 +186,8 @@ DebugTranslation::translateImpl(DILocalVariableAttr attr) {
-       llvmCtx, translate(attr.getScope()), getMDStringOrNull(attr.getName()),
-       translate(attr.getFile()), attr.getLine(), translate(attr.getType()),
-       attr.getArg(),
--      /*Flags=*/llvm::DINode::FlagZero, attr.getAlignInBits(),
-+      /*Flags=*/llvm::DINode::FlagZero, llvm::dwarf::DW_MSPACE_LLVM_none,
-+      attr.getAlignInBits(),
-       /*Annotations=*/nullptr);
- }
- 

pkgs/development/rocm-modules/6/llvm/stage-3/clang.nix

@@ -68,6 +68,6 @@ wrapCCWith rec {
 
     # GPU compilation uses builtin `lld`
     substituteInPlace $out/bin/{clang,clang++} \
-      --replace "-MM) dontLink=1 ;;" "-MM | --cuda-device-only) dontLink=1 ;;''\n--cuda-host-only | --cuda-compile-host-device) dontLink=0 ;;"
+      --replace-fail "-MM) dontLink=1 ;;" "-MM | --cuda-device-only) dontLink=1 ;;''\n--cuda-host-only | --cuda-compile-host-device) dontLink=0 ;;"
   '';
 }

pkgs/development/rocm-modules/6/llvm/stage-3/libclc.nix

@@ -21,13 +21,13 @@ in callPackage ../base.nix rec {
   # `clspv` tests fail, unresolved calls
   extraPostPatch = ''
     substituteInPlace CMakeLists.txt \
-      --replace "find_program( LLVM_CLANG clang PATHS \''${LLVM_BINDIR} NO_DEFAULT_PATH )" \
+      --replace-fail "find_program( LLVM_CLANG clang PATHS \''${LLVM_BINDIR} NO_DEFAULT_PATH )" \
         "find_program( LLVM_CLANG clang PATHS \"${clang}/bin\" NO_DEFAULT_PATH )" \
-      --replace "find_program( LLVM_SPIRV llvm-spirv PATHS \''${LLVM_BINDIR} NO_DEFAULT_PATH )" \
+      --replace-fail "find_program( LLVM_SPIRV llvm-spirv PATHS \''${LLVM_BINDIR} NO_DEFAULT_PATH )" \
         "find_program( LLVM_SPIRV llvm-spirv PATHS \"${spirv}/bin\" NO_DEFAULT_PATH )" \
-      --replace "  spirv-mesa3d-" "" \
-      --replace "  spirv64-mesa3d-" "" \
-      --replace "NOT \''${t} MATCHES" \
+      --replace-fail "  spirv-mesa3d-" "" \
+      --replace-fail "  spirv64-mesa3d-" "" \
+      --replace-fail "NOT \''${t} MATCHES" \
         "NOT \''${ARCH} STREQUAL \"clspv\" AND NOT \''${ARCH} STREQUAL \"clspv64\" AND NOT \''${t} MATCHES"
   '';
 

pkgs/development/rocm-modules/6/llvm/stage-3/mlir.nix

@@ -7,6 +7,7 @@
 , glslang
 , shaderc
 , lit
+, fetchpatch
 }:
 
 callPackage ../base.nix rec {
@@ -17,9 +18,12 @@ callPackage ../base.nix rec {
   targetDir = targetName;
 
   # Fix `DebugTranslation.cpp:139:10: error: no matching function for call to 'get'`
-  # We patch at a different source root, so we modify the patch and include it locally
-  # https://github.com/ROCm/llvm-project/commit/f1d1e10ec7e1061bf0b90abbc1e298d9438a5e74.patch
-  extraPatches = [ ./0000-mlir-fix-debugtranslation.patch ];
+  extraPatches = [
+    (fetchpatch {
+      url = "https://github.com/ROCm/llvm-project/commit/f1d1e10ec7e1061bf0b90abbc1e298d9438a5e74.patch";
+      hash = "sha256-3c91A9InMKxm+JcnWxoUeOU68y5I6w1AAXx6T9UByqI=";
+    })
+  ];
   extraNativeBuildInputs = [ clr ];
 
   extraBuildInputs = [
@@ -41,7 +45,7 @@ callPackage ../base.nix rec {
   extraPostPatch = ''
     # `add_library cannot create target "llvm_gtest" because an imported target with the same name already exists`
     substituteInPlace CMakeLists.txt \
-      --replace "EXISTS \''${UNITTEST_DIR}/googletest/include/gtest/gtest.h" "FALSE"
+      --replace-fail "EXISTS \''${UNITTEST_DIR}/googletest/include/gtest/gtest.h" "FALSE"
 
     # Mainly `No such file or directory`
     cat ${./1001-mlir-failing-tests.list} | xargs -d \\n rm

pkgs/development/rocm-modules/6/llvm/stage-3/polly.nix

@@ -11,7 +11,7 @@ callPackage ../base.nix rec {
   extraPostPatch = ''
     # `add_library cannot create target "llvm_gtest" because an imported target with the same name already exists`
     substituteInPlace CMakeLists.txt \
-      --replace "NOT TARGET gtest" "FALSE"
+      --replace-fail "NOT TARGET gtest" "FALSE"
   '';
 
   checkTargets = [ "check-${targetName}" ];