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
Merged

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

Show file tree
Hide file tree
Changes from 1 commit
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
46 changes: 46 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,7 @@
#include <gtest/gtest.h>
#include <limits>
#include <memory>
#include <utility>

#include <hip/hip_runtime.h>

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

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

using namespace TensileLite;
using namespace TensileLite::testing;
Expand Down Expand Up @@ -900,6 +902,32 @@ 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)
{
return checkSolution(solution, problem, false);
}

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 +1028,24 @@ 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);

ContractionSolution dynamicSolution;
initEquality512Solution(dynamicSolution, 4);
EXPECT_FALSE(iterator.accepts(dynamicSolution, problem))
<< "The explicit all-solutions client path must not launch an eight-queue "
"dynamic kernel on a six-XCD MI300A";

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

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