From e71edb799d6a8b281b73aca3d80d931c7a2baf54 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Thu, 10 Sep 2026 19:40:42 +0000 Subject: [PATCH 1/3] fix(tensilelite): align topology filtering with selection --- .../client/src/SolutionIterator.cpp | 23 ++--- .../tensilelite/src/ContractionSolution.cpp | 10 ++ .../tensilelite/tests/CuCount_test.cpp | 98 +++++++++++++++++-- 3 files changed, 110 insertions(+), 21 deletions(-) diff --git a/projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp b/projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp index c5a1cd169a50..efecc6107b99 100644 --- a/projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp +++ b/projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp @@ -104,20 +104,6 @@ namespace TensileLite return false; } - // The all-solutions benchmark path selects kernels by index instead - // 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. - if(!solution.streamKDynamicQueueSupported(problem, *m_hardware)) - { - 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); @@ -139,6 +125,15 @@ namespace TensileLite return false; } + // Explicit index selection bypasses SolutionLibrary; apply its + // StreamK topology predicate after the same problem/task checks. + if(!solution.streamKDynamicQueueSupported(problem, *m_hardware)) + { + if(isReportValid) + m_reporter->report(ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY"); + return false; + } + if(solution.requiredHostWorkspaceSizePerProblem == static_cast(-1)) { solution.requiredHostWorkspaceSizePerProblem diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index e42688a5256b..d314346b7008 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -4611,6 +4611,16 @@ namespace TensileLite if(!streamKDynamicQueueUnsupported(hardware)) return true; + // streamK5EffectiveDynamic() needs analytical hardware for AUTO mode. + // Unknown hardware is already classified as unsupported, so return the + // policy result instead of letting mode resolution throw. + auto const* hipAMDGPU = dynamic_cast(&hardware); + if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr) + { + warnStreamKDynamicQueueUnsupportedOnce(hardware); + return false; + } + // Runtime XCD count is not a power of two or does not equal the baked // per-XCD queue count. Only the dynamic-queue sub-path is affected: an // SK5 solution that resolves to the static (SK3) sub-path for this diff --git a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp index b25156b1bd9a..d50d0b12c5ab 100644 --- a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp @@ -5,7 +5,9 @@ #include #include #include +#include #include +#include #include @@ -932,6 +934,26 @@ namespace return nullptr; } }; + + class RecordingReporter : public Client::ResultReporter + { + public: + std::vector> reportedStrings; + + void reportValue_string(std::string const& key, std::string const& value) override + { + reportedStrings.emplace_back(key, value); + } + void reportValue_uint(std::string const&, uint64_t) override {} + void reportValue_int(std::string const&, int64_t) override {} + void reportValue_double(std::string const&, double) override {} + void reportValue_sizes(std::string const&, std::vector const&) override {} + void reportValue_vecOfSizes(std::string const&, + std::vector> const&) override + { + } + void finalizeReport() override {} + }; } // namespace TEST(StreamKDynamicQueueXcdGateTest, RejectsMi300aSixXcd) @@ -996,6 +1018,25 @@ TEST(StreamKDynamicQueueXcdGateTest, MissingAnalyticalHardwareIsUnsupported) << "SK3-static solution must remain selectable when NUM_XCD is unknown"; } +TEST(StreamKDynamicQueueXcdGateTest, Sk5AutoUnknownHardwareIsRejectedWithoutThrowing) +{ + hip::HipAMDGPU noAnalytical; + noAnalytical.processor = AMDGPU::Processor::gfx942; + noAnalytical.computeUnitCount = 304; + noAnalytical.deviceName = "test-gfx942-no-analytical"; + + ContractionSolution hybridSolution; + initEquality512Solution(hybridSolution, 5); + auto problem = makeGemmProblem(512, 512, 512); + problem.setParams().setStreamKTileSchedulingMode(2); // AUTO requires analytical hardware. + + bool supported = true; + EXPECT_NO_THROW( + supported = hybridSolution.streamKDynamicQueueSupported(problem, noAnalytical)); + EXPECT_FALSE(supported) + << "Unknown hardware must be rejected without entering the SK5 AUTO heuristic"; +} + // Selection-predicate contract: on MI300A (6 XCD) the dynamic-queue solution is // EXCLUDED from selection (supported == false) so a different solution serves // the GEMM, while on MI300X (8 XCD) the identical solution stays selectable. @@ -1037,13 +1078,7 @@ TEST(StreamKDynamicQueueXcdGateTest, ClientIteratorFiltersOnlyUnsupportedDynamic auto mi300a = std::make_shared(makeGfx942DeviceWithXcd(6)); TestSolutionIterator iterator(mi300a); auto problem = makeGemmProblem(512, 512, 512); - std::ostringstream reportOutput; - auto reporter = std::make_shared( - Client::LogLevel::Terse, - std::initializer_list{Client::ResultKey::Validation}, - reportOutput, - false, - false); + auto reporter = std::make_shared(); iterator.setReporter(reporter); ContractionSolution dynamicSolution; @@ -1051,13 +1086,62 @@ TEST(StreamKDynamicQueueXcdGateTest, ClientIteratorFiltersOnlyUnsupportedDynamic EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, true)) << "The explicit all-solutions client path must not launch an eight-queue " "dynamic kernel on a six-XCD MI300A"; + ASSERT_EQ(reporter->reportedStrings.size(), 1u); + EXPECT_EQ(reporter->reportedStrings[0].first, Client::ResultKey::Validation); + EXPECT_EQ(reporter->reportedStrings[0].second, "UNSUPPORTED_XCD_TOPOLOGY"); + + reporter->reportedStrings.clear(); EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, false)) << "Topology filtering must also apply during non-reporting prediction checks"; + EXPECT_TRUE(reporter->reportedStrings.empty()); ContractionSolution staticSolution; initEquality512Solution(staticSolution, 3); EXPECT_TRUE(iterator.accepts(staticSolution, problem, true)) << "The topology guard must retain static StreamK coverage on MI300A"; + EXPECT_TRUE(reporter->reportedStrings.empty()); +} + +TEST(StreamKDynamicQueueXcdGateTest, ClientIteratorReportsAssertsBeforeTopology) +{ + auto mi300a = std::make_shared(makeGfx942DeviceWithXcd(6)); + TestSolutionIterator iterator(mi300a); + auto problem = makeGemmProblem(512, 512, 512); + auto reporter = std::make_shared(); + iterator.setReporter(reporter); + + ContractionSolution dynamicSolution; + initEquality512Solution(dynamicSolution, 4); + dynamicSolution.problemPredicate + = std::make_shared>(); + + EXPECT_FALSE(iterator.accepts(dynamicSolution, problem, true)); + ASSERT_EQ(reporter->reportedStrings.size(), 1u); + EXPECT_EQ(reporter->reportedStrings[0].first, Client::ResultKey::Validation); + EXPECT_EQ(reporter->reportedStrings[0].second, "DID_NOT_SATISFY_ASSERTS"); +} + +TEST(StreamKDynamicQueueXcdGateTest, LogReporterTreatsUnsupportedTopologyAsTerse) +{ + std::ostringstream reportOutput; + Client::LogReporter reporter( + Client::LogLevel::Terse, + std::initializer_list{Client::ResultKey::Validation, + Client::ResultKey::TimeUS, + Client::ResultKey::SpeedGFlops}, + reportOutput, + false, + false); + + reporter.preProblem(nullptr); + reporter.preSolution(nullptr); + reporter.report(Client::ResultKey::TimeUS, 1.0); + reporter.report(Client::ResultKey::SpeedGFlops, 0.0); + reporter.report(Client::ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY"); + reporter.postSolution(); + reporter.postProblem(); + + EXPECT_NE(reportOutput.str().find("UNSUPPORTED_XCD_TOPOLOGY"), std::string::npos); } // =========================================================================== From 5602a27106d9d161cdf4d6b786a07bad2f20aaec Mon Sep 17 00:00:00 2001 From: David Dixon Date: Fri, 11 Sep 2026 20:26:28 +0000 Subject: [PATCH 2/3] fix(tensilelite): honour the resolved SK5 mode on unknown hardware 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/// includes. Co-Authored-By: Claude Opus 5 --- .../tensilelite/src/ContractionSolution.cpp | 10 -- .../tensilelite/tests/CuCount_test.cpp | 107 ++++++++++++++++-- 2 files changed, 99 insertions(+), 18 deletions(-) diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index d314346b7008..e42688a5256b 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -4611,16 +4611,6 @@ namespace TensileLite if(!streamKDynamicQueueUnsupported(hardware)) return true; - // streamK5EffectiveDynamic() needs analytical hardware for AUTO mode. - // Unknown hardware is already classified as unsupported, so return the - // policy result instead of letting mode resolution throw. - auto const* hipAMDGPU = dynamic_cast(&hardware); - if(hipAMDGPU == nullptr || hipAMDGPU->analyticalHardware == nullptr) - { - warnStreamKDynamicQueueUnsupportedOnce(hardware); - return false; - } - // Runtime XCD count is not a power of two or does not equal the baked // per-XCD queue count. Only the dynamic-queue sub-path is affected: an // SK5 solution that resolves to the static (SK3) sub-path for this diff --git a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp index d50d0b12c5ab..6bee2b6ecac3 100644 --- a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp @@ -2,9 +2,12 @@ // SPDX-License-Identifier: MIT #include +#include +#include #include #include #include +#include #include #include #include @@ -24,6 +27,7 @@ #include "FallbackTestUtils.hpp" #include "LogReporter.hpp" +#include "ResultReporter.hpp" #include "SolutionIterator.hpp" using namespace TensileLite; @@ -442,6 +446,43 @@ namespace return pack; } + // TENSILE_STREAMK5_FORCE_MODE short-circuits SK5 mode resolution, so any + // test that exercises resolution must pin it rather than inherit it from + // the environment or from a sibling test in this binary. + class ScopedStreamK5ForceMode + { + public: + explicit ScopedStreamK5ForceMode(const char* value) + { + const char* prior = std::getenv("TENSILE_STREAMK5_FORCE_MODE"); + m_had = (prior != nullptr); + if(m_had) + m_saved = prior; + + if(value) + setenv("TENSILE_STREAMK5_FORCE_MODE", value, /*overwrite=*/1); + else + unsetenv("TENSILE_STREAMK5_FORCE_MODE"); + Debug::Instance().reloadDebugBitsForTest(); + } + + ~ScopedStreamK5ForceMode() + { + if(m_had) + setenv("TENSILE_STREAMK5_FORCE_MODE", m_saved.c_str(), /*overwrite=*/1); + else + unsetenv("TENSILE_STREAMK5_FORCE_MODE"); + Debug::Instance().reloadDebugBitsForTest(); + } + + ScopedStreamK5ForceMode(ScopedStreamK5ForceMode const&) = delete; + ScopedStreamK5ForceMode& operator=(ScopedStreamK5ForceMode const&) = delete; + + private: + bool m_had = false; + std::string m_saved; + }; + void initEquality512Solution(ContractionSolution& solution, int streamK) { solution.sizeMapping.streamK = streamK; @@ -1018,8 +1059,60 @@ TEST(StreamKDynamicQueueXcdGateTest, MissingAnalyticalHardwareIsUnsupported) << "SK3-static solution must remain selectable when NUM_XCD is unknown"; } -TEST(StreamKDynamicQueueXcdGateTest, Sk5AutoUnknownHardwareIsRejectedWithoutThrowing) +// Unknown hardware and SK5: resolving the sub-mode is what decides whether the +// solution is selectable, so the three outcomes are tested separately. +// +// Policy: a mode that RESOLVES without analytical hardware is honoured -- +// explicitly static stays selectable, explicitly dynamic is excluded. A mode +// that cannot resolve (AUTO, or OFF with smCountTarget > 0, both of which enter +// the origami heuristic) throws rather than silently substituting a schedule. +TEST(StreamKDynamicQueueXcdGateTest, Sk5UnknownHardwareHonoursResolvedMode) +{ + ScopedStreamK5ForceMode noForce(nullptr); + ASSERT_EQ(Debug::Instance().streamK5ForceMode(), -1) + << "TENSILE_STREAMK5_FORCE_MODE must be unset or this test proves nothing"; + + hip::HipAMDGPU noAnalytical; + noAnalytical.processor = AMDGPU::Processor::gfx942; + noAnalytical.computeUnitCount = 304; + noAnalytical.deviceName = "test-gfx942-no-analytical"; + ASSERT_EQ(noAnalytical.analyticalHardware, nullptr); + + ContractionSolution hybridSolution; + initEquality512Solution(hybridSolution, 5); + + auto problemOff = makeGemmProblem(512, 512, 512); + problemOff.setParams().setStreamKTileSchedulingMode(0); + problemOff.setParams().setSmCountTarget(0); + EXPECT_TRUE(hybridSolution.streamKDynamicQueueSupported(problemOff, noAnalytical)) + << "SK5 resolving to the static sub-path must stay selectable when NUM_XCD is unknown"; + + auto problemOn = makeGemmProblem(512, 512, 512); + problemOn.setParams().setStreamKTileSchedulingMode(1); + EXPECT_FALSE(hybridSolution.streamKDynamicQueueSupported(problemOn, noAnalytical)) + << "SK5 resolving to the dynamic sub-path must be excluded when NUM_XCD is unknown"; + + auto problemAuto = makeGemmProblem(512, 512, 512); + problemAuto.setParams().setStreamKTileSchedulingMode(2); + EXPECT_THROW(hybridSolution.streamKDynamicQueueSupported(problemAuto, noAnalytical), + std::runtime_error) + << "AUTO cannot resolve without analytical hardware and must not pick a schedule silently"; + + auto problemOffHeuristic = makeGemmProblem(512, 512, 512); + problemOffHeuristic.setParams().setStreamKTileSchedulingMode(0); + problemOffHeuristic.setParams().setSmCountTarget(128); + EXPECT_THROW(hybridSolution.streamKDynamicQueueSupported(problemOffHeuristic, noAnalytical), + std::runtime_error) + << "OFF + smCountTarget>0 enters the same heuristic as AUTO and must behave identically"; +} + +TEST(StreamKDynamicQueueXcdGateTest, Sk5ForcedStaticStaysSelectableOnUnknownHardware) { + // TENSILE_STREAMK5_FORCE_MODE=0 short-circuits mode resolution to static + // before the heuristic is reached, so even AUTO stays selectable. + ScopedStreamK5ForceMode forceStatic("0"); + ASSERT_EQ(Debug::Instance().streamK5ForceMode(), 0); + hip::HipAMDGPU noAnalytical; noAnalytical.processor = AMDGPU::Processor::gfx942; noAnalytical.computeUnitCount = 304; @@ -1028,13 +1121,10 @@ TEST(StreamKDynamicQueueXcdGateTest, Sk5AutoUnknownHardwareIsRejectedWithoutThro ContractionSolution hybridSolution; initEquality512Solution(hybridSolution, 5); auto problem = makeGemmProblem(512, 512, 512); - problem.setParams().setStreamKTileSchedulingMode(2); // AUTO requires analytical hardware. + problem.setParams().setStreamKTileSchedulingMode(2); - bool supported = true; - EXPECT_NO_THROW( - supported = hybridSolution.streamKDynamicQueueSupported(problem, noAnalytical)); - EXPECT_FALSE(supported) - << "Unknown hardware must be rejected without entering the SK5 AUTO heuristic"; + EXPECT_TRUE(hybridSolution.streamKDynamicQueueSupported(problem, noAnalytical)) + << "Forced-static SK5 must stay selectable when NUM_XCD is unknown"; } // Selection-predicate contract: on MI300A (6 XCD) the dynamic-queue solution is @@ -1139,7 +1229,8 @@ TEST(StreamKDynamicQueueXcdGateTest, LogReporterTreatsUnsupportedTopologyAsTerse reporter.report(Client::ResultKey::SpeedGFlops, 0.0); reporter.report(Client::ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY"); reporter.postSolution(); - reporter.postProblem(); + // Deliberately no postProblem(): it printf()s a banner to stdout rather + // than to the captured stream, which would leak into this binary's output. EXPECT_NE(reportOutput.str().find("UNSUPPORTED_XCD_TOPOLOGY"), std::string::npos); } From cf40154acbd5d372e0024c11154cdd1df0c36517 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Fri, 11 Sep 2026 21:29:13 +0000 Subject: [PATCH 3/3] fix(tensilelite): link tensilelite-tests in the default shared build `TENSILELITE_BUILD_SHARED_LIBS` defaults to `ON`, and in that configuration `tensilelite-tests` failed to link. Two separate causes. Undefined LLVM YAML symbols. The serialization templates are instantiated in the test translation units, not only inside `tensilelite-host`. That library links `LLVMObjectYAML` `PRIVATE` and, when shared, is linked with `--exclude-libs,ALL` (`tensilelite/CMakeLists.txt:40`), so it exports none of those symbols. `RangeLibrary_test.cpp` referenced `llvm::yaml::ScalarTraits::input`, `ScalarValueTypeInfo::Get` and `KernelLanguageTypeInfo::Get` with nothing to resolve them against. Fixed by linking `LLVMObjectYAML` into the test target as well, and by adding `TENSILELITEHOST_EXPORT` to the enum-info `Get` overloads and to the `MLFeatures` / `MLPClassification` entry points the tests call. Duplicated `Debug` singleton. `Debug` inherits its instance accessor from `LazySingleton`, which holds the object in a function-local static of an inline template member. `VISIBILITY_INLINES_HIDDEN` gives that hidden visibility, so a shared `tensilelite-host` and the test binary each construct their own `Debug`. A test calling `reloadDebugBitsForTest()` would then mutate an instance the library never reads. `Debug::Instance()` is now defined out of line in `Debug.cpp` so both sides share one object. This failure predates the Stream-K commits that precede it on this branch. A shared build of `tensilelite-tests` at the branch point `267267524ae` emits the same 20 undefined symbols, and the files touched here are disjoint from those commits. Verified in `hipblaslt-ondemand:ubuntu24-gfx950-c865f12`: `tensilelite-tests` builds and runs in both the shared and static configurations, 416 passed / 3 skipped / 5 failed. The 5 failures are all `PciChipIdTest`, which call into HIP and get `hipErrorNoDevice` (100) because the container has no GPU. Co-Authored-By: Claude Opus 5 --- .../tensilelite/include/Tensile/Debug.hpp | 14 ++++++++++---- .../include/Tensile/KernelLanguageTypes.hpp | 11 +++++------ .../tensilelite/include/Tensile/MLFeatures.hpp | 11 ++++++----- .../include/Tensile/MLPClassification.hpp | 9 +++++---- .../include/Tensile/ScalarValueTypes.hpp | 11 +++++------ projects/hipblaslt/tensilelite/src/Debug.cpp | 13 ++++++++++--- .../hipblaslt/tensilelite/tests/CMakeLists.txt | 9 +++++++++ 7 files changed, 50 insertions(+), 28 deletions(-) diff --git a/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp b/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp index 944c2eaaaa2c..b5187de46082 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp @@ -47,6 +47,13 @@ namespace TensileLite class TENSILELITEHOST_EXPORT Debug : public LazySingleton { public: + // Defined out of line rather than inherited from LazySingleton. That base + // holds the instance in a function-local static of an inline template + // member, which VISIBILITY_INLINES_HIDDEN gives hidden visibility, so a + // shared tensilelite-host and its consumer each get their own Debug and + // one side's reloadDebugBitsForTest() cannot be seen by the other. + static Debug& Instance(); + bool printPropertyEvaluation() const; bool printPredicateEvaluation() const; bool printPredicateEvaluationVerbose() const; @@ -185,11 +192,10 @@ namespace TensileLite bool m_printMarker = false; bool m_disableStaggerU = false; // -1 = unset (use API attribute); 0 = force static SK3; 1 = force dynamic SK4 - int m_streamK5ForceMode = -1; - bool m_usePreciseSMTarget = false; - StringSet m_excludedFromGetAll; + int m_streamK5ForceMode = -1; + bool m_usePreciseSMTarget = false; + StringSet m_excludedFromGetAll; Debug(); }; } // namespace TensileLite - diff --git a/projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp b/projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp index a52f4831b022..a1901723a6cf 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp @@ -59,8 +59,8 @@ namespace TensileLite Count }; - TENSILELITEHOST_EXPORT std::string ToString(KernelLanguage d); - TENSILELITEHOST_EXPORT std::string TypeAbbrev(KernelLanguage d); + TENSILELITEHOST_EXPORT std::string ToString(KernelLanguage d); + TENSILELITEHOST_EXPORT std::string TypeAbbrev(KernelLanguage d); TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, KernelLanguage const& t); TENSILELITEHOST_EXPORT std::istream& operator>>(std::istream& stream, KernelLanguage& t); @@ -70,9 +70,9 @@ namespace TensileLite */ struct KernelLanguageTypeInfo { - static KernelLanguageTypeInfo const& Get(int index); - static KernelLanguageTypeInfo const& Get(KernelLanguage t); - static KernelLanguageTypeInfo const& Get(std::string const& str); + TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(int index); + TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(KernelLanguage t); + TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(std::string const& str); KernelLanguage m_kernelLanguage; std::string name; @@ -148,4 +148,3 @@ namespace std } }; } // namespace std - diff --git a/projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp b/projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp index 84dde1bdf58c..336751272956 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp @@ -59,11 +59,13 @@ namespace TensileLite float waveScale; }; - float tilesPerCU(ContractionProblemGemm const& problem, - CUGranularityScaleFactors const& cuFactors); + TENSILELITEHOST_EXPORT float tilesPerCU(ContractionProblemGemm const& problem, + CUGranularityScaleFactors const& cuFactors); - std::ostream& operator<<(std::ostream& stream, CUGranularityScaleFactors const& cugsf); - std::ostream& operator<<(std::ostream& stream, WaveGranularityScaleFactors const& wgsf); + TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, + CUGranularityScaleFactors const& cugsf); + TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, + WaveGranularityScaleFactors const& wgsf); /** * @brief A Property whose value is of type `float`. @@ -257,4 +259,3 @@ namespace TensileLite */ } // namespace MLFeatures } // namespace TensileLite - diff --git a/projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp b/projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp index ecea84559d6c..d798782cd388 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp @@ -73,7 +73,8 @@ namespace TensileLite struct DenseLayer { DenseLayer() = default; - DenseLayer(const std::vector& weights, const std::vector& bias); + TENSILELITEHOST_EXPORT DenseLayer(const std::vector& weights, + const std::vector& bias); std::vector operator()(const std::vector& F) const; @@ -101,9 +102,10 @@ namespace TensileLite { MLPNet() = default; - std::vector predict(std::vector const& probkey) const; + TENSILELITEHOST_EXPORT std::vector + predict(std::vector const& probkey) const; - bool valid(bool verbose = false) const; + TENSILELITEHOST_EXPORT bool valid(bool verbose = false) const; std::string description() const { @@ -117,4 +119,3 @@ namespace TensileLite } // namespace MLPClassification } // namespace TensileLite - diff --git a/projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp b/projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp index f83930268a2e..5a102256e441 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp @@ -61,8 +61,8 @@ namespace TensileLite Count }; - TENSILELITEHOST_EXPORT std::string ToString(ScalarValue d); - TENSILELITEHOST_EXPORT std::string TypeAbbrev(ScalarValue d); + TENSILELITEHOST_EXPORT std::string ToString(ScalarValue d); + TENSILELITEHOST_EXPORT std::string TypeAbbrev(ScalarValue d); TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, ScalarValue const& t); TENSILELITEHOST_EXPORT std::istream& operator>>(std::istream& stream, ScalarValue& t); @@ -72,9 +72,9 @@ namespace TensileLite */ struct ScalarValueTypeInfo { - static ScalarValueTypeInfo const& Get(int index); - static ScalarValueTypeInfo const& Get(ScalarValue t); - static ScalarValueTypeInfo const& Get(std::string const& str); + TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(int index); + TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(ScalarValue t); + TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(std::string const& str); ScalarValue m_value; std::string name; @@ -189,4 +189,3 @@ namespace std } }; } // namespace std - diff --git a/projects/hipblaslt/tensilelite/src/Debug.cpp b/projects/hipblaslt/tensilelite/src/Debug.cpp index 989b360d5062..d4e519e2aad7 100644 --- a/projects/hipblaslt/tensilelite/src/Debug.cpp +++ b/projects/hipblaslt/tensilelite/src/Debug.cpp @@ -205,6 +205,13 @@ namespace TensileLite m_excludedFromGetAll = excludedSet; } + Debug& Debug::Instance() + { + static Debug instance; + + return instance; + } + Debug::Debug() : m_value(DEBUG_SM) , m_value2(DEBUG_SM2) @@ -230,7 +237,7 @@ namespace TensileLite const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE"); if(sk5Force) { - char* end = nullptr; + char* end = nullptr; const long val = strtol(sk5Force, &end, 0); if(end != sk5Force && *end == '\0' && val >= -1 && val <= 1) m_streamK5ForceMode = static_cast(val); @@ -296,8 +303,8 @@ namespace TensileLite const char* db2 = std::getenv("TENSILE_DB2"); m_value2 = db2 ? static_cast(strtol(db2, nullptr, 0)) : DEBUG_SM2; - const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE"); - m_streamK5ForceMode = -1; + const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE"); + m_streamK5ForceMode = -1; if(sk5Force) { char* end = nullptr; diff --git a/projects/hipblaslt/tensilelite/tests/CMakeLists.txt b/projects/hipblaslt/tensilelite/tests/CMakeLists.txt index 83f7d60e1acc..734b727af49e 100644 --- a/projects/hipblaslt/tensilelite/tests/CMakeLists.txt +++ b/projects/hipblaslt/tensilelite/tests/CMakeLists.txt @@ -82,6 +82,15 @@ endif() target_link_libraries(tensilelite-tests PUBLIC GTest::gtest) +# In YAML builds the serialization templates are instantiated in these test +# translation units, not only inside tensilelite-host. tensilelite-host links +# LLVMObjectYAML PRIVATE and, when shared, is linked with --exclude-libs,ALL, +# so it does not re-export those symbols. Link them here as well. +if(HIPBLASLT_ENABLE_YAML) + target_link_libraries(tensilelite-tests PRIVATE LLVMObjectYAML) + target_include_directories(tensilelite-tests PRIVATE ${LLVM_INCLUDE_DIRS}) +endif() + if(HIPBLASLT_ENABLE_OPENMP) target_link_libraries(tensilelite-tests PRIVATE OpenMP::OpenMP_CXX) endif()