Repository navigation
fix(tensilelite): filter unsupported StreamK benchmark kernels - #11872
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11872 +/- ##
===========================================
+ Coverage 69.42% 69.43% +0.01%
===========================================
Files 2828 2828
Lines 470113 470119 +6
Branches 69417 69419 +2
===========================================
+ Hits 326343 326399 +56
+ Misses 120066 120006 -60
- Partials 23704 23714 +10
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
tony-davis
left a comment
tony-davis
left a comment
There was a problem hiding this comment.
Reviewed the fix and dug into the CI results directly on Math CI.
The change makes sense: the all-solutions benchmark path selects kernels by index and bypasses the normal library search, so it was skipping the streamKDynamicQueueSupported() predicate that gates dynamic StreamK kernels by XCD count. Adding that check explicitly in checkSolution() closes that gap, and it's a low-risk change since it just reuses an existing predicate and only affects the benchmark client. The new gtest (ClientIteratorFiltersOnlyUnsupportedDynamicQueue) covers it well: dynamic StreamK gets rejected on a 6-XCD MI300A, static StreamK still runs.
On CI: the preliminary(hipblaslt) build shows as failed overall, but I traced it stage by stage. The gfx942 Test stage (the lane this PR is actually about) passed. The failure is in the gfx12 Test stage, and it's hipErrorOutOfMemory from DataInitialization.cpp on 4 unrelated test configs, happening while pytest-xdist workers share 2 GPUs on that runner. That's a resource-contention flake, not something this diff could have caused (the change only touches StreamK/XCD topology logic).
Approving. Nice fix.
|
Overriding - this fixes a test flake but is being held up by other unrelated flakes. |
talumbau
left a comment
talumbau
left a comment
There was a problem hiding this comment.
Review of the StreamK topology filter for the benchmark client.
The approach is correct. The client selects kernels by index and does no SolutionLibrary search, so it did not apply the topology rule. This PR applies the predicate that the library already uses, and it changes an error result into a quiet skip. That is better than the old behavior, where solve() threw and exit-on-error could stop the run.
I have four items. One removes test coverage that the PR description says is present. Three are hardening.
- The new test does not check the report output. The
LogReporterbranch has no coverage. You can delete that branch and the test is still successful. - The new predicate call can throw an exception at a call site that
main.cppdoes not protect. - The order of the checks is not the same as
softwarePredicate, so a rejection reason can be wrong in the tuning log. - The test declares the stream after the object that points to it.
Please see the inline comments for the code and the suggested changes.
Items that I checked and that are not problems: the missing uniformSummationOrderSupported clause cannot matter here, because it gives true on its first line unless setUniformSummationOrder was called, and the only callers are in rocblaslt/tensile_host.cpp. The gate also does not lose tuning data or move the CSV columns: ProgressListener::preSolution still writes the column and the -1.0 value before runCurrentSolution() runs.
| // gfx942 MI300A reports six XCDs while these kernels bake eight | ||
| // per-XCD queues; launching one is invalid, but static StreamK and | ||
| // non-StreamK solutions in the same config remain runnable. | ||
| if(!solution.streamKDynamicQueueSupported(problem, *m_hardware)) |
There was a problem hiding this comment.
The new call can throw an exception at a call site that is not protected.
streamKDynamicQueueUnsupported() gives true when analyticalHardware is nullptr (ContractionSolution.cpp:145). Control then goes on to streamK5EffectiveDynamic() for an SK5 solution. That function throws at ContractionSolution.cpp:4568:
TENSILE_ASSERT_EXC(hipAMDGPU != nullptr);
TENSILE_ASSERT_EXC(hipAMDGPU->analyticalHardware != nullptr); // std::runtime_errorHipHardware.cpp:50 sets analyticalHardware only when origami knows the hardware. On other hardware the member stays nullptr. The test TriStateAutoRequiresAnalyticalHardware in this same file already asserts that throw.
client/main.cpp:1405 calls runCurrentSolution() in the condition of the if statement. The try block starts inside the body. Therefore the exception goes out of the loop and stops the full benchmark run. The old code never called streamK5EffectiveDynamic() from this path.
The conditions are narrow: an origami-unknown architecture, an SK5 solution, and AUTO mode or smCountTarget > 0.
Suggested change in ContractionSolution::streamKDynamicQueueSupported():
if(!streamKDynamicQueueUnsupported(hardware))
return true;
// Unknown hardware: streamK5EffectiveDynamic() needs analyticalHardware and
// throws without it. The policy for unknown hardware is UNSUPPORTED.
auto const* hipAMDGPU = dynamic_cast<hip::HipAMDGPU const*>(&hardware);
if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr)
{
warnStreamKDynamicQueueUnsupportedOnce(hardware);
return false;
}
const bool dynamicQueue = ...Why this is correct. The exception is possible only when analyticalHardware is nullptr. The guard stops this one condition, so the predicate always gives a bool. The answer does not change, because the comment in the same function already says that unknown hardware is UNSUPPORTED. The fix is in ContractionSolution, so it also protects the existing softwarePredicate call site.
| return false; | ||
| } | ||
|
|
||
| // The all-solutions benchmark path selects kernels by index instead |
There was a problem hiding this comment.
The order of the checks is not the same as the softwarePredicate that this comment names.
SolutionLibrary.hpp:91 uses the opposite order. It checks the asserts first and the topology rule third:
return (*solutions.problemPredicate)(problem) && (*solutions.taskPredicate)(task)
&& solutions.streamKDynamicQueueSupported(problem, hardware)
&& solutions.uniformSummationOrderSupported(problem, hardware);Here the new block is before the problemPredicate / taskPredicate block at line 124. There are two results:
- A solution that fails the asserts and the topology rule now reports
UNSUPPORTED_XCD_TOPOLOGY. The library reportsDID_NOT_SATISFY_ASSERTS. In a tuning log the user reads that the GPU part is wrong, but the real cause is a macro-tile or assert mismatch. This also uses the singlewarnStreamKDynamicQueueUnsupportedOncemessage on a solution that was never usable. - The predicate runs on solutions that the assert check rejects one line later. This makes the exception above possible for more solutions.
Suggested change — move the new block below the assert block:
if(!(*solution.problemPredicate)(problem) || !(*solution.taskPredicate)(task))
{
...
return false;
}
// Use the same clause order as softwarePredicate() in SolutionLibrary.hpp.
if(!solution.streamKDynamicQueueSupported(problem, *m_hardware))
{
if(isReportValid)
m_reporter->report(ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY");
return false;
}Why this is correct and safe. The two checks are independent. streamKDynamicQueueSupported reads only sizeMapping.streamK, the XCD counts, and the problem sizes. It does not read m_eligibleForPK. softwarePredicate proves this: it calls the predicate but never calls checkPersistentKernelEligibility (line 122 here is the only call site in the tree). Therefore the set of accepted solutions stays the same. Only the report value for a doubly-failing solution changes, and it changes to the library value. The new order is also less expensive.
Also: please make this comment shorter. The same MI300A text is in ContractionSolution.hpp:728, which this PR also edits, and again in the streamKDynamicQueueSupportedRef mirror in CuCount_test.cpp. Three copies become different with time. One line that points to the canonical text is enough.
|
|
||
| ContractionSolution dynamicSolution; | ||
| initEquality512Solution(dynamicSolution, 4); | ||
| EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, true)) |
There was a problem hiding this comment.
These two assertions check the same result, and the LogReporter half of the PR has no coverage.
report() writes only to a map in memory. Only CSVStackFile::writeRow writes to the stream, and only LogReporter::postSolution() calls writeRow. This test never calls postSolution(). Therefore reportOutput is always empty, and it is never read again in this file.
Two results:
- The
truecall and thefalsecall assert the same observable value. The test does not check the report. - If you delete the new
else if(value == "UNSUPPORTED_XCD_TOPOLOGY")branch inLogReporter.hpp:219, this test is still successful.
This matters because acceptValidation has no else branch and m_rowLevel (LogReporter.hpp:478) is set only in preSolution. A wrong or renamed string keeps the level of the previous row. The compiler shows no error and no test fails.
The PR description says the test "verifies both reporting and non-reporting checks". The report path is not checked.
Suggested change — record the values with a small double on the pure virtual hook ResultReporter::reportValue_string (ResultReporter.hpp:185):
class RecordingReporter : public Client::ResultReporter
{
public:
std::vector<std::pair<std::string, std::string>> reported;
void reportValue_string(std::string const& key, std::string const& value) override
{
reported.emplace_back(key, value);
}
// stub the other pure virtual functions
};auto reporter = std::make_shared<RecordingReporter>();
iterator.setReporter(reporter);
EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, /*reportResult=*/true));
ASSERT_EQ(reporter->reported.size(), 1u);
EXPECT_EQ(reporter->reported[0].second, "UNSUPPORTED_XCD_TOPOLOGY");
reporter->reported.clear();
EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, /*reportResult=*/false));
EXPECT_TRUE(reporter->reported.empty());Why this works. The full row path is difficult to drive, because postSolution() calls std::stof on an empty TimeUS field and throws. The virtual hook gives the value without that path. After the change the two calls have different results: if a person moves the guard under if(isReportValid), the second check fails; if a person changes the string in SolutionIterator.cpp:117, the first check fails.
| auto mi300a = std::make_shared<hip::HipAMDGPU>(makeGfx942DeviceWithXcd(6)); | ||
| TestSolutionIterator iterator(mi300a); | ||
| auto problem = makeGemmProblem(512, 512, 512); | ||
| std::ostringstream reportOutput; |
There was a problem hiding this comment.
The stream is destroyed before the object that points to it.
iterator is declared at line 1038 and reportOutput at line 1040. setReporter copies the shared_ptr into the iterator, so the iterator keeps the LogReporter alive. C++ destroys local objects in the opposite order of the declarations: reporter, then reportOutput, then problem, then iterator. The LogReporter is alive after the stream is gone.
LogReporter keeps std::ostream& m_stream (LogReporter.hpp:474), and its CSVStackFile keeps a pointer to the same object with a null_deleter.
There is no fault today, because the destructors are empty and nothing writes to the stream at destruction. But the fault becomes real if a person adds a flush operation to the destructor, and an ASAN build can show it now.
Suggested change — declare the stream first:
auto mi300a = std::make_shared<hip::HipAMDGPU>(makeGfx942DeviceWithXcd(6));
std::ostringstream reportOutput; // moved up
TestSolutionIterator iterator(mi300a);
auto problem = makeGemmProblem(512, 512, 512);Why this works. The destruction order becomes reporter, problem, iterator, reportOutput. The iterator releases the last reference to the LogReporter before the stream goes away. Therefore the stream is alive longer than each object that points to it, and it stays that way after any change to the destructor. The cost is one moved line.
If you take the change in the comment below, this problem also goes away, because that change removes the stream and the LogReporter from the test.
Separate remark on TestSolutionIterator (added above at line ~907): it writes stub bodies for five pure virtual functions only to reach the protected checkSolution. If a person adds a pure virtual function to RunListener or SolutionIterator, this StreamK test file fails to build for an unrelated reason.
TheRock Submodule Bump ActivityNewest first
|
Follow-up to review comments left by @talumbau on #11872 after it merged, carried by #11946. The previous guard rejected every SK5 solution as soon as a device had no analytical hardware, which is broader than the condition it was protecting against. Only the origami hybrid-mode heuristic needs analytical hardware; a mode that resolves without it was being discarded anyway. Drop the blanket early-out and let mode resolution decide. A mode that resolves -- explicit static, explicit dynamic, or a forced mode from TENSILE_STREAMK5_FORCE_MODE -- is honoured: static stays selectable, dynamic is excluded. A mode that cannot resolve (AUTO, or OFF with sm_count_target > 0, both of which enter the heuristic) throws rather than silently substituting a schedule the caller did not ask for. This also makes the grouped-GEMM workspace path, which reaches mode resolution before the topology guard runs, behave the same way as the single-GEMM path instead of diverging from it. This narrows the guard that #11946 added from the suggestion on SolutionIterator.cpp:114. That suggestion rested on "the answer does not change, because unknown hardware is already UNSUPPORTED", which does not hold for the modes that resolve without analytical hardware: they never reach the throwing heuristic, yet the guard rejected them too. Unknown analytical hardware only arises when origami reports the device as unsupported, and SK5 kernels are only built for supported architectures, so the throwing path is defensive rather than reachable in production. Tests: the old unknown-hardware test was vacuous whenever TENSILE_STREAMK5_FORCE_MODE was set in the environment, because that variable short-circuits mode resolution before the code under test runs. Pin it with a ScopedStreamK5ForceMode guard and assert the pinned value, then split the cases so each resolution outcome is checked on its own. Drop the postProblem() call from the LogReporter test, which printf()s a banner to stdout rather than to the captured stream, and add the missing ResultReporter.hpp/<cstdint>/<cstdlib>/<stdexcept> includes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
ISSUE ID : #11890
Motivation
The TensileLite benchmark client explicitly iterates generated GPU kernels to validate each candidate. Normal library selection rejects dynamic StreamK kernels when the device's runtime compute-die count does not match the queue count compiled into the kernel, but explicit benchmark iteration did not apply that check. On a six-XCD MI300A reported as gfx942, the client therefore attempted to launch kernels compiled with eight queues and failed the affected StreamK YAML tests.
Technical Details
streamKDynamicQueueSupportedpredicate before the benchmark client launches an explicitly selected solution.UNSUPPORTED_XCD_TOPOLOGYwithout counting them as execution failures. The existing runtime warning includes the detected XCD count and compiled queue count.No YAML tests or gfx942 architecture coverage are disabled.
Test Plan
tensilelite-testswithTENSILELITE_BUILD_TESTING=ON.tensilelite-tests --gtest_filter=StreamKDynamicQueueXcdGateTest.*.Test Result
StreamKDynamicQueueXcdGateTest.*: 8 passed.Submission Checklist
Risk level
Low (2/5): the change affects only benchmark-client candidate filtering. It uses the predicate already applied by normal library selection and does not change kernel generation or the hipBLASLt library interface.
Fixes #11890