[Alex-Vasile]

JIRA ID : AIHPBLAS-5257

Motivation

tensilelite-client, the C++ program that runs and validates generated GPU kernels, checks each solution against a CPU reference, the CPU computation of the expected result. For MX problems (microscaling formats, in which a block of values shares one scale) on gfx950, some solutions need their inputs regenerated with a different scale layout. ReferenceValidator recomputed the reference for every solution that could need new inputs, using a predicate (referenceNeedsPerSolutionRecompute()). DataInitialization, which owns the inputs, usually reused them from its cache. On subtile_mxfp4, 685 of 9,949 solutions regenerated anything, but the reference was recomputed for all 9,949, about 17 s of CPU time.

Technical Details

Production code, client only:

• DataInitialization keeps an input epoch, a counter that it increments wherever it rewrites the host inputs the reference reads. Those places are initializeCPUInputs, initTensorFromDefault, the grouped-GEMM MX path and a real regeneration in initOneMXSide. inputEpoch(solution) throws std::logic_error unless preSolution() already ran for that solution, so a listener-order mistake fails instead of giving a stale reference.
• ReferenceValidator stores the epoch each reference was computed with. In preSolution it recomputes, for any GEMM problem, only when the epoch moved. It no longer calls referenceNeedsPerSolutionRecompute(), which stays in the header until a later change removes it.
• A new client option --verify-reference-reuse, or TENSILELITE_VERIFY_REFERENCE_REUSE=1, recomputes anyway when the epoch did not move. It throws "Stale CPU reference" unless the result is bitwise identical, which catches a missed epoch increment.
• Timers (no-ops unless timing instrumentation is on) separate needed from redundant MX regeneration and reference recomputes. scripts/analyze_timing.py knows the new names.

Behavior change: on gfx950 MX runs, the reference is recomputed only for solutions whose inputs were rewritten. GPU work and results do not change.

Tests: DataInitializationInputEpoch.PerSolutionEpochRequiresPreSolution replaces the gfx950 predicate test. It runs only on gfx950 or gfx1250 hardware. ReferenceValidator_test.cpp calls DataInitialization::preSolution before the validator, as main does.

Test Plan

From projects/hipblaslt/tensilelite. C++ tests use one configured build: cmake --preset tensilelite -S .. -B <build> -DTENSILELITE_BUILD_TESTING=ON and cmake --build <build> --parallel.

• ctest --test-dir <build>/tensilelite/tests --output-on-failure -R 'DataInit|ReferenceValidator' on a gfx950 GPU checks the epoch rule and the validator's listener order.
• TENSILELITE_VERIFY_REFERENCE_REUSE=1 tox -e py3 -- tensilelite/Tests -m common -k 'subtile_mxfp8' on a gfx950 GPU checks that every reused reference equals a fresh one.
• pre-commit run --from-ref HEAD~1 --to-ref HEAD runs the repository hooks on this PR's files.

Test Result

Time saved: not measured in isolation. The closest evidence is the change's own commit on the original unsplit branch: before it, subtile_mxfp4 on gfx950 recomputed the CPU reference for all 9,949 solutions, about 17 s of CPU time, while only 685 of them had rewritten inputs. The gfx950 baseline was taken with this change already present, so the before/after full-run comparisons do not include it.

Other results, from the original unsplit branch on gfx950:

• subtile_mxfp4 in verify mode reused 9,264 references, each bitwise identical to a fresh recompute.
• The 33 gfx950 MX common configs passed with verification enabled.

Not yet run on this branch. Automated GitHub checks are pending.

Related

• Work tracking: AIHPBLAS-5257

Device / Architecture Coverage

Client host code (CPU reference scheduling); no kernel generation changes. The behavior change applies to MX problems on gfx950. gfx1250 MX runs also use the epoch. The gfx950 MX common configs need a passing CI run, ideally once with TENSILELITE_VERIFY_REFERENCE_REUSE=1. gfx1250 is uncovered.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Risk level

Moderate (3/5): validation is only as correct as the epoch increments. A host-input writer that misses one would let a solution be checked against a stale reference. Verify mode detects that, but CI does not enable it by default.

[therock-pr-bot]

❌ PR Check — Action Required

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	❌ Fail	Error: Check concluded with cancelled.
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

• ❌ pre-commit

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

🙋 Wish to Override Policy?

[therock-pr-bot]

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