Skip to content

rocmPackages_10: add composable_kernel - #26

Closed
randomizedcoder wants to merge 2 commits into
rocm10/08-sparse-solversfrom
rocm10/09-composable-kernel
Closed

randomizedcoder wants to merge 2 commits into
rocm10/08-sparse-solversfrom
rocm10/09-composable-kernel

Conversation

@randomizedcoder

@randomizedcoder randomizedcoder commented Oct 14, 2026 •

Copy link
Copy Markdown
Owner

Part 9/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: composable_kernel
  • 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
  • Convolution instances are compiled in 8 memory-bounded parts again, using ROCm 10's layout-specific OBJECT targets
  • The final derivation only archives their objects into device_conv_operations
  • Parts skip targets CMake didn't generate for the configuration, and fail on names the CK source no longer defines (this caught the removed streamk targets)

Testing

  • rocmPackages_10 evaluates: 58 derivations, 0 failures
  • Built on x86_64-linux with the default GPU targets, including the 8 convolution parts (peak memory ≤ ~30 GB each)
  • 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 #23
8 #25 sparse, solvers, rocALUTION, rocWMMA, hipfort #24
▶ 9 #26 composable_kernel (this PR) #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/09-composable-kernel branch from c632c92 to 5d3ae80 Compare October 15, 2026 23:02
@randomizedcoder
randomizedcoder force-pushed the rocm10/08-sparse-solvers branch 2 times, most recently from 44f8a64 to 235026f Compare October 17, 2026 02:50
@randomizedcoder
randomizedcoder force-pushed the rocm10/09-composable-kernel branch 2 times, most recently from 16bee19 to 8148431 Compare October 17, 2026 16:20
@randomizedcoder
randomizedcoder force-pushed the rocm10/08-sparse-solvers branch from 0914857 to 970adf2 Compare October 20, 2026 06:44
@randomizedcoder
randomizedcoder force-pushed the rocm10/09-composable-kernel branch from 8148431 to ce6e3df Compare October 20, 2026 06:44
@randomizedcoder
randomizedcoder force-pushed the rocm10/08-sparse-solvers branch from 970adf2 to d05854b Compare October 21, 2026 07:15
@randomizedcoder
randomizedcoder force-pushed the rocm10/09-composable-kernel branch 2 times, most recently from 3af6d0c to ddf83e6 Compare October 22, 2026 14:51
@randomizedcoder
randomizedcoder force-pushed the rocm10/08-sparse-solvers branch from d05854b to 6a96f46 Compare October 22, 2026 14:51
Verbatim copies of the rocmPackages (ROCm 7.2.4) expressions for
composable_kernel, 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)
Update the rocmPackages.composable_kernel copy from the previous commit
to ROCm 10.1 and add it to the rocmPackages_10 scope.

From "rocmPackages.composable_kernel: adapt to ROCm 10 conv targets and
clang 24":

ROCm 10 restructured CK's per-op conv instance libraries into OBJECT
libraries aggregated under a single STATIC target
device_conv_operations.

clang 24 forwards NIX_LDFLAGS to the HIP device linker, so every split
"parts" instance's device object is stamped with DT_RUNPATH=$out/lib.
That RUNPATH is baked into the --offload-compress'd .hip_fatbin (a CCOB
integrity-checked blob that cannot be scrubbed post-hoc without
corrupting the device code), and since those .o are copied into the
final CK derivation it becomes a disallowed cross-reference. The parts
only emit relocatable .o, so the self-rpath is useless here; strip
"-rpath $out/lib" from NIX_LDFLAGS (and any role-suffixed copies) in
preConfigure, keeping offload compression.

Also give fix-empty-offload-targets.diff a gfx90a fallback for an empty
offload-target list.

Revised during the split review:

Restore memory-bounded per-group parts for the convolution instances,
using ROCm 10's layout-specific OBJECT targets; the final build only
archives their objects into device_conv_operations. Parts skip targets
CMake did not generate for the configuration and fail on names the CK
source no longer defines. Drop the streamk targets, removed upstream.

Assisted-by: Claude Code (claude-opus-5-5)
@randomizedcoder

Copy link
Copy Markdown
Owner Author

Opened upstream as NixOS#579522.

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant