fix(hipblaslt): avoid regenerating unchanged MX inputs on every solution by minsukim-amd · Pull Request #11963 · ROCm/rocm-libraries

minsukim-amd · 2026-09-11T01:24:58Z

JIRA ID : AIHPBLAS-4461

Motivation

The tensilelite benchmark client regenerates the full host-side MX input set once per
solution when benchmarking a list of solutions against a single problem. The generated
bytes are identical every time.

MX input generation is expensive in absolute terms — anyone who has run an MX problem
through the benchmark clients has noticed that a single MX GEMM takes far longer to set up
than a comparable non-MX one. Per full pass the generator draws a random value and a
double-precision cos per element, packs sub-byte elements, computes a per-block E8M0
scale, and materialises a float reference copy. copySwizzledToGPUBuffer() then pads,
reshapes and permutes the whole scale tensor on the host.

Because that cost is paid per solution rather than per problem, benchmarking N solutions
against one MX problem costs roughly N full input generations. For a problem whose B tensor
is a couple of GiB, host-side input preparation dominates the run so completely that the
measured kernel time is a rounding error against it.

prepareGPUInputs() is called from the per-solution loop in main.cpp (resetInput is set
true and never cleared), and for MX problems none of the usual short-circuits apply:

The m_cachedGPUInputs fast path requires !m_problemDependentData && !needMXSwizzle.
needMXSwizzle is true for any problem with a non-zero MX block, and isMXProblem()
separately forces m_problemDependentData = true. Two independent reasons; no
command-line option reaches either.
The m_cpuPtrs.empty() fallback therefore fires on every call. m_cpuPtrs is populated
only under ReferenceValidator, so when validation is disabled it stays empty for the
lifetime of the process.
initializeMXData() has no memoization of its own.

Point 3 is the asymmetry with the non-MX path, which reaches initializeCPUInputs() on
every solution too — and does nothing, because it is guarded:

// Only update when the descriptor changed
if(p.second.initDescriptor[0] != tensors[i]) { ... }

That same idiom already exists inside initializeMXData, in its non-MX fallback
initTensorFromDefault. It was simply never applied to initOneMXSide. So a non-MX tensor
is generated once per problem while an MX tensor is generated once per solution.

Technical Details

Two files in the tensilelite benchmark client:

tensilelite/client/src/DataInitialization.cpp
tensilelite/client/include/DataInitialization.hpp

Two guards, both keyed on the tensor descriptors

initOneMXSide() returns early when the data and scale descriptors are unchanged,
mirroring the existing non-MX initDescriptor check.
copySwizzledToGPUBuffer() gains an else if branch that reuses the device buffer when
the scale descriptor it last swizzled is unchanged, instead of repeating the host
pad/reshape/permute. This is the MX counterpart of g_swizzleCache, which the MX branch
never reached. Guard 1 alone is not sufficient: with only the first guard applied, the
host-side permute still accounted for roughly two thirds of the remaining profile
samples.

The new state is a single member, m_mxSwizzledDescriptor, a tensor index → TensorDescriptor map placed alongside the other MX members.

Scoped to gfx1250

Both guards require m_mxScaleLayout == MXScaleLayout::GFX1250.
mxScaleLayoutForArchName() maps gfx1250 → GFX1250, gfx950 → GFX950, and gfx942 /
gfx90a → None, so this is a positive allowlist: gfx1250 (both ASIC revisions, which share
gcnArchName) changes behaviour, and every other architecture keeps the existing code
path until the change has been measured there.

This is keyed on the member m_mxScaleLayout — the device's layout — rather than the
per-call swizzleLayout parameter. swizzleLayout is MXScaleLayout::None unless
MXScaleFormat is set, so keying on it would silently disable the optimisation for the
common gfx1250 configuration.

MXScaleLayout::GFX950 is excluded deliberately, not merely by omission: its preswizzle
derives from m_currentSolution->sizeMapping.matrixInstruction, so its generated bytes
genuinely do vary per solution. gfx1250 derives its layout from problem geometry alone,
which is what makes descriptor equality a sufficient key there.

What is not changed

Generated data is identical — the same bytes are produced, they are simply not recomputed
when nothing they depend on has changed. No data-init mode, seed, scale format, or
validation behaviour is touched. Timed regions are unaffected: input preparation runs
outside the preEnqueues / postEnqueues bracket in BenchmarkTimer.

Related inefficiencies found but deliberately not addressed here

DataGenerator::packArray is a fully serial loop over every element (no OpenMP), on the
critical path for every sub-byte type.
DataGenerator caps itself at std::min(32, omp_get_max_threads()), with no override.
generateMXInput returns a full std::vector<float> reference copy that both tensilelite
call sites discard.
A HIP data generator (DataGeneratorGPU_impl.hpp) exists in the shared mxdatagenerator
library but is never called from hipBLASLt.

Each of these would further reduce MX setup cost for the benchmark clients generally, not
just for repeated solutions.

Test Plan

Direct phase measurement. Use the client's own --timing-instrumentation to record
the per-solution input-reset phase and the benchmark phase, before and after, on the same
problem and the same prebuilt library.
Kernel timings unchanged. Run the same solution indices from the same library under
the pre-change and post-change binaries and compare the reported kernel times, to confirm
the change is not visible in what the client measures.
Non-MX regression. Run a complete non-MX (bf16) benchmark sweep — a path the change
must not reach — and compare the selected kernel, its measured time, and its validation
status against the pre-change result.
End-to-end MX coverage. Re-run a full MX sweep across mxfp8, mxfp4 and nvfp4, for both
a compute-bound and a bandwidth-bound problem shape, and confirm the results are
numerically valid and consistent with the pre-change values rather than merely faster.
Compilation. Build the client with the change applied.

Test Result

Measured on gfx1250.

Per-solution input reset drops by ~550x, from tens of seconds to tens of
milliseconds, on an MX problem with a multi-GiB B tensor. The benchmark phase itself is
unchanged. What remains in a multi-solution run is the single unavoidable initial
preparation.
Kernel timings unchanged — maximum deviation 0.15% across the matched solutions, which
is within run-to-run noise for this client.
No non-MX regression. A full bf16 sweep selected the identical winning kernel, its
measured time differed by 0.08%, and it reported 0 validation failures.
End-to-end MX, six configurations (mxfp8, mxfp4, nvfp4 × compute-bound and
bandwidth-bound), all completing with 0 Check failed in output tensor across every
log:
- Two of the six had previously been run to completion under the unpatched client. Both
  produced results within ~1% of their pre-change values — the same order as the 0.08%
  seen on the deterministic non-MX re-run, and consistent with search-order variance
  rather than a measurement shift.
- Two of the six had never completed at all under the unpatched client; both now finish.
- Relative ordering across data types is preserved and physically sensible: for the
  fixed-byte bandwidth shape, throughput decreases monotonically as scale-block overhead
  increases, with 16-element scale blocks below 32-element ones.
Compiles clean with the project's own compile_commands.json flags.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

therock-pr-bot · 2026-09-11T01:29:19Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/hipblaslt/tensilelite/client/include/DataInitialization.hpp`, `projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp`; no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-09-11T01:29:21Z

🎉 All checks passed! This PR is ready for review.

codecov · 2026-09-11T02:43:57Z

Codecov Report

❌ Patch coverage is 0% with 19 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...aslt/tensilelite/client/src/DataInitialization.cpp	0.00%	19 Missing ⚠️

Additional details and impacted files

@@           Coverage Diff            @@
##           develop   #11963   +/-   ##
========================================
  Coverage    70.54%   70.55%           
========================================
  Files         2810     2810           
  Lines       462436   462506   +70     
  Branches     68077    68099   +22     
========================================
+ Hits        326219   326279   +60     
- Misses      112673   112684   +11     
+ Partials     23544    23543    -1

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`46.40% <0.00%> (-0.05%)`	⬇️
TensileLite-Unit	`76.06% <ø> (+0.03%)`	⬆️
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 76ae32e
hipBLASLt	`35.23% <ø> (ø)`
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 76ae32e
hipDNN	`87.02% <ø> (ø)`		Carriedforward from 76ae32e
hipFFT	`42.66% <ø> (ø)`		Carriedforward from 76ae32e
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 76ae32e
hipSOLVER	`68.92% <ø> (ø)`		Carriedforward from 76ae32e
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 76ae32e
rocBLAS	`48.31% <ø> (ø)`		Carriedforward from 76ae32e
rocFFT	`51.72% <ø> (ø)`		Carriedforward from 76ae32e
rocRAND	`56.90% <ø> (ø)`		Carriedforward from 76ae32e
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 76ae32e
rocSPARSE	`74.61% <ø> (ø)`		Carriedforward from 76ae32e
rocThrust	`91.60% <ø> (ø)`		Carriedforward from 76ae32e

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...aslt/tensilelite/client/src/DataInitialization.cpp	`0.00% <0.00%> (ø)`

... and 5 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

The client regenerates the full host-side MX input set once per solution. prepareGPUInputs() is called from the per-solution loop in main.cpp, and for MX problems none of the usual short-circuits apply: the m_cachedGPUInputs fast path requires !needMXSwizzle, and isMXProblem() forces m_problemDependentData, so initializeCPUInputs() runs every time. The non-MX path already avoids the redundant work with an initDescriptor check ("Only update when the descriptor changed"), and the same idiom exists in initTensorFromDefault inside initializeMXData, but it was never applied to initOneMXSide. So a non-MX tensor is generated once per problem while an MX tensor is generated once per solution. That is expensive because MX generation is not simply a larger version of the non-MX path: per full pass it draws a random value and a double-precision cos per element, packs sub-byte elements, computes a per-block E8M0 scale, and materialises a float reference copy. Measured on gfx1250, mxfp4, 16x65536x65536 (2 GiB B tensor), using the client's own --timing-instrumentation: gpu_input_reset 24386 ms per solution -> 44.5 ms (548x) benchmark_runs 24.5 ms per solution i.e. ~99.9% of tuning wall time was spent regenerating identical input data. Scoped to gfx1250 (MXScaleLayout::GFX1250), where the scale layout derives from problem geometry alone and the generated bytes therefore depend only on the descriptors. GFX950 preswizzles using the solution's matrix instruction and must keep regenerating; every other architecture keeps the previous behaviour until this has been measured there. Two guards, both keyed on the tensor descriptors: - initOneMXSide() returns early when the data and scale descriptors are unchanged. - copySwizzledToGPUBuffer() reuses the device buffer when the scale descriptor it last swizzled is unchanged, instead of repeating the host pad/reshape/permute. This mirrors g_swizzleCache on the non-MX side, which the MX branch never reached. Generated data is unchanged: the same bytes are produced, they are simply not recomputed when nothing they depend on has changed. Validation behaviour is untouched. Verified on gfx1250: - same mxfp4 solutions, before vs after: max kernel-time delta 0.15% (473.42 -> 474.09, 473.67 -> 473.50, 313.14 -> 312.66 us) - full bf16 tuning re-run (non-MX regression check): 12.6474 TB/s vs 12.6574 TB/s, identical winning kernel, 0 validation failures - mxfp8 MAF tuning job, same config before and after: GA fitness 3685000 vs 3648750 (-1.0%, independent-run variance), 4h22m -> 353s, 0 validation failures Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

aazz44ss

LGTM

Do you know why it regenerating unchanged MX inputs every solution?
Do we need to leave validation != 0 case unchanged?

Two follow-ups to the MX input memoization. Neither changes the cost of a benchmark sweep, which is the case #11963 optimised. 1. Key the scale-upload reuse on the layout in the buffer, not just the descriptor. The scale upload layout is chosen per solution from problemType.mxScaleFormat; that is why the per-solution upload exists at all (main.cpp sets resetInput = true so the upload follows preSolution() and can read the picked solution's format). A NoSwizzle solution writes the canonical layout into the same gpuInput.valid buffer the swizzle path uses, so keying reuse on the descriptor alone let a later swizzled solution reuse a buffer holding the wrong layout. Every non-NoSwizzle format takes the same fallback swizzle, so the key only needs to record whether the buffer holds the canonical layout. That also lets the first, pre-solution call (m_currentSolution == nullptr, format sentinel -1) populate an entry the first real solution can reuse. Each solution carries its own problemType and MXScaleFormat is a settable solution parameter, so values can differ within a run. In practice it is set globally and is uniform, which is why this did not surface in testing. 2. Disable the reuse when validation is enabled. Re-uploading inputs for every solution also refreshes them. With the reuse in place that no longer happens, so a kernel writing outside its buffers could cause a later solution to be reported as failing. Gate both guards on m_elementsToValidate == 0. Benchmark sweeps run with validation off and are unaffected. Measured on gfx1250, mxfp4, 12 solutions from one library: validation off merged 484 ms total, 44 ms per reset this 484 ms total, 44 ms per reset (identical) validation on reuse disabled, inputs refreshed per solution as before Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

Follow-up to #11963. The reuse added there keyed only on the tensor descriptor, but the buffer it reuses can hold either of two layouts. The scale upload layout is chosen per solution from problemType.mxScaleFormat, which is why the per-solution upload exists at all: main.cpp sets resetInput = true so the upload happens after preSolution() and can read the picked solution's format. A NoSwizzle solution writes the canonical layout into the same gpuInput.valid buffer the swizzle path uses, so with a descriptor-only key a later swizzled solution could reuse a buffer holding the wrong layout, producing incorrect results with no validation failure during a sweep. Every non-NoSwizzle format takes the same fallback swizzle, so the key only needs to record whether the buffer currently holds the canonical layout. That also lets the first, pre-solution call (m_currentSolution == nullptr) populate an entry the first real solution can reuse, so no extra regeneration is introduced. Each solution carries its own problemType and MXScaleFormat is a settable solution parameter, so values can differ within a run. In practice it is set globally and is uniform, which is why this did not surface in testing. Measured on gfx1250, mxfp4, 12 solutions from one library, per-solution input reset: validation off #11963 484 ms total / 44 ms each this 484 ms total / 44 ms each validation on #11963 484 ms total / 44 ms each this 484 ms total / 44 ms each, results still PASSED Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

minsukim-amd · 2026-09-11T20:24:31Z

LGTM

Do you know why it regenerating unchanged MX inputs every solution? Do we need to leave validation != 0 case unchanged?

Thanks Henry for the review! To your questions,

Q1. Do you know why it is regenerating unchanged MX inputs every solution?

Yes. It is a side effect of a deliberate change, not an accident in itself.

main.cpp sets resetInput = true unconditionally (commit e9f7fc5231, "Extend MX
BufferLoad + NoSwizzle to all transposes on gfx1250"). That was intentional: the MX scale
upload layout is chosen per solution from problemType.mxScaleFormat, so the upload has
to happen after preSolution() in order to read the picked solution's format.

What rides along with it is the expensive part. prepareGPUInputs() reaches
initializeCPUInputs() on every solution, and for MX none of the usual short-circuits
apply:

the m_cachedGPUInputs fast path requires !needMXSwizzle, which is false for any
problem with a non-zero MX block;
isMXProblem() separately forces m_problemDependentData = true;
so the m_cpuPtrs.empty() fallback fires on every call.

The non-MX path reaches initializeCPUInputs() just as often and does nothing, because it
is guarded:

// Only update when the descriptor changed
if(p.second.initDescriptor[0] != tensors[i]) { ... }

That idiom is even present inside initializeMXData, in its non-MX fallback
initTensorFromDefault — it was simply never applied to initOneMXSide. So a non-MX tensor
is generated once per problem, while an MX tensor is generated once per solution.

This PR adds that missing guard. The per-solution upload behaviour introduced by
e9f7fc5231 is unchanged; only the redundant regeneration is skipped.

Q2. Do we need to leave the `validation != 0` case unchanged?

I thought so initially and implemented exactly that, then measured it and backed it out —
the data does not support special-casing validation.

The concern is fair: re-uploading inputs on every solution also refreshes them, so reuse
removes that refresh. But no datatype gates input reuse on validation. The non-MX
initDescriptor guard is unconditional, and has been for a long time.

Measured on gfx1250, per-solution input reset, 12 solutions from one library:

	validate = 0	validate = -1
bf16	23 ms	23 ms
mxfp4 (this PR)	44 ms	44 ms
mxfp4 (with a validation gate)	44 ms	~12,000 ms

Gating on validation would make MX correctness runs roughly 500x more expensive per solution
than bf16 ones, for a protection bf16 has never had.

It is also worth being precise about what the refresh actually buys. The CPU reference is
computed from the host copy, so if a kernel wrote outside its buffers the GPU result
would diverge and the next solution would report a spurious FAIL. The exposure is a
misattributed failure, not a missed one — a debuggability property rather than a correctness
hole. If we do want that property for MX later, the right implementation is to cache the
swizzled host buffer and re-upload it cheaply, which is effectively what the non-MX path
already does.

So: no validation gate.

Your question did surface a real defect in this PR, though. The reuse keyed only on the
tensor descriptor, while the buffer can hold either the canonical or the swizzled
layout, depending on the picked solution's mxScaleFormat. A NoSwizzle solution can
overwrite the buffer, and a later swizzled solution with an unchanged descriptor would then
reuse the wrong layout — giving incorrect results with no validation failure, since sweeps
run with validation off.

Fixed in #11982, at no measured cost:

	#11963	#11982
validation off	484 ms total / 44 ms per reset	484 ms total / 44 ms per reset
validation on	484 ms total / 44 ms per reset	484 ms total / 44 ms per reset

Validation reports PASSED unchanged.

assistant-librarian · 2026-09-14T20:12:30Z

TheRock Submodule Bump Activity

Newest first

2026-09-14 — Included in TheRock via ROCm/TheRock#8188.

JIRA ID : AIHPBLAS-4756 ## Motivation This change adds regression coverage for the MX swizzle cache behavior introduced and refined around PR #11963. The key issue is that MX scale reuse is not uniform across architectures: - on gfx950, the host-side preswizzled MX scale can be solution-dependent because the swizzled layout is tied to the selected matrix-instruction shape - on gfx1250, the swizzled MX scale layout is derived from problem geometry, so unchanged descriptors can safely reuse previously generated swizzled device data That difference means stale reuse bugs are easy to introduce if the host cache state is not kept truthful. In this path, correctness depends on the interaction between canonical host buffers, preswizzled device buffers, per-solution reinitialization, and the `m_mxSwizzledDescriptor` cache tag. This PR adds direct regression coverage for those rules and makes the control flow easier to reason about. This work builds upon PR #11963 ## Technical Details This PR strengthens both the implementation clarity and the regression net around MX scale swizzle reuse in TensileLite. In `tensilelite/client/src/DataInitialization.cpp`, the existing MX initialization and upload flow is now explicitly described at the decision points that matter for correctness: - The `initializeMXData()` path distinguishes between gfx950 and gfx1250 reuse semantics. - For gfx1250, if the MX data and scale descriptors are unchanged, the generated swizzled scale bytes are fully determined by problem geometry, so the existing `gpuInput.valid` swizzled buffer remains reusable. - For gfx950, descriptor equality alone is not enough because the preswizzled layout also depends on the selected solution’s matrix-instruction shape. Reuse is therefore only safe when the cached swizzled descriptor still matches the current scale descriptor and the preswizzled state has been explicitly revalidated. - The `m_mxSwizzledDescriptor` cache is treated as truth metadata, not just an optimization hint. - When the path falls back to canonical scale upload, the descriptor cache entry is erased before upload so later calls cannot incorrectly assume that `gpuInput.valid` still contains a valid swizzled layout. - When a new swizzled MX scale is generated and copied into `gpuInput.valid`, the descriptor cache is repopulated with the exact scale descriptor so future reuse checks can skip unnecessary host-side swizzling. - When fallback host-side swizzle generation is required, stale cache metadata is cleared before regenerating device state. - The `copySwizzledToGPUBuffer()` path now has regression coverage for the core branch behavior: - if a kernel requests canonical MX scales (`problemType.mxScaleFormat == 0`), the code uploads `cpuInput.valid` and invalidates any cached swizzled descriptor for that tensor - if a preswizzled scale was already produced into `gpuInput.valid`, the device buffer is reused directly - if the cache records that `gpuInput.valid` already contains the correct swizzled bytes for the current descriptor, host-side re-swizzling is skipped - otherwise the code falls back to rebuilding the swizzled layout from canonical host data In `tensilelite/tests/DataInit_test.cpp`, this PR adds runtime-targeted coverage for the architecture split: - `Gfx950RequiresPerSolutionRecomputeForSwizzledMX` - constructs a real `ClientProblemFactory` + `DataInitialization` fixture using valid client-style arguments - verifies that on gfx950, `referenceNeedsPerSolutionRecompute(...)` is false for a canonical solution (`mxScaleFormat == 0`) and true for a swizzled solution (`mxScaleFormat == 1`) - this directly covers the solution-dependent host preswizzle behavior added for gfx950 - `Gfx1250ReusesPreswizzledMXAcrossSolutions` - uses the same production fixture path - verifies that on gfx1250, a swizzled MX solution does not require per-solution recompute, which preserves the original descriptor-based reuse behavior - this ensures the gfx1250 optimization remains tested even though the original behavior previously had no dedicated regression coverage The reasoning for these branches and cache transitions is also captured inline in source comments at the changed control-flow points, so future edits do not have to reconstruct the design intent from behavior alone. ## Test Plan 1. Rebuild the TensileLite C++ test binary with tests enabled: - `cmake --preset tensilelite -S .. -B my-custom-build -DTENSILELITE_BUILD_TESTING=ON -DHIPBLASLT_ENABLE_YAML=OFF` - `cmake --build my-custom-build --parallel --target tensilelite-tests` 2. Run the focused MX runtime regression tests: - `./my-custom-build/tensilelite/tests/tensilelite-tests --gtest_filter='DataInitializationReferenceRecompute.*'` 3. Re-run the targeted Tensile config tests covering the MX YAML surfaces: - `python -m pytest Tensile/Tests/common/test_config.py -q -k 'sk_mx32f4_quick.yaml or subtile_mxfp8_tail_loop.yaml or mxfp4_mxfp4_fp32_tn_act.yaml or sk_mx32f4_force_dp_only.yaml or sk_dynamic_mxfp4.yaml or subtile_mxfp8_tail_loop_smoke.yaml or sk_mxfp4_atomic_variants.yaml or sk_mxfp4_problem_sizes.yaml or mx32f4_tn.yaml or sk_mxfp4_debug_modes.yaml or sk_mxfp4_tree_reduction.yaml or sk_mxfp4_comprehensive.yaml or sk_mxfp4_full_coverage.yaml or sk_mxfp4_basic_variants.yaml or sk_mxf4gemm_pap.yaml or subtile_mxfp4.yaml or mxfp4_mxfp4_bf16_tn_act.yaml'` ## Test Result - `tensilelite-tests` built successfully from the custom build tree with: - `TENSILELITE_BUILD_TESTING=ON` - `HIPBLASLT_ENABLE_YAML=OFF` - Focused MX runtime gtests: - `DataInitializationReferenceRecompute.Gfx950RequiresPerSolutionRecomputeForSwizzledMX` passed on gfx950 - Targeted Tensile config tests: - `17 passed, 398 deselected, 15 warnings` - warnings were `PytestUnknownMarkWarning` for existing custom markers and did not indicate behavioral failures ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

…ning The benchmark client's prepareGPUInputs() copies all GPU tensors on every kernel evaluation for MX problems because two structural flags (m_problemDependentData, needMXSwizzle) prevent the cached-input fast path. After #11963 made the upstream MX data generation a per-problem no-op on gfx1250, those copies are redundant: every evaluation copies identical bytes. Unify the non-MX and gfx1250-MX fast paths into a single inputsUnchanged condition. Non-MX evaluates the same sub-expression as before. gfx1250 MX now also returns cached inputs, since the descriptor caches added in #11963 guarantee the GPU buffers are unchanged between solutions. gfx950 is excluded because its preswizzled MX scale layout depends on the solution's matrix-instruction shape. Measured on gfx1250, mxfp4, [1382, 37324, 1, 65536]: per-eval wall-clock 3.9 s -> 0.09 s (43x) total tuning time ~6 h -> ~8 min (~45x) Co-authored-by: Cursor <cursoragent@cursor.com>

minsukim-amd requested a review from a team as a code owner September 11, 2026 01:24

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Sep 11, 2026

therock-pr-bot Bot added the Not ready to Review label Sep 11, 2026

minsukim-amd force-pushed the users/minsukim/mx_overhead_fix_20260910 branch from e0a5aef to 0d3bbd3 Compare September 11, 2026 01:41

minsukim-amd changed the title ~~Avoid regenerating unchanged MX inputs on every solution~~ fix(hipblaslt): avoid regenerating unchanged MX inputs on every solution Sep 11, 2026

assistant-librarian Bot added the organization: ROCm label Sep 11, 2026

therock-pr-bot Bot removed the Not ready to Review label Sep 11, 2026

minsukim-amd force-pushed the users/minsukim/mx_overhead_fix_20260910 branch from 0d3bbd3 to 06e3d77 Compare September 11, 2026 05:10

minsukim-amd requested review from aazz44ss and jichangjichang September 11, 2026 05:29

aazz44ss approved these changes Sep 11, 2026

View reviewed changes

minsukim-amd merged commit f502f0a into develop Sep 11, 2026
112 of 119 checks passed

minsukim-amd deleted the users/minsukim/mx_overhead_fix_20260910 branch September 11, 2026 13:40

minsukim-amd mentioned this pull request Sep 11, 2026

Harden the MX input reuse added in #11963 #11982

Open

1 task done

amd-ysanchez mentioned this pull request Sep 16, 2026

Enable fix from PR #11963 for gfx950 #12187

Merged

1 task done

pkamd mentioned this pull request Oct 26, 2026

perf(hipblaslt): skip redundant GPU tensor copies for gfx1250 MX tuning #14425

Open

2 of 4 tasks

Navigation Menu

Repository navigation

fix(hipblaslt): avoid regenerating unchanged MX inputs on every solution - #11963

fix(hipblaslt): avoid regenerating unchanged MX inputs on every solution#11963
minsukim-amd merged 1 commit into
developfrom
users/minsukim/mx_overhead_fix_20260910

minsukim-amd commented Sep 11, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 11, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 11, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Sep 11, 2026 •

edited

Loading

Uh oh!

aazz44ss left a comment

Uh oh!

Uh oh!

minsukim-amd commented Sep 11, 2026

Uh oh!

assistant-librarian Bot commented Sep 14, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

minsukim-amd commented Sep 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Sep 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Sep 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

aazz44ss left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

minsukim-amd commented Sep 11, 2026

Q1. Do you know why it is regenerating unchanged MX inputs every solution?

Q2. Do we need to leave the validation != 0 case unchanged?

Uh oh!

assistant-librarian Bot commented Sep 14, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

minsukim-amd commented Sep 11, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 11, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 11, 2026 •

edited

Loading

codecov Bot commented Sep 11, 2026 •

edited

Loading

Q2. Do we need to leave the `validation != 0` case unchanged?