diff --git a/pkgs/by-name/ll/llama-cpp/package.nix b/pkgs/by-name/ll/llama-cpp/package.nix index c3ae8e1dc0c7..890c3ce7abea 100644 --- a/pkgs/by-name/ll/llama-cpp/package.nix +++ b/pkgs/by-name/ll/llama-cpp/package.nix @@ -12,7 +12,7 @@ rocmSupport ? config.rocmSupport, rocmPackages ? { }, - rocmGpuTargets ? builtins.concatStringsSep ";" rocmPackages.clr.gpuTargets, + rocmGpuTargets ? rocmPackages.clr.localGpuTargets or rocmPackages.clr.gpuTargets, openclSupport ? false, clblast, @@ -150,7 +150,7 @@ effectiveStdenv.mkDerivation (finalAttrs: { ] ++ optionals rocmSupport [ (cmakeFeature "CMAKE_HIP_COMPILER" "${rocmPackages.clr.hipClangPath}/clang++") - (cmakeFeature "CMAKE_HIP_ARCHITECTURES" rocmGpuTargets) + (cmakeFeature "CMAKE_HIP_ARCHITECTURES" (builtins.concatStringsSep ";" rocmGpuTargets)) ] ++ optionals metalSupport [ (cmakeFeature "CMAKE_C_FLAGS" "-D__ARM_FEATURE_DOTPROD=1") diff --git a/pkgs/by-name/ma/magma/package.nix b/pkgs/by-name/ma/magma/package.nix index 9c5d762fd8eb..06fbe93a003d 100644 --- a/pkgs/by-name/ma/magma/package.nix +++ b/pkgs/by-name/ma/magma/package.nix @@ -170,6 +170,10 @@ stdenv.mkDerivation (finalAttrs: { ] ); + env.CFLAGS = "-DADD_" + lib.optionalString rocmSupport " -fopenmp"; + env.CXXFLAGS = finalAttrs.env.CFLAGS; + env.FFLAGS = "-DADD_"; + cmakeFlags = [ (strings.cmakeFeature "GPU_TARGET" gpuTargetString) (strings.cmakeBool "MAGMA_ENABLE_CUDA" cudaSupport) @@ -179,9 +183,6 @@ stdenv.mkDerivation (finalAttrs: { # otherwise not be set in NVCC_FLAGS or DEVCCFLAGS (which we cannot modify). # See https://github.com/NixOS/nixpkgs/issues/281656#issuecomment-1902931289 (strings.cmakeBool "USE_FORTRAN" true) - (strings.cmakeFeature "CMAKE_C_FLAGS" "-DADD_") - (strings.cmakeFeature "CMAKE_CXX_FLAGS" "-DADD_") - (strings.cmakeFeature "FORTRAN_CONVENTION" "-DADD_") ] ++ lists.optionals cudaSupport [ (strings.cmakeFeature "CMAKE_CUDA_ARCHITECTURES" cudaArchitecturesString) @@ -192,8 +193,8 @@ stdenv.mkDerivation (finalAttrs: { # Can't easily apply the PR as a patch because we rely on the tarball with pregenerated # hipified files ∴ fetchpatch of the PR will apply cleanly but fail to build (strings.cmakeFeature "ROCM_CORE" "${rocmPackages.clr}") - (strings.cmakeFeature "CMAKE_C_COMPILER" "${rocmPackages.clr}/bin/hipcc") - (strings.cmakeFeature "CMAKE_CXX_COMPILER" "${rocmPackages.clr}/bin/hipcc") + (strings.cmakeFeature "CMAKE_C_COMPILER" "${rocmPackages.clr}/bin/clang") + (strings.cmakeFeature "CMAKE_CXX_COMPILER" "${rocmPackages.clr}/bin/clang++") ]; # Magma doesn't have a test suite we can easily run, just loose executables, all of which require a GPU. diff --git a/pkgs/by-name/uc/ucx/package.nix b/pkgs/by-name/uc/ucx/package.nix index 393842b7b07b..b3a10fc8454e 100644 --- a/pkgs/by-name/uc/ucx/package.nix +++ b/pkgs/by-name/uc/ucx/package.nix @@ -32,11 +32,8 @@ let name = "rocm"; paths = rocmList; }; - - # rocm build fails with gcc stdenv due to unrecognised arg parallel-jobs - stdenv' = if enableRocm then rocmPackages.stdenv else stdenv; in -stdenv'.mkDerivation (finalAttrs: { +stdenv.mkDerivation (finalAttrs: { pname = "ucx"; version = "1.19.0"; diff --git a/pkgs/development/python-modules/torch/source/default.nix b/pkgs/development/python-modules/torch/source/default.nix index aa6de6c0a2aa..d9b25ed09bb9 100644 --- a/pkgs/development/python-modules/torch/source/default.nix +++ b/pkgs/development/python-modules/torch/source/default.nix @@ -187,49 +187,52 @@ let # TODO: Retest after ROCm 6.4 or torch 2.8 "gfx1010" "gfx1012" - - # Strix Halo seems to be broken as well, see - # https://github.com/NixOS/nixpkgs/pull/440359. - "gfx1151" - ] (rocmPackages.clr.localGpuTargets or rocmPackages.clr.gpuTargets) + ] rocmPackages.clr.localGpuTargets or rocmPackages.clr.gpuTargets else throw "No GPU targets specified" ); + # Use vendored CK as header only dep if rocmPackages' CK doesn't properly support targets + vendorComposableKernel = rocmSupport && !rocmPackages.composable_kernel.anyMfmaTarget; + rocmtoolkit_joined = symlinkJoin { name = "rocm-merged"; - paths = with rocmPackages; [ - rocm-core - clr - rccl - miopen - aotriton - composable_kernel - rocrand - rocblas - rocsparse - hipsparse - rocthrust - rocprim - hipcub - roctracer - rocfft - rocsolver - hipfft - hiprand - hipsolver - hipblas-common - hipblas - hipblaslt - rocminfo - rocm-comgr - rocm-device-libs - rocm-runtime - rocm-smi - clr.icd - hipify - ]; + paths = + with rocmPackages; + [ + rocm-core + clr + rccl + miopen + aotriton + rocrand + rocblas + rocsparse + hipsparse + rocthrust + rocprim + hipcub + roctracer + rocfft + rocsolver + hipfft + hiprand + hipsolver + hipblas-common + hipblas + hipblaslt + rocminfo + rocm-comgr + rocm-device-libs + rocm-runtime + rocm-smi + clr.icd + hipify + ] + ++ lib.optionals (!vendorComposableKernel) [ + composable_kernel + ]; # Fix `setuptools` not being found postBuild = '' @@ -373,8 +376,9 @@ buildPythonPackage rec { # Doesn't pick up the environment variable? substituteInPlace third_party/kineto/libkineto/CMakeLists.txt \ --replace-fail "\''$ENV{ROCM_SOURCE_DIR}" "${rocmtoolkit_joined}" - - # Use composable kernel as dependency, rather than built-in third-party + '' + # When possible, composable kernel as dependency, rather than built-in third-party + + lib.optionalString (rocmSupport && !vendorComposableKernel) '' substituteInPlace aten/src/ATen/CMakeLists.txt \ --replace-fail "list(APPEND ATen_HIP_INCLUDE \''${CMAKE_CURRENT_SOURCE_DIR}/../../../third_party/composable_kernel/include)" "" \ --replace-fail "list(APPEND ATen_HIP_INCLUDE \''${CMAKE_CURRENT_SOURCE_DIR}/../../../third_party/composable_kernel/library/include)" "" @@ -412,7 +416,7 @@ buildPythonPackage rec { export ROCM_PATH=${rocmtoolkit_joined} export ROCM_SOURCE_DIR=${rocmtoolkit_joined} export PYTORCH_ROCM_ARCH="${gpuTargetString}" - export CMAKE_CXX_FLAGS="-I${rocmtoolkit_joined}/include -I${rocmtoolkit_joined}/include/rocblas" + export CMAKE_CXX_FLAGS="-I${rocmtoolkit_joined}/include" python tools/amd_build/build_amd.py ''; diff --git a/pkgs/development/rocm-modules/6/amdsmi/default.nix b/pkgs/development/rocm-modules/6/amdsmi/default.nix index 2b4db7be7aa6..6f26889ab01b 100644 --- a/pkgs/development/rocm-modules/6/amdsmi/default.nix +++ b/pkgs/development/rocm-modules/6/amdsmi/default.nix @@ -15,24 +15,26 @@ let esmi_ib_src = fetchFromGitHub { owner = "amd"; repo = "esmi_ib_library"; - rev = "esmi_pkg_ver-3.0.3"; - hash = "sha256-q0w5c5c+CpXkklmSyfzc+sbkt4cHNxscGJA3AXwvHxQ="; + rev = "esmi_pkg_ver-4.2"; + hash = "sha256-czF9ezkAO0PuDkXh8y639AcOZH+KVcWiXPX74H5W/nw="; }; in stdenv.mkDerivation (finalAttrs: { pname = "amdsmi"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "rocm"; repo = "amdsmi"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-hrPqd4ZWqzTG7JRfVwc1SZx6TNS0Q/LFg8yDxrM3mPo="; + hash = "sha256-9O29O4HGkQxFDglAhHKv5KWA7p97RwMGG2x/fkOS2jE="; }; postPatch = '' substituteInPlace goamdsmi_shim/CMakeLists.txt \ --replace-fail "amd_smi)" ${"'"}''${AMD_SMI_TARGET})' \ --replace-fail 'target_link_libraries(''${GOAMDSMI_SHIM_TARGET} -L' '#' + substituteInPlace CMakeLists.txt \ + --replace-fail "if(NOT latest_esmi_tag STREQUAL current_esmi_tag)" "if(OFF)" # Manually unpack esmi_ib_src and add amd_hsmp.h so execute-process git clone doesn't run cp -rf --no-preserve=mode ${esmi_ib_src} ./esmi_ib_library @@ -41,10 +43,10 @@ stdenv.mkDerivation (finalAttrs: { ''; patches = [ - # Fix ld.lld undefined reference: drmGetVersion (fetchpatch { - url = "https://github.com/ROCm/amdsmi/commit/c3864bf6171970d86dc50fd23f06377736823997.patch"; - hash = "sha256-zRG1tBD8sIQCWdKfCbXC/Z/6d6NTrRYvRpddPWdM4j8="; + name = "esmi-to-tag-4.2.patch"; + url = "https://github.com/ROCm/amdsmi/commit/49aa2af045a4bc688e6f3ee0545f12afc45c1efe.patch"; + hash = "sha256-5dH9N4m+2mJIGVEB86SvdK3uAYyGFTfbCBJ8e09iQ3w="; }) ]; diff --git a/pkgs/development/rocm-modules/6/aotriton/default.nix b/pkgs/development/rocm-modules/6/aotriton/default.nix index 26b22e61d8c1..99d0c1d46bda 100644 --- a/pkgs/development/rocm-modules/6/aotriton/default.nix +++ b/pkgs/development/rocm-modules/6/aotriton/default.nix @@ -23,17 +23,21 @@ let "gfx90a" "gfx942" "gfx950" + # some gfx1100 kernels fail with error: branch size exceeds simm16 + # but build proceeds and those ops will fallback so it's ok "gfx1100" "gfx1151" "gfx1150" "gfx1201" "gfx1200" ] gpuTargets; - supportedTargets' = lib.concatStringsSep ";" supportedTargets; anySupportedTargets = supportedTargets != [ ]; + # Pick a single arbitrary target to speed up shim build when we can't support our target + supportedTargets' = + if anySupportedTargets then lib.concatStringsSep ";" supportedTargets else "gfx1200"; in stdenv.mkDerivation (finalAttrs: { - pname = "aotriton"; + pname = "aotriton${lib.optionalString (!anySupportedTargets) "-shim"}"; version = "0.10b"; src = fetchFromGitHub { @@ -132,13 +136,16 @@ stdenv.mkDerivation (finalAttrs: { (lib.cmakeBool "AOTRITON_NOIMAGE_MODE" (!anySupportedTargets)) # Use preinstalled triton from our python's site-packages (lib.cmakeBool "AOTRITON_INHERIT_SYSTEM_SITE_TRITON" true) + # FP32 kernels are optional, turn them off to speed up builds and save space + # Perf sensitive code should be using BF16 or F16 + (lib.cmakeBool "AOTRITON_ENABLE_FP32_INPUTS" false) + # Avoid kernels being skipped if build host is overloaded + (lib.cmakeFeature "AOTRITON_GPU_BUILD_TIMEOUT" "0") # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 (lib.cmakeFeature "CMAKE_INSTALL_BINDIR" "bin") (lib.cmakeFeature "CMAKE_INSTALL_LIBDIR" "lib") (lib.cmakeFeature "CMAKE_INSTALL_INCLUDEDIR" "include") - ] - ++ lib.optionals anySupportedTargets [ # Note: build will warn "AMDGPU_TARGETS was not set, and system GPU detection was unsuccsesful." # but this can safely be ignored, aotriton uses a different approach to pass targets (lib.cmakeFeature "AOTRITON_TARGET_ARCH" supportedTargets') diff --git a/pkgs/development/rocm-modules/6/aqlprofile/default.nix b/pkgs/development/rocm-modules/6/aqlprofile/default.nix new file mode 100644 index 000000000000..e1f492515922 --- /dev/null +++ b/pkgs/development/rocm-modules/6/aqlprofile/default.nix @@ -0,0 +1,36 @@ +{ + lib, + stdenv, + clr, + cmake, + fetchFromGitHub, +}: + +stdenv.mkDerivation { + pname = "aqlprofile"; + version = "6.4.3"; + + src = fetchFromGitHub { + # TODO: Will move to rocm-systems repo and have proper tags in 7.x + # pinned to oddly named tag for now + owner = "ROCm"; + repo = "aqlprofile"; + tag = "rocm-42"; + hash = "sha256-avL78ZfB+rJ1TYaejSUzU6i5L9JeMawMwIxaTQINQdE="; + }; + + env.CXXFLAGS = "-DROCP_LD_AQLPROFILE=1"; + + nativeBuildInputs = [ + cmake + clr + ]; + + meta = with lib; { + description = "AQLPROFILE library for AMD HSA runtime API extension support"; + homepage = "https://github.com/ROCm/aqlprofile/"; + license = with licenses; [ mit ]; + teams = [ teams.rocm ]; + platforms = platforms.linux; + }; +} diff --git a/pkgs/development/rocm-modules/6/clr/default.nix b/pkgs/development/rocm-modules/6/clr/default.nix index 11a8cd244743..15c742fc4b70 100644 --- a/pkgs/development/rocm-modules/6/clr/default.nix +++ b/pkgs/development/rocm-modules/6/clr/default.nix @@ -47,26 +47,29 @@ let "--set ROCM_PATH $out" ]; amdclang = writeShellScriptBin "amdclang" '' - exec clang "$@" + exec ${hipClang}/bin/clang "$@" ''; amdclangxx = writeShellScriptBin "amdclang++" '' - exec clang++ "$@" + exec ${hipClang}/bin/clang++ "$@" ''; in stdenv.mkDerivation (finalAttrs: { pname = "clr"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" "icd" ]; + __structuredAttrs = true; + strictDeps = true; + src = fetchFromGitHub { owner = "ROCm"; repo = "clr"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-4qjfnn0kto2sNaSumXxHRHFrf3a3RZILOdhVSxkEs1I="; + hash = "sha256-DOAAuC9TN1//v56GXyUMJwQHgOuctC+WsC5agrgL+QM="; }; nativeBuildInputs = [ @@ -99,8 +102,10 @@ stdenv.mkDerivation (finalAttrs: { rocminfo ]; + cmakeBuildType = "RelWithDebInfo"; + separateDebugInfo = true; + cmakeFlags = [ - "-DCMAKE_BUILD_TYPE=Release" "-DCMAKE_POLICY_DEFAULT_CMP0072=NEW" # Prefer newer OpenGL libraries "-DCLR_BUILD_HIP=ON" "-DCLR_BUILD_OCL=ON" @@ -124,33 +129,27 @@ stdenv.mkDerivation (finalAttrs: { patches = [ ./cmake-find-x11-libgl.patch - (fetchpatch { - # Fix handling of old fatbin version https://github.com/ROCm/clr/issues/99 - sha256 = "sha256-CK/QwgWJQEruiG4DqetF9YM0VEWpSiUMxAf1gGdJkuA="; - url = "https://src.fedoraproject.org/rpms/rocclr/raw/rawhide/f/0001-handle-v1-of-compressed-fatbins.patch"; + # [PATCH] improve rocclr isa compatibility check + sha256 = "sha256-oj1loBEuqzuMihOKoN0wR92Wo25AshN5MpBuTq/9TMw="; + url = "https://github.com/GZGavinZhao/clr/commit/f675b9b46d9f7bb8e003f4f47f616fa86a0b7a5e.patch"; }) (fetchpatch { - # improve rocclr isa compatibility check - sha256 = "sha256-wUrhpYN68AbEXeFU5f366C6peqHyq25kujJXY/bBJMs="; - url = "https://github.com/GZGavinZhao/clr/commit/22c17a0ac09c6b77866febf366591f669a1ed133.patch"; - }) - (fetchpatch { - # [PATCH] Improve hipamd compat check - sha256 = "sha256-uZQ8rMrWH61CCbxwLqQGggDmXFmYTi6x8OcgYPrZRC8="; - url = "https://github.com/GZGavinZhao/clr/commit/63c6ee630966744d4199fdfb854e98d2da9e1122.patch"; - }) - (fetchpatch { - # [PATCH] SWDEV-504340 - Move cast of cl_mem inside the condition - # Fixes crash due to UB in KernelBlitManager::setArgument - sha256 = "sha256-nL4CZ7EOXqsTVUtYhuu9DLOMpnMeMRUhkhylEQLTg9I="; - url = "https://github.com/ROCm/clr/commit/fa63919a6339ea2a61111981ba2362c97fbdf743.patch"; + # [PATCH] improve hipamd isa compatibility check + sha256 = "sha256-E3ERoVjUVWCiYHuE1GaVY5jMrAVx3B1cAVHM4/HPuaQ="; + url = "https://github.com/GZGavinZhao/clr/commit/aec0fc56ee2d10a2bc269c418fa847da2ee9969a.patch"; }) (fetchpatch { # [PATCH] SWDEV-507104 - Removes alignment requirement for Semaphore class to resolve runtime misaligned memory issues sha256 = "sha256-nStJ22B/CM0fzQTvYjbHDbQt0GlE8DXxVK+UDU9BAx4="; url = "https://github.com/ROCm/clr/commit/21d764518363d74187deaef2e66c1a127bc5aa64.patch"; }) + (fetchpatch { + # CMake 4 compat + # [PATCH] SWDEV-509213 - make cmake_minimum_required consistent across clr + url = "https://github.com/ROCm/clr/commit/fcaefe97b862afe12aaac0147f1004e6dc595fce.patch"; + hash = "sha256-hRZXbASbIOOETe+T4mDyyiRWLXd6RDKRieN2ns1w/rs="; + }) ]; postPatch = '' @@ -207,6 +206,7 @@ stdenv.mkDerivation (finalAttrs: { ln -s ${rocm-core}/.info/ $out/.info ln -s ${hipClang} $out/llvm + ln -s ${hipClang}/bin/{ld.lld,lld,clang-offload-bundler,llvm-objcopy,clang,clang++} $out/bin/ ''; disallowedRequisites = [ @@ -217,28 +217,33 @@ stdenv.mkDerivation (finalAttrs: { # All known and valid general GPU targets # We cannot use this for each ROCm library, as each defines their own supported targets # See: https://github.com/ROCm/ROCm/blob/77cbac4abab13046ee93d8b5bf410684caf91145/README.md#library-target-matrix - # Generic targets are not yet available in rocm-6.3.1 llvm gpuTargets = lib.forEach [ - # "9-generic" + # "9-generic" # can handle all Vega variants "900" # MI25, Vega 56/64 - "906" # MI50/60, Radeon VII - "908" # MI100 - "90a" # MI210 / MI250 - # "9-4-generic" + # "902" # Vega 8 + # "909" # Renoir Vega APU + # "90c" # Renoir Vega APU + # Past this point cards need their own kernels for perf despite gfx9-generic compat + "906" # MI50/60, Radeon VII - adds dot product & mixed precision FMA ops + "908" # MI100 - adds MFMA (matrix fused multiply-add) ops + "90a" # MI210/MI250 - additional MFMA variants + # "9-4-generic" - since only 942 is valid for 6.4 target it directly # 940/1 - never released publicly, maybe HPE cray specific MI3xx? - "942" # MI300 - # "10-1-generic" + "942" # MI300A/X, MI325X + # "950" # MI350X TODO: Expected in ROCm 7.x + # "10-1-generic" # fine for all RDNA1 cards "1010" - "1012" # "10-3-generic" "1030" # W6800, various Radeon cards - # "11-generic" + # "11-generic" # will handle 7600, hopefully ryzen AI series iGPUs "1100" "1101" "1102" + # 7.x "1150" "1151" # Strix Halo - "1200" # RX 9070 - "1201" # RX 9070 XT + # "12-generic" + "1200" # RX 9060 + "1201" # RX 9070 + XT ] (target: "gfx${target}"); inherit hipClangPath; diff --git a/pkgs/development/rocm-modules/6/composable_kernel/avoid-extra-host-compile.patch b/pkgs/development/rocm-modules/6/composable_kernel/avoid-extra-host-compile.patch new file mode 100644 index 000000000000..b1ba9929b64f --- /dev/null +++ b/pkgs/development/rocm-modules/6/composable_kernel/avoid-extra-host-compile.patch @@ -0,0 +1,40 @@ +From 6ac72ec84269737626b1f5e43e64729f0922d182 Mon Sep 17 00:00:00 2001 +From: "Ding, Yi" +Date: Wed, 9 Jul 2025 03:12:39 +0000 +Subject: [PATCH] Avoid compile kernel in host pass + +--- + include/ck_tile/host/kernel_launch.hpp | 9 +++++++-- + 1 file changed, 7 insertions(+), 2 deletions(-) + +diff --git a/include/ck_tile/host/kernel_launch.hpp b/include/ck_tile/host/kernel_launch.hpp +index 9770e99738..f6ccb6968b 100644 +--- a/include/ck_tile/host/kernel_launch.hpp ++++ b/include/ck_tile/host/kernel_launch.hpp +@@ -4,11 +4,12 @@ + #pragma once + + #include "ck_tile/core/config.hpp" +-#include "ck_tile/host/stream_config.hpp" ++#include "ck_tile/core/utility/ignore.hpp" + #include "ck_tile/host/hip_check_error.hpp" ++#include "ck_tile/host/stream_config.hpp" + #include "ck_tile/host/timer.hpp" +-#include + #include ++#include + + namespace ck_tile { + +@@ -24,7 +25,11 @@ __launch_bounds__(MaxThreadPerBlock, MinBlockPerCu) + #endif + __global__ void kentry(Args... args) + { ++#if defined(__HIP_DEVICE_COMPILE__) + Kernel{}(args...); ++#else ++ (..., (ignore = args, 0)); ++#endif + } + + // diff --git a/pkgs/development/rocm-modules/6/composable_kernel/base.nix b/pkgs/development/rocm-modules/6/composable_kernel/base.nix index 6856b3378576..a4f59c2e1d5f 100644 --- a/pkgs/development/rocm-modules/6/composable_kernel/base.nix +++ b/pkgs/development/rocm-modules/6/composable_kernel/base.nix @@ -32,6 +32,10 @@ ), }: +# TODO: in 7.x CK is likely to gain support for +# a) miopen kernel only build (MIOPEN_REQ_LIBS_ONLY) +# b) header only build (useful for torch) https://github.com/ROCm/composable_kernel/issues/2030 +# that will likely allow us to get rid of this complicated split part build! stdenv.mkDerivation (finalAttrs: { preBuild = '' echo "This derivation isn't intended to be built directly and only exists to be overridden and built in chunks"; @@ -39,9 +43,7 @@ stdenv.mkDerivation (finalAttrs: { ''; pname = "composable_kernel_base"; - # Picked this version over 6.3 because much easier to get to build - # and it matches the version torch 2.6 wants - version = "6.4.0-unstable-2024-12-20"; + version = "6.4-unstable-2025-05-22"; outputs = [ "out" @@ -56,8 +58,9 @@ stdenv.mkDerivation (finalAttrs: { src = fetchFromGitHub { owner = "ROCm"; repo = "composable_kernel"; - rev = "07339c738396ebeae57374771ded4dcf11bddf1e"; - hash = "sha256-EvEBxlOpQ71BF57VW79WBo/cdxAwTKFXFMiYKyGyyEs="; + # Using a dev snapshot, trying to get MIOpen to work + rev = "bc2551ac3b27edc31f20863e3a873508fb73aad2"; + hash = "sha256-bfmwbgR1ya+zkME3wOyaZX/e+1+ie0sSlugK/kozLsI="; }; nativeBuildInputs = [ @@ -100,7 +103,9 @@ stdenv.mkDerivation (finalAttrs: { "-DCMAKE_INSTALL_LIBDIR=lib" "-DCMAKE_INSTALL_INCLUDEDIR=include" "-DBUILD_DEV=OFF" + "-DBUILD_MHA_LIB=ON" "-DROCM_PATH=${clr}" + "-DENABLE_CLANG_CPP_CHECKS=OFF" "-DCMAKE_HIP_COMPILER_ROCM_ROOT=${clr}" # FP8 can build for 908/90a but very slow build @@ -118,6 +123,11 @@ stdenv.mkDerivation (finalAttrs: { "-DGOOGLETEST_DIR=${gtest.src}" # Custom linker names ]; + patches = [ + # Significant build performance improvement + ./avoid-extra-host-compile.patch + ]; + # No flags to build selectively it seems... postPatch = # Reduce configure time by preventing thousands of clang-tidy targets being added @@ -128,6 +138,14 @@ stdenv.mkDerivation (finalAttrs: { --replace-fail clang_tidy_check '#clang_tidy_check' substituteInPlace CMakeLists.txt \ --replace-fail "add_subdirectory(profiler)" "" + substituteInPlace cmake/EnableCompilerWarnings.cmake \ + --replace-fail "-Werror" "" + + # Apply equivalent change to https://github.com/ROCm/composable_kernel/pull/2564 + # TODO: Remove after ROCm 7.1 + find include/ck/tensor_operation/ -type f -name "*.hpp" -exec sed -i \ + -e 's/!defined(__HIP_DEVICE_COMPILE__)/false/g' \ + {} + '' # Optionally remove tests + lib.optionalString (!buildTests) '' @@ -152,20 +170,29 @@ stdenv.mkDerivation (finalAttrs: { mv $out/bin/example_* $example/bin ''; - passthru.updateScript = rocmUpdateScript { - name = finalAttrs.pname; - inherit (finalAttrs.src) owner; - inherit (finalAttrs.src) repo; + passthru = { + inherit gpuTargets; + updateScript = rocmUpdateScript { + name = finalAttrs.pname; + inherit (finalAttrs.src) owner; + inherit (finalAttrs.src) repo; + }; + anyGfx9Target = lib.lists.any (lib.strings.hasPrefix "gfx9") gpuTargets; + anyMfmaTarget = + (lib.lists.intersectLists gpuTargets [ + "gfx908" + "gfx90a" + "gfx942" + "gfx950" + ]) != [ ]; }; - passthru.anyGfx9Target = lib.lists.any (lib.strings.hasPrefix "gfx9") gpuTargets; - meta = with lib; { description = "Performance portable programming model for machine learning tensor operators"; homepage = "https://github.com/ROCm/composable_kernel"; license = with licenses; [ mit ]; teams = [ teams.rocm ]; platforms = platforms.linux; - broken = true; + broken = true; # this base package shouldn't be built directly }; }) diff --git a/pkgs/development/rocm-modules/6/composable_kernel/ck4inductor.nix b/pkgs/development/rocm-modules/6/composable_kernel/ck4inductor.nix index f7068c547955..c60552c490a2 100644 --- a/pkgs/development/rocm-modules/6/composable_kernel/ck4inductor.nix +++ b/pkgs/development/rocm-modules/6/composable_kernel/ck4inductor.nix @@ -14,7 +14,7 @@ buildPythonPackage { setuptools setuptools-scm ]; - version = "6.4.0"; + version = "6.4.3"; inherit (composable_kernel) src; pythonImportsCheck = [ "ck4inductor" diff --git a/pkgs/development/rocm-modules/6/composable_kernel/default.nix b/pkgs/development/rocm-modules/6/composable_kernel/default.nix index 0ec02b962cc1..42c634f0aa94 100644 --- a/pkgs/development/rocm-modules/6/composable_kernel/default.nix +++ b/pkgs/development/rocm-modules/6/composable_kernel/default.nix @@ -17,61 +17,90 @@ let targets = [ "device_mha_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; extraCmakeFlags = [ "-DHIP_CLANG_NUM_PARALLEL_JOBS=2" ]; }; gemm_multiply_multiply = { targets = [ "device_gemm_multiply_multiply_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; extraCmakeFlags = [ "-DHIP_CLANG_NUM_PARALLEL_JOBS=2" ]; + onlyFor = [ + "gfx942" + "gfx950" + ]; }; - grouped_conv = { + gemm_multiply_multiply_wp = { + targets = [ + "device_gemm_multiply_multiply_wp_instance" + ]; + extraCmakeFlags = [ "-DHIP_CLANG_NUM_PARALLEL_JOBS=2" ]; + onlyFor = [ + "gfx942" + "gfx950" + ]; + }; + grouped_conv_bwd = { targets = [ "device_grouped_conv1d_bwd_weight_instance" "device_grouped_conv2d_bwd_data_instance" "device_grouped_conv2d_bwd_weight_instance" + ]; + }; + grouped_conv_fwd = { + targets = [ "device_grouped_conv1d_fwd_instance" "device_grouped_conv2d_fwd_instance" + "device_grouped_conv2d_fwd_bias_relu_instance" "device_grouped_conv2d_fwd_dynamic_op_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; - grouped_conv_bwd_3d = { + grouped_conv_bwd_3d1 = { targets = [ "device_grouped_conv3d_bwd_data_instance" "device_grouped_conv3d_bwd_data_bilinear_instance" "device_grouped_conv3d_bwd_data_scale_instance" + ]; + }; + grouped_conv_bwd_3d2 = { + targets = [ "device_grouped_conv3d_bwd_weight_instance" "device_grouped_conv3d_bwd_weight_bilinear_instance" "device_grouped_conv3d_bwd_weight_scale_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; - grouped_conv_fwd_3d = { + grouped_conv_fwd_3d1 = { targets = [ "device_grouped_conv3d_fwd_instance" + "device_grouped_conv3d_fwd_bias_relu_instance" "device_grouped_conv3d_fwd_bilinear_instance" "device_grouped_conv3d_fwd_convinvscale_instance" "device_grouped_conv3d_fwd_convscale_instance" "device_grouped_conv3d_fwd_convscale_add_instance" + ]; + }; + grouped_conv_fwd_3d2 = { + targets = [ "device_grouped_conv3d_fwd_convscale_relu_instance" "device_grouped_conv3d_fwd_dynamic_op_instance" "device_grouped_conv3d_fwd_scale_instance" "device_grouped_conv3d_fwd_scaleadd_ab_instance" "device_grouped_conv3d_fwd_scaleadd_scaleadd_relu_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; - batched_gemm = { + batched_gemm1 = { targets = [ "device_batched_gemm_instance" + "device_batched_gemm_b_scale_instance" + "device_batched_gemm_multi_d_instance" "device_batched_gemm_add_relu_gemm_add_instance" "device_batched_gemm_bias_permute_instance" "device_batched_gemm_gemm_instance" "device_batched_gemm_reduce_instance" "device_batched_gemm_softmax_gemm_instance" + ]; + }; + batched_gemm2 = { + targets = [ "device_batched_gemm_softmax_gemm_permute_instance" "device_grouped_gemm_instance" "device_grouped_gemm_bias_instance" @@ -80,27 +109,33 @@ let "device_grouped_gemm_fixed_nk_multi_abd_instance" "device_grouped_gemm_tile_loop_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; - gemm_universal = { + gemm_universal1 = { targets = [ "device_gemm_universal_instance" "device_gemm_universal_batched_instance" + ]; + }; + gemm_universal2 = { + targets = [ "device_gemm_universal_reduce_instance" "device_gemm_universal_streamk_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; - extraCmakeFlags = [ "-DHIP_CLANG_NUM_PARALLEL_JOBS=2" ]; }; - gemm_other = { + gemm_other1 = { targets = [ "device_gemm_instance" + "device_gemm_b_scale_instance" "device_gemm_ab_scale_instance" "device_gemm_add_instance" "device_gemm_add_add_fastgelu_instance" "device_gemm_add_fastgelu_instance" "device_gemm_add_multiply_instance" "device_gemm_add_relu_instance" + ]; + }; + gemm_other2 = { + targets = [ "device_gemm_add_relu_add_layernorm_instance" "device_gemm_add_silu_instance" "device_gemm_bias_add_reduce_instance" @@ -112,7 +147,6 @@ let "device_gemm_splitk_instance" "device_gemm_streamk_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; conv = { targets = [ @@ -123,7 +157,6 @@ let "device_conv2d_fwd_bias_relu_add_instance" "device_conv3d_bwd_data_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; pool = { targets = [ @@ -141,11 +174,7 @@ let "device_contraction_scale_instance" "device_elementwise_instance" "device_elementwise_normalization_instance" - "device_normalization_bwd_data_instance" - "device_normalization_bwd_gamma_beta_instance" - "device_normalization_fwd_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; other2 = { targets = [ @@ -154,10 +183,16 @@ let "device_permute_scale_instance" "device_quantization_instance" "device_reduce_instance" + ]; + }; + other3 = { + targets = [ + "device_normalization_bwd_data_instance" + "device_normalization_bwd_gamma_beta_instance" + "device_normalization_fwd_instance" "device_softmax_instance" "device_transpose_instance" ]; - requiredSystemFeatures = [ "big-parallel" ]; }; }; tensorOpBuilder = @@ -165,51 +200,59 @@ let part, targets, extraCmakeFlags ? [ ], - requiredSystemFeatures ? [ ], + requiredSystemFeatures ? [ "big-parallel" ], + onlyFor ? [ ], }: - composable_kernel_base.overrideAttrs (old: { - inherit requiredSystemFeatures; - pname = "composable_kernel${clr.gpuArchSuffix}-${part}"; - makeTargets = targets; - preBuild = '' - echo "Building ${part}" - makeFlagsArray+=($makeTargets) - substituteInPlace Makefile \ - --replace-fail '.NOTPARALLEL:' "" - ''; + let + supported = + onlyFor == [ ] || (lib.lists.intersectLists composable_kernel_base.gpuTargets onlyFor) != [ ]; + in + if supported then + (composable_kernel_base.overrideAttrs (old: { + inherit requiredSystemFeatures; + pname = "composable_kernel${clr.gpuArchSuffix}-${part}"; + makeTargets = targets; + preBuild = '' + echo "Building ${part}" + makeFlagsArray+=($makeTargets) + substituteInPlace $(find ./ -name "Makefile" -type f) \ + --replace-fail '.NOTPARALLEL:' '.UNUSED_NOTPARALLEL:' + ''; - # Compile parallelism adjusted based on available RAM - # Never uses less than NIX_BUILD_CORES/4, never uses more than NIX_BUILD_CORES - # CK uses an unusually high amount of memory per core in the build step - # Nix/nixpkgs doesn't really have any infra to tell it that this build is unusually memory hungry - # So, bodge. Otherwise you end up having to build all of ROCm with a low core limit when - # it's only this package that has trouble. - preConfigure = old.preConfigure or "" + '' - MEM_GB_TOTAL=$(awk '/MemTotal/ { printf "%d \n", $2/1024/1024 }' /proc/meminfo) - MEM_GB_AVAILABLE=$(awk '/MemAvailable/ { printf "%d \n", $2/1024/1024 }' /proc/meminfo) - APPX_GB=$((MEM_GB_AVAILABLE > MEM_GB_TOTAL ? MEM_GB_TOTAL : MEM_GB_AVAILABLE)) - MAX_CORES=$((1 + APPX_GB/3)) - MAX_CORES=$((MAX_CORES < NIX_BUILD_CORES/3 ? NIX_BUILD_CORES/3 : MAX_CORES)) - export NIX_BUILD_CORES="$((NIX_BUILD_CORES > MAX_CORES ? MAX_CORES : NIX_BUILD_CORES))" - echo "Picked new core limit NIX_BUILD_CORES=$NIX_BUILD_CORES based on available mem: $APPX_GB GB" - cmakeFlagsArray+=( - "-DCK_PARALLEL_COMPILE_JOBS=$NIX_BUILD_CORES" - ) - ''; - cmakeFlags = old.cmakeFlags ++ extraCmakeFlags; - # Early exit after build phase with success, skips fixups etc - # Will get copied back into /build of the final CK - postBuild = '' - find . -name "*.o" -type f | while read -r file; do - mkdir -p "$out/$(dirname "$file")" - cp --reflink=auto "$file" "$out/$file" - done - exit 0 - ''; - meta = old.meta // { - broken = false; - }; - }); + # Compile parallelism adjusted based on available RAM + # Never uses less than NIX_BUILD_CORES/4, never uses more than NIX_BUILD_CORES + # CK uses an unusually high amount of memory per core in the build step + # Nix/nixpkgs doesn't really have any infra to tell it that this build is unusually memory hungry + # So, bodge. Otherwise you end up having to build all of ROCm with a low core limit when + # it's only this package that has trouble. + preConfigure = old.preConfigure or "" + '' + MEM_GB_TOTAL=$(awk '/MemTotal/ { printf "%d \n", $2/1024/1024 }' /proc/meminfo) + MEM_GB_AVAILABLE=$(awk '/MemAvailable/ { printf "%d \n", $2/1024/1024 }' /proc/meminfo) + APPX_GB=$((MEM_GB_AVAILABLE > MEM_GB_TOTAL ? MEM_GB_TOTAL : MEM_GB_AVAILABLE)) + MAX_CORES=$((1 + APPX_GB/3)) + MAX_CORES=$((MAX_CORES < NIX_BUILD_CORES/3 ? NIX_BUILD_CORES/3 : MAX_CORES)) + export NIX_BUILD_CORES="$((NIX_BUILD_CORES > MAX_CORES ? MAX_CORES : NIX_BUILD_CORES))" + echo "Picked new core limit NIX_BUILD_CORES=$NIX_BUILD_CORES based on available mem: $APPX_GB GB" + cmakeFlagsArray+=( + "-DCK_PARALLEL_COMPILE_JOBS=$NIX_BUILD_CORES" + ) + ''; + cmakeFlags = old.cmakeFlags ++ extraCmakeFlags; + # Early exit after build phase with success, skips fixups etc + # Will get copied back into /build of the final CK + postBuild = '' + find . -name "*.o" -type f | while read -r file; do + mkdir -p "$out/$(dirname "$file")" + cp --reflink=auto "$file" "$out/$file" + done + exit 0 + ''; + meta = old.meta // { + broken = false; + }; + })) + else + null; composable_kernel_parts = builtins.mapAttrs ( part: targets: tensorOpBuilder (targets // { inherit part; }) ) parts; @@ -218,8 +261,8 @@ in composable_kernel_base.overrideAttrs ( finalAttrs: old: { pname = "composable_kernel${clr.gpuArchSuffix}"; - parts_dirs = builtins.attrValues composable_kernel_parts; - disallowedReferences = builtins.attrValues composable_kernel_parts; + parts_dirs = builtins.filter (x: x != null) (builtins.attrValues composable_kernel_parts); + disallowedReferences = builtins.filter (x: x != null) (builtins.attrValues composable_kernel_parts); preBuild = '' for dir in $parts_dirs; do find "$dir" -type f -name "*.o" | while read -r file; do @@ -240,7 +283,7 @@ composable_kernel_base.overrideAttrs ( parts = composable_kernel_parts; }; meta = old.meta // { - # Builds which don't don't target any gfx9 cause cmake errors in dependent projects + # Builds without any gfx9 fail broken = !finalAttrs.passthru.anyGfx9Target; }; } diff --git a/pkgs/development/rocm-modules/6/default.nix b/pkgs/development/rocm-modules/6/default.nix index 81954e88bb25..7c259148e23d 100644 --- a/pkgs/development/rocm-modules/6/default.nix +++ b/pkgs/development/rocm-modules/6/default.nix @@ -6,15 +6,12 @@ recurseIntoAttrs, symlinkJoin, fetchFromGitHub, - ffmpeg_4, boost179, opencv, libjpeg_turbo, python3Packages, - triton-llvm, openmpi, stdenv, - rocmGpuArches ? [ ], }: let @@ -22,14 +19,16 @@ let self: let inherit (self) llvm; + origStdenv = stdenv; pyPackages = python3Packages; openmpi-orig = openmpi; + rocmClangStdenv = llvm.rocmClangStdenv; in { - inherit rocmGpuArches; + inherit rocmClangStdenv; + stdenv = rocmClangStdenv; buildTests = false; buildBenchmarks = false; - stdenv = llvm.rocmClangStdenv; rocmPath = self.callPackage ./rocm-path { }; rocmUpdateScript = self.callPackage ./update.nix { }; @@ -37,38 +36,43 @@ let ## ROCm ## llvm = recurseIntoAttrs ( callPackage ./llvm/default.nix { - inherit (self) rocm-device-libs rocm-runtime; + # rocm-device-libs is used for .src only + # otherwise would cause infinite recursion + inherit (self) rocm-device-libs; } ); inherit (self.llvm) rocm-merged-llvm clang openmp; - rocm-core = self.callPackage ./rocm-core { }; + rocm-core = self.callPackage ./rocm-core { stdenv = origStdenv; }; + + rocm-cmake = self.callPackage ./rocm-cmake { stdenv = origStdenv; }; + + rocm-device-libs = self.callPackage ./rocm-device-libs { + stdenv = origStdenv; + inherit (llvm) rocm-merged-llvm; + }; + + rocm-runtime = self.callPackage ./rocm-runtime { + stdenv = origStdenv; + inherit (llvm) rocm-merged-llvm; + }; + + rocm-comgr = self.callPackage ./rocm-comgr { + stdenv = origStdenv; + inherit (llvm) rocm-merged-llvm; + }; + + rocminfo = self.callPackage ./rocminfo { stdenv = origStdenv; }; + amdsmi = pyPackages.callPackage ./amdsmi { inherit (self) rocmUpdateScript; }; - rocm-cmake = self.callPackage ./rocm-cmake { }; - rocm-smi = pyPackages.callPackage ./rocm-smi { inherit (self) rocmUpdateScript; }; - rocm-device-libs = self.callPackage ./rocm-device-libs { - inherit (llvm) rocm-merged-llvm; - }; - - rocm-runtime = self.callPackage ./rocm-runtime { - inherit (llvm) rocm-merged-llvm; - }; - - rocm-comgr = self.callPackage ./rocm-comgr { - inherit (llvm) rocm-merged-llvm; - }; - - rocminfo = self.callPackage ./rocminfo { }; - - # Unfree - hsa-amd-aqlprofile-bin = self.callPackage ./hsa-amd-aqlprofile-bin { }; + aqlprofile = self.callPackage ./aqlprofile { }; rdc = self.callPackage ./rdc { }; @@ -76,17 +80,18 @@ let hip-common = self.callPackage ./hip-common { }; - # Eventually will be in the LLVM repo hipcc = self.callPackage ./hipcc { + stdenv = origStdenv; inherit (llvm) rocm-merged-llvm; }; # Replaces hip, opencl-runtime, and rocclr clr = self.callPackage ./clr { }; - aotriton = self.callPackage ./aotriton { inherit stdenv; }; + aotriton = self.callPackage ./aotriton { stdenv = origStdenv; }; hipify = self.callPackage ./hipify { + stdenv = origStdenv; inherit (llvm) clang rocm-merged-llvm @@ -103,7 +108,6 @@ let inherit (llvm) clang; }; - # Needs GCC roctracer = self.callPackage ./roctracer { }; rocgdb = self.callPackage ./rocgdb { }; @@ -128,13 +132,6 @@ let rccl = self.callPackage ./rccl { }; - # RCCL with sanitizers and tests - # Can't have with sanitizer build as dep of other packages without - # runtime crashes due to ASAN not loading first - rccl-tests = self.callPackage ./rccl { - buildTests = true; - }; - hipcub = self.callPackage ./hipcub { }; hipsparse = self.callPackage ./hipsparse { }; @@ -194,7 +191,7 @@ let miopen-hip = self.miopen; - migraphx = self.callPackage ./migraphx { }; + migraphx = self.callPackage ./migraphx { stdenv = origStdenv; }; rpp = self.callPackage ./rpp { }; @@ -214,10 +211,8 @@ let }; mivisionx = self.callPackage ./mivisionx { + stdenv = origStdenv; opencv = opencv.override { enablePython = true; }; - # TODO: Remove this pin in ROCm 6.4+ - # FFMPEG support was improved in https://github.com/ROCm/MIVisionX/pull/1460 - ffmpeg = ffmpeg_4; # Unfortunately, rocAL needs a custom libjpeg-turbo until further notice # See: https://github.com/ROCm/MIVisionX/issues/1051 libjpeg_turbo = libjpeg_turbo.overrideAttrs { @@ -275,7 +270,7 @@ let rocm-developer-tools = symlinkJoin { name = "rocm-developer-tools-meta"; paths = [ - hsa-amd-aqlprofile-bin + aqlprofile rocm-core rocr-debug-agent roctracer @@ -298,7 +293,6 @@ let name = "rocm-ml-libraries-meta"; paths = [ llvm.clang - llvm.mlir llvm.openmp rocm-core miopen-hip @@ -357,7 +351,6 @@ let paths = [ rocm-core llvm.clang - llvm.mlir llvm.openmp # openmp-extras-devel (https://github.com/ROCm/aomp) rocm-language-runtime ]; @@ -389,7 +382,6 @@ let hipify rocm-cmake llvm.clang - llvm.mlir llvm.openmp rocm-runtime rocm-hip-runtime @@ -440,6 +432,10 @@ let }; } // lib.optionalAttrs config.allowAliases { + hsa-amd-aqlprofile-bin = lib.warn '' + 'hsa-amd-aqlprofile-bin' has been replaced by 'aqlprofile'. + '' self.aqlprofile; # Added 2025-08-27 + triton = throw '' 'rocmPackages.triton' has been removed. Please use python3Packages.triton ''; # Added 2025-08-24 diff --git a/pkgs/development/rocm-modules/6/half/default.nix b/pkgs/development/rocm-modules/6/half/default.nix index 953152b37e0f..429439591ffd 100644 --- a/pkgs/development/rocm-modules/6/half/default.nix +++ b/pkgs/development/rocm-modules/6/half/default.nix @@ -9,7 +9,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "half"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; diff --git a/pkgs/development/rocm-modules/6/hip-common/default.nix b/pkgs/development/rocm-modules/6/hip-common/default.nix index a169c97cae54..1c967cf065b8 100644 --- a/pkgs/development/rocm-modules/6/hip-common/default.nix +++ b/pkgs/development/rocm-modules/6/hip-common/default.nix @@ -7,13 +7,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "hip-common"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "HIP"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-eKDbuG57KCg/nZpy+RcDsjyw0pP68qO9K/3ZpSPuMNw="; + hash = "sha256-B4Gc119iff3ak9tmpz3rUJBtCk5T1AA8z67K9PshTLQ="; }; dontConfigure = true; diff --git a/pkgs/development/rocm-modules/6/hipblas-common/default.nix b/pkgs/development/rocm-modules/6/hipblas-common/default.nix index 626c36833d63..8763e916c034 100644 --- a/pkgs/development/rocm-modules/6/hipblas-common/default.nix +++ b/pkgs/development/rocm-modules/6/hipblas-common/default.nix @@ -8,13 +8,13 @@ }: stdenv.mkDerivation (finalAttrs: { pname = "hipblas-common"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "hipBLAS-common"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-tvNz4ymQ1y3YSUQxAtNu2who79QzSKR+3JEevr+GDWo="; + hash = "sha256-eTwoAXH2HGdSAOLTZHJUFHF+c2wWHixqeMqr60KxJrc="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hipblas/default.nix b/pkgs/development/rocm-modules/6/hipblas/default.nix index ca236f1b85ff..0622f6f0b263 100644 --- a/pkgs/development/rocm-modules/6/hipblas/default.nix +++ b/pkgs/development/rocm-modules/6/hipblas/default.nix @@ -18,12 +18,14 @@ buildTests ? false, buildBenchmarks ? false, buildSamples ? false, + # for passthru.tests + hipblas, }: # Can also use cuBLAS stdenv.mkDerivation (finalAttrs: { pname = "hipblas"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -42,15 +44,16 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipBLAS"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-Rz1KAhBUbvErHTF2PM1AkVhqo4OHldfSNMSpp5Tx9yk="; + hash = "sha256-lQv8Ik6+0ldqyeJ05CSGB0309nIpzlRL3CRYeQxVfd0="; }; patches = [ - # https://github.com/ROCm/hipBLAS/pull/952 (fetchpatch { - name = "transitively-depend-hipblas-common.patch"; - url = "https://github.com/ROCm/hipBLAS/commit/54220fdaebf0fb4fd0921ee9e418ace5b143ec8f.patch"; - hash = "sha256-MFEhv8Bkrd2zD0FFIDg9oJzO7ztdyMAF+R9oYA0rmwQ="; + # Subject: [PATCH] Add gfx1150, gfx1150, gfx1200, gfx1201 support (#1055) + # This was merged to release/rocm-rel-6.4 but AMD forgot to tag it for 6.4.3 + name = "release-6.4-arch-extra.patch"; + url = "https://github.com/ROCm/hipBLAS/commit/0100b32ccff9a0f12134694315b4e44884e25a8e.patch"; + hash = "sha256-BmktlLJpYaTcogHzEKpZdCnksIIysEO47WMezXoxvCs="; }) ]; @@ -82,15 +85,13 @@ stdenv.mkDerivation (finalAttrs: { ]; cmakeFlags = [ - "-DCMAKE_BUILD_TYPE=Release" - "-DCMAKE_CXX_COMPILER=${lib.getExe' clr "hipcc"}" - # Upstream is migrating to amdclang++, it is likely this will be correct in next version bump - #"-DCMAKE_CXX_COMPILER=${lib.getBin clr}/bin/amdclang++" + "-DCMAKE_CXX_COMPILER=${lib.getExe' clr "amdclang++"}" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" "-DCMAKE_INSTALL_LIBDIR=lib" "-DCMAKE_INSTALL_INCLUDEDIR=include" + "-DBUILD_WITH_SOLVER=ON" "-DAMDGPU_TARGETS=${rocblas.amdgpu_targets}" ] ++ lib.optionals buildTests [ @@ -125,6 +126,11 @@ stdenv.mkDerivation (finalAttrs: { inherit (finalAttrs.src) owner; inherit (finalAttrs.src) repo; }; + passthru.tests.hipblas-tested = hipblas.override { + buildTests = true; + buildBenchmarks = true; + buildSamples = true; + }; meta = with lib; { description = "ROCm BLAS marshalling library"; diff --git a/pkgs/development/rocm-modules/6/hipblaslt/default.nix b/pkgs/development/rocm-modules/6/hipblaslt/default.nix index cbcca61af76f..47bfd00909fc 100644 --- a/pkgs/development/rocm-modules/6/hipblaslt/default.nix +++ b/pkgs/development/rocm-modules/6/hipblaslt/default.nix @@ -1,232 +1,229 @@ { lib, stdenv, - fetchpatch, fetchFromGitHub, + writableTmpDirAsHomeHook, cmake, rocm-cmake, + rocm-smi, + pkg-config, clr, gfortran, gtest, - msgpack, + boost, + msgpack-cxx, + amd-blis, libxml2, python3, python3Packages, openmp, hipblas-common, - tensile, lapack-reference, ncurses, + ninja, libffi, zlib, zstd, rocmUpdateScript, buildTests ? false, - buildBenchmarks ? false, buildSamples ? false, # hipblaslt supports only devices with MFMA or WMMA - # WMMA on gfx1100 may be broken - # MFMA on MI100 may be broken - # MI200/MI300 known to work - gpuTargets ? ( - clr.localGpuTargets or [ - # "gfx908" FIXME: confirm MFMA on MI100 works - "gfx90a" - "gfx942" - # "gfx1100" FIXME: confirm WMMA targets work - ] - ), + gpuTargets ? (clr.localGpuTargets or clr.gpuTargets), }: -stdenv.mkDerivation ( - finalAttrs: - let - supportsTargetArches = - (builtins.any (lib.strings.hasPrefix "gfx9") gpuTargets) - || (builtins.any (lib.strings.hasPrefix "gfx11") gpuTargets); - tensile' = (tensile.override { isTensileLite = true; }).overrideAttrs { - inherit (finalAttrs) src; - sourceRoot = "${finalAttrs.src.name}/tensilelite"; - }; - py = python3.withPackages (ps: [ - ps.pyyaml - ps.setuptools - ps.packaging - ]); - gpuTargets' = lib.optionalString supportsTargetArches (lib.concatStringsSep ";" gpuTargets); - compiler = "amdclang++"; - cFlags = "-O3 -I${msgpack}/include"; # FIXME: cmake files need patched to include this properly - in - { - pname = "hipblaslt${clr.gpuArchSuffix}"; - version = "6.3.3"; - - src = fetchFromGitHub { - owner = "ROCm"; - repo = "hipBLASLt"; - rev = "rocm-${finalAttrs.version}"; - hash = "sha256-ozfHwsxcczzYXN9SIkyfRvdtaCqlDN4bh3UHZNS2oVQ="; - }; - env.CXX = compiler; - env.CFLAGS = cFlags; - env.CXXFLAGS = cFlags; - env.ROCM_PATH = "${clr}"; - env.TENSILE_ROCM_ASSEMBLER_PATH = lib.getExe' clr "amdclang++"; - env.TENSILE_GEN_ASSEMBLY_TOOLCHAIN = lib.getExe' clr "amdclang++"; - # Some tensile scripts look for this as an env var rather than a cmake flag - env.CMAKE_CXX_COMPILER = lib.getExe' clr "amdclang++"; - requiredSystemFeatures = [ "big-parallel" ]; - - outputs = [ - "out" +let + # hipblaslt is extremely particular about what it will build with + # so intersect with a known supported list and use only those + supportedTargets = ( + lib.lists.intersectLists gpuTargets [ + "gfx908" + "gfx90a" + "gfx942" + "gfx950" + "gfx1100" + "gfx1101" + # 7.x "gfx1150" + "gfx1151" + "gfx1200" + "gfx1201" ] - ++ lib.optionals buildTests [ - "test" - ] - ++ lib.optionals buildBenchmarks [ - "benchmark" - ] - ++ lib.optionals buildSamples [ - "sample" - ]; + ); + supportsTargetArches = supportedTargets != [ ]; + py = python3.withPackages (ps: [ + ps.pyyaml + ps.setuptools + ps.packaging + ps.nanobind + ps.joblib + ps.msgpack + ]); + # workaround: build for one working target if no targets are supported + # a few CXX files are still build for the device + gpuTargets' = + if supportsTargetArches then (lib.concatStringsSep ";" supportedTargets) else "gfx1200"; + compiler = "amdclang++"; + # no-switch due to spammy warnings on some cases with fixme messages + # FIXME(LunNova@): cmake files need patched to include this properly or + # maybe we improve the toolchain to use config files + assemble a sysroot + # so system wide include assumptions work + cFlags = "-Wno-switch -fopenmp -I${lib.getDev zstd}/include -I${amd-blis}/include/blis/ -I${lib.getDev msgpack-cxx}/include"; +in +stdenv.mkDerivation (finalAttrs: { + pname = "hipblaslt${clr.gpuArchSuffix}"; + version = "6.5-unstable-2025-08-21"; - postPatch = '' - mkdir -p build/Tensile/library - # git isn't needed and we have no .git - substituteInPlace cmake/Dependencies.cmake \ - --replace-fail "find_package(Git REQUIRED)" "" - substituteInPlace CMakeLists.txt \ - --replace-fail "include(virtualenv)" "" \ - --replace-fail "virtualenv_install(\''${Tensile_TEST_LOCAL_PATH})" "" \ - --replace-fail "virtualenv_install(\''${CMAKE_SOURCE_DIR}/tensilelite)" "" \ - --replace-fail 'find_package(Tensile 4.33.0 EXACT REQUIRED HIP LLVM OpenMP PATHS "''${INSTALLED_TENSILE_PATH}")' "find_package(Tensile)" \ - --replace-fail 'Tensile_CPU_THREADS ""' 'Tensile_CPU_THREADS "$ENV{NIX_BUILD_CORES}"' - # FIXME: TensileCreateExtOpLibraries build failure due to unsupported null operand - # Working around for now by disabling the ExtOp libs - substituteInPlace library/src/amd_detail/rocblaslt/src/CMakeLists.txt \ - --replace-fail 'TensileCreateExtOpLibraries("' '# skipping TensileCreateExtOpLibraries' - substituteInPlace library/src/amd_detail/rocblaslt/src/kernels/compile_code_object.sh \ - --replace-fail '${"\${rocm_path}"}/bin/' "" - ''; + src = fetchFromGitHub { + owner = "ROCm"; + repo = "rocm-libraries"; + rev = "a676499add42941ff6af1e8d3f0504416dac7429"; + hash = "sha256-zIYdHFbHyP2V6dkx6Ueb6NBqWu8tJji2hSWF9zWEJa4="; + sparseCheckout = [ "projects/hipblaslt" ]; + }; + sourceRoot = "${finalAttrs.src.name}/projects/hipblaslt"; + env.CXX = compiler; + env.CFLAGS = cFlags; + env.CXXFLAGS = cFlags; + env.ROCM_PATH = "${clr}"; + env.TENSILE_ROCM_ASSEMBLER_PATH = lib.getExe' clr "amdclang++"; + env.TENSILE_GEN_ASSEMBLY_TOOLCHAIN = lib.getExe' clr "amdclang++"; + requiredSystemFeatures = [ "big-parallel" ]; - # Apply patches to allow building without a target arch if we need to do that - patches = lib.optionals (!supportsTargetArches) [ - # Add ability to build without specitying any arch. - (fetchpatch { - sha256 = "sha256-VW3bPzmQvfo8+iKsVfpn4sbqAe41fLzCEUfBh9JxVyk="; - url = "https://raw.githubusercontent.com/gentoo/gentoo/refs/heads/master/sci-libs/hipBLASLt/files/hipBLASLt-6.1.1-no-arch.patch"; - }) - # Followup to above patch for 6.3.x - (fetchpatch { - sha256 = "sha256-GCsrne6BiWzwj8TMAfFuaYz1Pij97hoCc6E3qJhWb10="; - url = "https://raw.githubusercontent.com/gentoo/gentoo/refs/heads/master/sci-libs/hipBLASLt/files/hipBLASLt-6.3.0-no-arch-extra.patch"; - }) - ]; + __structuredAttrs = true; + strictDeps = true; - doCheck = false; - doInstallCheck = false; + outputs = [ + "out" + # benchmarks are non-optional + "benchmark" + ] + ++ lib.optionals buildTests [ + "test" + ] + ++ lib.optionals buildSamples [ + "sample" + ]; - nativeBuildInputs = [ - cmake - rocm-cmake - py - clr - gfortran - # need make to get streaming console output so nix knows build is still running - # so deliberately not using ninja - ]; + patches = [ + # Upstream issue requesting properly specifying + # parallel-jobs for these invocations + # https://github.com/ROCm/rocm-libraries/issues/1242 + ./parallel-buildSourceCodeObjectFile.diff + # Support loading zstd compressed .dat files, required to keep output under + # hydra size limit + ./messagepack-compression-support.patch + ]; - buildInputs = [ - hipblas-common - tensile' - openmp - libffi - ncurses + postPatch = '' + # git isn't needed and we have no .git + substituteInPlace cmake/dependencies.cmake \ + --replace-fail "find_package(Git REQUIRED)" "" + substituteInPlace CMakeLists.txt \ + --replace-fail " LANGUAGES CXX" " LANGUAGES CXX C ASM" + ''; - # Tensile deps - not optional, building without tensile isn't actually supported - msgpack # FIXME: not included in cmake! - libxml2 - python3Packages.msgpack - python3Packages.joblib - zlib - zstd - ] - ++ lib.optionals buildTests [ - gtest - ] - ++ lib.optionals (buildTests || buildBenchmarks) [ - lapack-reference - ]; + doCheck = false; + doInstallCheck = false; - cmakeFlags = [ - "-Wno-dev" - "-DCMAKE_BUILD_TYPE=Release" - "-DCMAKE_VERBOSE_MAKEFILE=ON" - "-DVIRTUALENV_PYTHON_EXENAME=${lib.getExe py}" - "-DTENSILE_USE_HIP=ON" - "-DTENSILE_BUILD_CLIENT=OFF" - "-DTENSILE_USE_FLOAT16_BUILTIN=ON" - "-DCMAKE_CXX_COMPILER=${compiler}" - # Manually define CMAKE_INSTALL_ - # See: https://github.com/NixOS/nixpkgs/pull/197838 - "-DCMAKE_INSTALL_BINDIR=bin" - "-DCMAKE_INSTALL_LIBDIR=lib" - "-DCMAKE_INSTALL_INCLUDEDIR=include" - "-DHIPBLASLT_ENABLE_MARKER=Off" - # FIXME what are the implications of hardcoding this? - "-DTensile_CODE_OBJECT_VERSION=V5" - "-DTensile_COMPILER=${compiler}" - "-DAMDGPU_TARGETS=${gpuTargets'}" - "-DGPU_TARGETS=${gpuTargets'}" - "-DTensile_LIBRARY_FORMAT=msgpack" - ] - ++ lib.optionals (!supportsTargetArches) [ - "-DBUILD_WITH_TENSILE=OFF" - ] - ++ lib.optionals buildTests [ - "-DBUILD_CLIENTS_TESTS=ON" - ] - ++ lib.optionals buildBenchmarks [ - "-DBUILD_CLIENTS_BENCHMARKS=ON" - ] - ++ lib.optionals buildSamples [ - "-DBUILD_CLIENTS_SAMPLES=ON" - ]; + nativeBuildInputs = [ + cmake + rocm-cmake + py + clr + gfortran + pkg-config + ninja + rocm-smi + ]; - postInstall = - lib.optionalString buildTests '' + buildInputs = [ + clr + rocm-cmake + hipblas-common + amd-blis + rocm-smi + openmp + libffi + ncurses + lapack-reference + + # Tensile deps - not optional, building without tensile isn't actually supported + msgpack-cxx + libxml2 + python3Packages.msgpack + python3Packages.joblib + zlib + zstd + ] + ++ lib.optionals buildTests [ + gtest + ]; + + cmakeFlags = [ + (lib.cmakeFeature "Boost_INCLUDE_DIR" "${lib.getDev boost}/include") # msgpack FindBoost fails to find boost + (lib.cmakeFeature "GPU_TARGETS" gpuTargets') + (lib.cmakeBool "BUILD_TESTING" buildTests) + (lib.cmakeBool "HIPBLASLT_ENABLE_BLIS" true) + (lib.cmakeBool "HIPBLASLT_BUILD_TESTING" buildTests) + (lib.cmakeBool "HIPBLASLT_ENABLE_SAMPLES" buildSamples) + (lib.cmakeBool "HIPBLASLT_ENABLE_DEVICE" supportsTargetArches) + # FIXME: Enable for ROCm 7.x + (lib.cmakeBool "HIPBLASLT_ENABLE_ROCROLLER" false) + "-DCMAKE_C_COMPILER=amdclang" + "-DCMAKE_HIP_COMPILER=${compiler}" + "-DCMAKE_CXX_COMPILER=${compiler}" + "-DROCM_FOUND=ON" # hipblaslt tries to download rocm-cmake if this isn't set + "-DBLIS_ROOT=${amd-blis}" + "-DBLIS_LIB=${amd-blis}/lib/libblis-mt.so" + "-DBLIS_INCLUDE_DIR=${amd-blis}/include/blis/" + "-DBLA_PREFER_PKGCONFIG=ON" + "-DFETCHCONTENT_SOURCE_DIR_NANOBIND=${python3Packages.nanobind.src}" + # Manually define CMAKE_INSTALL_ + # See: https://github.com/NixOS/nixpkgs/pull/197838 + "-DCMAKE_INSTALL_BINDIR=bin" + "-DCMAKE_INSTALL_LIBDIR=lib" + "-DCMAKE_INSTALL_INCLUDEDIR=include" + "-DHIPBLASLT_ENABLE_MARKER=Off" + ]; + + postInstall = + # Compress msgpack .dat files to stay under hydra output size limit + # Relies on messagepack-compression-support.patch + '' + for file in $out/lib/hipblaslt/library/*.dat; do + zstd -19 --long -f "$file" -o "$file.tmp" && mv "$file.tmp" "$file" + done + '' + # Move binaries to appropriate outputs and delete leftover /bin + + '' + mkdir -p $benchmark/bin + mv $out/bin/hipblaslt-{api-overhead,sequence,bench*} $out/bin/*.yaml $out/bin/*.py $benchmark/bin + ${lib.optionalString buildTests '' mkdir -p $test/bin mv $out/bin/hipblas-test $test/bin - '' - + lib.optionalString buildBenchmarks '' - mkdir -p $benchmark/bin - mv $out/bin/hipblas-bench $benchmark/bin - '' - + lib.optionalString buildSamples '' + ''} + ${lib.optionalString buildSamples '' mkdir -p $sample/bin mv $out/bin/example-* $sample/bin - '' - + lib.optionalString (buildTests || buildBenchmarks || buildSamples) '' - rmdir $out/bin - ''; - # If this is false there are no kernels in the output lib - # and it's useless at runtime - # so if it's an optional dep it's best to not depend on it - # Some packages like torch need hipblaslt to compile - # and are fine ignoring it at runtime if it's not supported - # so we have to support building an empty hipblaslt - passthru.supportsTargetArches = supportsTargetArches; - passthru.updateScript = rocmUpdateScript { - name = finalAttrs.pname; - inherit (finalAttrs.src) owner repo; - }; - passthru.tensilelite = tensile'; - meta = with lib; { - description = "Library that provides general matrix-matrix operations with a flexible API"; - homepage = "https://github.com/ROCm/hipBLASlt"; - license = with licenses; [ mit ]; - teams = [ teams.rocm ]; - platforms = platforms.linux; - }; - } -) + ''} + rmdir $out/bin + ''; + + # If this is false there are no kernels in the output lib + # supporting the target device + # so if it's an optional dep it's best to not depend on it + # Some packages like torch need hipblaslt to compile + # and are fine ignoring it at runtime if it's not supported + # so we have to support building an empty hipblaslt + passthru.supportsTargetArches = supportsTargetArches; + passthru.updateScript = rocmUpdateScript { + name = finalAttrs.pname; + inherit (finalAttrs.src) owner repo; + }; + meta = with lib; { + description = "Library that provides general matrix-matrix operations with a flexible API"; + homepage = "https://github.com/ROCm/hipBLASlt"; + license = with licenses; [ mit ]; + teams = [ teams.rocm ]; + platforms = platforms.linux; + }; +}) diff --git a/pkgs/development/rocm-modules/6/hipblaslt/messagepack-compression-support.patch b/pkgs/development/rocm-modules/6/hipblaslt/messagepack-compression-support.patch new file mode 100644 index 000000000000..ace6b2b728ad --- /dev/null +++ b/pkgs/development/rocm-modules/6/hipblaslt/messagepack-compression-support.patch @@ -0,0 +1,56 @@ +diff --git a/Tensile/Source/lib/source/msgpack/MessagePack.cpp b/Tensile/Source/lib/source/msgpack/MessagePack.cpp +index de97929c..dbc397e0 100644 +--- a/tensilelite/src/msgpack/MessagePack.cpp ++++ b/tensilelite/src/msgpack/MessagePack.cpp +@@ -28,6 +28,8 @@ + + #include + ++#include ++ + #include + + namespace Tensile +@@ -86,6 +88,34 @@ namespace Tensile + return nullptr; + } + ++ // Check if the file is zstd compressed ++ char magic[4]; ++ in.read(magic, 4); ++ bool isCompressed = (in.gcount() == 4 && magic[0] == '\x28' && magic[1] == '\xB5' && magic[2] == '\x2F' && magic[3] == '\xFD'); ++ // Reset file pointer to the beginning ++ in.seekg(0, std::ios::beg); ++ ++ if (isCompressed) { ++ // Decompress zstd file ++ std::vector compressedData((std::istreambuf_iterator(in)), std::istreambuf_iterator()); ++ ++ size_t decompressedSize = ZSTD_getFrameContentSize(compressedData.data(), compressedData.size()); ++ if (decompressedSize == ZSTD_CONTENTSIZE_ERROR || decompressedSize == ZSTD_CONTENTSIZE_UNKNOWN) { ++ if(Debug::Instance().printDataInit()) ++ std::cout << "Error: Unable to determine decompressed size for " << filename << std::endl; ++ return false; ++ } ++ ++ std::vector decompressedData(decompressedSize); ++ size_t dSize = ZSTD_decompress(decompressedData.data(), decompressedSize, compressedData.data(), compressedData.size()); ++ if (ZSTD_isError(dSize)) { ++ if(Debug::Instance().printDataInit()) ++ std::cout << "Error: ZSTD decompression failed for " << filename << std::endl; ++ return false; ++ } ++ ++ msgpack::unpack(result, decompressedData.data(), dSize); ++ } else { + msgpack::unpacker unp; + bool finished_parsing; + constexpr size_t buffer_size = 1 << 19; +@@ -109,6 +139,7 @@ namespace Tensile + + return nullptr; + } ++ } + } + catch(std::runtime_error const& exc) + { diff --git a/pkgs/development/rocm-modules/6/hipblaslt/parallel-buildSourceCodeObjectFile.diff b/pkgs/development/rocm-modules/6/hipblaslt/parallel-buildSourceCodeObjectFile.diff new file mode 100644 index 000000000000..019e7a6123b1 --- /dev/null +++ b/pkgs/development/rocm-modules/6/hipblaslt/parallel-buildSourceCodeObjectFile.diff @@ -0,0 +1,12 @@ +diff --git a/tensilelite/Tensile/Toolchain/Source.py b/tensilelite/Tensile/Toolchain/Source.py +index c9862e6c..dfa5ba40 100644 +--- a/tensilelite/Tensile/Toolchain/Source.py ++++ b/tensilelite/Tensile/Toolchain/Source.py +@@ -102,6 +102,7 @@ def buildSourceCodeObjectFiles( + coPaths= [] + + objPath = str(tmpObjDir / objFilename) ++ compiler.default_args += ["-parallel-jobs=8"] + compiler(str(includeDir), cmdlineArchs, str(kernelPath), objPath) + + for target in bundler.targets(objPath): diff --git a/pkgs/development/rocm-modules/6/hipcub/default.nix b/pkgs/development/rocm-modules/6/hipcub/default.nix index c559760738aa..137b7e76f9c4 100644 --- a/pkgs/development/rocm-modules/6/hipcub/default.nix +++ b/pkgs/development/rocm-modules/6/hipcub/default.nix @@ -17,7 +17,7 @@ # CUB can also be used as a backend instead of rocPRIM. stdenv.mkDerivation (finalAttrs: { pname = "hipcub"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -33,7 +33,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipCUB"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-uECOQWG9C64tg5YZdm9/3+fZXaZVGslu8vElK3m23GY="; + hash = "sha256-yRJxhYFZYiKNR2xrn5fif/+vjHKKcKdn0JKPi972g+0="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hipfft/default.nix b/pkgs/development/rocm-modules/6/hipfft/default.nix index 21507e774c6d..1002b0b3d16f 100644 --- a/pkgs/development/rocm-modules/6/hipfft/default.nix +++ b/pkgs/development/rocm-modules/6/hipfft/default.nix @@ -22,7 +22,7 @@ # Can also use cuFFT stdenv.mkDerivation (finalAttrs: { pname = "hipfft"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -41,7 +41,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipFFT"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-Jq/YHEtOo7a0/Ki7gxZATKmSqPU6cyLf5gx3A4MAZNw="; + hash = "sha256-4W93OOKTqNteoQ4GKycr06cjvGy5NF7RR08F+rfn+0o="; fetchSubmodules = true; }; diff --git a/pkgs/development/rocm-modules/6/hipfort/default.nix b/pkgs/development/rocm-modules/6/hipfort/default.nix index d3916bf6d20c..e4c60cade6b7 100644 --- a/pkgs/development/rocm-modules/6/hipfort/default.nix +++ b/pkgs/development/rocm-modules/6/hipfort/default.nix @@ -10,13 +10,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "hipfort"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "hipfort"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-V5XDNM0bYHKnpkcnaDyxIS1zwsgaByJj+znFxJ6VxR0="; + hash = "sha256-Nks1+0X8bLtZ9HqZXJOtrAWQlJquMH+feuu1stf/9Vo="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hipify/default.nix b/pkgs/development/rocm-modules/6/hipify/default.nix index 8747dba16b96..1fe6e969a776 100644 --- a/pkgs/development/rocm-modules/6/hipify/default.nix +++ b/pkgs/development/rocm-modules/6/hipify/default.nix @@ -14,13 +14,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "hipify"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "HIPIFY"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-o/1LNsNtAyQcSug1gf7ujGNRRbvC33kwldrJKZi2LA0="; + hash = "sha256-uj25WmGCpwouS1yzW9Oil5Vyrbyj5yRITvWF9WaGozM="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hiprand/default.nix b/pkgs/development/rocm-modules/6/hiprand/default.nix index d39e01ac2a78..d22dd2ee6708 100644 --- a/pkgs/development/rocm-modules/6/hiprand/default.nix +++ b/pkgs/development/rocm-modules/6/hiprand/default.nix @@ -14,7 +14,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "hiprand"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -27,7 +27,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipRAND"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-TVc+qFwRiS5tAo1OKI1Wu5hadlwPZmSVZ9SvVvH1w7Y="; + hash = "sha256-f/AWDV7vkjt8BnjhLwT2m1Y6dlZy054+z6J0UW9Glg8="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hipsolver/default.nix b/pkgs/development/rocm-modules/6/hipsolver/default.nix index d87c58eb3dbf..17b888d2a269 100644 --- a/pkgs/development/rocm-modules/6/hipsolver/default.nix +++ b/pkgs/development/rocm-modules/6/hipsolver/default.nix @@ -21,7 +21,7 @@ # Can also use cuSOLVER stdenv.mkDerivation (finalAttrs: { pname = "hipsolver"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -40,7 +40,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipSOLVER"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-ZQUKU3L4DgZ5zM7pCYEix0ulRkl78x/5wJnyCndTAwk="; + hash = "sha256-YP88fSM27Do0/tZ676Pvv2tr+lnlQa4vr3UnVNaVpLA="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hipsparse/default.nix b/pkgs/development/rocm-modules/6/hipsparse/default.nix index 0a1795d3ab16..f10280335ba0 100644 --- a/pkgs/development/rocm-modules/6/hipsparse/default.nix +++ b/pkgs/development/rocm-modules/6/hipsparse/default.nix @@ -20,7 +20,7 @@ # This can also use cuSPARSE as a backend instead of rocSPARSE stdenv.mkDerivation (finalAttrs: { pname = "hipsparse"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -36,7 +36,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "hipSPARSE"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-3a7fKpYyiqG3aGOg7YrTHmKoH4rgTVLD16DvrZ3YY1g="; + hash = "sha256-fbh9fKlzxuIBTeCV/bEQbUS3lO6O3KoGF7/tTqRaCpE="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/default.nix b/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/default.nix deleted file mode 100644 index 68d7a337f89d..000000000000 --- a/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/default.nix +++ /dev/null @@ -1,53 +0,0 @@ -{ - lib, - stdenv, - fetchurl, - callPackage, - dpkg, - rocm-core, -}: - -stdenv.mkDerivation (finalAttrs: { - pname = "hsa-amd-aqlprofile-bin"; - version = "6.3.3"; - - src = - let - inherit (finalAttrs) version; - patch = rocm-core.ROCM_LIBPATCH_VERSION; - majorMinor = lib.versions.majorMinor version; - poolVersion = if majorMinor + ".0" == version then majorMinor else version; - incremental = "74"; - osRelease = "22.04"; - in - fetchurl { - url = "https://repo.radeon.com/rocm/apt/${poolVersion}/pool/main/h/hsa-amd-aqlprofile/hsa-amd-aqlprofile_1.0.0.${patch}-${incremental}~${osRelease}_amd64.deb"; - hash = "sha256-Lo6gU9ywkujtsKvnOAwL3L8qQNPwjjm0Pm4OyzoUYao="; - }; - - nativeBuildInputs = [ dpkg ]; - dontPatch = true; - dontConfigure = true; - dontBuild = true; - - installPhase = '' - runHook preInstall - - mkdir -p $out - cp -a opt/rocm-${finalAttrs.version}*/* $out - chmod +x $out/lib/libhsa-amd-aqlprofile64.so.1.* - chmod +x $out/lib/hsa-amd-aqlprofile/librocprofv2_att.so - - runHook postInstall - ''; - - passthru.updateScript = (callPackage ./update.nix { }) { inherit (finalAttrs) version; }; - - meta = with lib; { - description = "AQLPROFILE library for AMD HSA runtime API extension support"; - homepage = "https://rocm.docs.amd.com/en/latest/"; - license = with licenses; [ unfree ]; - teams = [ teams.rocm ]; - platforms = platforms.linux; - }; -}) diff --git a/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/update.nix b/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/update.nix deleted file mode 100644 index ed834c84a678..000000000000 --- a/pkgs/development/rocm-modules/6/hsa-amd-aqlprofile-bin/update.nix +++ /dev/null @@ -1,61 +0,0 @@ -{ - lib, - writeScript, -}: - -{ version }: - -let - prefix = "hsa-amd-aqlprofile"; - extVersion = lib.strings.concatStrings ( - lib.strings.intersperse "0" (lib.versions.splitVersion version) - ); - major = lib.versions.major version; - minor = lib.versions.minor version; - patch = lib.versions.patch version; - - updateScript = writeScript "update.sh" '' - #!/usr/bin/env nix-shell - #!nix-shell -i bash -p curl common-updater-scripts - apt="https://repo.radeon.com/rocm/apt" - pool="pool/main/h/${prefix}/" - url="$apt/latest/$pool" - res="$(curl -sL "$url")" - deb="${prefix}$(echo "$res" | grep -o -P "(?<=href=\"${prefix}).*(?=\">)" | tail -1)" - patch="${patch}" - - # Try up to 10 patch versions - for i in {1..10}; do - ((patch++)) - extVersion="$(echo "$deb" | grep -o -P "(?<=\.....).*(?=\..*-)")" - - if (( ''${#extVersion} == 6 )) && (( $extVersion <= ${extVersion} )); then - url="https://repo.radeon.com/rocm/apt/${major}.${minor}.$patch/pool/main/h/${prefix}/" - res="$(curl -sL "$url")" - deb="${prefix}$(echo "$res" | grep -o -P "(?<=href=\"${prefix}).*(?=\">)" | tail -1)" - else - break - fi - done - - extVersion="$(echo $deb | grep -o -P "(?<=\.....).*(?=\..*-)")" - version="$(echo $extVersion | sed "s/0/./1" | sed "s/0/./1")" - IFS='.' read -a version_arr <<< "$version" - - if (( ''${version_arr[0]} > 6 )); then - echo "'rocmPackages_6.${prefix}-bin' is already at it's maximum allowed version.''\nAny further upgrades should go into 'rocmPackages_X.${prefix}-bin'." 1>&2 - exit 1 - fi - - if (( ''${#extVersion} == 6 )); then - repoVersion="$version" - - if (( ''${version:4:1} == 0 )); then - repoVersion=''${version:0:3} - fi - - update-source-version rocmPackages_6.${prefix}-bin "$version" "" "$apt/$repoVersion/$pool$deb" --ignore-same-hash - fi - ''; -in -[ updateScript ] diff --git a/pkgs/development/rocm-modules/6/llvm/default.nix b/pkgs/development/rocm-modules/6/llvm/default.nix index c8fb82fa0235..8dc0456bf366 100644 --- a/pkgs/development/rocm-modules/6/llvm/default.nix +++ b/pkgs/development/rocm-modules/6/llvm/default.nix @@ -1,16 +1,17 @@ { lib, stdenv, - llvmPackages_18, + # default LLVM version is used as stdenv to build our toolchain + llvmPackages, + # LLVM version closest to ROCm fork to override + llvmPackages_20, overrideCC, rocm-device-libs, - rocm-runtime, fetchFromGitHub, runCommand, symlinkJoin, rdfind, wrapBintoolsWith, - emptyDirectory, zstd, zlib, gcc-unwrapped, @@ -21,33 +22,49 @@ removeReferencesTo, fetchpatch, # Build compilers and stdenv suitable for profiling - # compressed line tables (-g1 -gz) and - # frame pointers for sampling profilers (-fno-omit-frame-pointer -momit-leaf-frame-pointer) - # TODO: Should also apply to downstream packages which use rocmClangStdenv + # leaving compressed line tables (-g1 -gz) unstripped + # TODO: Should also apply to downstream packages which use rocmClangStdenv? profilableStdenv ? false, -}: + # Whether to use LTO when building the ROCm toolchain + # Slows down this toolchain's build, for typical ROCm usecase + # time saved building composable_kernel and other heavy packages + # will outweight that. ~3-4% speedup multiplied by thousands + # of corehours. + withLto ? true, + # whether rocm stdenv uses libcxx (clang c++ stdlib) instead of gcc stdlibc++ + withLibcxx ? false, +}@args: let - llvmPackagesNoBintools = llvmPackages_18.override { + version = "6.4.3"; + # major version of this should be the clang version ROCm forked from + rocmLlvmVersion = "19.0.0-rocm"; + # llvmPackages_base version should be close to rocmLlvmVersion, + # may be one off because AMD backports a lot and the +1 patches + # may be easier to get to build + llvmPackages_base = llvmPackages_20; + llvmPackagesNoBintools = llvmPackages.override { bootBintools = null; bootBintoolsNoLibc = null; }; - useLibcxx = false; # whether rocm stdenv uses libcxx (clang c++ stdlib) instead of gcc stdlibc++ - llvmStdenv = overrideCC llvmPackagesNoBintools.libcxxStdenv llvmPackagesNoBintools.clangUseLLVM; - llvmLibstdcxxStdenv = overrideCC llvmPackagesNoBintools.stdenv ( - llvmPackagesNoBintools.libstdcxxClang.override { - inherit (llvmPackages_18) bintools; - } - ); - stdenvToBuildRocmLlvm = if useLibcxx then llvmStdenv else llvmLibstdcxxStdenv; + stdenvToBuildRocmLlvm = + if withLibcxx then + overrideCC llvmPackagesNoBintools.libcxxStdenv llvmPackagesNoBintools.clangUseLLVM + else + # oddly fuse-ld=lld fails without this override + overrideCC llvmPackagesNoBintools.stdenv ( + llvmPackagesNoBintools.libstdcxxClang.override { + inherit (llvmPackages) bintools; + } + ); + gcc-include = runCommand "gcc-include" { } '' mkdir -p $out ln -s ${gcc-unwrapped}/include/ $out/ ln -s ${gcc-unwrapped}/lib/ $out/ ''; - # A prefix for use as the GCC prefix when building rocmcxx disallowedRefsForToolchain = [ stdenv.cc stdenv.cc.cc @@ -57,40 +74,47 @@ let stdenvToBuildRocmLlvm.cc stdenvToBuildRocmLlvm.cc.cc ]; - gcc-prefix = - let - gccPrefixPaths = [ - gcc-unwrapped - gcc-unwrapped.lib - glibc.dev - ]; - in - symlinkJoin { - name = "gcc-prefix"; - paths = gccPrefixPaths ++ [ - glibc - ]; - disallowedRequisites = gccPrefixPaths; - postBuild = '' - rm -rf $out/{bin,libexec,nix-support,lib64,share,etc} - rm $out/lib/gcc/x86_64-unknown-linux-gnu/*/plugin/include/auto-host.h + # A prefix for use as the GCC prefix when building rocmcxx + gcc-prefix-headers = symlinkJoin { + name = "gcc-prefix-headers"; + paths = [ + glibc.dev + gcc-unwrapped.out + ]; + disallowedRequisites = [ + glibc.dev + gcc-unwrapped.out + ]; + postBuild = '' + rm -rf $out/{bin,libexec,nix-support,lib64,share,etc} + rm $out/lib/gcc/x86_64-unknown-linux-gnu/*/plugin/include/auto-host.h - mkdir /build/tmpout - mv $out/* /build/tmpout - cp -Lr --no-preserve=mode /build/tmpout/* $out/ - set -x - versionedIncludePath="$(echo $out/include/c++/*/)" - mv $versionedIncludePath/* $out/include/c++/ - rm -rf $versionedIncludePath/ - - find $out/lib -type f -exec ${removeReferencesTo}/bin/remove-references-to -t ${gcc-unwrapped.lib} {} + - - ln -s $out $out/x86_64-unknown-linux-gnu - ''; - }; - version = "6.3.1"; - # major version of this should be the clang version ROCm forked from - rocmLlvmVersion = "18.0.0-${llvmSrc.rev}"; + mkdir /build/tmpout + mv $out/* /build/tmpout + cp -Lr --no-preserve=mode /build/tmpout/* $out/ + set -x + versionedIncludePath="$(echo $out/include/c++/*/)" + mv $versionedIncludePath/* $out/include/c++/ + rm -rf $versionedIncludePath/ + ''; + }; + gcc-prefix = symlinkJoin { + name = "gcc-prefix"; + paths = [ + gcc-prefix-headers + glibc + gcc-unwrapped.lib + ]; + disallowedRequisites = [ + glibc.dev + gcc-unwrapped.out + ]; + postBuild = '' + rm -rf $out/{bin,libexec,nix-support,lib64,share,etc} + rm $out/lib/ld-linux-x86-64.so.2 + ln -s $out $out/x86_64-unknown-linux-gnu + ''; + }; usefulOutputs = drv: builtins.filter (x: x != null) [ @@ -100,27 +124,24 @@ let ]; listUsefulOutputs = builtins.concatMap usefulOutputs; llvmSrc = fetchFromGitHub { - # Performance improvements cherry-picked on top of rocm-6.3.x - # most importantly, amdgpu-early-alwaysinline memory usage fix - owner = "LunNova"; - repo = "llvm-project-rocm"; - rev = "4182046534deb851753f0d962146e5176f648893"; - hash = "sha256-sPmYi1WiiAqnRnHVNba2nPUxGflBC01FWCTNLPlYF9c="; + owner = "ROCm"; + repo = "llvm-project"; + rev = "rocm-${version}"; + hash = "sha256-12ftH5fMPAsbcEBmhADwW1YY/Yxo/MAK1FafKczITMg="; }; - llvmSrcFixed = llvmSrc; llvmMajorVersion = lib.versions.major rocmLlvmVersion; # An llvmPackages (pkgs/development/compilers/llvm/) built from ROCm LLVM's source tree # optionally using LLVM libcxx - llvmPackagesRocm = llvmPackages_18.override (_old: { - stdenv = stdenvToBuildRocmLlvm; # old.stdenv #llvmPackagesNoBintools.libcxxStdenv; + llvmPackagesRocm = llvmPackages_base.override (_old: { + stdenv = stdenvToBuildRocmLlvm; # not setting gitRelease = because that causes patch selection logic to use git patches - # ROCm LLVM is closer to 18 official + # ROCm LLVM is closer to 20 official # gitRelease = {}; officialRelease = null; officialRelease = { }; # Set but empty because we're overriding everything from it. version = rocmLlvmVersion; - src = llvmSrcFixed; - monorepoSrc = llvmSrcFixed; + src = llvmSrc; + monorepoSrc = llvmSrc; doCheck = false; }); sysrootCompiler = @@ -158,82 +179,112 @@ let sed -i 's|set(CLANG_INCLUDE_DIRS.*$|set(CLANG_INCLUDE_DIRS "${cc.dev}/include")|g' $out/lib/cmake/clang/ClangConfig.cmake ${lib.getExe rdfind} -makesymlinks true $out/ # create links *within* the sysroot to save space ''; - findClangNostdlibincPatch = + # Removes patches which either aren't desired, or don't apply against ROCm LLVM + removeInapplicablePatches = x: ( (lib.strings.hasSuffix "add-nostdlibinc-flag.patch" (builtins.baseNameOf x)) || (lib.strings.hasSuffix "clang-at-least-16-LLVMgold-path.patch" (builtins.baseNameOf x)) ); + tablegenUsage = x: !(lib.strings.hasInfix "llvm-tblgen" x); llvmTargetsFlag = "-DLLVM_TARGETS_TO_BUILD=AMDGPU;${ { "x86_64" = "X86"; "aarch64" = "AArch64"; } - .${llvmStdenv.targetPlatform.parsed.cpu.name} - or (throw "Unsupported CPU architecture: ${llvmStdenv.targetPlatform.parsed.cpu.name}") + .${stdenv.targetPlatform.parsed.cpu.name} + or (throw "Unsupported CPU architecture: ${stdenv.targetPlatform.parsed.cpu.name}") }"; - # -ffat-lto-objects = emit LTO object files that are compatible with non-LTO-supporting builds too - # FatLTO objects are a special type of fat object file that contain LTO compatible IR in addition to generated object code, - # instead of containing object code for multiple target architectures. This allows users to defer the choice of whether to - # use LTO or not to link-time, and has been a feature available in other compilers, like GCC, for some time. + llvmMeta = { + # TODO(@LunNova): it would be nice to support aarch64 for rocmPackages + platforms = [ "x86_64-linux" ]; + }; + # TODO(@LunNova): Some of this might be worth supporting in llvmPackages, dropping from here + commonCmakeFlags = [ + llvmTargetsFlag + # Compression support is required for compressed offload kernels + # Set FORCE_ON so that failure to find the compression libs will be a build error + (lib.cmakeFeature "LLVM_ENABLE_ZSTD" "FORCE_ON") + (lib.cmakeFeature "LLVM_ENABLE_ZLIB" "FORCE_ON") + # required for threaded ThinLTO to work + (lib.cmakeBool "LLVM_ENABLE_THREADS" true) + # LLVM tries to call git to embed VCS info if FORCE_VC_ aren't set + (lib.cmakeFeature "LLVM_FORCE_VC_REVISION" "rocm-${version}") + (lib.cmakeFeature "LLVM_FORCE_VC_REPOSITORY" "https://github.com/ROCm/llvm-project") + (lib.cmakeFeature "LLVM_VERSION_SUFFIX" "") + (lib.cmakeBool "LLVM_ENABLE_LIBCXX" withLibcxx) + (lib.cmakeFeature "CLANG_DEFAULT_CXX_STDLIB" (if withLibcxx then "libc++" else "libstdc++")) + (lib.cmakeFeature "CLANG_VENDOR" "nixpkgs-AMD") + (lib.cmakeFeature "CLANG_REPOSITORY_STRING" "https://github.com/ROCm/llvm-project/tree/rocm-${version}") + ] + ++ lib.optionals withLibcxx [ + (lib.cmakeFeature "CLANG_DEFAULT_RTLIB" "compiler-rt") + ] + ++ lib.optionals withLto [ + (lib.cmakeBool "CMAKE_INTERPROCEDURAL_OPTIMIZATION" true) + (lib.cmakeBool "LLVM_ENABLE_FATLTO" false) + ] + ++ lib.optionals (withLto && stdenvToBuildRocmLlvm.cc.isClang) [ + (lib.cmakeFeature "LLVM_ENABLE_LTO" "FULL") + (lib.cmakeFeature "LLVM_USE_LINKER" "lld") + ]; - tablegenUsage = x: !(lib.strings.hasInfix "llvm-tblgen" x); - addGccLtoCmakeFlags = !llvmPackagesRocm.stdenv.cc.isClang; - llvmExtraCflags = - "-O3 -DNDEBUG -march=skylake -mtune=znver3" - + (lib.optionalString addGccLtoCmakeFlags " -D_GLIBCXX_USE_CXX11_ABI=0 -flto -ffat-lto-objects -flto-compression-level=19 -Wl,-flto") - + (lib.optionalString llvmPackagesRocm.stdenv.cc.isClang " -flto=thin -ffat-lto-objects") - + (lib.optionalString profilableStdenv " -fno-omit-frame-pointer -momit-leaf-frame-pointer -gz -g1"); + llvmExtraCflags = lib.concatStringsSep " " ( + lib.optionals (stdenv.hostPlatform.system == "x86_64-linux") [ + # Unprincipled decision to build x86_64 ROCm clang for at least skylake and tune for zen3+ + # In practice building the ROCm package set with anything earlier than zen3 is annoying + # and earlier than skylake is implausible due to too few cores and too little RAM + # Speeds up composable_kernel builds by ~4% + # If this causes trouble in practice we can drop this. Set since 2025-03-24. + "-march=skylake" + "-mtune=znver3" + ] + ++ lib.optionals profilableStdenv [ + # compressed line only debug info for profiling + "-gz" + "-g1" + ] + ); in rec { inherit (llvmPackagesRocm) libunwind; inherit (llvmPackagesRocm) libcxx; + inherit args; # Pass through original attrs for debugging where non-overridden llvm/clang is getting used # llvm-orig = llvmPackagesRocm.llvm; # nix why-depends --derivation .#rocmPackages.clr .#rocmPackages.llvm.llvm-orig # clang-orig = llvmPackagesRocm.clang; # nix why-depends --derivation .#rocmPackages.clr .#rocmPackages.llvm.clang-orig - llvm = (llvmPackagesRocm.llvm.override { ninja = emptyDirectory; }).overrideAttrs (old: { + llvm = llvmPackagesRocm.llvm.overrideAttrs (old: { + patches = old.patches ++ [ + (fetchpatch { + # fix compile error in tools/gold/gold-plugin.cpp + name = "gold-plugin-fix.patch"; + url = "https://github.com/llvm/llvm-project/commit/b0baa1d8bd68a2ce2f7c5f2b62333e410e9122a1.patch"; + hash = "sha256-yly93PvGIXOnFeDGZ2W+W6SyhdWFM6iwA+qOeaptrh0="; + relative = "llvm"; + }) + (fetchpatch { + # fix tools/llvm-exegesis/X86/latency/ failing with glibc 2.4+ + name = "exegesis-latency-glibc-fix.patch"; + sha256 = "sha256-CjKxQlYwHXTM0mVnv8k/ssg5OXuKpJxRvBZGXjrFZAg="; + url = "https://github.com/llvm/llvm-project/commit/1e8df9e85a1ff213e5868bd822877695f27504ad.patch"; + relative = "llvm"; + }) + ./perf-increase-namestring-size.patch + # TODO: consider reapplying "Don't include aliases in RegisterClassInfo::IgnoreCSRForAllocOrder" + # it was reverted as it's a pessimization for non-GPU archs, but this compiler + # is used mostly for amdgpu + ]; dontStrip = profilableStdenv; + hardeningDisable = [ "all" ]; nativeBuildInputs = old.nativeBuildInputs ++ [ removeReferencesTo ]; buildInputs = old.buildInputs ++ [ zstd zlib ]; - env.NIX_BUILD_ID_STYLE = "fast"; - postPatch = '' - ${old.postPatch or ""} - patchShebangs lib/OffloadArch/make_generated_offload_arch_h.sh - ''; - LDFLAGS = "-Wl,--build-id=sha1,--icf=all,--compress-debug-sections=zlib"; - cmakeFlags = - (builtins.filter tablegenUsage old.cmakeFlags) - ++ [ - llvmTargetsFlag - "-DCMAKE_BUILD_TYPE=Release" - "-DLLVM_ENABLE_ZSTD=FORCE_ON" - "-DLLVM_ENABLE_ZLIB=FORCE_ON" - "-DLLVM_ENABLE_THREADS=ON" - "-DLLVM_ENABLE_LTO=Thin" - "-DLLVM_USE_LINKER=lld" - (lib.cmakeBool "LLVM_ENABLE_LIBCXX" useLibcxx) - "-DCLANG_DEFAULT_CXX_STDLIB=${if useLibcxx then "libc++" else "libstdc++"}" - ] - ++ lib.optionals addGccLtoCmakeFlags [ - "-DCMAKE_AR=${gcc-unwrapped}/bin/gcc-ar" - "-DCMAKE_RANLIB=${gcc-unwrapped}/bin/gcc-ranlib" - "-DCMAKE_NM=${gcc-unwrapped}/bin/gcc-nm" - ] - ++ lib.optionals useLibcxx [ - "-DLLVM_ENABLE_LTO=Thin" - "-DLLVM_USE_LINKER=lld" - "-DLLVM_ENABLE_LIBCXX=ON" - ]; - preConfigure = '' - ${old.preConfigure or ""} - cmakeFlagsArray+=( - '-DCMAKE_C_FLAGS_RELEASE=${llvmExtraCflags}' - '-DCMAKE_CXX_FLAGS_RELEASE=${llvmExtraCflags}' - ) - ''; + env = (old.env or { }) // { + NIX_CFLAGS_COMPILE = "${(old.env or { }).NIX_CFLAGS_COMPILE or ""} ${llvmExtraCflags}"; + }; + cmakeFlags = (builtins.filter tablegenUsage old.cmakeFlags) ++ commonCmakeFlags; # Ensure we don't leak refs to compiler that was used to bootstrap this LLVM disallowedReferences = (old.disallowedReferences or [ ]) ++ disallowedRefsForToolchain; postFixup = '' @@ -243,48 +294,26 @@ rec { find $lib -type f -exec remove-references-to -t ${stdenvToBuildRocmLlvm.cc} {} + find $lib -type f -exec remove-references-to -t ${stdenv.cc.bintools} {} + ''; + meta = old.meta // llvmMeta; }); lld = (llvmPackagesRocm.lld.override { libllvm = llvm; - ninja = emptyDirectory; }).overrideAttrs (old: { - patches = builtins.filter ( - x: !(lib.strings.hasSuffix "more-openbsd-program-headers.patch" (builtins.baseNameOf x)) - ) old.patches; dontStrip = profilableStdenv; + hardeningDisable = [ "all" ]; nativeBuildInputs = old.nativeBuildInputs ++ [ - llvmPackagesNoBintools.lld removeReferencesTo ]; buildInputs = old.buildInputs ++ [ zstd zlib ]; - env.NIX_BUILD_ID_STYLE = "fast"; - LDFLAGS = "-Wl,--build-id=sha1,--icf=all,--compress-debug-sections=zlib"; - cmakeFlags = - (builtins.filter tablegenUsage old.cmakeFlags) - ++ [ - llvmTargetsFlag - "-DCMAKE_BUILD_TYPE=Release" - "-DLLVM_ENABLE_ZSTD=FORCE_ON" - "-DLLVM_ENABLE_ZLIB=FORCE_ON" - "-DLLVM_ENABLE_THREADS=ON" - "-DLLVM_ENABLE_LTO=Thin" - "-DLLVM_USE_LINKER=lld" - (lib.cmakeBool "LLVM_ENABLE_LIBCXX" useLibcxx) - "-DCLANG_DEFAULT_CXX_STDLIB=${if useLibcxx then "libc++" else "libstdc++"}" - ] - ++ lib.optionals addGccLtoCmakeFlags [ - "-DCMAKE_AR=${gcc-unwrapped}/bin/gcc-ar" - "-DCMAKE_RANLIB=${gcc-unwrapped}/bin/gcc-ranlib" - "-DCMAKE_NM=${gcc-unwrapped}/bin/gcc-nm" - ] - ++ lib.optionals useLibcxx [ - "-DLLVM_ENABLE_LIBCXX=ON" - ]; + env = (old.env or { }) // { + NIX_CFLAGS_COMPILE = "${(old.env or { }).NIX_CFLAGS_COMPILE or ""} ${llvmExtraCflags}"; + }; + cmakeFlags = (builtins.filter tablegenUsage old.cmakeFlags) ++ commonCmakeFlags; # Ensure we don't leak refs to compiler that was used to bootstrap this LLVM disallowedReferences = (old.disallowedReferences or [ ]) ++ disallowedRefsForToolchain; postFixup = '' @@ -292,36 +321,40 @@ rec { find $lib -type f -exec remove-references-to -t ${stdenv.cc.cc} {} + find $lib -type f -exec remove-references-to -t ${stdenv.cc.bintools} {} + ''; - preConfigure = '' - ${old.preConfigure or ""} - cmakeFlagsArray+=( - '-DCMAKE_C_FLAGS_RELEASE=${llvmExtraCflags}' - '-DCMAKE_CXX_FLAGS_RELEASE=${llvmExtraCflags}' - ) - ''; + meta = old.meta // llvmMeta; }); clang-unwrapped = ( (llvmPackagesRocm.clang-unwrapped.override { libllvm = llvm; - ninja = emptyDirectory; }).overrideAttrs ( old: let - filteredPatches = builtins.filter (x: !(findClangNostdlibincPatch x)) old.patches; + filteredPatches = builtins.filter (x: !(removeInapplicablePatches x)) old.patches; in { - meta.platforms = [ - "x86_64-linux" - ]; - pname = "${old.pname}-rocm"; - patches = filteredPatches ++ [ + passthru = old.passthru // { + inherit gcc-prefix; + }; + patches = [ + (fetchpatch { + # [PATCH] [clang] Install scan-build-py into plain "lib" directory + # Backported so 19/clang/gnu-install-dirs patch applies to AMD's LLVM fork + hash = "sha256-bOqAjBwRKcERpQkiBpuojGs6ddd5Ht3zL5l3TuJK2w8="; + url = "https://github.com/llvm/llvm-project/commit/816fde1cbb700ebcc8b3df81fb93d675c04c12cd.patch"; + relative = "clang"; + }) + ] + ++ filteredPatches + ++ [ # Never add FHS include paths ./clang-bodge-ignore-systemwide-incls.diff # Prevents builds timing out if a single compiler invocation is very slow but # per-arch jobs are completing by ensuring there's terminal output ./clang-log-jobs.diff + ./opt-offload-compress-on-by-default.patch + ./perf-shorten-gcclib-include-paths.patch (fetchpatch { # [ClangOffloadBundler]: Add GetBundleIDsInFile to OffloadBundler sha256 = "sha256-G/mzUdFfrJ2bLJgo4+mBcR6Ox7xGhWu5X+XxT4kH2c8="; @@ -334,65 +367,41 @@ rec { libllvmLibdir = "${llvm.lib}/lib"; }) ]; + hardeningDisable = [ "all" ]; nativeBuildInputs = old.nativeBuildInputs ++ [ - llvmPackagesNoBintools.lld removeReferencesTo ]; buildInputs = old.buildInputs ++ [ zstd zlib ]; - dontStrip = profilableStdenv; - LDFLAGS = "-Wl,--build-id=sha1,--icf=all,--compress-debug-sections=zlib"; env = (old.env or { }) // { - NIX_BUILD_ID_STYLE = "fast"; + NIX_CFLAGS_COMPILE = "${(old.env or { }).NIX_CFLAGS_COMPILE or ""} ${llvmExtraCflags}"; }; + dontStrip = profilableStdenv; # Ensure we don't leak refs to compiler that was used to bootstrap this LLVM disallowedReferences = (old.disallowedReferences or [ ]) ++ disallowedRefsForToolchain; # Enable structured attrs for separateDebugInfo, because it is required with disallowedReferences set __structuredAttrs = true; - requiredSystemFeatures = (old.requiredSystemFeatures or [ ]) ++ [ "big-parallel" ]; # https://github.com/llvm/llvm-project/blob/6976deebafa8e7de993ce159aa6b82c0e7089313/clang/cmake/caches/DistributionExample-stage2.cmake#L9-L11 cmakeFlags = (builtins.filter tablegenUsage old.cmakeFlags) - ++ [ - llvmTargetsFlag - "-DCMAKE_BUILD_TYPE=Release" - "-DLLVM_ENABLE_ZSTD=FORCE_ON" - "-DLLVM_ENABLE_ZLIB=FORCE_ON" - "-DLLVM_ENABLE_THREADS=ON" - "-DLLVM_ENABLE_LTO=Thin" - "-DLLVM_USE_LINKER=lld" - (lib.cmakeBool "LLVM_ENABLE_LIBCXX" useLibcxx) - "-DCLANG_DEFAULT_CXX_STDLIB=${if useLibcxx then "libc++" else "libstdc++"}" - ] - ++ lib.optionals addGccLtoCmakeFlags [ - "-DCMAKE_AR=${gcc-unwrapped}/bin/gcc-ar" - "-DCMAKE_RANLIB=${gcc-unwrapped}/bin/gcc-ranlib" - "-DCMAKE_NM=${gcc-unwrapped}/bin/gcc-nm" - ] - ++ lib.optionals useLibcxx [ - "-DLLVM_ENABLE_LTO=Thin" - "-DLLVM_ENABLE_LIBCXX=ON" - "-DLLVM_USE_LINKER=lld" - "-DCLANG_DEFAULT_RTLIB=compiler-rt" - ] - ++ lib.optionals (!useLibcxx) [ + ++ commonCmakeFlags + ++ lib.optionals (!withLibcxx) [ # FIXME: Config file in rocmcxx instead of GCC_INSTALL_PREFIX? + # Expected to be fully removed eventually + "-DUSE_DEPRECATED_GCC_INSTALL_PREFIX=ON" "-DGCC_INSTALL_PREFIX=${gcc-prefix}" ]; postFixup = (old.postFixup or "") + '' + find $lib -type f -exec remove-references-to -t ${stdenvToBuildRocmLlvm} {} + find $lib -type f -exec remove-references-to -t ${stdenvToBuildRocmLlvm.cc} {} + + find $lib -type f -exec remove-references-to -t ${stdenvToBuildRocmLlvm.cc.cc} {} + find $lib -type f -exec remove-references-to -t ${stdenv.cc} {} + find $lib -type f -exec remove-references-to -t ${stdenv.cc.cc} {} + find $lib -type f -exec remove-references-to -t ${stdenv.cc.bintools} {} + ''; - preConfigure = (old.preConfigure or "") + '' - cmakeFlagsArray+=( - '-DCMAKE_C_FLAGS_RELEASE=${llvmExtraCflags}' - '-DCMAKE_CXX_FLAGS_RELEASE=${llvmExtraCflags}' - ) - ''; + meta = old.meta // llvmMeta; } ) ) @@ -409,11 +418,12 @@ rec { clang-unwrapped bintools compiler-rt + openmp ] - ++ (lib.optionals useLibcxx [ + ++ (lib.optionals withLibcxx [ libcxx ]) - ++ (lib.optionals (!useLibcxx) [ + ++ (lib.optionals (!withLibcxx) [ gcc-include glibc glibc.dev @@ -438,7 +448,14 @@ rec { hash = "sha256-Db+L1HFMWVj4CrofsGbn5lnMoCzEcU+7q12KKFb17/g="; relative = "compiler-rt"; }) + (fetchpatch { + # Fixes fortify hardening compile error related to openat usage + hash = "sha256-pgpN1q1vIQrPXHPxNSZ6zfgV2EflHO5Amzl+2BDjXbs="; + url = "https://github.com/llvm/llvm-project/commit/155b7a12820ec45095988b6aa6e057afaf2bc892.patch"; + relative = "compiler-rt"; + }) ]; + meta = old.meta // llvmMeta; }); compiler-rt = compiler-rt-libc; bintools = wrapBintoolsWith { @@ -450,6 +467,7 @@ rec { clang = rocmcxx; # Emulate a monolithic ROCm LLVM build to support building ROCm's in-tree LLVM projects + # TODO(@LunNova): destroy this rocm-merged-llvm = symlinkJoin { name = "rocm-llvm-merge"; paths = [ @@ -464,7 +482,7 @@ rec { compiler-rt.dev rocmcxx ] - ++ lib.optionals useLibcxx [ + ++ lib.optionals withLibcxx [ libcxx libcxx.out libcxx.dev @@ -489,37 +507,32 @@ rec { }; rocmClangStdenv = overrideCC ( - if useLibcxx then llvmPackagesRocm.libcxxStdenv else llvmPackagesRocm.stdenv + if withLibcxx then llvmPackagesRocm.libcxxStdenv else llvmPackagesRocm.stdenv ) clang; # Projects openmp = (llvmPackagesRocm.openmp.override { - stdenv = rocmClangStdenv; - llvm = rocm-merged-llvm; - targetLlvm = rocm-merged-llvm; - clang-unwrapped = clang; + llvm = llvm; + targetLlvm = llvm; + clang-unwrapped = clang-unwrapped; }).overrideAttrs (old: { disallowedReferences = (old.disallowedReferences or [ ]) ++ disallowedRefsForToolchain; - nativeBuildInputs = (old.nativeBuildInputs or [ ]) ++ [ removeReferencesTo ]; + nativeBuildInputs = (old.nativeBuildInputs or [ ]) ++ [ + removeReferencesTo + ]; cmakeFlags = old.cmakeFlags + ++ commonCmakeFlags ++ [ "-DDEVICELIBS_ROOT=${rocm-device-libs.src}" - # OMPD support is broken in ROCm 6.3. Haven't investigated why. + # OMPD support is broken in ROCm 6.3+ Haven't investigated why. "-DLIBOMP_OMPD_SUPPORT:BOOL=FALSE" "-DLIBOMP_OMPD_GDB_SUPPORT:BOOL=FALSE" - ] - ++ lib.optionals addGccLtoCmakeFlags [ - "-DCMAKE_AR=${gcc-unwrapped}/bin/gcc-ar" - "-DCMAKE_RANLIB=${gcc-unwrapped}/bin/gcc-ranlib" ]; - env.LLVM = "${rocm-merged-llvm}"; - env.LLVM_DIR = "${rocm-merged-llvm}"; buildInputs = old.buildInputs ++ [ - rocm-device-libs - rocm-runtime + clang-unwrapped zlib zstd libxml2 diff --git a/pkgs/development/rocm-modules/6/llvm/opt-offload-compress-on-by-default.patch b/pkgs/development/rocm-modules/6/llvm/opt-offload-compress-on-by-default.patch new file mode 100644 index 000000000000..5b43528c3271 --- /dev/null +++ b/pkgs/development/rocm-modules/6/llvm/opt-offload-compress-on-by-default.patch @@ -0,0 +1,22 @@ +From 762511ad9c9cac7769a795cc4df888a1559dca02 Mon Sep 17 00:00:00 2001 +From: Luna Nova +Date: Mon, 16 Dec 2024 14:56:57 -0800 +Subject: [PATCH] [AMDGPU] CommonArgs: Enable offload compression by default + +--- + clang/lib/Driver/ToolChains/CommonArgs.cpp | 2 +- + 1 file changed, 1 insertion(+), 1 deletion(-) + +diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp +index 33f12fb2d075d..ba35f65210c7a 100644 +--- a/lib/Driver/ToolChains/CommonArgs.cpp ++++ b/lib/Driver/ToolChains/CommonArgs.cpp +@@ -2914,7 +2914,7 @@ void tools::addHIPRuntimeLibArgs(const ToolChain &TC, Compilation &C, + void tools::addOffloadCompressArgs(const llvm::opt::ArgList &TCArgs, + llvm::opt::ArgStringList &CmdArgs) { + if (TCArgs.hasFlag(options::OPT_offload_compress, +- options::OPT_no_offload_compress, false)) ++ options::OPT_no_offload_compress, true)) + CmdArgs.push_back("-compress"); + if (TCArgs.hasArg(options::OPT_v)) + CmdArgs.push_back("-verbose"); diff --git a/pkgs/development/rocm-modules/6/llvm/perf-increase-namestring-size.patch b/pkgs/development/rocm-modules/6/llvm/perf-increase-namestring-size.patch new file mode 100644 index 000000000000..84046a588630 --- /dev/null +++ b/pkgs/development/rocm-modules/6/llvm/perf-increase-namestring-size.patch @@ -0,0 +1,28 @@ +From a57eb548d46a37e4aed33eefccc9e6374b579ff9 Mon Sep 17 00:00:00 2001 +From: Luna Nova +Date: Mon, 16 Dec 2024 13:58:54 -0800 +Subject: [PATCH] [AMDGPU] TargetMachine: increase NameString size to 256 + +128 is too small a default for compilers used to build +AMD's rocBLAS, composable_kernel and similar libs because +they use very long symbol names such as +device_gemm_xdl_universal_bf16_bf16_bf16_mk_nk_mn_comp_default_instance + +This fixes spending ~10% of compile time in malloc inside this function. +--- + llvm/lib/Target/TargetMachine.cpp | 2 +- + 1 file changed, 1 insertion(+), 1 deletion(-) + +diff --git a/llvm/lib/Target/TargetMachine.cpp b/llvm/lib/Target/TargetMachine.cpp +index abd0fdf2390c0..86ce468b75643 100644 +--- a/lib/Target/TargetMachine.cpp ++++ b/lib/Target/TargetMachine.cpp +@@ -270,7 +270,7 @@ MCSymbol *TargetMachine::getSymbol(const GlobalValue *GV) const { + if (MCSymbol *TargetSymbol = TLOF->getTargetSymbol(GV, *this)) + return TargetSymbol; + +- SmallString<128> NameStr; ++ SmallString<256> NameStr; + getNameWithPrefix(NameStr, GV, TLOF->getMangler()); + return TLOF->getContext().getOrCreateSymbol(NameStr); + } diff --git a/pkgs/development/rocm-modules/6/llvm/perf-shorten-gcclib-include-paths.patch b/pkgs/development/rocm-modules/6/llvm/perf-shorten-gcclib-include-paths.patch new file mode 100644 index 000000000000..a973851725d8 --- /dev/null +++ b/pkgs/development/rocm-modules/6/llvm/perf-shorten-gcclib-include-paths.patch @@ -0,0 +1,57 @@ +From ef6c5b353861be727c98f1319c81d0c6b609d644 Mon Sep 17 00:00:00 2001 +From: Luna Nova +Date: Tue, 17 Dec 2024 04:29:11 -0800 +Subject: [PATCH] HACK: Get canonical GCC include path so doesn't have + ../../../../ + +This allows more of the strings used in compilation to fit inside +fixed size stack allocated buffers instead of spilling into the heap + +--- + clang/lib/Driver/ToolChains/Gnu.cpp | 16 ++++++++++------ + 1 file changed, 10 insertions(+), 6 deletions(-) + +diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp +index af9fd46f0ce7b..a63a7a93f6a78 100644 +--- a/lib/Driver/ToolChains/Gnu.cpp ++++ b/lib/Driver/ToolChains/Gnu.cpp +@@ -3394,29 +3394,33 @@ bool Generic_GCC::addLibStdCXXIncludePaths(Twine IncludeDir, StringRef Triple, + if (!getVFS().exists(IncludeDir)) + return false; + ++ SmallString<260> CanonicalIncludeDir; ++ if (getVFS().getRealPath(IncludeDir, CanonicalIncludeDir)) ++ return false; ++ + // Debian native gcc uses g++-multiarch-incdir.diff which uses + // include/x86_64-linux-gnu/c++/10$IncludeSuffix instead of + // include/c++/10/x86_64-linux-gnu$IncludeSuffix. +- std::string Dir = IncludeDir.str(); + StringRef Include = +- llvm::sys::path::parent_path(llvm::sys::path::parent_path(Dir)); ++ llvm::sys::path::parent_path(llvm::sys::path::parent_path(CanonicalIncludeDir)); + std::string Path = +- (Include + "/" + Triple + Dir.substr(Include.size()) + IncludeSuffix) ++ (Include + "/" + Triple + CanonicalIncludeDir.substr(Include.size()) + IncludeSuffix) + .str(); + if (DetectDebian && !getVFS().exists(Path)) + return false; + + // GPLUSPLUS_INCLUDE_DIR +- addSystemInclude(DriverArgs, CC1Args, IncludeDir); ++ addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir); + // GPLUSPLUS_TOOL_INCLUDE_DIR. If Triple is not empty, add a target-dependent + // include directory. + if (DetectDebian) + addSystemInclude(DriverArgs, CC1Args, Path); + else if (!Triple.empty()) + addSystemInclude(DriverArgs, CC1Args, +- IncludeDir + "/" + Triple + IncludeSuffix); ++ CanonicalIncludeDir + "/" + Triple + IncludeSuffix); + // GPLUSPLUS_BACKWARD_INCLUDE_DIR +- addSystemInclude(DriverArgs, CC1Args, IncludeDir + "/backward"); ++ if (getVFS().exists(CanonicalIncludeDir + "/backward")) ++ addSystemInclude(DriverArgs, CC1Args, CanonicalIncludeDir + "/backward"); + return true; + } + diff --git a/pkgs/development/rocm-modules/6/migraphx/default.nix b/pkgs/development/rocm-modules/6/migraphx/default.nix index b1b88b1c170c..646bddf0fdf3 100644 --- a/pkgs/development/rocm-modules/6/migraphx/default.nix +++ b/pkgs/development/rocm-modules/6/migraphx/default.nix @@ -18,7 +18,8 @@ abseil-cpp, half, nlohmann_json, - msgpack, + boost, + msgpack-cxx, sqlite, oneDNN, blaze, @@ -28,6 +29,7 @@ docutils, ghostscript, python3Packages, + writableTmpDirAsHomeHook, buildDocs ? false, buildTests ? false, gpuTargets ? clr.gpuTargets, @@ -63,7 +65,7 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "migraphx"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -79,9 +81,13 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "AMDMIGraphX"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-h9cTbrMwHeRGVJS/uHQnCXplNcrBqxbhwz2AcAEso0M="; + hash = "sha256-8iOBoRBygTvn9eX5f9cG0kBHKgKSeflqHkV6Qwh/ruA="; }; + patches = [ + ./msgpack-6-compat.patch + ]; + nativeBuildInputs = [ pkg-config cmake @@ -97,6 +103,7 @@ stdenv.mkDerivation (finalAttrs: { ghostscript python3Packages.sphinx-rtd-theme python3Packages.breathe + writableTmpDirAsHomeHook ]; buildInputs = [ @@ -110,7 +117,8 @@ stdenv.mkDerivation (finalAttrs: { protobuf half nlohmann_json - msgpack + boost + msgpack-cxx sqlite oneDNN' blaze @@ -136,9 +144,6 @@ stdenv.mkDerivation (finalAttrs: { # migraphxs relies on miopen which relies on current composable_kernel # impossible to build with this ON; we can't link both of them even if we package both "-DMIGRAPHX_USE_COMPOSABLEKERNEL=OFF" - "-DOpenMP_C_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_CXX_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_omp_LIBRARY=${openmp}/lib" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" @@ -166,7 +171,6 @@ stdenv.mkDerivation (finalAttrs: { # Unfortunately, it seems like we have to call make on this manually preInstall = lib.optionalString buildDocs '' - export HOME=$(mktemp -d) make -j$NIX_BUILD_CORES doc cd ../doc/pdf make -j$NIX_BUILD_CORES diff --git a/pkgs/development/rocm-modules/6/migraphx/msgpack-6-compat.patch b/pkgs/development/rocm-modules/6/migraphx/msgpack-6-compat.patch new file mode 100644 index 000000000000..5c4890a02ce2 --- /dev/null +++ b/pkgs/development/rocm-modules/6/migraphx/msgpack-6-compat.patch @@ -0,0 +1,29 @@ +diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt +index 1dba2c8d..7795ba78 100644 +--- a/src/CMakeLists.txt ++++ b/src/CMakeLists.txt +@@ -358,13 +358,19 @@ target_link_libraries(migraphx PRIVATE nlohmann_json::nlohmann_json) + find_package(SQLite3 REQUIRED) + target_link_libraries(migraphx PRIVATE SQLite::SQLite3) + +-find_package(msgpackc-cxx QUIET) +-if(NOT msgpackc-cxx_FOUND) +- find_package(msgpack REQUIRED) ++# See: https://github.com/msgpack/msgpack-c/wiki/Q%26A#how-to-support-both-msgpack-c-c-version-5x-and-6x- ++# Prefer 6.x (msgpack-cxx) ++find_package(msgpack-cxx) ++if(msgpack-cxx_FOUND) ++ message(STATUS "Found msgpack-cxx (>=6.x)") ++else() ++ find_package(msgpackc-cxx REQUIRED NAMES msgpackc-cxx msgpack) ++ message(STATUS "Found msgpackc-cxx (<=5.x)") ++ add_library(msgpack-cxx ALIAS msgpackc-cxx) + endif() +-target_link_libraries(migraphx PRIVATE msgpackc-cxx) ++target_link_libraries(migraphx PRIVATE msgpack-cxx) + # Make this available to the tests +-target_link_libraries(migraphx INTERFACE $) ++target_link_libraries(migraphx INTERFACE $) + + add_library(migraphx_all_targets INTERFACE) + diff --git a/pkgs/development/rocm-modules/6/miopen/default.nix b/pkgs/development/rocm-modules/6/miopen/default.nix index 3a24c267ffaf..80b60f69a8f1 100644 --- a/pkgs/development/rocm-modules/6/miopen/default.nix +++ b/pkgs/development/rocm-modules/6/miopen/default.nix @@ -34,41 +34,50 @@ rocm-comgr, roctracer, python3Packages, - # FIXME: should be able to use all clr targets - gpuTargets ? [ + gpuTargets ? clr.localGpuTargets or clr.gpuTargets, + buildDocs ? false, # Needs internet because of rocm-docs-core + buildTests ? false, + withComposableKernel ? composable_kernel.anyMfmaTarget, +}: + +let + # FIXME: cmake files need patched to include this properly + cFlags = "-Wno-documentation-pedantic --offload-compress -I${hipblas-common}/include -I${hipblas}/include -I${roctracer}/include -I${nlohmann_json}/include -I${sqlite.dev}/include -I${rocrand}/include"; + version = "6.4.3"; + + # Targets outside this list will get + # error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD' + supportedTargets = lib.intersectLists [ "gfx900" "gfx906" "gfx908" "gfx90a" "gfx942" + "gfx950" "gfx1030" + "gfx1031" "gfx1100" "gfx1101" "gfx1102" - ], # clr.gpuTargets - buildDocs ? false, # Needs internet because of rocm-docs-core - buildTests ? false, - withComposableKernel ? composable_kernel.anyGfx9Target, -}: - -let - # FIXME: cmake files need patched to include this properly - cFlags = "-O3 -DNDEBUG -Wno-documentation-pedantic --offload-compress -I${hipblas-common}/include -I${hipblas}/include -I${roctracer}/include -I${nlohmann_json}/include -I${sqlite.dev}/include -I${rocrand}/include"; - version = "6.3.3"; + "gfx1150" + "gfx1151" + "gfx1200" + "gfx1201" + ] gpuTargets; src = fetchFromGitHub { owner = "ROCm"; repo = "MIOpen"; rev = "rocm-${version}"; - hash = "sha256-rX+BE6wBDMnLyc6eai3bDVvmfahomDO0s10n6HhWu7c="; + hash = "sha256-DEcVj2vOwIYYyNKEKFqZ0fb9o+/QRpwiSksxwnmgEMc="; fetchLFS = true; fetchSubmodules = true; # WORKAROUND: .lfsconfig is incorrectly set to exclude everything upstream leaveDotGit = true; + # FIXME: if someone can reduce the level of awful here that would be really nice postFetch = '' export HOME=$(mktemp -d) cd $out - set -x git remote add origin $url git fetch origin +refs/tags/rocm-${version}:refs/tags/rocm-${version} git clean -fdx @@ -77,8 +86,8 @@ let rm .lfsconfig git lfs install git lfs track "*.kdb.bz2" - GIT_TRACE=1 git lfs fetch --include="src/kernels/**" - GIT_TRACE=1 git lfs pull --include="src/kernels/**" + git lfs fetch --include="src/kernels/**" + git lfs pull --include="src/kernels/**" git lfs checkout rm -rf .git @@ -132,7 +141,6 @@ stdenv.mkDerivation (finalAttrs: { # Find zstd and add to target. Mainly for torch. patches = [ ./skip-preexisting-dbs.patch - ./fix-isnan.patch # https://github.com/ROCm/MIOpen/pull/3448 (fetchpatch { url = "https://github.com/ROCm/MIOpen/commit/e608b4325646afeabb5e52846997b926d2019d19.patch"; hash = "sha256-oxa3qlIC2bzbwGxrQOZXoY/S7CpLsMrnWRB7Og0tk0M="; @@ -205,9 +213,9 @@ stdenv.mkDerivation (finalAttrs: { ]; cmakeFlags = [ - "-DAMDGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" - "-DGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" - "-DGPU_ARCHS=${lib.concatStringsSep ";" gpuTargets}" + "-DAMDGPU_TARGETS=${lib.concatStringsSep ";" supportedTargets}" + "-DGPU_TARGETS=${lib.concatStringsSep ";" supportedTargets}" + "-DGPU_ARCHS=${lib.concatStringsSep ";" supportedTargets}" "-DMIOPEN_USE_SQLITE_PERFDB=ON" "-DCMAKE_VERBOSE_MAKEFILE=ON" "-DCMAKE_MODULE_PATH=${clr}/hip/cmake" @@ -271,7 +279,7 @@ stdenv.mkDerivation (finalAttrs: { ''; postInstall = '' - rm $out/bin/install_precompiled_kernels.sh + rm $out/libexec/miopen/install_precompiled_kernels.sh ln -sf ${gfx900} $out/share/miopen/db/gfx900.kdb ln -sf ${gfx906} $out/share/miopen/db/gfx906.kdb ln -sf ${gfx908} $out/share/miopen/db/gfx908.kdb diff --git a/pkgs/development/rocm-modules/6/miopen/fix-isnan.patch b/pkgs/development/rocm-modules/6/miopen/fix-isnan.patch deleted file mode 100644 index 71c140d57c8e..000000000000 --- a/pkgs/development/rocm-modules/6/miopen/fix-isnan.patch +++ /dev/null @@ -1,31 +0,0 @@ -From 17f67e0aa31cd2f1c1cb012d3858abf6956acc72 Mon Sep 17 00:00:00 2001 -From: "Sv. Lockal" -Date: Tue, 24 Dec 2024 14:43:10 +0000 -Subject: [PATCH] Fix missing isnan definition on libstdc++ >=14 systems - -Closes #3441 ---- - driver/reducecalculation_driver.hpp | 3 ++- - 1 file changed, 2 insertions(+), 1 deletion(-) - -diff --git a/driver/reducecalculation_driver.hpp b/driver/reducecalculation_driver.hpp -index 8226b3c953..2001969509 100644 ---- a/driver/reducecalculation_driver.hpp -+++ b/driver/reducecalculation_driver.hpp -@@ -33,6 +33,7 @@ - #include "random.hpp" - #include - #include -+#include - #include - #include - #include -@@ -77,7 +78,7 @@ int32_t mloReduceCalculationForwardRunHost(miopenTensorDescriptor_t inputDesc, - for(size_t i = 0; i < reduce_size; ++i) - { - Tcheck val = static_cast(input[input_idx]); -- if(nanPropagation && isnan(val)) -+ if(nanPropagation && std::isnan(val)) - { - val = 0.0f; - } diff --git a/pkgs/development/rocm-modules/6/mivisionx/default.nix b/pkgs/development/rocm-modules/6/mivisionx/default.nix index 0008b69f404a..324f8d884acf 100644 --- a/pkgs/development/rocm-modules/6/mivisionx/default.nix +++ b/pkgs/development/rocm-modules/6/mivisionx/default.nix @@ -5,7 +5,6 @@ rocmUpdateScript, cmake, rocm-cmake, - rocm-device-libs, clr, pkg-config, rpp, @@ -42,13 +41,13 @@ stdenv.mkDerivation (finalAttrs: { "cpu" ); - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "MIVisionX"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-SisCbUDCAiWQ1Ue7qrtoT6vO/1ztzqji+3cJD6MXUNw="; + hash = "sha256-07MivgCYmKLnhGDjOYsFBfwIxEoQLYNoRbOo3MPpVzE="; }; patches = [ @@ -58,19 +57,18 @@ stdenv.mkDerivation (finalAttrs: { nativeBuildInputs = [ cmake rocm-cmake - clr pkg-config ] + ++ lib.optionals (!useOpenCL && !useCPU) [ + clr + ] ++ lib.optionals buildDocs [ rocm-docs-core python3Packages.python ]; buildInputs = [ - miopen - migraphx rpp - rocblas openmp half protobuf @@ -83,28 +81,31 @@ stdenv.mkDerivation (finalAttrs: { rapidjson python3Packages.pybind11 python3Packages.numpy - python3Packages.torchWithRocm + ] + ++ lib.optionals (!useOpenCL && !useCPU) [ + miopen + rocblas + migraphx ]; cmakeFlags = [ - "-DROCM_PATH=${clr}" - "-DAMDRPP_PATH=${rpp}" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" "-DCMAKE_INSTALL_LIBDIR=lib" "-DCMAKE_INSTALL_INCLUDEDIR=include" "-DCMAKE_INSTALL_PREFIX_PYTHON=lib" - "-DOpenMP_C_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_CXX_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_omp_LIBRARY=${openmp}/lib" # "-DAMD_FP16_SUPPORT=ON" `error: typedef redefinition with different types ('__half' vs 'half_float::half')` ] ++ lib.optionals (gpuTargets != [ ]) [ "-DAMDGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" ] ++ lib.optionals (!useOpenCL && !useCPU) [ + "-DROCM_PATH=${clr}" + "-DAMDRPP_PATH=${rpp}" "-DBACKEND=HIP" + "-DCMAKE_C_COMPILER=hipcc" + "-DCMAKE_CXX_COMPILER=hipcc" ] ++ lib.optionals (useOpenCL && !useCPU) [ "-DBACKEND=OCL" @@ -114,17 +115,17 @@ stdenv.mkDerivation (finalAttrs: { ]; postPatch = '' - # We need to not use hipcc and define the CXXFLAGS manually due to `undefined hidden symbol: tensorflow:: ...` - export CXXFLAGS+=" --rocm-path=${clr} --rocm-device-lib-path=${rocm-device-libs}/amdgcn/bitcode" - # Properly find miopen, fix ffmpeg version detection - substituteInPlace amd_openvx_extensions/CMakeLists.txt \ - --replace-fail "miopen PATHS \''${ROCM_PATH} QUIET" "miopen PATHS ${miopen} QUIET" \ - --replace-fail "\''${ROCM_PATH}/include/miopen/config.h" "${miopen}/include/miopen/config.h" - # Properly find turbojpeg substituteInPlace cmake/FindTurboJpeg.cmake \ --replace-fail "\''${TURBO_JPEG_PATH}/include" "${libjpeg_turbo.dev}/include" \ --replace-fail "\''${TURBO_JPEG_PATH}/lib" "${libjpeg_turbo.out}/lib" + + ${lib.optionalString (!useOpenCL && !useCPU) '' + # Properly find miopen + substituteInPlace amd_openvx_extensions/CMakeLists.txt \ + --replace-fail "miopen PATHS \''${ROCM_PATH} QUIET" "miopen PATHS ${miopen} QUIET" \ + --replace-fail "\''${ROCM_PATH}/include/miopen/config.h" "${miopen}/include/miopen/config.h" + ''} ''; postBuild = lib.optionalString buildDocs '' diff --git a/pkgs/development/rocm-modules/6/rccl/default.nix b/pkgs/development/rocm-modules/6/rccl/default.nix index b4c24708eece..32c096c475e5 100644 --- a/pkgs/development/rocm-modules/6/rccl/default.nix +++ b/pkgs/development/rocm-modules/6/rccl/default.nix @@ -7,6 +7,7 @@ rocm-cmake, rocm-smi, rocm-core, + pkg-config, clr, mscclpp, perl, @@ -14,11 +15,14 @@ python3, gtest, chrpath, + roctracer, rocprofiler, rocprofiler-register, autoPatchelfHook, buildTests ? false, gpuTargets ? (clr.localGpuTargets or [ ]), + # for passthru.tests + rccl, }: let @@ -35,7 +39,7 @@ in # infiniband ib_peer_mem support isn't in the mainline kernel but is carried by some distros stdenv.mkDerivation (finalAttrs: { pname = "rccl${clr.gpuArchSuffix}"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -48,13 +52,14 @@ stdenv.mkDerivation (finalAttrs: { ./fix-mainline-support-and-ub.diff ./enable-mscclpp-on-all-gfx9.diff ./rccl-test-missing-iomanip.diff + ./fix_hw_reg_hw_id_gt_gfx10.patch ]; src = fetchFromGitHub { owner = "ROCm"; repo = "rccl"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-998tDiC0Qp9hhcXtFpiCWqwdKPVT2vNp0GU/rng03Bw="; + hash = "sha256-XpD+UjgdbAoGYK5UvvTX3f8rny4tiEDH/vYoCdZhtjo="; }; nativeBuildInputs = [ @@ -64,12 +69,14 @@ stdenv.mkDerivation (finalAttrs: { perl hipify python3 + pkg-config autoPatchelfHook # ASAN doesn't add rpath without this ]; buildInputs = [ rocm-smi gtest + roctracer rocprofiler rocprofiler-register mscclpp @@ -88,8 +95,9 @@ stdenv.mkDerivation (finalAttrs: { "-DROCM_VERSION=${rocm-core.ROCM_LIBPATCH_VERSION}" "-DBUILD_BFD=OFF" # Can't get it to detect bfd.h "-DENABLE_MSCCL_KERNEL=ON" - "-DENABLE_MSCCLPP=ON" - "-DMSCCLPP_ROOT=${mscclpp}" + # FIXME: this is still running a download because if(NOT mscclpp_nccl_FOUND) is commented out T_T + "-DENABLE_MSCCLPP=OFF" + #"-DMSCCLPP_ROOT=${mscclpp}" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" @@ -107,11 +115,15 @@ stdenv.mkDerivation (finalAttrs: { # -O2 and -fno-strict-aliasing due to UB issues in RCCL :c # Reported upstream - env.CFLAGS = "-I${clr}/include -O2 -fno-strict-aliasing ${san}-fno-omit-frame-pointer -momit-leaf-frame-pointer"; - env.CXXFLAGS = "-I${clr}/include -O2 -fno-strict-aliasing ${san}-fno-omit-frame-pointer -momit-leaf-frame-pointer"; + env.CFLAGS = "-I${clr}/include -I${roctracer}/include -O2 -fno-strict-aliasing ${san}-fno-omit-frame-pointer -momit-leaf-frame-pointer"; + env.CXXFLAGS = "-I${clr}/include -I${roctracer}/include -O2 -fno-strict-aliasing ${san}-fno-omit-frame-pointer -momit-leaf-frame-pointer"; env.LDFLAGS = "${san}"; postPatch = '' patchShebangs src tools + substituteInPlace CMakeLists.txt \ + --replace-fail '${"\${HOST_OS_ID}"}' '"ubuntu"' \ + --replace-fail 'target_include_directories(rccl PRIVATE ''${ROCM_SMI_INCLUDE_DIR})' \ + 'target_include_directories(rccl PRIVATE ''${ROCM_SMI_INCLUDE_DIRS})' ''; postInstall = @@ -130,6 +142,12 @@ stdenv.mkDerivation (finalAttrs: { inherit (finalAttrs.src) repo; }; + # This package with sanitizers + manual integration test binaries built + # must be ran manually + passthru.tests.rccl = rccl.override { + buildTests = true; + }; + meta = with lib; { description = "ROCm communication collectives library"; homepage = "https://github.com/ROCm/rccl"; diff --git a/pkgs/development/rocm-modules/6/rccl/fix-mainline-support-and-ub.diff b/pkgs/development/rocm-modules/6/rccl/fix-mainline-support-and-ub.diff index 15efac82c5ab..e65b23940c4c 100644 --- a/pkgs/development/rocm-modules/6/rccl/fix-mainline-support-and-ub.diff +++ b/pkgs/development/rocm-modules/6/rccl/fix-mainline-support-and-ub.diff @@ -92,13 +92,7 @@ index 6d77784..49762d3 100644 } if (moduleLoaded == -1) { -@@ -586,13 +586,14 @@ ncclResult_t ncclIbGdrSupport() { - // or created under a different path like `/sys/kernel/` or `/sys/` (depending on your ib_peer_mem module) - const char* memory_peers_paths[] = {"/sys/kernel/mm/memory_peers/amdkfd/version", - "/sys/kernel/memory_peers/amdkfd/version", -- "/sys/memory_peers/amdkfd/version"}; -+ "/sys/memory_peers/amdkfd/version", -+ NULL}; +@@ -586,9 +586,9 @@ ncclResult_t ncclIbGdrSupport() { int i = 0; while (memory_peers_paths[i]) { diff --git a/pkgs/development/rocm-modules/6/rccl/fix_hw_reg_hw_id_gt_gfx10.patch b/pkgs/development/rocm-modules/6/rccl/fix_hw_reg_hw_id_gt_gfx10.patch new file mode 100644 index 000000000000..c69000814a8a --- /dev/null +++ b/pkgs/development/rocm-modules/6/rccl/fix_hw_reg_hw_id_gt_gfx10.patch @@ -0,0 +1,33 @@ +fix: __trace_hwreg should use HW_REG_HW_ID1 for all gfx10/11/12 +diff --git a/src/device/common.h b/src/device/common.h +index c6c61021..742885f4 100644 +--- a/src/device/common.h ++++ b/src/device/common.h +@@ -26,8 +26,9 @@ + { __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST); } + #endif + +-#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1200__) || defined(__gfx1201__) +-#define __trace_hwreg() ++#if defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__) ++#define __trace_hwreg() \ ++ asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s" (collTrace->data_0)); + #else + #define __trace_hwreg() \ + asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (collTrace->data_0)); +diff --git a/tools/JitterBench/Common.hpp b/tools/JitterBench/Common.hpp +index bad12a1b..b59856a8 100644 +--- a/tools/JitterBench/Common.hpp ++++ b/tools/JitterBench/Common.hpp +@@ -43,9 +43,9 @@ THE SOFTWARE. + #endif + + // Macro for collecting HW_REG_HW_ID +-#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__NVCC__) ++#if defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__) + #define GetHwId(val) \ +- val = 0 ++ asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s" (val)); + #else + #define GetHwId(val) \ + asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (val)); diff --git a/pkgs/development/rocm-modules/6/rdc/default.nix b/pkgs/development/rocm-modules/6/rdc/default.nix index 36f85f90c3dc..315e922feb5a 100644 --- a/pkgs/development/rocm-modules/6/rdc/default.nix +++ b/pkgs/development/rocm-modules/6/rdc/default.nix @@ -48,7 +48,7 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rdc"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -64,7 +64,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rdc"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-s/31b8/Kn5l1QJ941UMSB8SCzpvODsPfOLMmEBUYYmY="; + hash = "sha256-zILZPW9Lx5T+cMDqTg/zWy3ro+Nypzc9bDNTupZjt4s="; }; nativeBuildInputs = [ @@ -95,6 +95,7 @@ stdenv.mkDerivation (finalAttrs: { cmakeFlags = [ "-DCMAKE_VERBOSE_MAKEFILE=OFF" "-DRDC_INSTALL_PREFIX=${placeholder "out"}" + "-DBUILD_RVS=OFF" # TODO: Needs RVS package "-DBUILD_ROCRTEST=ON" "-DRSMI_INC_DIR=${rocm-smi}/include" "-DRSMI_LIB_DIR=${rocm-smi}/lib" diff --git a/pkgs/development/rocm-modules/6/rocalution/default.nix b/pkgs/development/rocm-modules/6/rocalution/default.nix index c0b7222747fd..0fbf4025bf3f 100644 --- a/pkgs/development/rocm-modules/6/rocalution/default.nix +++ b/pkgs/development/rocm-modules/6/rocalution/default.nix @@ -23,7 +23,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocalution"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -42,7 +42,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocALUTION"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-xdZ3HUiRGsreHfJH8RgL/s3jGyC5ABmBKcEfgtqWg8Y="; + hash = "sha256-bZx1Cc2jcIfysohKCKzj5mowM3IeCelRhVaBU73KnTo="; }; nativeBuildInputs = [ @@ -65,11 +65,7 @@ stdenv.mkDerivation (finalAttrs: { gtest ]; - CXXFLAGS = "-I${openmp.dev}/include"; cmakeFlags = [ - "-DOpenMP_C_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_CXX_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_omp_LIBRARY=${openmp}/lib" "-DROCM_PATH=${clr}" "-DHIP_ROOT_DIR=${clr}" "-DSUPPORT_HIP=ON" diff --git a/pkgs/development/rocm-modules/6/rocblas/default.nix b/pkgs/development/rocm-modules/6/rocblas/default.nix index b155bfebcc86..768ab094b38f 100644 --- a/pkgs/development/rocm-modules/6/rocblas/default.nix +++ b/pkgs/development/rocm-modules/6/rocblas/default.nix @@ -4,12 +4,14 @@ fetchFromGitHub, fetchpatch, rocmUpdateScript, + writableTmpDirAsHomeHook, cmake, rocm-cmake, clr, python3, tensile, - msgpack, + boost, + msgpack-cxx, libxml2, gtest, gfortran, @@ -17,46 +19,19 @@ git, amd-blis, zstd, + roctracer, hipblas-common, hipblaslt, python3Packages, rocm-smi, + pkg-config, buildTensile ? true, buildTests ? true, buildBenchmarks ? true, - # https://github.com/ROCm/Tensile/issues/1757 - # Allows gfx101* users to use rocBLAS normally. - # Turn the below two values to `true` after the fix has been cherry-picked - # into a release. Just backporting that single fix is not enough because it - # depends on some previous commits. tensileSepArch ? true, tensileLazyLib ? true, withHipBlasLt ? true, - # `gfx940`, `gfx941` are not present in this list because they are early - # engineering samples, and all final MI300 hardware are `gfx942`: - # https://github.com/NixOS/nixpkgs/pull/298388#issuecomment-2032791130 - # - # `gfx1012` is not present in this list because the ISA compatibility patches - # would force all `gfx101*` GPUs to run as `gfx1010`, so `gfx101*` GPUs will - # always try to use `gfx1010` code objects, hence building for `gfx1012` is - # useless: https://github.com/NixOS/nixpkgs/pull/298388#issuecomment-2076327152 - gpuTargets ? ( - clr.localGpuTargets or [ - "gfx900" - "gfx906" - "gfx908" - "gfx90a" - "gfx942" - "gfx1010" - "gfx1030" - "gfx1100" - "gfx1101" - "gfx1102" - "gfx1151" - "gfx1200" - "gfx1201" - ] - ), + gpuTargets ? (clr.localGpuTargets or clr.gpuTargets), }: let @@ -64,17 +39,13 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rocblas${clr.gpuArchSuffix}"; - version = "6.3.3"; - - outputs = [ - "out" - ]; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocBLAS"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-IYcrVcGH4yZDkFZeNOJPfG0qsPS/WiH0fTSUSdo1BH4="; + hash = "sha256-FCzo/BOk4xLEFkdOdqcCXh4a9t3/OIIBEy8oz6oOMWg="; }; nativeBuildInputs = [ @@ -83,6 +54,7 @@ stdenv.mkDerivation (finalAttrs: { rocm-cmake clr git + pkg-config ] ++ lib.optionals buildTensile [ tensile @@ -91,13 +63,16 @@ stdenv.mkDerivation (finalAttrs: { buildInputs = [ python3 hipblas-common + roctracer + openmp + amd-blis ] ++ lib.optionals withHipBlasLt [ hipblaslt ] ++ lib.optionals buildTensile [ zstd - msgpack + msgpack-cxx libxml2 python3Packages.msgpack python3Packages.zstandard @@ -107,27 +82,19 @@ stdenv.mkDerivation (finalAttrs: { ] ++ lib.optionals (buildTests || buildBenchmarks) [ gfortran - openmp - amd-blis rocm-smi ] ++ lib.optionals (buildTensile || buildTests || buildBenchmarks) [ python3Packages.pyyaml ]; - dontStrip = true; - env.CXXFLAGS = - "-O3 -DNDEBUG -I${hipblas-common}/include" - + lib.optionalString (buildTests || buildBenchmarks) " -I${amd-blis}/include/blis"; - # Fails to link tests if we don't add amd-blis libs - env.LDFLAGS = lib.optionalString ( - buildTests || buildBenchmarks - ) "-Wl,--as-needed -L${amd-blis}/lib -lblis-mt -lcblas"; + env.CXXFLAGS = "-fopenmp -I${lib.getDev boost}/include -I${hipblas-common}/include -I${roctracer}/include"; + # Fails to link tests with undefined symbol: cblas_* + env.LDFLAGS = lib.optionalString (buildTests || buildBenchmarks) "-Wl,--as-needed -lcblas"; env.TENSILE_ROCM_ASSEMBLER_PATH = "${stdenv.cc}/bin/clang++"; cmakeFlags = [ - (lib.cmakeFeature "CMAKE_BUILD_TYPE" "Release") - (lib.cmakeBool "CMAKE_VERBOSE_MAKEFILE" true) + (lib.cmakeFeature "Boost_INCLUDE_DIR" "${lib.getDev boost}/include") # msgpack FindBoost fails to find boost (lib.cmakeFeature "CMAKE_EXECUTE_PROCESS_COMMAND_ECHO" "STDERR") (lib.cmakeFeature "CMAKE_Fortran_COMPILER" "${lib.getBin gfortran}/bin/gfortran") (lib.cmakeFeature "CMAKE_Fortran_COMPILER_AR" "${lib.getBin gfortran}/bin/ar") @@ -144,8 +111,8 @@ stdenv.mkDerivation (finalAttrs: { (lib.cmakeBool "BUILD_CLIENTS_BENCHMARKS" buildBenchmarks) (lib.cmakeBool "BUILD_CLIENTS_SAMPLES" buildBenchmarks) (lib.cmakeBool "BUILD_OFFLOAD_COMPRESS" true) - # Temporarily set variables to work around upstream CMakeLists issue - # Can be removed once https://github.com/ROCm/rocm-cmake/issues/121 is fixed + # # Temporarily set variables to work around upstream CMakeLists issue + # # Can be removed once https://github.com/ROCm/rocm-cmake/issues/121 is fixed "-DCMAKE_INSTALL_BINDIR=bin" "-DCMAKE_INSTALL_INCLUDEDIR=include" "-DCMAKE_INSTALL_LIBDIR=lib" @@ -153,6 +120,9 @@ stdenv.mkDerivation (finalAttrs: { ++ lib.optionals buildTensile [ "-DCPACK_SET_DESTDIR=OFF" "-DLINK_BLIS=ON" + "-DBLIS_LIB=${amd-blis}/lib/libblis-mt.so" + "-DBLIS_INCLUDE_DIR=${amd-blis}/include/blis/" + "-DBLA_PREFER_PKGCONFIG=ON" "-DTensile_CODE_OBJECT_VERSION=default" "-DTensile_LOGIC=asm_full" "-DTensile_LIBRARY_FORMAT=msgpack" @@ -169,6 +139,7 @@ stdenv.mkDerivation (finalAttrs: { url = "https://github.com/GZGavinZhao/rocBLAS/commit/89b75ff9cc731f71f370fad90517395e117b03bb.patch"; hash = "sha256-W/ohOOyNCcYYLOiQlPzsrTlNtCBdJpKVxO8s+4G7sjo="; }) + ./hiplaslt-unstable-compat.patch ]; # Pass $NIX_BUILD_CORES to Tensile @@ -176,7 +147,8 @@ stdenv.mkDerivation (finalAttrs: { substituteInPlace cmake/build-options.cmake \ --replace-fail 'Tensile_CPU_THREADS ""' 'Tensile_CPU_THREADS "$ENV{NIX_BUILD_CORES}"' substituteInPlace CMakeLists.txt \ - --replace-fail "4.42.0" "4.43.0" + --replace-fail "4.43.0" "4.44.0" \ + --replace-fail '0.10' '1.0' ''; passthru.updateScript = rocmUpdateScript { diff --git a/pkgs/development/rocm-modules/6/rocblas/hiplaslt-unstable-compat.patch b/pkgs/development/rocm-modules/6/rocblas/hiplaslt-unstable-compat.patch new file mode 100644 index 000000000000..4d8c885c8d39 --- /dev/null +++ b/pkgs/development/rocm-modules/6/rocblas/hiplaslt-unstable-compat.patch @@ -0,0 +1,99 @@ +diff --git a/library/src/hipblaslt_host.cpp b/library/src/hipblaslt_host.cpp +index 8080070c..97d5216e 100644 +--- a/library/src/hipblaslt_host.cpp ++++ b/library/src/hipblaslt_host.cpp +@@ -155,22 +155,22 @@ namespace + hipblaslt_compute_type); + + hipblaslt_ext::GemmProblemType problemType; +- problemType.op_a = (hipblasOperation_t)prob.trans_a; +- problemType.op_b = (hipblasOperation_t)prob.trans_b; +- problemType.type_a = hipblaslt_datatype; +- problemType.type_b = hipblaslt_datatype; +- problemType.type_c = hipblaslt_datatype; +- problemType.type_d = hipblaslt_datatype; +- problemType.type_compute = hipblaslt_compute_type; ++ problemType.setOpA((hipblasOperation_t)prob.trans_a); ++ problemType.setOpB((hipblasOperation_t)prob.trans_b); ++ problemType.setTypeA(hipblaslt_datatype); ++ problemType.setTypeB(hipblaslt_datatype); ++ problemType.setTypeC(hipblaslt_datatype); ++ problemType.setTypeD(hipblaslt_datatype); ++ problemType.setTypeCompute(hipblaslt_compute_type); + + hipblaslt_ext::GemmEpilogue epilogue; + hipblaslt_ext::GemmInputs inputs; +- inputs.a = (void*)(prob.A + prob.buffer_offset_a); +- inputs.b = (void*)(prob.B + prob.buffer_offset_b); +- inputs.c = (void*)(prob.C + prob.buffer_offset_c); +- inputs.d = (void*)(prob.D + prob.buffer_offset_d); +- inputs.alpha = (void*)prob.alpha; +- inputs.beta = (void*)prob.beta; ++ inputs.setA((void*)(prob.A + prob.buffer_offset_a)); ++ inputs.setB((void*)(prob.B + prob.buffer_offset_b)); ++ inputs.setC((void*)(prob.C + prob.buffer_offset_c)); ++ inputs.setD((void*)(prob.D + prob.buffer_offset_d)); ++ inputs.setAlpha((void*)prob.alpha); ++ inputs.setBeta((void*)prob.beta); + + gemm.setProblem(prob.m, + prob.n, +@@ -214,13 +214,13 @@ namespace + hipblaslt_compute_type); + + hipblaslt_ext::GemmProblemType problemType; +- problemType.op_a = (hipblasOperation_t)prob.trans_a; +- problemType.op_b = (hipblasOperation_t)prob.trans_b; +- problemType.type_a = hipblaslt_datatype; +- problemType.type_b = hipblaslt_datatype; +- problemType.type_c = hipblaslt_datatype; +- problemType.type_d = hipblaslt_datatype; +- problemType.type_compute = hipblaslt_compute_type; ++ problemType.setOpA((hipblasOperation_t)prob.trans_a); ++ problemType.setOpB((hipblasOperation_t)prob.trans_b); ++ problemType.setTypeA(hipblaslt_datatype); ++ problemType.setTypeB(hipblaslt_datatype); ++ problemType.setTypeC(hipblaslt_datatype); ++ problemType.setTypeD(hipblaslt_datatype); ++ problemType.setTypeCompute(hipblaslt_compute_type); + + std::vector Ms(prob.batch_count); + std::vector Ns(prob.batch_count); +@@ -251,12 +251,12 @@ namespace + stridecs[batch] = prob.batch_stride_c; + strideds[batch] = prob.batch_stride_d; + batch_counts[batch] = 1; +- inputs[batch].a = (void*)(prob.batch_A[batch] + prob.buffer_offset_a); +- inputs[batch].b = (void*)(prob.batch_B[batch] + prob.buffer_offset_b); +- inputs[batch].c = (void*)(prob.batch_C[batch] + prob.buffer_offset_c); +- inputs[batch].d = (void*)(prob.batch_D[batch] + prob.buffer_offset_d); +- inputs[batch].alpha = (void*)prob.alpha; +- inputs[batch].beta = (void*)prob.beta; ++ inputs[batch].setA((void*)(prob.batch_A[batch] + prob.buffer_offset_a)); ++ inputs[batch].setB((void*)(prob.batch_B[batch] + prob.buffer_offset_b)); ++ inputs[batch].setC((void*)(prob.batch_C[batch] + prob.buffer_offset_c)); ++ inputs[batch].setD((void*)(prob.batch_D[batch] + prob.buffer_offset_d)); ++ inputs[batch].setAlpha((void*)prob.alpha); ++ inputs[batch].setBeta((void*)prob.beta); + } + + gemm.setProblem(Ms, +diff --git a/library/src/tensile_host.cpp b/library/src/tensile_host.cpp +index 1b1289f3..ed463725 100644 +--- a/library/src/tensile_host.cpp ++++ b/library/src/tensile_host.cpp +@@ -271,14 +271,6 @@ namespace + { + return Tensile::LazyLoadingInit::gfx90a; + } +- else if(deviceString.find("gfx940") != std::string::npos) +- { +- return Tensile::LazyLoadingInit::gfx940; +- } +- else if(deviceString.find("gfx941") != std::string::npos) +- { +- return Tensile::LazyLoadingInit::gfx941; +- } + else if(deviceString.find("gfx942") != std::string::npos) + { + return Tensile::LazyLoadingInit::gfx942; diff --git a/pkgs/development/rocm-modules/6/rocdbgapi/default.nix b/pkgs/development/rocm-modules/6/rocdbgapi/default.nix index cd07af5d4cdf..0140218222a7 100644 --- a/pkgs/development/rocm-modules/6/rocdbgapi/default.nix +++ b/pkgs/development/rocm-modules/6/rocdbgapi/default.nix @@ -12,6 +12,7 @@ texliveSmall, doxygen, graphviz, + writableTmpDirAsHomeHook, buildDocs ? true, }: @@ -52,7 +53,7 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rocdbgapi"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -65,7 +66,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "ROCdbgapi"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-6itfBrWVspobU47aiJAOQoxT8chwrq9scRn0or3bXto="; + hash = "sha256-Rr8+SNeFps0rjk4Jn2+rFmtRJfL42l0tNOz13oZQy+I="; }; # FIXME: remove once https://github.com/doxygen/doxygen/issues/11634 is resolved @@ -84,6 +85,7 @@ stdenv.mkDerivation (finalAttrs: { git ] ++ lib.optionals buildDocs [ + writableTmpDirAsHomeHook latex doxygen graphviz @@ -106,7 +108,6 @@ stdenv.mkDerivation (finalAttrs: { # Unfortunately, it seems like we have to call make on this manually postBuild = lib.optionalString buildDocs '' - export HOME=$(mktemp -d) make -j$NIX_BUILD_CORES doc ''; diff --git a/pkgs/development/rocm-modules/6/rocfft/default.nix b/pkgs/development/rocm-modules/6/rocfft/default.nix index dc64f6bf3e93..0f948250297c 100644 --- a/pkgs/development/rocm-modules/6/rocfft/default.nix +++ b/pkgs/development/rocm-modules/6/rocfft/default.nix @@ -14,18 +14,19 @@ gtest, openmp, rocrand, + hiprand, gpuTargets ? clr.localGpuTargets or clr.gpuTargets, }: stdenv.mkDerivation (finalAttrs: { pname = "rocfft${clr.gpuArchSuffix}"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocFFT"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-RrxdwZ64uC7lQzyJI1eGHX2dmRnW8TfNThnuvuz5XWo="; + hash = "sha256-yaOjBF2aJkCBlxkydyOsrfT4lNZ0BVkS2jJC0fEiBug="; }; nativeBuildInputs = [ @@ -35,14 +36,23 @@ stdenv.mkDerivation (finalAttrs: { rocm-cmake ]; - # FIXME: rocfft_aot_helper runs at the end of the build and has a risk of timing it out - # due to a long period with no terminal output - buildInputs = [ sqlite ]; + buildInputs = [ + sqlite + hiprand + ]; + + patches = [ + # Fixes build timeout due to no log output during rocfft_aot step + ./log-every-n-aot-jobs.patch + ]; cmakeFlags = [ - "-DCMAKE_C_COMPILER=hipcc" - "-DCMAKE_CXX_COMPILER=hipcc" "-DSQLITE_USE_SYSTEM_PACKAGE=ON" + "-DHIP_PLATFORM=amd" + "-DBUILD_CLIENTS=OFF" + "-DBUILD_SHARED_LIBS=ON" + "-DUSE_HIPRAND=ON" + "-DROCFFT_KERNEL_CACHE_ENABLE=ON" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" @@ -50,7 +60,7 @@ stdenv.mkDerivation (finalAttrs: { "-DCMAKE_INSTALL_INCLUDEDIR=include" ] ++ lib.optionals (gpuTargets != [ ]) [ - "-DAMDGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" + "-DGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" ]; passthru = { @@ -74,11 +84,7 @@ stdenv.mkDerivation (finalAttrs: { gtest openmp rocrand - ]; - - cmakeFlags = [ - "-DCMAKE_C_COMPILER=hipcc" - "-DCMAKE_CXX_COMPILER=hipcc" + hiprand ]; postInstall = '' @@ -112,11 +118,6 @@ stdenv.mkDerivation (finalAttrs: { rocrand ]; - cmakeFlags = [ - "-DCMAKE_C_COMPILER=hipcc" - "-DCMAKE_CXX_COMPILER=hipcc" - ]; - postInstall = '' cp -a ../../../scripts/perf "$out/bin" ''; @@ -141,11 +142,6 @@ stdenv.mkDerivation (finalAttrs: { rocrand ]; - cmakeFlags = [ - "-DCMAKE_C_COMPILER=hipcc" - "-DCMAKE_CXX_COMPILER=hipcc" - ]; - installPhase = '' runHook preInstall mkdir "$out" diff --git a/pkgs/development/rocm-modules/6/rocfft/log-every-n-aot-jobs.patch b/pkgs/development/rocm-modules/6/rocfft/log-every-n-aot-jobs.patch new file mode 100644 index 000000000000..266f495d3e7c --- /dev/null +++ b/pkgs/development/rocm-modules/6/rocfft/log-every-n-aot-jobs.patch @@ -0,0 +1,35 @@ +diff --git a/library/src/rocfft_aot_helper.cpp b/library/src/rocfft_aot_helper.cpp +index f0a889f4..452eb37f 100644 +--- a/library/src/rocfft_aot_helper.cpp ++++ b/library/src/rocfft_aot_helper.cpp +@@ -771,26 +771,22 @@ int main(int argc, char** argv) + for(size_t i = 0; i < NUM_THREADS; ++i) + { + threads.emplace_back([&queue, &gpu_archs]() { ++ int compile_count = 0; + while(true) + { + auto item = queue.pop(); + if(item.kernel_name.empty()) + break; + ++ if(++compile_count % 16 == 0) ++ std::cerr << "rocfft_aot_helper processing " << item.kernel_name << std::endl << std::flush; + for(const auto& gpu_arch : gpu_archs) + { +- if(item.sol_arch_name.empty()) ++ if(item.sol_arch_name.empty() || gpu_arch.find(item.sol_arch_name) != std::string::npos) + { + RTCCache::cached_compile( + item.kernel_name, gpu_arch, item.generate_src, generator_sum()); + } +- else if(gpu_arch.find(item.sol_arch_name) != std::string::npos) +- { +- // std::cout << "arch: " << gpu_arch +- // << ", solution-kernel: " << item.kernel_name << std::endl; +- RTCCache::cached_compile( +- item.kernel_name, gpu_arch, item.generate_src, generator_sum()); +- } + } + } + }); diff --git a/pkgs/development/rocm-modules/6/rocgdb/default.nix b/pkgs/development/rocm-modules/6/rocgdb/default.nix index 8e009e583b08..96909655c5c2 100644 --- a/pkgs/development/rocm-modules/6/rocgdb/default.nix +++ b/pkgs/development/rocm-modules/6/rocgdb/default.nix @@ -23,13 +23,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocgdb"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "ROCgdb"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-Z+uk+ViLXgk5hXrIhVHRY0Kly7mktYms7M3o9Tmxv8s="; + hash = "sha256-evDWg2w2FHv6OU5BQOCAXTlDm7JpwdJ3Wh5a2i5r1gQ="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/rocm-cmake/default.nix b/pkgs/development/rocm-modules/6/rocm-cmake/default.nix index 23a48dcd55b2..8ccc8ff182ad 100644 --- a/pkgs/development/rocm-modules/6/rocm-cmake/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-cmake/default.nix @@ -9,13 +9,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocm-cmake"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocm-cmake"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-U4vGkH2iUlNJkqiNmVuFianD4WR9yuGvZsYG58smg0k="; + hash = "sha256-wAipNWAB66YNf7exLSNPAzg3NgkGD9LPKfKiulL5yak="; }; nativeBuildInputs = [ cmake ]; diff --git a/pkgs/development/rocm-modules/6/rocm-comgr/default.nix b/pkgs/development/rocm-modules/6/rocm-comgr/default.nix index 58b35a1a7621..de0b2434cfac 100644 --- a/pkgs/development/rocm-modules/6/rocm-comgr/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-comgr/default.nix @@ -37,8 +37,8 @@ stdenv.mkDerivation (finalAttrs: { }) #[Comgr] Extend ISA compatibility for CCOB (fetchpatch { - sha256 = "sha256-6Rwz12Lk4R2JK3olii3cr2Zd0ZLYe7VSpK1YRCOsJWY="; - url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/2d8c459a4d4c0567a7a275b4b54560d88e5c6919.patch"; + sha256 = "sha256-PCi0QHLiEQCTIYRtSSbhOjXANJ3zC3VLdMED1BEfQeg="; + url = "https://github.com/GZGavinZhao/rocm-llvm-project/commit/fa80abb77d5ae6f8d89ab956e7ebda9c802a804f.patch"; relative = "amd/comgr"; }) ]; diff --git a/pkgs/development/rocm-modules/6/rocm-core/default.nix b/pkgs/development/rocm-modules/6/rocm-core/default.nix index 8f5100ec0000..6c0e178aad86 100644 --- a/pkgs/development/rocm-modules/6/rocm-core/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-core/default.nix @@ -17,13 +17,13 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rocm-core"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocm-core"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-UDnPGvgwzwv49CzF+Kt0v95CsxS33BZeqNcKw1K6jRI="; + hash = "sha256-Z21k+0vi/P35WTHGAdfUAX7/jVv+d9g5YBo+HTi1Mpk="; }; patches = [ diff --git a/pkgs/development/rocm-modules/6/rocm-docs-core/default.nix b/pkgs/development/rocm-modules/6/rocm-docs-core/default.nix index c86fa10cdff4..c8329621d5f6 100644 --- a/pkgs/development/rocm-modules/6/rocm-docs-core/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-docs-core/default.nix @@ -24,14 +24,14 @@ # FIXME: Move to rocmPackages_common buildPythonPackage rec { pname = "rocm-docs-core"; - version = "1.17.0"; + version = "1.23.0"; format = "pyproject"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocm-docs-core"; rev = "v${version}"; - hash = "sha256-fGRJyQq0Eook1Dc9Qy+dehQ5BVNX+6pkkFN9adb21Eo="; + hash = "sha256-5Qh83eJ9vju/uUb9gGA4B5Bh1WZCygIRbSnaEZzIdbw="; }; buildInputs = [ setuptools ]; diff --git a/pkgs/development/rocm-modules/6/rocm-runtime/default.nix b/pkgs/development/rocm-modules/6/rocm-runtime/default.nix index c40f270699ce..2bc797759af3 100644 --- a/pkgs/development/rocm-modules/6/rocm-runtime/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-runtime/default.nix @@ -19,17 +19,18 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocm-runtime"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "ROCR-Runtime"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-du20+5VNYgwchGO7W7FIVebBqLPtfSBnmPVbPpgEZjo="; + hash = "sha256-zs0nydwYUY+8uiPyJxgTfAiV7spUMbESb0jUUvFf+AU="; }; - env.CFLAGS = "-I${numactl.dev}/include -I${elfutils.dev}/include -w"; - env.CXXFLAGS = "-I${numactl.dev}/include -I${elfutils.dev}/include -w"; + cmakeBuildType = "RelWithDebInfo"; + separateDebugInfo = true; + __structuredAttrs = true; nativeBuildInputs = [ pkg-config @@ -56,17 +57,17 @@ stdenv.mkDerivation (finalAttrs: { ]; patches = [ + (fetchpatch { + # rocr: Extend HIP ISA compatibility check + sha256 = "sha256-8r2Lb5lBfFaZC3knCxfXGcnkzNv6JxOKyJn2rD5gus4="; + url = "https://github.com/GZGavinZhao/ROCR-Runtime/commit/7c63e7185d8fcf08537a278908946145f6231121.patch"; + }) # Patches for UB at runtime https://github.com/ROCm/ROCR-Runtime/issues/272 (fetchpatch { # [PATCH] hsa-runtime: set underlying type of hsa_region_info_t and hsa_amd_region_info_t to int url = "https://github.com/ROCm/ROCR-Runtime/commit/39a6a168fa07e289a10f6e20e6ead4e303e99ba0.patch"; hash = "sha256-CshJJDvII1nNyNmt+YjwMwfBHUTlrdsxkhwfgBwO+WE="; }) - (fetchpatch { - # [PATCH] rocr: refactor of runtime.cpp based on Coverity - url = "https://github.com/ROCm/ROCR-Runtime/commit/441bd9fe6c7bdb5c4c31f71524ed642786bc923e.patch"; - hash = "sha256-7bQXxGkipzgT2aXRxCuh3Sfmo/zc/IOmA0x1zB+fMb0="; - }) (fetchpatch { # [PATCH] queues: fix UB due to 1 << 31 url = "https://github.com/ROCm/ROCR-Runtime/commit/9b8a0f5dbee1903fa990a7d8accc1c5fbc549636.patch"; @@ -82,6 +83,9 @@ stdenv.mkDerivation (finalAttrs: { url = "https://github.com/ROCm/ROCR-Runtime/commit/41bfc66aef437a5b349f71105fa4b907cc7e17d5.patch"; hash = "sha256-A7VhPR3eSsmjq2cTBSjBIz9i//WiNjoXm0EsRKtF+ns="; }) + # This causes a circular dependency, aqlprofile relies on hsa-runtime64 + # which is part of rocm-runtime + # Worked around by having rocprofiler load aqlprofile directly ./remove-hsa-aqlprofile-dep.patch ]; diff --git a/pkgs/development/rocm-modules/6/rocm-smi/default.nix b/pkgs/development/rocm-modules/6/rocm-smi/default.nix index 981c3ca6f892..f146d6f9fb43 100644 --- a/pkgs/development/rocm-modules/6/rocm-smi/default.nix +++ b/pkgs/development/rocm-modules/6/rocm-smi/default.nix @@ -3,26 +3,35 @@ stdenv, fetchFromGitHub, rocmUpdateScript, + pkg-config, + libdrm, cmake, wrapPython, }: stdenv.mkDerivation (finalAttrs: { pname = "rocm-smi"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocm_smi_lib"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-j9pkyUt+p6IkhawIhiTymqDBydxXZunxmdyCyRN0RxE="; + hash = "sha256-yJ3Bf+tM39JWbY+A0NlpHNkvythdAdz6ZVp1AvLcXhk="; }; - patches = [ ./cmake.patch ]; + patches = [ + ./cmake.patch + ]; + + propagatedBuildInputs = [ + libdrm + ]; nativeBuildInputs = [ cmake wrapPython + pkg-config ]; cmakeFlags = [ @@ -33,11 +42,21 @@ stdenv.mkDerivation (finalAttrs: { "-DCMAKE_INSTALL_INCLUDEDIR=include" ]; - postInstall = '' - wrapPythonProgramsIn $out - mv $out/libexec/rocm_smi/.rsmiBindingsInit.py-wrapped $out/libexec/rocm_smi/rsmiBindingsInit.py - mv $out/libexec/rocm_smi/.rsmiBindings.py-wrapped $out/libexec/rocm_smi/rsmiBindings.py - ''; + postInstall = + # wrap python programs, but undo two that need to be importable at that path + '' + wrapPythonProgramsIn $out + mv $out/libexec/rocm_smi/.rsmiBindingsInit.py-wrapped $out/libexec/rocm_smi/rsmiBindingsInit.py + mv $out/libexec/rocm_smi/.rsmiBindings.py-wrapped $out/libexec/rocm_smi/rsmiBindings.py + '' + # workaround: propagate libdrm/ manually + # rocmcxx doesn't automatically add buildInputs to isystem include path like + # wrapper based toolchains, cmake files often don't find_package(rocm-smi) so + # can't rely on cmake propagated interface + # upstream have been shipping libdrm copied into /opt/rocm + + '' + ln -s ${libdrm.dev}/include/libdrm/ $out/include/ + ''; passthru.updateScript = rocmUpdateScript { name = finalAttrs.pname; diff --git a/pkgs/development/rocm-modules/6/rocminfo/default.nix b/pkgs/development/rocm-modules/6/rocminfo/default.nix index 18166bc780df..e3e76b4fd35f 100644 --- a/pkgs/development/rocm-modules/6/rocminfo/default.nix +++ b/pkgs/development/rocm-modules/6/rocminfo/default.nix @@ -9,25 +9,17 @@ busybox, python3, gnugrep, - clr, # Only for localGpuTargets - # rocminfo requires that the calling user have a password and be in - # the video group. If we let rocm_agent_enumerator rely upon - # rocminfo's output, then it, too, has those requirements. Instead, - # we can specify the GPU targets for this system (e.g. "gfx803" for - # Polaris) such that no system call is needed for downstream - # compilers to determine the desired target. - defaultTargets ? (clr.localGpuTargets or [ ]), }: stdenv.mkDerivation (finalAttrs: { - version = "6.3.3"; + version = "6.4.3"; pname = "rocminfo"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocminfo"; rev = "rocm-${finalAttrs.version}"; - sha256 = "sha256-fQPtO5TNbCbaZZ7VtGkkqng5QZ+FcScdh1opWr5YkLU="; + sha256 = "sha256-YscZ5sFsLOVBg98w2X6vTzniTvl9NfCkIE+HAH6vv5Y="; }; strictDeps = true; @@ -46,10 +38,6 @@ stdenv.mkDerivation (finalAttrs: { sed 's,lsmod | grep ,${busybox}/bin/lsmod | ${gnugrep}/bin/grep ,' -i rocminfo.cc ''; - postInstall = lib.optionalString (defaultTargets != [ ]) '' - echo '${lib.concatStringsSep "\n" defaultTargets}' > $out/bin/target.lst - ''; - passthru.updateScript = rocmUpdateScript { name = finalAttrs.pname; inherit (finalAttrs.src) owner; diff --git a/pkgs/development/rocm-modules/6/rocmlir/default.nix b/pkgs/development/rocm-modules/6/rocmlir/default.nix index 480b9d2fa48d..281b75f76814 100644 --- a/pkgs/development/rocm-modules/6/rocmlir/default.nix +++ b/pkgs/development/rocm-modules/6/rocmlir/default.nix @@ -5,7 +5,6 @@ rocmUpdateScript, cmake, rocm-cmake, - rocminfo, clr, git, libxml2, @@ -41,7 +40,7 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rocmlir${suffix}"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -54,13 +53,12 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocMLIR"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-0SQ6uLDRfVfdCX+8a7D6pu6dYlFvX0HFzCDEvlKYfak="; + hash = "sha256-p/gvr1Z6yZtO5N+ecSouXiCrf520jt1HMOy/tohUHfI="; }; nativeBuildInputs = [ cmake rocm-cmake - clr python3Packages.python python3Packages.tomli ]; @@ -77,10 +75,6 @@ stdenv.mkDerivation (finalAttrs: { ncurses ]; - patches = [ - ./initparamdata-sort-const.patch - ]; - cmakeFlags = [ "-DLLVM_TARGETS_TO_BUILD=AMDGPU;${llvmNativeTarget}" "-DCMAKE_BUILD_TYPE=Release" @@ -109,12 +103,8 @@ stdenv.mkDerivation (finalAttrs: { substituteInPlace mlir/lib/Analysis/BufferDependencyAnalysis.cpp \ --replace-fail "enum EffectType { read, write, unknown };" "enum class EffectType { read, write, unknown };" - # remove when no longer required - substituteInPlace mlir/test/{e2e/generateE2ETest.py,fusion/e2e/generate-fusion-tests.py} \ - --replace-fail "\"/opt/rocm/bin" "\"${rocminfo}/bin" - substituteInPlace mlir/utils/performance/common/CMakeLists.txt \ - --replace-fail "/opt/rocm" "${clr}" + --replace-fail " PATHS /opt/rocm" "" ''; dontBuild = true; diff --git a/pkgs/development/rocm-modules/6/rocmlir/initparamdata-sort-const.patch b/pkgs/development/rocm-modules/6/rocmlir/initparamdata-sort-const.patch deleted file mode 100644 index 37f521f6e02d..000000000000 --- a/pkgs/development/rocm-modules/6/rocmlir/initparamdata-sort-const.patch +++ /dev/null @@ -1,13 +0,0 @@ -diff --git a/mlir/include/mlir/Dialect/Rock/Tuning/GridwiseGemmParams.h b/mlir/include/mlir/Dialect/Rock/Tuning/GridwiseGemmParams.h -index 3f5ee596819a..590d53788822 100644 ---- a/mlir/include/mlir/Dialect/Rock/Tuning/GridwiseGemmParams.h -+++ b/mlir/include/mlir/Dialect/Rock/Tuning/GridwiseGemmParams.h -@@ -209,7 +209,7 @@ private: - size_t original_pos; - int64_t padding_amount; - -- bool operator<(const InitParamData &rhs) { -+ bool operator<(const InitParamData &rhs) const { - if (this->padding_amount < rhs.padding_amount) { - return true; - } else if (this->padding_amount == rhs.padding_amount) { diff --git a/pkgs/development/rocm-modules/6/rocprim/default.nix b/pkgs/development/rocm-modules/6/rocprim/default.nix index 1b92663403c4..567fe7b6ae97 100644 --- a/pkgs/development/rocm-modules/6/rocprim/default.nix +++ b/pkgs/development/rocm-modules/6/rocprim/default.nix @@ -15,7 +15,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocprim"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocPRIM"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-0aHxpBuYIYhI2UER45YhHHL5YcxA+XeXoihcUs2AmCo="; + hash = "sha256-lH4MlBEkVJocq1VliGBtb7VvWfS6p/uIKWR239fSjRY="; }; nativeBuildInputs = [ @@ -49,7 +49,7 @@ stdenv.mkDerivation (finalAttrs: { ]; cmakeFlags = [ - "-DCMAKE_BUILD_TYPE=Release" + "-DHIP_CXX_COMPILER=amdclang++" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" diff --git a/pkgs/development/rocm-modules/6/rocprofiler-register/default.nix b/pkgs/development/rocm-modules/6/rocprofiler-register/default.nix index 64ef999a506c..4b580f47674c 100644 --- a/pkgs/development/rocm-modules/6/rocprofiler-register/default.nix +++ b/pkgs/development/rocm-modules/6/rocprofiler-register/default.nix @@ -18,13 +18,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocprofiler-register"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocprofiler-register"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-UZsCiGnudsbL1v5lKBx7Vz3/HRnGn4f86Pd+qu3ryh0="; + hash = "sha256-HaN4XMHuCRDfKOpfuZ2SkOEQfAZKouh6luqbtATUYm0="; fetchSubmodules = true; }; diff --git a/pkgs/development/rocm-modules/6/rocprofiler/0000-dont-install-tests-hsaco.patch b/pkgs/development/rocm-modules/6/rocprofiler/0000-dont-install-tests-hsaco.patch deleted file mode 100644 index 70550d0119a9..000000000000 --- a/pkgs/development/rocm-modules/6/rocprofiler/0000-dont-install-tests-hsaco.patch +++ /dev/null @@ -1,15 +0,0 @@ -diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt -index 46efbd5..ca2cc3b 100644 ---- a/test/CMakeLists.txt -+++ b/test/CMakeLists.txt -@@ -127,10 +127,6 @@ function(generate_hsaco TARGET_ID INPUT_FILE OUTPUT_FILE) - DEPENDS ${INPUT_FILE} clang - COMMENT "Building ${OUTPUT_FILE}..." - VERBATIM) -- install( -- FILES ${PROJECT_BINARY_DIR}/${OUTPUT_FILE} -- DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/tests-v1 -- COMPONENT tests) - set(HSACO_TARGET_LIST - ${HSACO_TARGET_LIST} ${PROJECT_BINARY_DIR}/${OUTPUT_FILE} - PARENT_SCOPE) diff --git a/pkgs/development/rocm-modules/6/rocprofiler/default.nix b/pkgs/development/rocm-modules/6/rocprofiler/default.nix index 06612a65d591..341aae4d1e95 100644 --- a/pkgs/development/rocm-modules/6/rocprofiler/default.nix +++ b/pkgs/development/rocm-modules/6/rocprofiler/default.nix @@ -7,6 +7,7 @@ cmake, clang, clr, + aqlprofile, rocm-core, rocm-runtime, rocm-device-libs, @@ -44,22 +45,16 @@ let in stdenv.mkDerivation (finalAttrs: { pname = "rocprofiler"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocprofiler"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-x6DVt1logBE8aNnuwukQhsv/vRqkJALcfAF+6yEQuIk="; + hash = "sha256-CgW8foM4W3K19kUK/l8IsH2Q9DHi/z88viXTxhNqlHQ="; fetchSubmodules = true; }; - patches = [ - # These just simply won't build - ./0000-dont-install-tests-hsaco.patch - ./optional-aql-in-cmake.patch - ]; - nativeBuildInputs = [ cmake clang @@ -80,12 +75,11 @@ stdenv.mkDerivation (finalAttrs: { mpi systemd gtest + aqlprofile ]; propagatedBuildInputs = [ rocmtoolkit-merged ]; - # HACK: allow building without aqlprofile, probably explodes at runtime if use profiling - env.LDFLAGS = "-z nodefs -Wl,-undefined,dynamic_lookup,--unresolved-symbols=ignore-all"; #HACK: rocprofiler's cmake doesn't add these deps properly env.CXXFLAGS = "-I${libpciaccess}/include -I${numactl.dev}/include -I${rocmtoolkit-merged}/include -I${elfutils.dev}/include -w"; @@ -95,9 +89,6 @@ stdenv.mkDerivation (finalAttrs: { "-DGPU_TARGETS=${lib.concatStringsSep ";" gpuTargets}" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 - "-DBUILD_TEST=OFF" - "-DROCPROFILER_BUILD_TESTS=0" - "-DROCPROFILER_BUILD_SAMPLES=0" "-DCMAKE_INSTALL_BINDIR=bin" "-DCMAKE_INSTALL_LIBDIR=lib" "-DCMAKE_INSTALL_INCLUDEDIR=include" @@ -123,6 +114,14 @@ stdenv.mkDerivation (finalAttrs: { postInstall = '' # Why do these have the executable bit set? chmod -x $out/libexec/rocprofiler/counters/*.xml + # rocprof shell script wants to find it in the same bin dir, easiest to symlink in + ln -s ${clr}/bin/rocm_agent_enumerator $out/bin/rocm_agent_enumerator + ''; + + postFixup = '' + patchelf $out/lib/*.so \ + --add-rpath ${aqlprofile}/lib \ + --add-needed libhsa-amd-aqlprofile64.so ''; passthru.updateScript = rocmUpdateScript { diff --git a/pkgs/development/rocm-modules/6/rocprofiler/optional-aql-in-cmake.patch b/pkgs/development/rocm-modules/6/rocprofiler/optional-aql-in-cmake.patch deleted file mode 100644 index 83bf88fb1e7e..000000000000 --- a/pkgs/development/rocm-modules/6/rocprofiler/optional-aql-in-cmake.patch +++ /dev/null @@ -1,147 +0,0 @@ -From https://raw.githubusercontent.com/AphidGit/rocm_compile/refs/heads/main/rocprofiler.patch -diff --git a/cmake_modules/rocprofiler_env.cmake b/cmake_modules/rocprofiler_env.cmake -index 7b7c472..0aba3ed 100644 ---- a/cmake_modules/rocprofiler_env.cmake -+++ b/cmake_modules/rocprofiler_env.cmake -@@ -36,6 +36,7 @@ if(ROCPROFILER_DEBUG_TRACE) - target_compile_definitions(rocprofiler-build-flags INTERFACE DEBUG_TRACE=1) - endif() - -+set(ROCPROFILER_LD_AQLPROFILE false) - # Enable direct loading of AQL-profile HSA extension - if(ROCPROFILER_LD_AQLPROFILE) - target_compile_definitions(rocprofiler-build-flags INTERFACE ROCP_LD_AQLPROFILE=1) -@@ -80,9 +81,3 @@ if("${ROCM_ROOT_DIR}" STREQUAL "") - message(FATAL_ERROR "ROCM_ROOT_DIR is not found.") - endif() - --find_library( -- HSA_AMD_AQLPROFILE_LIBRARY -- NAMES hsa-amd-aqlprofile64 -- HINTS ${CMAKE_PREFIX_PATH} -- PATHS ${ROCM_ROOT_DIR} -- PATH_SUFFIXES lib REQUIRED) -diff --git a/src/api/CMakeLists.txt b/src/api/CMakeLists.txt -index 61782f0..16c83bf 100644 ---- a/src/api/CMakeLists.txt -+++ b/src/api/CMakeLists.txt -@@ -51,15 +51,6 @@ find_file( - NO_DEFAULT_PATH REQUIRED) - get_filename_component(HSA_RUNTIME_INC_PATH ${HSA_H} DIRECTORY) - --find_library( -- AQLPROFILE_LIB "libhsa-amd-aqlprofile64.so" -- HINTS ${CMAKE_PREFIX_PATH} -- PATHS ${ROCM_PATH} -- PATH_SUFFIXES lib) -- --if(NOT AQLPROFILE_LIB) -- message(FATAL_ERROR "AQL_PROFILE not installed. Please install hsa-amd-aqlprofile!") --endif() - - # ######################################################################################## - # Adding Old Library Files -@@ -247,7 +238,7 @@ target_include_directories( - PUBLIC $ - PRIVATE ${LIB_DIR} ${ROOT_DIR} ${PROJECT_SOURCE_DIR}/include/rocprofiler) - target_link_libraries( -- ${ROCPROFILER_TARGET} PRIVATE ${AQLPROFILE_LIB} hsa-runtime64::hsa-runtime64 c stdc++ -+ ${ROCPROFILER_TARGET} PRIVATE hsa-runtime64::hsa-runtime64 c stdc++ - dl rocprofiler::build-flags rocprofiler::memcheck) - - get_target_property(ROCPROFILER_LIBRARY_V1_NAME ${ROCPROFILER_TARGET} NAME) -@@ -325,8 +316,7 @@ target_link_options( - -Wl,--no-undefined) - target_link_libraries( - rocprofiler-v2 -- PRIVATE ${AQLPROFILE_LIB} -- hsa-runtime64::hsa-runtime64 -+ PRIVATE hsa-runtime64::hsa-runtime64 - Threads::Threads - atomic - numa -diff --git a/src/util/hsa_rsrc_factory.cpp b/src/util/hsa_rsrc_factory.cpp -index 2c47186..6b39634 100644 ---- a/src/util/hsa_rsrc_factory.cpp -+++ b/src/util/hsa_rsrc_factory.cpp -@@ -155,17 +155,6 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize - if (kern_arg_pool_ == nullptr) - CHECK_STATUS("Kern-arg memory pool is not found", HSA_STATUS_ERROR); - -- // Get AqlProfile API table -- aqlprofile_api_ = {}; --#ifdef ROCP_LD_AQLPROFILE -- status = LoadAqlProfileLib(&aqlprofile_api_); --#else -- status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, -- hsa_ven_amd_aqlprofile_VERSION_MAJOR, -- sizeof(aqlprofile_api_), &aqlprofile_api_); --#endif -- CHECK_STATUS("aqlprofile API table load failed", status); -- - // Get Loader API table - loader_api_ = {}; - status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, -diff --git a/test/util/hsa_rsrc_factory.cpp b/test/util/hsa_rsrc_factory.cpp -index 0a44d18..fab5b75 100644 ---- a/test/util/hsa_rsrc_factory.cpp -+++ b/test/util/hsa_rsrc_factory.cpp -@@ -137,17 +137,6 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize - if (cpu_pool_ == NULL) CHECK_STATUS("CPU memory pool is not found", HSA_STATUS_ERROR); - if (kern_arg_pool_ == NULL) CHECK_STATUS("Kern-arg memory pool is not found", HSA_STATUS_ERROR); - -- // Get AqlProfile API table -- aqlprofile_api_ = {0}; --#ifdef ROCP_LD_AQLPROFILE -- status = LoadAqlProfileLib(&aqlprofile_api_); --#else -- status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, -- hsa_ven_amd_aqlprofile_VERSION_MAJOR, -- sizeof(aqlprofile_api_), &aqlprofile_api_); --#endif -- CHECK_STATUS("aqlprofile API table load failed", status); -- - // Get Loader API table - loader_api_ = {0}; - status = hsa_api_.hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, -diff --git a/tests-v2/unittests/core/CMakeLists.txt b/tests-v2/unittests/core/CMakeLists.txt -index 107cb51..0f6d4bf 100644 ---- a/tests-v2/unittests/core/CMakeLists.txt -+++ b/tests-v2/unittests/core/CMakeLists.txt -@@ -235,8 +235,7 @@ set_target_properties(runCoreUnitTests PROPERTIES - INSTALL_RPATH "${CMAKE_INSTALL_PREFIX}/${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/tests") - target_link_libraries( - runCoreUnitTests -- PRIVATE ${AQLPROFILE_LIB} -- test_hsatool_library -+ PRIVATE test_hsatool_library - hsa-runtime64::hsa-runtime64 - Threads::Threads - GTest::gtest GTest::gtest_main -@@ -285,4 +284,4 @@ endif() - # for the *_FilePlugin tests - if(NOT EXISTS "${PROJECT_BINARY_DIR}/test-output") - file(MAKE_DIRECTORY "${PROJECT_BINARY_DIR}/test-output") --endif() -\ No newline at end of file -+endif() -diff --git a/tests-v2/unittests/profiler/CMakeLists.txt b/tests-v2/unittests/profiler/CMakeLists.txt -index 53180d5..0c4d4a7 100644 ---- a/tests-v2/unittests/profiler/CMakeLists.txt -+++ b/tests-v2/unittests/profiler/CMakeLists.txt -@@ -122,7 +122,7 @@ target_compile_definitions( - PRIVATE PROF_API_IMPL HIP_PROF_HIP_API_STRING=1 __HIP_PLATFORM_AMD__=1) - - target_link_libraries( -- runUnitTests PRIVATE rocprofiler-v2 ${AQLPROFILE_LIB} hsa-runtime64::hsa-runtime64 -+ runUnitTests PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64 - GTest::gtest GTest::gtest_main stdc++fs ${PCIACCESS_LIBRARIES} dw elf c dl) - - add_dependencies(tests runUnitTests) -@@ -158,4 +158,4 @@ endif() - # for the *_FilePlugin tests - if(NOT EXISTS "${PROJECT_BINARY_DIR}/test-output") - file(MAKE_DIRECTORY "${PROJECT_BINARY_DIR}/test-output") --endif() -\ No newline at end of file -+endif() diff --git a/pkgs/development/rocm-modules/6/rocr-debug-agent/default.nix b/pkgs/development/rocm-modules/6/rocr-debug-agent/default.nix index 7fbf0a04da09..2b1501dca51a 100644 --- a/pkgs/development/rocm-modules/6/rocr-debug-agent/default.nix +++ b/pkgs/development/rocm-modules/6/rocr-debug-agent/default.nix @@ -12,13 +12,13 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocr-debug-agent"; - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rocr_debug_agent"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-HYag5/E72hopDhS9EVcdyGgSvzbCMzKqLC+SIS28Y9M="; + hash = "sha256-otoxZ2NHkPDIFhvn4/nvaQ/W4LF38Nx9MZ9IYEf1DyY="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/rocrand/default.nix b/pkgs/development/rocm-modules/6/rocrand/default.nix index d58e596d97a8..f0219f66955b 100644 --- a/pkgs/development/rocm-modules/6/rocrand/default.nix +++ b/pkgs/development/rocm-modules/6/rocrand/default.nix @@ -15,7 +15,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocrand${clr.gpuArchSuffix}"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocRAND"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-rrRLPqEw39M+6dtPW8DcnQiSZNwxWNINJ1wjU098Vkk="; + hash = "sha256-pZAwmsWup0byGxJ5ADbztco8svBpIjBWITjs+OgKvQc="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/rocsolver/default.nix b/pkgs/development/rocm-modules/6/rocsolver/default.nix index f5daea4de5db..c055e5abf735 100644 --- a/pkgs/development/rocm-modules/6/rocsolver/default.nix +++ b/pkgs/development/rocm-modules/6/rocsolver/default.nix @@ -36,7 +36,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocsolver${clr.gpuArchSuffix}"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -52,7 +52,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocSOLVER"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-+sGU+0CB48iolJSyYo+xH36q5LCUp+nKtOYbguzMuhg="; + hash = "sha256-JthNY5rkrrm9bf1fgkO9HnARdX6H0kiF9EW1jMbBmj4="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/rocsparse/default.nix b/pkgs/development/rocm-modules/6/rocsparse/default.nix index ee7fd4d96f15..5b17880b2ea0 100644 --- a/pkgs/development/rocm-modules/6/rocsparse/default.nix +++ b/pkgs/development/rocm-modules/6/rocsparse/default.nix @@ -20,7 +20,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocsparse${clr.gpuArchSuffix}"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -36,7 +36,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocSPARSE"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-6Cut5rbyqKFzHaXfJZGApyY9Mj1Zq/+U8MkXgy4X4Pw="; + hash = "sha256-PrLyqHMAJYACkBMz97iBfwCaQ9Kf1IBL7IEf1IF1/m0="; }; nativeBuildInputs = [ @@ -59,7 +59,7 @@ stdenv.mkDerivation (finalAttrs: { ]; cmakeFlags = [ - "-DCMAKE_BUILD_TYPE=Release" + "-DCMAKE_CXX_COMPILER=amdclang++" # Manually define CMAKE_INSTALL_ # See: https://github.com/NixOS/nixpkgs/pull/197838 "-DCMAKE_INSTALL_BINDIR=bin" diff --git a/pkgs/development/rocm-modules/6/rocthrust/default.nix b/pkgs/development/rocm-modules/6/rocthrust/default.nix index 0a5254f60b95..bf334571c68f 100644 --- a/pkgs/development/rocm-modules/6/rocthrust/default.nix +++ b/pkgs/development/rocm-modules/6/rocthrust/default.nix @@ -15,7 +15,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocthrust"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -31,7 +31,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocThrust"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-c1+hqP/LipaQ2/lPJo79YBd9H0n0Y7yHkxe0/INE14s="; + hash = "sha256-IfMBVISClD1dk7FnAakP2GIpyZFrCnAloFRTaNdSKyw="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/roctracer/default.nix b/pkgs/development/rocm-modules/6/roctracer/default.nix index e45b340c9a63..f560e4a235ec 100644 --- a/pkgs/development/rocm-modules/6/roctracer/default.nix +++ b/pkgs/development/rocm-modules/6/roctracer/default.nix @@ -19,7 +19,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "roctracer"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -35,7 +35,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "roctracer"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-GhnF7rqNLQLLB7nzIp0xNqyqBOwj9ZJ+hzzj1EAaXWU="; + hash = "sha256-Dwk5cBZLysmsVA2kwpQM0FQt2KXOGcaZcAw/d8VUaXw="; }; nativeBuildInputs = [ diff --git a/pkgs/development/rocm-modules/6/rocwmma/default.nix b/pkgs/development/rocm-modules/6/rocwmma/default.nix index 718f45b18519..bfed05aaf72b 100644 --- a/pkgs/development/rocm-modules/6/rocwmma/default.nix +++ b/pkgs/development/rocm-modules/6/rocwmma/default.nix @@ -19,7 +19,7 @@ stdenv.mkDerivation (finalAttrs: { pname = "rocwmma"; - version = "6.3.3"; + version = "6.4.3"; outputs = [ "out" @@ -38,7 +38,7 @@ stdenv.mkDerivation (finalAttrs: { owner = "ROCm"; repo = "rocWMMA"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-kih3hn6QhcMmyj9n8f8eO+RIgKQgWKIuzg8fb0eoRPE="; + hash = "sha256-fjyxMrzt74rE7Gf4v4WawYltuw1fvahwZUpauMIE3qc="; }; patches = lib.optionals (buildTests || buildBenchmarks) [ @@ -61,9 +61,6 @@ stdenv.mkDerivation (finalAttrs: { ]; cmakeFlags = [ - "-DOpenMP_C_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_CXX_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_omp_LIBRARY=${openmp}/lib" "-DROCWMMA_BUILD_TESTS=${if buildTests || buildBenchmarks then "ON" else "OFF"}" "-DROCWMMA_BUILD_SAMPLES=${if buildSamples then "ON" else "OFF"}" # Manually define CMAKE_INSTALL_ diff --git a/pkgs/development/rocm-modules/6/rpp/default.nix b/pkgs/development/rocm-modules/6/rpp/default.nix index 4185c91e079f..41cbf54089e0 100644 --- a/pkgs/development/rocm-modules/6/rpp/default.nix +++ b/pkgs/development/rocm-modules/6/rpp/default.nix @@ -29,13 +29,13 @@ stdenv.mkDerivation (finalAttrs: { "cpu" ); - version = "6.3.3"; + version = "6.4.3"; src = fetchFromGitHub { owner = "ROCm"; repo = "rpp"; rev = "rocm-${finalAttrs.version}"; - hash = "sha256-METwagek17/DdZGaOTQqvyU6xGt7OBMLHk4YM4KmgtA="; + hash = "sha256-rccVjSrOVIe4ZDtloCoCCI3u9UIcUqdirHIzS7ffAas="; }; nativeBuildInputs = [ @@ -54,12 +54,7 @@ stdenv.mkDerivation (finalAttrs: { boost ]; - CFLAGS = "-I${openmp.dev}/include"; - CXXFLAGS = "-I${openmp.dev}/include"; cmakeFlags = [ - "-DOpenMP_C_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_CXX_INCLUDE_DIR=${openmp.dev}/include" - "-DOpenMP_omp_LIBRARY=${openmp}/lib" "-DROCM_PATH=${clr}" ] ++ lib.optionals (gpuTargets != [ ]) [ diff --git a/pkgs/development/rocm-modules/6/tensile/default.nix b/pkgs/development/rocm-modules/6/tensile/default.nix index 9446452692f2..d54c5dd92365 100644 --- a/pkgs/development/rocm-modules/6/tensile/default.nix +++ b/pkgs/development/rocm-modules/6/tensile/default.nix @@ -6,55 +6,43 @@ buildPythonPackage, pytestCheckHook, setuptools, + distro, pyyaml, msgpack, - simplejson, - ujson, - orjson, pandas, joblib, filelock, clr, rich, - isTensileLite ? false, }: buildPythonPackage rec { - pname = if isTensileLite then "tensilelite" else "tensile"; - # Using a specific commit which has code object compression support from after the 6.3 release + pname = "tensile"; + # Using a specific commit which has compression support from after the 6.4 release # Without compression packages are too large for hydra - version = "6.3-unstable-2024-12-10"; + version = "6.4-unstable-2025-06-12"; format = "pyproject"; src = fetchFromGitHub { owner = "ROCm"; repo = "Tensile"; - rev = "1752af518190500891a865379a4569b8abf6ba01"; - hash = "sha256-Wvz4PVs//3Ox7ykZHpjPzOVwlyATyc+MmVVenfTzWK4="; + rev = "1ce87a9fe73610ffb962082f0a882360cd39b103"; + hash = "sha256-qIuoIbmridy1HQVV10qPTzbccuxNJPsOvePaQQnClZc="; }; # TODO: It should be possible to run asm caps test ONCE for all supported arches # We currently disable the test because it's slow and runs each time tensile launches - - postPatch = - lib.optionalString (!isTensileLite) '' - if grep -F .SafeLoader Tensile/LibraryIO.py; then - substituteInPlace Tensile/LibraryIO.py \ - --replace-fail "yaml.SafeLoader" "yaml.CSafeLoader" - fi - # See TODO above about asm caps test - substituteInPlace Tensile/Common.py \ - --replace-fail 'if globalParameters["AssemblerPath"] is not None:' "if False:" - '' - + '' - # Add an assert that the fallback 9,0,0 is supported before setting the kernel to it - # If it's not detected as supported we have an issue with compiler paths or the compiler is broken - # and it's better to stop immediately - substituteInPlace Tensile/KernelWriter.py \ - --replace-fail '= (9,0,0)' '= (9,0,0);assert(globalParameters["AsmCaps"][(9,0,0)]["SupportedISA"])' - find . -type f -iname "*.sh" -exec chmod +x {} \; - patchShebangs Tensile - ''; + postPatch = '' + substituteInPlace Tensile/Common.py \ + --replace-fail 'if globalParameters["AssemblerPath"] is not None:' "if False:" + # Add an assert that the fallback 9,0,0 is supported before setting the kernel to it + # If it's not detected as supported we have an issue with compiler paths or the compiler is broken + # and it's better to stop immediately + substituteInPlace Tensile/KernelWriter.py \ + --replace-fail '= (9,0,0)' '= (9,0,0);assert(globalParameters["AsmCaps"][(9,0,0)]["SupportedISA"])' + find . -type f -iname "*.sh" -exec chmod +x {} \; + patchShebangs Tensile + ''; buildInputs = [ setuptools ]; @@ -63,27 +51,19 @@ buildPythonPackage rec { msgpack pandas joblib - ] - ++ lib.optionals (!isTensileLite) [ + distro rich - ] - ++ lib.optionals isTensileLite [ - simplejson - ujson - orjson ]; - patches = - lib.optional (!isTensileLite) ./tensile-solutionstructs-perf-fix.diff - ++ lib.optional (!isTensileLite) ./tensile-create-library-dont-copy-twice.diff - ++ lib.optional (!isTensileLite) (fetchpatch { + patches = [ + ./tensile-solutionstructs-perf-fix.diff + ./tensile-create-library-dont-copy-twice.diff + (fetchpatch { # [PATCH] Extend Tensile HIP ISA compatibility sha256 = "sha256-d+fVf/vz+sxGqJ96vuxe0jRMgbC5K6j5FQ5SJ1e3Sl8="; url = "https://github.com/GZGavinZhao/Tensile/commit/855cb15839849addb0816a6dde45772034a3e41f.patch"; }) - ++ lib.optional isTensileLite ./tensilelite-create-library-dont-copy-twice.diff - ++ lib.optional isTensileLite ./tensilelite-gen_assembly-venv-err-handling.diff - ++ lib.optional isTensileLite ./tensilelite-compression.diff; + ]; doCheck = false; # Too many errors, not sure how to set this up properly diff --git a/pkgs/development/rocm-modules/6/tensile/tensilelite-compression.diff b/pkgs/development/rocm-modules/6/tensile/tensilelite-compression.diff deleted file mode 100644 index bfc0146efe3d..000000000000 --- a/pkgs/development/rocm-modules/6/tensile/tensilelite-compression.diff +++ /dev/null @@ -1,345 +0,0 @@ -diff --git a/Tensile/TensileCreateLibrary.py b/Tensile/TensileCreateLibrary.py -index b8cea84558..1bc24bd1dd 100644 ---- a/Tensile/TensileCreateLibrary.py -+++ b/Tensile/TensileCreateLibrary.py -@@ -41,6 +41,7 @@ - from .SolutionLibrary import MasterSolutionLibrary - from .SolutionStructs import Solution - from .CustomYamlLoader import load_logic_gfx_arch -+from .Utilities.Profile import profile - - import argparse - import collections -@@ -1233,7 +1234,7 @@ def validateLibrary(masterLibraries: MasterSolutionLibrary, - ################################################################################ - # Tensile Create Library - ################################################################################ --@timing -+@profile - def TensileCreateLibrary(): - print1("") - print1(HR) -@@ -1558,7 +1559,6 @@ def param(key, value): - - print1("# Check if generated files exists.") - -- @timing - def checkFileExistence(files): - for filePath in files: - if not os.path.exists(filePath): -diff --git a/Tensile/Utilities/Profile.py b/Tensile/Utilities/Profile.py -new file mode 100644 -index 0000000000..cc3c7eb44c ---- /dev/null -+++ b/Tensile/Utilities/Profile.py -@@ -0,0 +1,77 @@ -+################################################################################ -+# -+# Copyright (C) 2016-2024 Advanced Micro Devices, Inc. All rights reserved. -+# -+# Permission is hereby granted, free of charge, to any person obtaining a copy -+# of this software and associated documentation files (the "Software"), to deal -+# in the Software without restriction, including without limitation the rights -+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -+# copies of the Software, and to permit persons to whom the Software is -+# furnished to do so, subject to the following conditions: -+# -+# The above copyright notice and this permission notice shall be included in -+# all copies or substantial portions of the Software. -+# -+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -+# SOFTWARE. -+# -+################################################################################ -+ -+import cProfile -+import pstats -+import os -+ -+from pathlib import Path -+from datetime import datetime, timezone -+from typing import Callable, Tuple -+ -+PROFILE_ENV_VAR = "TENSILE_PROFILE" -+ -+def profile(func: Callable) -> Callable: -+ """Profiling decorator. -+ -+ Add ``@profile`` to mark a function for profiling; set the environment variable -+ TENSILE_PROFILE=ON to enable profiling decorated functions. -+ """ -+ if not envVariableIsSet(PROFILE_ENV_VAR): -+ return func -+ def wrapper(*args, **kwargs): -+ path, filename = initProfileArtifacts(func.__name__) -+ -+ prof = cProfile.Profile() -+ output = prof.runcall(func, *args, **kwargs) -+ result = pstats.Stats(prof) -+ result.sort_stats(pstats.SortKey.TIME) -+ result.dump_stats(path/filename) -+ -+ return output -+ return wrapper -+ -+def envVariableIsSet(varName: str) -> bool: -+ """Checks if the provided environment variable is set to "ON", "TRUE", or "1" -+ Args: -+ varName: Environment variable name. -+ Returns: -+ True if the environment variable is set, otherwise False. -+ """ -+ value = os.environ.get(varName, "").upper() -+ return True if value in ["ON", "TRUE", "1"] else False -+ -+def initProfileArtifacts(funcName: str) -> Tuple[Path, str]: -+ """Initializes filenames and paths for profiling artifacts based on the current datetime -+ Args: -+ funcName: The name of the function being profiled, nominally passed via func.__name__ -+ Returns: -+ A tuple (path, filename) where the path is the artifact directory and filename is -+ a .prof file with the profiling results. -+ """ -+ dt = datetime.now(timezone.utc) -+ filename = f"{funcName}-{dt.strftime('%Y-%m-%dT%H-%M-%SZ')}.prof" -+ path = Path().cwd()/f"profiling-results-{dt.strftime('%Y-%m-%d')}" -+ path.mkdir(exist_ok=True) -+ return path, filename - -diff --git a/Tensile/TensileCreateLibrary.py b/Tensile/TensileCreateLibrary.py -index e62b0072df..2c843ba936 100644 ---- a/Tensile/TensileCreateLibrary.py -+++ b/Tensile/TensileCreateLibrary.py -@@ -56,7 +56,7 @@ - import sys - from timeit import default_timer as timer - from pathlib import Path --from typing import Sequence, List -+from typing import Sequence, List, Union - - def timing(func): - def wrapper(*args, **kwargs): -@@ -90,87 +90,142 @@ def processKernelSource(kernel, kernelWriterAssembly, ti): - - return (err, src, header, kernelName, filename) - --def getAssemblyCodeObjectFiles(kernels, kernelWriterAssembly, outputPath): -- destDir = ensurePath(os.path.join(outputPath, 'library')) -- asmDir = kernelWriterAssembly.getAssemblyDirectory() -- archs = collections.defaultdict(list) -+def linkIntoCodeObject( -+ objFiles: List[str], coPathDest: Union[Path, str], kernelWriterAssembly: KernelWriterAssembly -+): -+ """Links object files into a code object file. - -- for k in filter(lambda k: k['KernelLanguage'] == 'Assembly', kernels): -- archs[tuple(k['ISA'])].append(k) -+ Args: -+ objectFiles: A list of object files to be linked. -+ coPathDest: The destination path for the code object file. -+ kernelWriterAssembly: An instance of KernelWriterAssembly to get link arguments. - -- coFiles = [] -+ Raises: -+ RuntimeError: If linker invocation fails. -+ """ -+ if os.name == "nt": -+ # On Windows, the objectFiles list command line (including spaces) -+ # exceeds the limit of 8191 characters, so using response file -+ -+ responseFile = os.path.join('/tmp', 'clangArgs.txt') -+ with open(responseFile, 'wt') as file: -+ file.write(" ".join(objFiles)) -+ file.flush() -+ -+ args = [globalParameters['AssemblerPath'], '-target', 'amdgcn-amd-amdhsa', '-o', coFileRaw, '@clangArgs.txt'] -+ subprocess.check_call(args, cwd=asmDir) -+ else: -+ numObjFiles = len(objFiles) -+ maxObjFiles = 10000 -+ -+ if numObjFiles > maxObjFiles: -+ batchedObjFiles = [objFiles[i:i+maxObjFiles] for i in range(0, numObjFiles, maxObjFiles)] -+ batchSize = int(math.ceil(numObjFiles / maxObjFiles)) -+ -+ newObjFiles = [str(coPathDest) + "." + str(i) for i in range(0, batchSize)] -+ newObjFilesOutput = [] -+ -+ for batch, filename in zip(batchedObjFiles, newObjFiles): -+ if len(batch) > 1: -+ args = [globalParameters["ROCmLdPath"], "-r"] + batch + [ "-o", filename] -+ print2(f"Linking object files into fewer object files: {' '.join(args)}") -+ subprocess.check_call(args) -+ newObjFilesOutput.append(filename) -+ else: -+ newObjFilesOutput.append(batchedObjFiles[0]) -+ -+ args = kernelWriterAssembly.getLinkCodeObjectArgs(newObjFilesOutput, str(coPathDest)) -+ print2(f"Linking object files into code object: {' '.join(args)}") -+ subprocess.check_call(args) -+ else: -+ args = kernelWriterAssembly.getLinkCodeObjectArgs(objFiles, str(coPathDest)) -+ print2(f"Linking object files into code object: {' '.join(args)}") -+ subprocess.check_call(args) -+ -+ -+def compressCodeObject( -+ coPathSrc: Union[Path, str], coPathDest: Union[Path, str], gfx: str, bundler: str -+): -+ """Compresses a code object file using the provided bundler. -+ -+ Args: -+ coPathSrc: The source path of the code object file to be compressed. -+ coPathDest: The destination path for the compressed code object file. -+ gfx: The target GPU architecture. -+ bundler: The path to the Clang Offload Bundler executable. -+ -+ Raises: -+ RuntimeError: If compressing the code object file fails. -+ """ -+ args = [ -+ bundler, -+ "--compress", -+ "--type=o", -+ "--bundle-align=4096", -+ f"--targets=host-x86_64-unknown-linux,hipv4-amdgcn-amd-amdhsa--{gfx}", -+ "--input=/dev/null", -+ f"--input={str(coPathSrc)}", -+ f"--output={str(coPathDest)}", -+ ] -+ -+ print2(f"Bundling/compressing code objects: {' '.join(args)}") -+ try: -+ out = subprocess.check_output(args, stderr=subprocess.STDOUT) -+ print2(f"Output: {out}") -+ except subprocess.CalledProcessError as err: -+ raise RuntimeError( -+ f"Error compressing code object via bundling: {err.output}\nFailed command: {' '.join(args)}" -+ ) -+ -+def buildAssemblyCodeObjectFiles(kernels, kernelWriterAssembly, outputPath): -+ -+ isAsm = lambda k: k["KernelLanguage"] == "Assembly" -+ -+ extObj = ".o" -+ extCo = ".co" -+ extCoRaw = ".co.raw" - -- for arch, archKernels in archs.items(): -+ destDir = Path(ensurePath(os.path.join(outputPath, 'library'))) -+ asmDir = Path(kernelWriterAssembly.getAssemblyDirectory()) -+ -+ archKernelMap = collections.defaultdict(list) -+ for k in filter(isAsm, kernels): -+ archKernelMap[tuple(k['ISA'])].append(k) -+ -+ coFiles = [] -+ for arch, archKernels in archKernelMap.items(): - if len(archKernels) == 0: - continue - -- archName = getGfxName(arch) -+ gfx = getGfxName(arch) - - if globalParameters["MergeFiles"] or globalParameters["NumMergedFiles"] > 1 or globalParameters["LazyLibraryLoading"]: -- objectFiles = [kernelWriterAssembly.getKernelFileBase(k) + '.o' for k in archKernels if 'codeObjectFile' not in k] -+ objectFiles = [str(asmDir / (kernelWriterAssembly.getKernelFileBase(k) + extObj)) for k in archKernels if 'codeObjectFile' not in k] - -- #Group kernels from placeholder libraries - coFileMap = collections.defaultdict(list) -+ - if len(objectFiles): -- coFileMap[os.path.join(destDir, "TensileLibrary_"+archName+".co")] = objectFiles -+ coFileMap[asmDir / ("TensileLibrary_"+ gfx + extCoRaw)] = objectFiles - - for kernel in archKernels: - coName = kernel.get("codeObjectFile", None) - if coName: -- coFileMap[os.path.join(destDir, coName+".co")] += [kernelWriterAssembly.getKernelFileBase(kernel) + '.o'] -+ coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (kernelWriterAssembly.getKernelFileBase(kernel) + extObj))) - -- for coFile, objectFiles in coFileMap.items(): -- if os.name == "nt": -- # On Windows, the objectFiles list command line (including spaces) -- # exceeds the limit of 8191 characters, so using response file -+ for coFileRaw, objFiles in coFileMap.items(): - -- responseArgs = objectFiles -- responseFile = os.path.join(asmDir, 'clangArgs.txt') -- with open(responseFile, 'wt') as file: -- file.write( " ".join(responseArgs) ) -- file.flush() -- -- args = [globalParameters['AssemblerPath'], '-target', 'amdgcn-amd-amdhsa', '-o', coFile, '@clangArgs.txt'] -- subprocess.check_call(args, cwd=asmDir) -- else: -- numOfObjectFiles = len(objectFiles) -- splitFiles = 10000 -- if numOfObjectFiles > splitFiles: -- slicedObjectFilesList = [objectFiles[x:x+splitFiles] for x in range(0, numOfObjectFiles, splitFiles)] -- objectFileBasename = os.path.split(coFile)[-1].split('.')[0] -- numOfOneSliceOfObjectFiles = int(math.ceil(numOfObjectFiles / splitFiles)) -- newObjectFiles = [ objectFileBasename + "_" + str(i) + ".o" for i in range(0, numOfOneSliceOfObjectFiles)] -- newObjectFilesOutput = [] -- for slicedObjectFiles, objectFile in zip(slicedObjectFilesList, newObjectFiles): -- if len(slicedObjectFiles) > 1: -- args = [globalParameters["ROCmLdPath"], "-r"] + slicedObjectFiles + [ "-o", objectFile ] -- if globalParameters["PrintCodeCommands"]: -- print(asmDir) -- print(' '.join(args)) -- subprocess.check_call(args, cwd=asmDir) -- newObjectFilesOutput.append(objectFile) -- else: -- newObjectFilesOutput.append(slicedObjectFiles[0]) -- args = kernelWriterAssembly.getLinkCodeObjectArgs(newObjectFilesOutput, coFile) -- if globalParameters["PrintCodeCommands"]: -- print(asmDir) -- print(' '.join(args)) -- subprocess.check_call(args, cwd=asmDir) -- else: -- args = kernelWriterAssembly.getLinkCodeObjectArgs(objectFiles, coFile) -- if globalParameters["PrintCodeCommands"]: -- print(asmDir) -- print(' '.join(args)) -- subprocess.check_call(args, cwd=asmDir) -+ linkIntoCodeObject(objFiles, coFileRaw, kernelWriterAssembly) -+ coFile = destDir / coFileRaw.name.replace(extCoRaw, extCo) -+ compressCodeObject(coFileRaw, coFile, gfx, globalParameters["ClangOffloadBundlerPath"]) - - coFiles.append(coFile) - else: - # no mergefiles - def newCoFileName(kName): - if globalParameters["PackageLibrary"]: -- return os.path.join(destDir, archName, kName + '.co') -+ return os.path.join(destDir, gfx, kName + '.co') - else: -- return os.path.join(destDir, kName + '_' + archName + '.co') -+ return os.path.join(destDir, kName + '_' + gfx + '.co') - - def orgCoFileName(kName): - return os.path.join(asmDir, kName + '.co') -@@ -179,6 +234,8 @@ def orgCoFileName(kName): - map(lambda k: kernelWriterAssembly.getKernelFileBase(k), archKernels)), "Copying code objects"): - shutil.copyfile(src, dst) - coFiles.append(dst) -+ printWarning("Code object files are not compressed in `--no-merge-files` build mode.") -+ - return coFiles - - def which(p): -@@ -645,7 +702,7 @@ def success(kernel): - - if not globalParameters["GenerateSourcesAndExit"]: - codeObjectFiles += buildSourceCodeObjectFiles(CxxCompiler, kernelFiles, outputPath) -- codeObjectFiles += getAssemblyCodeObjectFiles(kernelsToBuild, kernelWriterAssembly, outputPath) -+ codeObjectFiles += buildAssemblyCodeObjectFiles(kernelsToBuild, kernelWriterAssembly, outputPath) - - Common.popWorkingPath() # build_tmp - Common.popWorkingPath() # workingDir - diff --git a/pkgs/development/rocm-modules/6/tensile/tensilelite-create-library-dont-copy-twice.diff b/pkgs/development/rocm-modules/6/tensile/tensilelite-create-library-dont-copy-twice.diff deleted file mode 100644 index c630803c191f..000000000000 --- a/pkgs/development/rocm-modules/6/tensile/tensilelite-create-library-dont-copy-twice.diff +++ /dev/null @@ -1,37 +0,0 @@ -diff --git a/Tensile/TensileCreateLibrary.py b/Tensile/TensileCreateLibrary.py -index 2b9da394..b001fa7c 100644 ---- a/Tensile/TensileCreateLibrary.py -+++ b/Tensile/TensileCreateLibrary.py -@@ -808,10 +808,13 @@ def copyStaticFiles(outputPath=None): - "ReductionTemplate.h", - "memory_gfx.h" ] - -+ import filecmp - for fileName in libraryStaticFiles: -- # copy file -- shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ -- outputPath ) -+ src = os.path.join(globalParameters["SourcePath"], fileName) -+ dst = os.path.join(outputPath, os.path.basename(src)) -+ # no need to copy twice if it has already been copied -+ if not os.path.isfile(dst) or not filecmp.cmp(src, dst): -+ shutil.copyfile(src, dst) - - return libraryStaticFiles - -@@ -1417,9 +1420,13 @@ def TensileCreateLibrary(): - writeCMake(outputPath, solutionFiles, sourceKernelFiles, staticFiles, masterLibraries) - - # Make sure to copy the library static files. -+ import filecmp - for fileName in staticFiles: -- shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ -- outputPath ) -+ src = os.path.join(globalParameters["SourcePath"], fileName) -+ dst = os.path.join(outputPath, os.path.basename(src)) -+ # no need to copy twice if it has already been copied -+ if not os.path.isfile(dst) or not filecmp.cmp(src, dst): -+ shutil.copyfile(src, dst) - - # write solutions and kernels - codeObjectFiles = writeSolutionsAndKernels(outputPath, CxxCompiler, None, solutions, diff --git a/pkgs/development/rocm-modules/6/tensile/tensilelite-gen_assembly-venv-err-handling.diff b/pkgs/development/rocm-modules/6/tensile/tensilelite-gen_assembly-venv-err-handling.diff deleted file mode 100644 index 0667599a5839..000000000000 --- a/pkgs/development/rocm-modules/6/tensile/tensilelite-gen_assembly-venv-err-handling.diff +++ /dev/null @@ -1,36 +0,0 @@ -diff --git a/Tensile/Ops/gen_assembly.sh b/Tensile/Ops/gen_assembly.sh -index 0b21b6c6..609f1dd1 100755 ---- a/Tensile/Ops/gen_assembly.sh -+++ b/Tensile/Ops/gen_assembly.sh -@@ -23,6 +23,8 @@ - # - ################################################################################ - -+set -x -+ - archStr=$1 - dst=$2 - venv=$3 -@@ -35,7 +37,13 @@ fi - - toolchain=${rocm_path}/llvm/bin/clang++ - --. ${venv}/bin/activate -+if ! [ -z ${TENSILE_GEN_ASSEMBLY_TOOLCHAIN+x} ]; then -+ toolchain="${TENSILE_GEN_ASSEMBLY_TOOLCHAIN}" -+fi -+ -+if [ -f ${venv}/bin/activate ]; then -+ . ${venv}/bin/activate -+fi - - IFS=';' read -r -a archs <<< "$archStr" - -@@ -77,4 +85,6 @@ for arch in "${archs[@]}"; do - python3 ./ExtOpCreateLibrary.py --src=$dst --co=$dst/extop_$arch.co --output=$dst --arch=$arch - done - --deactivate -+if [ -f ${venv}/bin/activate ]; then -+ deactivate -+fi