Repository navigation
rocmPackages: 7.2.3 -> 10.1.0 - #573885
randomizedcoder wants to merge 36 commits into
Conversation
Introduce a per-architecture Tensile "shard" split for rocBLAS: base.nix is the raw single-/multi-arch builder (the former default.nix), and default.nix now assembles one joined output from per-arch shards. For a multi-arch build the host librocblas.so + headers + cmake are built once for the head arch, and every other arch's self-contained lib/rocblas/library/<gfx>/ Tensile directory is built as a separate minimal shard and merged in postInstall. rocBLAS builds Tensile with SEPARATE_ARCHITECTURES + LAZY_LIBRARY_LOADING, so in a shared build the kernels are not linked into librocblas.so; each per-arch dir is loaded by path at runtime, making the join a plain per-arch-dir merge with no relink. Single-arch scopes (e.g. rocmPackages.gfx906.*) have no tail archs and return the base build verbatim, so their derivation stays byte-identical to before the split. The merge strips the shard store path that Tensile bakes into a per-arch code object's ELF metadata (remove-references-to). That reference is not load-bearing -- the unsplit multi-arch build self-references $out in the same spot and runs fine -- and disallowedReferences guards against a regression. Validated on a gfx1102 (Radeon Pro W7500): a [gfx906-head, gfx1102-shard] join runs a correct rocBLAS SGEMM on the gfx1102 hardware, loading the merged gfx1102 Tensile directory from the joined output. Assisted-by: Claude Code (Claude Opus 4.8, Claude Fable 5.1)
LLVM 23 removed the legacy standalone OpenMP build (cmake <llvm>/openmp now FATAL_ERRORs). Build OpenMP through the runtimes entry point (cmake <llvm>/runtimes -DLLVM_ENABLE_RUNTIMES=openmp) for release_version >= 23, keeping the standalone layout for older LLVMs. Needed by the ROCm 10 (LLVM 24) toolchain. The runtimes-only attributes are added with lib.optionalAttrs so that the derivations of older LLVMs are unchanged (an unconditional empty preConfigure would re-key llvmPackages.openmp on every platform). Assisted-by: Claude Code (Claude Opus 4.8, Claude Fable 5.1)
There was a problem hiding this comment.
Thanks for contributing to the documentation.
Follow the documentation style guide, especially:
- Show, don’t tell. Start with a minimal working example. Explain after the code.
- Avoid meta-commentary. Don’t write "This section explains how to…".
- Use present tense, imperative mood, and active voice. For example: "Run the command," not "The user should run the following command."
- Avoid vague references like "this" when the referent is unclear. Write "This command creates a folder," not "This will create a folder."
- Be direct. Avoid "should", "might", "typically", and "usually".
- Cut filler words like "simply", "just", "easily", and "basically". Use "to" instead of "in order to".
For larger changes, such as adding or removing whole sections, ask the NixOS documentation team for review.
|
|
||
| Same fix as https://github.com/ROCm/amdsmi/pull/165, redone for ROCm 10, | ||
| where the bundled header moved from include/amd_smi/impl/amdgpu_drm.h to | ||
| include/libdrm/drm_mode.h. |
There was a problem hiding this comment.
https://github.com/ROCm/rocm-systems/blob/therock-10.1/projects/rocm-smi-lib/include/rocm_smi/amdgpu_drm.h#L1300 has the PR referenced struct.
https://github.com/ROCm/rocm-systems/blob/therock-10.1/projects/amdsmi/include/libdrm/drm_mode.h#L828 has a different struct from the referenced PR with the same name.
~all of this explanation is just a little off. Validate these patches changes are correct and fix the explanation.
There was a problem hiding this comment.
The old patch may still be applicable but rewritten to the updated rcom_smi file. That would be for the rocm-smi package and is outside this pr.
| # into hardware.graphics.extraPackages, so the vendor file would collide | ||
| # with the one $icd provides below. Remove them; rmdir fails loudly if a | ||
| # future release adds anything else under $out/etc. | ||
| rm $out/etc/OpenCL/vendors/amdocl64.icd $out/etc/ld.so.conf.d/10-rocm-opencl.conf |
There was a problem hiding this comment.
To facilitate maintenance update the comment to point at the upstream issue and drop most of the slop. opencl cmake code that introduces the files: https://github.com/ROCm/rocm-systems/blob/therock-10.1/projects/clr/opencl/CMakeLists.txt#L124-L143. The change was introduced more that a year ago.
There was a problem hiding this comment.
Extracted this change to #574457.
As a note: the comment added here is redundant with the commit message, and is too much about the change rather than the content. The purpose is also fairly apparently without the comment.
AMDMIGraphX rocm-10.1 requires ROCm/rocmlirTriton@release/rocm-rel-10.1 instead of ROCm/rocMLIR. Taken from NixOS#573885: pybind11, host target in LLVM_TARGETS_TO_BUILD, MLIR_INCLUDE_TESTS off and a writable HOME for the bundled Triton.
clang 24 forwards NIX_LDFLAGS to the HIP device linker, so each part's compressed fatbin carries -rpath $out/lib and the final CK fails disallowedReferences. Taken from NixOS#573885.
|
Thanks @anund for the careful review and all the time you put into it, and thanks @LunNova and @amarshall for the feedback too. You were right that this was too much for one PR. I've split it up as you suggested:
Each PR has a table showing where it fits in the series, and your review comments are addressed in the PR that touches that code. It took a few more days than I expected, mostly compile time: a full rebuild of the ROCm set and its main consumers is about 200 derivations and takes around two days on one 24-core machine. The final head was rebuilt from the toolchain up. That includes torch for Python 3.13 and 3.14, vllm, ollama, llama-cpp, blender, onnxruntime and magma. It's been tested on two different AMD GPUs, a Radeon Pro W7500 (gfx1102, RDNA3) and an MI50 (gfx906, Vega 20). Both pass the same checks: rocminfo and amd-smi, a hipcc kernel, torch matmul and an MIOpen convolution, Thanks also for the pointer to ROCm/rocm-systems#7789. It's the same OMPT deadlock. Until it lands, the series carries a small rocprofiler-sdk patch instead of the torch workaround. I'll close this PR now that the series is open. |
|
If you're going to launder my work you could be a little more up front about it. |
|
@anund - Sorry, I certainly mean no offense. I'm trying to address your feedback as best I can. I've broken the big PR up as you suggested into lots of smaller PRs. If that's not the direction you want to go, then I'd be very happy to collaborate on an alternative strategy. My objective is to get to ROCM 10.1, because it has a lot of bug fixes, including multiple ones I need |
Description of changes
This PR brings
rocmPackagesfrom ROCm 7.2.3 to ROCm 10.1.0, AMD's current release. It is a long way forward in one step: AMD skipped the 8.x and 9.x series, moved its sources into the TheRock monorepos and switched to an LLVM 24 based compiler, and the nixpkgs ROCm set has been held at 7.2.3 through all of it. At 10.1.0 the full ROCm library matrix (default 16-arch config) and the whole consumer fleet (torch, vllm, ollama, llama-cpp, magma, blender, onnxruntime, …) build green, and llama-cpp, ollama and torch have been run on real RDNA3 and CDNA hardware. The two things reviewers will want to weigh (the rebuild scope, and a ROCm LLVM 24 built on the releasedllvmPackages_23base because nixpkgs has nollvmPackages_24yet) are set out further down.ROCm 10.x is built from the TheRock monorepos (
rocm-systems,rocm-libraries,llvm-project) at the two-componenttherock-10.1tags. This is an in-place bump of the single ROCm set (as with every prior bump), so all consumers follow it andrelease-26.05remains the compatibility channel at ROCm 7.2.x.Highlights
rocmLlvmVersion = "24.0.0-rocm", built from ROCm's own LLVM source on the releasedllvmPackages_23base; the base only supplies the shared packaging and the bootstrap compiler).llvmPackages.openmpbuilds via theruntimesentry point for LLVM >= 23 (the standalone build was removed upstream).rev = "therock-${lib.versions.majorMinor version}"withversion = "10.1.0". hipfort uses its own three-componentrocm-10.1.0tag; rocprof-trace-decoder (upstream version 0.2.2) is fetched fromrocm-systemsattherock-10.1explicitly, since its tag cannot be derived from its version; rocmlir tracks AMDMIGraphX'srocm-10.1rocmlirTriton commit.rocm-systems/rocm-librariesvia sparseprojects/<name>+shared.rocblas_base+ join), so adding or changing one arch rebuilds only that arch's kernels.libamd_smi.so.27, libnl/libmnl); MIOpen 3.6.0 (OpenCL backend removed, C++20,MIOPEN_STANDALONE_BUILD=OFF, hipconv disabled, see below); hipBLASLt 1.4.1 (shared/stinkytofu,amd_smi, nanobind,CMAKE_SUPPRESS_REGENERATION); rocprofiler-sdk patch stack dropped as upstreamed; composable_kernel conv-target restructure.-Lfor the target-id probe; composable_kernel scrubs the device-object self-rpath clang 24 bakes in; amdsmi system-python sitelib configure guard; clang-linker-wrapper needs-lzstdon the default link path.import torchhang); ollama (ROCm >= 8 maps onto the upstreamrocm_v7_2backend label); blender HIP soname lookup; opensplat/torchcodec libdrm. whisper-cpp needs no change since it now takes its backends from llama-cpp.GPU support note
composable_kernel, the MIOpen CK solvers, and vllm no longer support gfx900/gfx906 on ROCm 10. rocBLAS/hipBLAS-only consumers (llama-cpp, ollama) still build and run for those targets. Documented in the 26.11 release notes.
MIOpen is built with
-DMIOPEN_USE_HIPCONV=OFF. hipconv serves only cdna3/4/5 (gfx942/gfx950+) and its cdna4direct_wgradkernels do not compile under clang 24 (constexpr values used in device code rejected as host references). I have no cdna3/4 hardware to validate a source patch, so the subsystem is disabled;conv_hipconv.cppself-guards on the macro so the path is cleanly absent. Restoring it is a follow-up.Known follow-ups
rocm_v7_2label remap.rocm-toolchainis a symlinkJoin of clang with glibc, so with-fopenmpthe clang driver's own-I <prefix>/../include(added to findomp.h) turns the libc headers into user headers; with glibc 2.44 that makes clang 24 rejectmm_malloc.h'sposix_memalignredeclaration. rpp carries a one-line-isystemworkaround; a toolchain-level fix belongs in a follow-up because it re-keys the whole set.Testing
Builds (x86_64-linux, default full 16-arch config): on the current head (the
llvmPackages_23base), the whole ROCm set plus the consumer fleet builds green: all 152 derivations ofrocmPackages.rocm-tests.tests.rocmPackagesDerivations(everyrocmPackagesattribute, full matrix, including rocBLAS and its 15 Tensile shards, hipBLASLt, composable_kernel, MIOpen, MIGraphX, rocSOLVER/hipSOLVER, hipSPARSELt, rccl, rocprofiler-sdk, rdc, rpp, mivisionx) and torch 2.13.0 (py3.13 + py3.14), vllm 0.29.0, ollama 0.40.2, llama-cpp, onnxruntime 1.27.1, magma, aotriton 0.12b, amd-aiter, bitsandbytes, torchcodec, blender, whisper-cpp, stable-diffusion-cpp, adaptivecpp, btop-rocm, zluda, opensplat, ctranslate2. The built toolchain reportsHIP version: 7.16.0andclang version 24.0.0 (https://github.com/ROCm/llvm-project/tree/therock-10.1). Before that, the series had been rebased onto current master and checked derivation by derivation: the rebase changed no build input of the ROCm set apart from migraphx'sonnxPython dependency (upstream bump 1.22.0 → 1.23.2) and rocm-smi, which now uses the upstream-merged packaging.Hardware inference:
llama-benchtg64), ollama 58 tok/s (llama3.2 3B, 29/29 layers offloaded,library=ROCm compute=gfx1102), torch 2.13.0 (py3.14)is_available()true, fp32 2048³ matmul matches the CPU result to 5e-4.llama-benchtg64, pp128 ~3.2k tok/s), all 99 layers offloaded,found 1 ROCm devices … gfx908:sramecc+:xnack-. The previous head (llvmPackages_gitbase) measured ~180 tok/s on the same model.Not tested: full Hydra matrix; CDNA2/3/4 (gfx90a/gfx942/gfx950); vLLM under real load on a supported arch. Independent builds and smoke tests on other hardware are very welcome.
Things done
passthru.tests:rocmPackages.rocm-tests.tests.rocmPackagesDerivations(everyrocmPackagesattribute, full 16-arch matrix) built green; see Testing.nixpkgs-review:nixpkgs-review rev -b master(master at 30ca09c as the base,rocmSupport = true) on 16 key attrs: clr, hipcc, rocminfo, rocm-smi, amdsmi, rocblas, hipblas, hipblaslt, composable_kernel, miopen, rccl, rocprofiler-sdk, llama-cpp, ollama, python3Packages.torch, onnxruntime. 16 built, 0 failed; the result paths are the same store paths as the full-set build above.rl-2611entry for the ROCm 10.1 update, the ABI/API breakage and the gfx900/gfx906 support change.Assisted-by:trailer, and this description is disclosed at the end.Additional checks:
rocmPackages: 7.2.3 -> 7.2.4rocmPackages: 7.2.3 -> 7.2.4 #572764 andollama: 0.40.2 -> 0.40.3merged). The conflicts were the 7.2.4 version and hash lines against this series' 10.1.0 lines; every ROCm derivation hash is unchanged from the verified builds, ollama was rebuilt at 0.40.3. nixfmt clean.clr(the 2026-10-14$out/etcfix re-keyed every ROCm derivation; the results above are from the pre-fixclr, which differs only by two files under$out/etc; see Build history)rocmPackagesevaluates (every attr with a ROCm 10.1 release is at 10.1; documented pins with no 10.x upstream: mivisionx 7.2.3, rocm-bandwidth-test 6.3.3, mscclpp, hiprt)config.allowAliases = false(the nixpkgs CI setting):rocmPackages,magma,pkgsRocm.ollamallvmPackages_23.openmp(whose standalone build is broken on master) changes:llvmPackages.openmpfor LLVM < 23 is derivation-identical to master on linux and darwinRebuild scope, target branch and LLVM base
Every package in
rocmPackageschanges version, so every ROCm consumer (torch, vllm, llama-cpp, ollama, magma, blender, onnxruntime, …) rebuilds. The fork PR's CI eval compare (randomizedcoder#9) on this head reportsRebuild: linux 105, darwin 2(labels10.rebuild-linux: 101-500,10.rebuild-darwin: 1-10); the two darwin rebuilds arellvmPackages_23.openmp, whose standalone build is broken on master and which this PR fixes. That is under the 500-rebuild rule of thumb in CONTRIBUTING for targetingstaging, somasteris a valid target; the individual builds are long (see the scale note below), which is a Hydra-capacity question rather than a rebuild-count one. This PR is fork→fork for review; the upstream target branch is the reviewers' call.ROCm 10.1's bundled compiler is LLVM 24, and nixpkgs has no released
llvmPackages_24. The ROCmllvmmodule builds LLVM 24 from ROCm's ownllvm-projectsource and uses a nixpkgs LLVM package set only as the base to override: the base supplies the shared packaging incompilers/llvm/commonplus the bootstrap compiler and bintools, andversion = "24.0.0-rocm"selects the LLVM 24 patch set. The first version of this PR usedllvmPackages_git(the24.0.0-gitsnapshot) as that base, matching the ROCm major. That does not pass nixpkgs CI:llvmPackages_gitis an alias reserved for integration testing, and its package-set splice points at the alias, so the set cannot be evaluated withconfig.allowAliases = false, which is what CI uses. The PR now uses the releasedllvmPackages_23as the base, one major behind, which is fine because the base's own libllvm is only used to bootstrap (LTO is off for ROCm's LLVM). The ROCm clang still reports24.0.0 (ROCm/llvm-project therock-10.1). WhenllvmPackages_24is released, the base is a one-line change. Alternative if reviewers prefer a toolchain whose base matches its major: ship ROCm 10.0.0 first, which is LLVM 23 onllvmPackages_23; I have a proven-green 10.0.0 series and can repost it as a separate PR.Build history: context for the amount of building behind this PR
This bump was built twice: first to ROCm 10.0.0 (LLVM 23), then retargeted in place to 10.1.0 (LLVM 24) when AMD released 10.1 mid-effort. The packaging debugging is version-independent, so the 10.0 pass validated the approach against a released LLVM before taking on the unreleased-LLVM-24 question. Below is the dated record from the working log, condensed. Times are US/Pacific unless marked UTC. 68 distinct build breakages were catalogued and fixed along the way (C1–C68 in the log); only the ones that changed the shape of the PR are called out here.
Hardware used
ll2Pass 1: 7.2.3 → 10.0.0 (LLVM 23 on
llvmPackages_23)37710547. Same day: ROCm LLVM 23 toolchain built (libllvm passed its full lit suite, 65,969 tests); 7 nixpkgs backport patches dropped as upstreamed; device-libs, comgr, hipcc, rocm-cmake, half, rocm-core, hip-common, rocm-runtime, rocminfo, clr, rocm-smi and amdsmi 27 all built.$HOMEhelper-cache crash, a postInstall bin move), so each cost most of a day. Green on attempt 7, 2026-09-23 ~01:10.device_conv_operationsaggregate. MIOpen 3.6.0 green on attempt 8 (09:32 → 11:10) after six sequential fixes.find_package(rocm_smi), CK FMHA generator shebang, aotriton adapter header, then torch 2.13 pinning aotriton 0.12b (bumped; needed the git-tree-sha and aiter-vendoring fixes already used by nixpkgs PR #562330). Attempt 5 died with simultaneous SIGBUS inld.lldwhile hipBLASLt and rocBLAS compiled concurrently (resource exhaustion, not a packaging defect). Attempt 6, narrowed to gfx906, compiled and installed 2026-09-26.import torchhang root-caused to the OpenMP tool interface (OMPT) and fixed. torch computes on a real MI50 (gfx906); transformer inference ran; ollama 75 tok/s on the MI50. vLLM (py3.14) built onl2after three OMPT deadlock fixes in its build pipeline. vLLM engine start then proved impossible on gfx906: Triton's AMD backend has no gfx906 target (upstream limitation, documented). rocBLAS per-arch Tensile split drafted.nix-eval-jobsgate passed. Full-matrix heavy-hitter build launched.l2GPU swapped MI50 → MI100 (gfx908). Heavy-hitter full-matrix build (CK, CK conv ops, MIOpen, rocBLAS shards, rccl) finished 16:32 → 21:52, 0 failures.l2hung with a single python process at 99 % for four days.l2warm-rebooted, MI100 recovered, llama-cpp smoke on gfx908 passed (~180 tok/s). Hung torch build killed; GC had reclaimed the gfx908 artifacts. Commit series structured (30 commits). Full-matrix torch launched. onnxruntime + MIGraphX + obs-backgroundremoval chain built.rocm_add_version_resourcestub.llvmPackages_2323.0 → 23.1, changing every ROCm store path, so the toolchain was rebuilt and re-verified before pushing.Pass 2: 10.0.0 → 10.1.0 (LLVM 24 on
llvmPackages_git)rocm-10-10.0-greenon the proven 10.0 head. Gate build ofllvmPackages_git(24.0.0 snapshot) and then the ROCm clang on top of it:clang --versionreports24.0.0 (ROCm/llvm-project therock-10.1)the same evening. LLVM patch stack re-audited against thetherock-10.1source: four more patches dropped as upstreamed. Mass bump: 44 files to 10.1.0, 45 sources rehashed (42 changed, 3 identical between tags). rocmlir moved to therocmlirTritoncommit pinned by AMDMIGraphX atrocm-10.1.rocm_check_target_idsprobe. Root cause: LLVM 24'sclang-linker-wrapperlinks-lzstdwhen it bundles compressed offload images, and nothing supplied it inside the sandbox (it passed innix developonly because the impure shell had zstd onLIBRARY_PATH). Fixed in the toolchain's clang config, which re-keyed every downstream path: fleet relaunched acrossl+l2. rocprofiler-sdk: six patches confirmed upstream and dropped. rocm-comgr ISA-compat patch rebased for LLVM 24. rocmlir-rock green after a four-fix rebase.disallowedReferences: two new-in-10.1 conv instances embed the build part's own store path inside.hip_fatbin(survivesstrip). Root-caused and scrubbed. rocBLAS postPatch fixed for a path the 10.1 source no longer contains. MIOpen's newhipconvsubsystem (cdna3/4/5 only) does not compile under clang 24; disabled, MIOpen and migraphx green. Core library fleet 15/15 green. Consumer fleet launched 23:21 in two waves; wave A (8 torch-independent consumers) green the same night after a whisper-cpp OpenMP link fix. A finish-time estimate was written down at 23:58 to be scored later.rocprof-trace-decoder/hipfortmoved to their 10.1 sources.rocm-10.1, which closed randomizedcoder#6; randomizedcoder#9 opened with the same content. Rebased onto current master and audited derivation by derivation (see Testing).Pass 3: LLVM base
llvmPackages_git→llvmPackages_23(CI fix)rocmPackagesdoes not evaluate withallowAliases = falsebecausellvmPackages_gitis a testing-only alias whose package-set splice resolves through the alias. Root-caused, thellvmmodule moved to the releasedllvmPackages_23base (ROCm's LLVM 24 source and patch set unchanged), aliases-off evaluation verified for the ROCm set, magma andpkgsRocm.ollama. Every ROCm store path re-keyed, so the in-flight consumer rebuild was stopped, the toolchain + clr gate launched on the new base, the branch rebased and force-pushed, and the full set plus consumer fleet (152 derivations) relaunched acrossl+l2. CI's eval compare on the new head then showed every aarch64-darwin package rebuilding: the openmp commit set an unconditional (empty)preConfigure, which re-keyedllvmPackages.openmpon every platform. Made conditional withlib.optionalAttrs; openmp 21 is byte-identical to master again on linux and darwin, and the ROCm derivations are unchanged, so the running rebuild is unaffected.half.hpp:HALF::halfwas linked PRIVATE although the public header includes it; exported PUBLIC) and rpp-cpu (clang 24mm_malloc.hvs glibc 2.44posix_memalignconflict, root-caused to the-fopenmpinclude path described under Known follow-ups;-isystemworkaround).llvmPackages_23base (152 derivations acrossl+l2) finished green in three runs: run 1 lost hipBLASLt after its 18 h Tensile phase when a store garbage collection on the remote builder removed one of its inputs mid-build (an operational mistake, not a packaging failure); run 2 rebuilt hipBLASLt in 5.5 h uncontended and everything downstream except py3.13 torch, which lost its output registration to the same class of remote-builder cleanup; run 3 rebuilt that one package. Nothing in the series changed between the runs. gfx1102 and gfx908 hardware smoke re-run with the resulting binaries (numbers under Testing); the gfx908 run waited for a power cycle of the MI100 host, whose card had dropped off the PCIe bus again (a recurring fault on that box, not a build result).rocmPackages: 7.2.3 -> 7.2.4(#572764) andollama: 0.40.2 -> 0.40.3merged. All 54 conflicts were the 7.2.4 version/hash lines against this series' 10.1.0 lines, resolved to 10.1.0; upstream's 7.2.4-only rocprofiler-sdk patch files were dropped (that stack is upstreamed in therock-10.1). Every ROCm derivation hash is unchanged from the verified builds; ollama rebuilt at 0.40.3 and smoke-tested on gfx1102. Opened upstream.clr: dropped the distro-styleetc/OpenCL/vendors/amdocl64.icdandetc/ld.so.conf.d/10-rocm-opencl.confthat CLR 10.1 newly installs into$out(7.2.3 did not); they collided with theicdoutput whenhardware.amdgpu.opencl.enableput both outputs intohardware.graphics.extraPackages(buildEnvconflict atnixos-rebuild). Found deploying to the two test hosts. This re-keys every derivation that depends onclr: the full-set build, nixpkgs-review and hardware results above were obtained on the pre-fixclr, whose outputs differ from the fixed one only by those two files under$out/etc(verified by comparing file lists).clritself rebuilt and checked; the full-set rebuild on the fixedclris the next step.Scale, for calibration. A full uncontended pass of the ROCm library matrix on the 24-thread builder is roughly a day: hipBLASLt ~8 h, composable_kernel ~9 h, rocBLAS ~2.5 h, rocsparse ~3 h, rocsolver ~2 h, MIOpen ~1.5 h, hipSPARSELt ~50 min, the rest minutes each. The consumer fleet on top is another ~15 h with torch and aotriton as the long pole. Because Nix has no intra-derivation resume, a late failure in a Tensile package costs the whole build; that is where most of the 20 calendar days went.
Assisted-by: Claude Code (Claude Opus 4.8, Claude Fable 5.1)
🤖 Generated with Claude Code