Skip to content

rocmPackages_10: add Tensile and rocBLAS - #24

Closed
randomizedcoder wants to merge 3 commits into
rocm10/06-hipblasltfrom
rocm10/07-tensile-rocblas
Closed

rocmPackages_10: add Tensile and rocBLAS#24
randomizedcoder wants to merge 3 commits into
rocm10/06-hipblasltfrom
rocm10/07-tensile-rocblas

Conversation

@randomizedcoder

@randomizedcoder randomizedcoder commented Oct 14, 2026 •
edited
Loading

Copy link
Copy Markdown
Owner

Part 7/14. Part of splitting upstream PR 573885 (rocmPackages: 7.2.4 -> 10.1) into smaller PRs, as asked in its review.

  • Adds to rocmPackages_10: tensile, rocblas
  • Target: master. No existing derivation changes: rocmPackages_10 is not recursed into

How to review

  • First commit: copies the 7.2.4 expressions unchanged into rocm-modules/10/
  • Then one commit per package: the 10.1 changes, with the reasons in the commit message

Notes

  • Split out on reviewer request
  • Tensile kernels are built as one shard per GPU arch, then joined into rocblas
  • rocblas_base is marked broken (like composable_kernel_base), so Hydra never builds the all-arch base directly
  • Overrides such as rocblas.override { buildTests = false; } are passed on to base.nix
  • Sharding applies only to separate-architecture, lazy-loaded Tensile builds

Testing

  • rocmPackages_10 evaluates: 47 derivations, 0 failures
  • Built on x86_64-linux with the default GPU targets, including the 15 per-arch Tensile shards
  • nixpkgs-review: no rebuilds in the default package set
  • Hardware smoke tests: see rocmPackages: 7.2.4 -> 10.1 #31

Series

Preparatory PRs, each against master: #10, #11, #15, #16, #17.

Part PR Contents Base
1 #18 rocmPackages_10: init master
2 #19 llvm (LLVM 24) #18
3 #20 runtime #19
4 #21 profiling and debugging tools, RCCL #20
5 #22 rocPRIM, rocRAND, rocFFT #21
6 #23 hipBLASLt #22
▶ 7 #24 Tensile and rocBLAS (this PR) #23
8 #25 sparse, solvers, rocALUTION, rocWMMA, hipfort #24
9 #26 composable_kernel #25
10 #27 rocMLIR #26
11 #28 MIOpen #27
12 #29 MIGraphX, RPP, MIVisionX #28
13 #30 remaining packages #29
14 #31 rocmPackages: 7.2.4 -> 10.1 (the switch) #30

🤖 Generated with Claude Code

@randomizedcoder
randomizedcoder force-pushed the rocm10/07-tensile-rocblas branch from 0f84a57 to 6b000b8 Compare October 15, 2026 23:02
@randomizedcoder
randomizedcoder force-pushed the rocm10/06-hipblaslt branch 2 times, most recently from 202550d to 742a184 Compare October 17, 2026 02:50
@randomizedcoder
randomizedcoder force-pushed the rocm10/07-tensile-rocblas branch from 6b000b8 to 6b95357 Compare October 17, 2026 02:50
@randomizedcoder
randomizedcoder force-pushed the rocm10/07-tensile-rocblas branch 3 times, most recently from e278b88 to c13eb13 Compare October 21, 2026 07:15
@randomizedcoder
randomizedcoder force-pushed the rocm10/07-tensile-rocblas branch from c13eb13 to e268c5f Compare October 22, 2026 14:51
@randomizedcoder
rocmPackages_10: copy Tensile and rocBLAS from rocmPackages
b043375 
Verbatim copies of the rocmPackages (ROCm 7.2.4) expressions for
tensile, rocblas, not yet part of the rocmPackages_10 scope. The
following commits update each one to ROCm 10.1, so their diffs read
against 7.2.4.

Assisted-by: Claude Code (claude-opus-5-5)
@randomizedcoder
rocmPackages_10.tensile: init at 10.1
baf7dc9 
Update the rocmPackages.tensile copy from the previous commit to ROCm
10.1 and add it to the rocmPackages_10 scope. Source/version bump only.

Assisted-by: Claude Code (claude-opus-5-5)
@randomizedcoder
rocmPackages_10.rocblas: init at 10.1
c0ec6f6 
Update the rocmPackages.rocblas copy from the previous commit to ROCm
10.1 and add it to the rocmPackages_10 scope.

From "rocmPackages.rocblas: split Tensile kernels into per-arch shards":

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.

From "rocmPackages.rocblas: 7.2.4 -> 10.1.0":

- Clients switched rocm_smi -> amd_smi (find_package(amd_smi CONFIG
  REQUIRED)); add amdsmi for the test/benchmark builds.
- Drop the GZGavinZhao "Extend rocBLAS HIP ISA compatibility" patch:
  library/src/handle.cpp now has explicit Processor::gfx906/gfx90c
  cases, so with gfx906 in gpuTargets Tensile builds real gfx906
  kernels.
- Give cblas an explicit -L. rocBLAS compiles with the unwrapped
  rocm-toolchain clang, so buildInput lib dirs are not injected into the
  clang-linker-wrapper host-link path. With LLVM 24 the rocm-cmake
  rocm_check_target_ids probe (run before find_package(BLAS)) does a
  full --hip-link, and CMake seeds CMAKE_EXE_LINKER_FLAGS from $LDFLAGS,
  so a bare `-lcblas` makes that probe fail, emptying the
  supported-target list and aborting with "Unsupported target". Add
  amd-blis's libdir.
- Drop the testing_logging.hpp /usr/bin/diff substitution; the source no
  longer shells out to it.
- ROCm 10 installs a bin/rocblas/ CTest scaffolding dir alongside the
  client yaml/txt, so remove the whole bin/ tree from the library
  output.

Revised during the split review:

Forward package overrides to base.nix and read the effective settings
from its passthru. Shard only builds whose Tensile libraries are
separate-architecture and lazily loaded; use the base build otherwise.
Omit -lzstd and the Tensile reference removal when Tensile is disabled.

Assisted-by: Claude Code (claude-opus-5-5)
@randomizedcoder
randomizedcoder force-pushed the rocm10/07-tensile-rocblas branch from e268c5f to c0ec6f6 Compare October 27, 2026 04:12
This was referenced Oct 27, 2026
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
Closed
@randomizedcoder

randomizedcoder commented Oct 27, 2026

Copy link
Copy Markdown
Owner Author

Opened upstream as NixOS#579520.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@randomizedcoder