Skip to content

fix(tensilelite): filter unsupported StreamK benchmark kernels #11872

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
davidd-amd merged 2 commits into from
Sep 10, 2026
+79 −1
Merged

fix(tensilelite): filter unsupported StreamK benchmark kernels #11872

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
44e042f
fix(tensilelite): filter unsupported StreamK benchmark kernels
davidd-amd Sep 9, 2026
f00216d
test(tensilelite): cover topology rejection reporting
davidd-amd Sep 9, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions projects/hipblaslt/tensilelite/client/include/LogReporter.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -216,6 +216,8 @@ namespace TensileLite
m_rowLevel = LogLevel::Terse;
else if(value == "DID_NOT_SATISFY_ASSERTS")
m_rowLevel = LogLevel::Terse;
else if(value == "UNSUPPORTED_XCD_TOPOLOGY")
m_rowLevel = LogLevel::Terse;
else if(value == "INVALID")
{
m_rowLevel = LogLevel::Error;
Expand Down
14 changes: 14 additions & 0 deletions projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -104,6 +104,20 @@ namespace TensileLite
return false;
}

// The all-solutions benchmark path selects kernels by index instead

@talumbau talumbau Sep 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 reports DID_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 single warnStreamKDynamicQueueUnsupportedOnce message 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.

// of going through a SolutionLibrary search. Apply the dynamic
// StreamK topology predicate explicitly so this path has the same
// support boundary as normal library selection. In particular, a
// 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.
Comment thread
davidd-amd marked this conversation as resolved.
Resolved
if(!solution.streamKDynamicQueueSupported(problem, *m_hardware))

@talumbau talumbau Sep 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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_error

HipHardware.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.

{
if(isReportValid)
m_reporter->report(ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY");
return false;
}

// Test if the persistent kernel is eligible for the current hw and solution
problem.checkPersistentKernelEligibility(solution, *m_hardware);
Task task(*m_hardware, problem, solution);
Expand Down
4 changes: 3 additions & 1 deletion projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -732,7 +732,9 @@ namespace TensileLite
// dynamic-queue path but the hardware's NUM_XCD is not a power of two
// (e.g. MI300A = 6), so the solution is EXCLUDED from selection rather
// than silently degraded to tree reduction. All other solutions return
// true. Wired into softwarePredicate() (SolutionLibrary.hpp).
// true. Wired into softwarePredicate() (SolutionLibrary.hpp) and the
// client's explicit all-solutions iterator, which otherwise bypasses
// library selection.
bool streamKDynamicQueueSupported(Problem const& problem,
Hardware const& hardware) const;
// Selection-time filter for uniform summation order. Resolves StreamK /
Expand Down
60 changes: 60 additions & 0 deletions projects/hipblaslt/tensilelite/tests/CuCount_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,8 @@
#include <gtest/gtest.h>
#include <limits>
#include <memory>
#include <sstream>
#include <utility>

#include <hip/hip_runtime.h>

Expand All @@ -19,6 +21,8 @@
#include <origami/streamk.hpp>

#include "FallbackTestUtils.hpp"
#include "LogReporter.hpp"
#include "SolutionIterator.hpp"

using namespace TensileLite;
using namespace TensileLite::testing;
Expand Down Expand Up @@ -900,6 +904,34 @@ namespace
makeGfx942HardwareWithXcd(numXCD));
return device;
}

class TestSolutionIterator : public Client::SolutionIterator
{
public:
explicit TestSolutionIterator(std::shared_ptr<Hardware> hardware)
: SolutionIterator(nullptr, std::move(hardware), false)
{
}

bool accepts(ContractionSolution& solution,
ContractionProblemGemm& problem,
bool reportResult)
{
return checkSolution(solution, problem, reportResult);
}

void postProblem() override {}
void preSolution(ContractionSolution* const) override {}
void postSolution() override {}
bool moreSolutionsInProblem() const override
{
return false;
}
std::shared_ptr<ContractionSolution> getSolution() override
{
return nullptr;
}
};
} // namespace

TEST(StreamKDynamicQueueXcdGateTest, RejectsMi300aSixXcd)
Expand Down Expand Up @@ -1000,6 +1032,34 @@ TEST(StreamKDynamicQueueXcdGateTest, KeepsNonDynamicQueueSolutionsOnMi300a)
<< "SK3-static solution must remain selectable on MI300A";
}

TEST(StreamKDynamicQueueXcdGateTest, ClientIteratorFiltersOnlyUnsupportedDynamicQueue)
{
auto mi300a = std::make_shared<hip::HipAMDGPU>(makeGfx942DeviceWithXcd(6));
TestSolutionIterator iterator(mi300a);
auto problem = makeGemmProblem(512, 512, 512);
std::ostringstream reportOutput;

@talumbau talumbau Sep 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

auto reporter = std::make_shared<Client::LogReporter>(
Client::LogLevel::Terse,
std::initializer_list<std::string>{Client::ResultKey::Validation},
reportOutput,
false,
false);
iterator.setReporter(reporter);

ContractionSolution dynamicSolution;
initEquality512Solution(dynamicSolution, 4);
EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, true))

@talumbau talumbau Sep 10, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 true call and the false call assert the same observable value. The test does not check the report.
  • If you delete the new else if(value == "UNSUPPORTED_XCD_TOPOLOGY") branch in LogReporter.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.

<< "The explicit all-solutions client path must not launch an eight-queue "
"dynamic kernel on a six-XCD MI300A";
EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, false))
<< "Topology filtering must also apply during non-reporting prediction checks";

ContractionSolution staticSolution;
initEquality512Solution(staticSolution, 3);
EXPECT_TRUE(iterator.accepts(staticSolution, problem, true))
<< "The topology guard must retain static StreamK coverage on MI300A";
}

// ===========================================================================
// StreamKFlagBound -- the bound getSKGrid puts on a Stream-K grid.
//
Expand Down
Loading