rocmPackages.llvm: compress outputs of clang-offload-bundler

This patches the clang-offload-bundler tool to add a compression option
from a more recent version of clang.
This compression option reduces the size of ROCm's fat binaries.
Those binaries contain .hip_fatbin sections with GPU-specific code,
for each target.
Compression is automatically turned on for all produced outputs
via a wrapper, because it's difficult to identify all the places
where the -compression argument would be needed.

Once upsteam introduces handeling for this argument,
we should drop the wrapper again.
This transistion will create inconsistsency, but I do not think that
it will impact any actual users and it's what's practical to implement.

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

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"
@@ -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-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/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 = [