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/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() diff --git a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp index b25156b1bd9a..6bee2b6ecac3 100644 --- a/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/CuCount_test.cpp @@ -2,10 +2,15 @@ // SPDX-License-Identifier: MIT #include +#include +#include #include #include #include +#include +#include #include +#include #include @@ -22,6 +27,7 @@ #include "FallbackTestUtils.hpp" #include "LogReporter.hpp" +#include "ResultReporter.hpp" #include "SolutionIterator.hpp" using namespace TensileLite; @@ -440,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; @@ -932,6 +975,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 +1059,74 @@ TEST(StreamKDynamicQueueXcdGateTest, MissingAnalyticalHardwareIsUnsupported) << "SK3-static solution must remain selectable when NUM_XCD is unknown"; } +// 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; + noAnalytical.deviceName = "test-gfx942-no-analytical"; + + ContractionSolution hybridSolution; + initEquality512Solution(hybridSolution, 5); + auto problem = makeGemmProblem(512, 512, 512); + problem.setParams().setStreamKTileSchedulingMode(2); + + 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 // EXCLUDED from selection (supported == false) so a different solution serves // the GEMM, while on MI300X (8 XCD) the identical solution stays selectable. @@ -1037,13 +1168,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 +1176,63 @@ 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(); + // 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); } // ===========================================================================