From a52534d14ddf30f05237cff77f19b704bce2bc7b Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Fri, 25 Sep 2026 05:04:51 +0000 Subject: [PATCH 1/5] fix(hipblaslt): return an error when a matmul throws after choosing its algo runContractionProblem's exception handlers left status as it was. For a call that passed no algo, status already held getBestSolutions' success, so an exception thrown while solving or launching returned success for a GEMM that never ran. Co-authored-by: Cursor --- projects/hipblaslt/CHANGELOG.md | 1 + .../library/src/amd_detail/rocblaslt/src/tensile_host.cpp | 4 ++++ 2 files changed, 5 insertions(+) diff --git a/projects/hipblaslt/CHANGELOG.md b/projects/hipblaslt/CHANGELOG.md index e38003776195..14fe975534cd 100644 --- a/projects/hipblaslt/CHANGELOG.md +++ b/projects/hipblaslt/CHANGELOG.md @@ -88,6 +88,7 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip * A tuning file row cut short, for example when the disk fills while `hipblaslt-bench` appends it, could be replayed with a truncated `solution_index`. A row whose cells do not line up with its header is now ignored. * `GemmInstance::getSolutionName()` in `hipblaslt_ext` crashed for a RocRoller solution. It looked the encoded RocRoller index up in the Tensile library and dereferenced the null result. It now returns the RocRoller short name, as `hipblaslt_ext::getSolutionNameFromAlgo()` already did. * The logger's destructor called `close()` on the log file, which throws out of a destructor and terminates the process at exit when the log file has failed, for example on a full disk. Reproducible with `HIPBLASLT_LOG_LEVEL=4` and a failing `HIPBLASLT_LOG_FILE`. +* `hipblasLtMatmul` with no `algo` could return success for a GEMM that never launched. Its status already held the result of choosing the algorithm, and an exception thrown after that returned it unchanged instead of an error. ## hipBLASLt 1.4.1 for ROCm 7.14 diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 3fe9def63b6b..53d8ce67ad5a 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -3891,10 +3891,14 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle catch(const std::exception& e) { log_error(__func__, e.what()); + // status may already hold getBestSolutions' success for a call that + // passed no algo. + status = rocblaslt_status_internal_error; } catch(...) { log_error(__func__, "unknown exception"); + status = rocblaslt_status_internal_error; } return status; From 6fb6475b706bd1b3745d7722f1437bd99821a1a4 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Fri, 25 Sep 2026 05:20:24 +0000 Subject: [PATCH 2/5] feat(hipblaslt): replay a managed tuning cache (HIPBLASLT_TUNING_MODE=cache) HIPBLASLT_TUNING_MODE=cache with HIPBLASLT_TUNING_CACHE_PATH replays a tuning file through the store, validating each entry as the override file does. - Mode and path are read once per process through TuningModeConfig, and a process in a secure execution context ignores both. With a mode set only the cache is consulted; HIPBLASLT_TUNING_OVERRIDE_FILE is ignored. - One launch rule on both paths: the heuristic hands out the cached entry first, an explicit algo launches as given, and a matmul with no algo launches the first entry that still validates and supports the call. - A startup line and a closing summary report what loaded and what was served, without a log level; per-problem hit, miss and invalidation lines appear at info level, once per key. - Test hooks report which solution a matmul launched and the tallies. - Docs: cache section in the offline tuning how-to, the variables in the environment reference, and corrections to the logging docs. Co-authored-by: Cursor --- projects/hipblaslt/CHANGELOG.md | 1 + .../clients/tests/src/tuning_cache_test.cpp | 567 +++++++++++++++++- .../clients/tests/src/tuning_store_test.cpp | 70 +++ .../how-to-use-hipblaslt-offline-tuning.rst | 65 ++ .../docs/how-to/use-logging-heuristics.rst | 8 +- .../docs/reference/env-variables.rst | 42 +- .../rocblaslt/src/TuningCacheStore.cpp | 23 + .../rocblaslt/src/UserDrivenTuningParser.cpp | 330 +++++++++- .../src/include/TuningCacheStore.hpp | 31 + .../src/include/UserDrivenTuningParser.hpp | 141 +++++ .../rocblaslt/src/include/tensile_host.hpp | 30 + .../rocblaslt/src/include/utility.hpp | 40 ++ .../rocblaslt/src/rocblaslt_auxiliary.cpp | 225 +++++-- .../amd_detail/rocblaslt/src/tensile_host.cpp | 99 ++- 14 files changed, 1593 insertions(+), 79 deletions(-) diff --git a/projects/hipblaslt/CHANGELOG.md b/projects/hipblaslt/CHANGELOG.md index 14fe975534cd..1ac0f818a1fa 100644 --- a/projects/hipblaslt/CHANGELOG.md +++ b/projects/hipblaslt/CHANGELOG.md @@ -23,6 +23,7 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip * Complex CGEMM/ZGEMM support for gfx1250. * `TENSILE_FIXED_WGMXCCSPLITK` environment variable to override the split-K work-group XCC mapping factor for StreamK GEMMs. * `HIPBLASLT_MATRIX_LAYOUT_OFFSET` matrix-layout attribute for 64-bit element offsets into sub-matrices in General Batched GEMM (`batch_mode=1`), along with `hipblaslt-bench` `batch_offset_a/b/c/d` arguments; nonzero offsets require `HIPBLASLT_BATCH_MODE_POINTER_ARRAY` and are rejected for sub-byte MX types (`HIP_R_6F_E2M3`, `HIP_R_6F_E3M2`, `HIP_R_4F_E2M1`) with `HIPBLAS_STATUS_NOT_SUPPORTED`. +* Opt-in runtime tuning cache, off by default. `HIPBLASLT_TUNING_MODE=cache` with `HIPBLASLT_TUNING_CACHE_PATH=` replays a tuning file, validating each entry as `HIPBLASLT_TUNING_OVERRIDE_FILE` does. Unlike the override file, it also serves `hipblasLtMatmul` calls that pass no `algo`, and it reports what it loaded and a closing summary of how many problems it served without needing a log level. An explicit `algo` always launches as given. With a tuning mode set, `HIPBLASLT_TUNING_OVERRIDE_FILE` is ignored. A process in a secure execution context ignores both variables. ### Changed diff --git a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp index c3c2014bd964..cb2382c5542c 100644 --- a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp +++ b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp @@ -4,12 +4,13 @@ // Tuning file tests. // // Each case writes a tuning file the way hipblaslt-bench does, points -// HIPBLASLT_TUNING_OVERRIDE_FILE at it, and asserts on the solution index the -// heuristic hands back, which is the entry replay chose. Rows are built from -// solutions the heuristic itself offers for the problem, so every recorded -// index is real and only the recorded name or build decides whether it is used. +// HIPBLASLT_TUNING_OVERRIDE_FILE or cache mode at it, and asserts on the +// solution index the heuristic hands back or the matmul launches, and on the +// library's own counters. Rows are built from solutions the heuristic itself +// offers for the problem, so every recorded index is real and only the +// recorded name or build decides whether it is used. // -// The override variable is read on first use and each file is loaded once per +// The tuning variables are read on first use and each file is loaded once per // process, so each case starts from hipblaslt_tuning_reset_for_test(). #include @@ -20,9 +21,12 @@ #include #include #include +#include #include +#include #include #include +#include #include #include #include @@ -33,7 +37,9 @@ #include #endif -extern "C" void hipblaslt_tuning_reset_for_test(); +extern "C" void hipblaslt_tuning_reset_for_test(); +extern "C" int hipblaslt_tuning_last_launch_for_test(); +extern "C" size_t hipblaslt_tuning_stats_for_test(uint64_t* values, size_t count); #ifdef WIN32 static int setenv(const char* name, const char* value, int overwrite) @@ -387,6 +393,212 @@ namespace return ok; } + /** The library's tallies, in the order hipblaslt_tuning_stats_for_test documents. */ + struct Counters + { + uint64_t loaded = 0, hits = 0, misses = 0, invalidated = 0; + uint64_t shapes = 0, matched = 0, fellback = 0; + }; + + Counters counters() + { + uint64_t v[7] = {}; + const size_t known = hipblaslt_tuning_stats_for_test(v, 7); + EXPECT_GE(known, 7u) << "the library reports fewer tallies than this reads"; + + Counters c; + c.loaded = v[0]; + c.hits = v[1]; + c.misses = v[2]; + c.invalidated = v[3]; + c.shapes = v[4]; + c.matched = v[5]; + c.fellback = v[6]; + return c; + } + + uint16_t halfBits(float value) + { + const _Float16 h = static_cast<_Float16>(value); + uint16_t bits; + std::memcpy(&bits, &h, sizeof(bits)); + return bits; + } + + /** Where launchGemm takes the algo it hands hipblasLtMatmul from. */ + enum class AlgoFrom + { + Heuristic, // the heuristic's first result, queried just before the call + Index, // a given solution index, the way a caller reuses one algo + Null, // no algo, which leaves the choice to the library + }; + + /** + * One matmul of the problem with separate C and D, reporting the solution + * the library actually launched. + * + * verifyProduct fills A and B with ones and checks every element of D is k, + * which proves the launched kernel ran on the caller's buffers. + */ + bool launchGemm(AlgoFrom from, int index, int* launched, bool verifyProduct = false) + { + *launched = -1; + + hipblasLtHandle_t handle = nullptr; + if(hipblasLtCreate(&handle) != HIPBLAS_STATUS_SUCCESS) + return false; + + void* dA = nullptr; + void* dB = nullptr; + void* dC = nullptr; + void* dD = nullptr; + void* dWs = nullptr; + + bool ok = hipMalloc(&dA, kM * kK * sizeof(uint16_t)) == hipSuccess + && hipMalloc(&dB, kK * kN * sizeof(uint16_t)) == hipSuccess + && hipMalloc(&dC, kM * kN * sizeof(uint16_t)) == hipSuccess + && hipMalloc(&dD, kM * kN * sizeof(uint16_t)) == hipSuccess + && hipMalloc(&dWs, kWorkspaceBytes) == hipSuccess; + + if(ok) + { + const uint16_t fill = verifyProduct ? halfBits(1.0f) : 0; + const std::vector a(static_cast(kM * kK), fill); + const std::vector b(static_cast(kK * kN), fill); + ok = hipMemcpy(dA, a.data(), a.size() * sizeof(uint16_t), hipMemcpyHostToDevice) + == hipSuccess + && hipMemcpy(dB, b.data(), b.size() * sizeof(uint16_t), hipMemcpyHostToDevice) + == hipSuccess + && hipMemset(dC, 0, kM * kN * sizeof(uint16_t)) == hipSuccess + && hipMemset(dD, 0, kM * kN * sizeof(uint16_t)) == hipSuccess; + } + + { + Problem problem; + ok = ok && problem.create(); + + const float alpha = 1.0f; + const float beta = 0.0f; + hipblasLtMatmulAlgo_t algo = {}; + const hipblasLtMatmulAlgo_t* algoPtr = nullptr; + + if(ok && from == AlgoFrom::Heuristic) + { + hipblasLtMatmulHeuristicResult_t heuristic[1]; + int returned = 0; + ok = hipblasLtMatmulAlgoGetHeuristic(handle, + problem.desc, + problem.layoutA, + problem.layoutB, + problem.layoutC, + problem.layoutC, + problem.pref, + 1, + heuristic, + &returned) + == HIPBLAS_STATUS_SUCCESS + && returned > 0; + algo = heuristic[0].algo; + algoPtr = &algo; + } + else if(ok && from == AlgoFrom::Index) + { + std::vector indexes{index}; + std::vector results; + ok = hipblaslt_ext::getAlgosFromIndex(handle, indexes, results) + == HIPBLAS_STATUS_SUCCESS + && !results.empty(); + if(ok) + { + algo = results[0].algo; + algoPtr = &algo; + } + } + + if(ok) + { + ok = hipblasLtMatmul(handle, + problem.desc, + &alpha, + dA, + problem.layoutA, + dB, + problem.layoutB, + &beta, + dC, + problem.layoutC, + dD, + problem.layoutC, + algoPtr, + dWs, + kWorkspaceBytes, + nullptr) + == HIPBLAS_STATUS_SUCCESS + && hipDeviceSynchronize() == hipSuccess; + *launched = hipblaslt_tuning_last_launch_for_test(); + } + } + + if(ok && verifyProduct) + { + std::vector out(static_cast(kM * kN), 0); + ok = hipMemcpy(out.data(), dD, out.size() * sizeof(uint16_t), hipMemcpyDeviceToHost) + == hipSuccess; + + const uint16_t expected = halfBits(static_cast(kK)); + for(size_t i = 0; ok && i < out.size(); i++) + ok = out[i] == expected; + } + + static_cast(hipFree(dWs)); + static_cast(hipFree(dD)); + static_cast(hipFree(dC)); + static_cast(hipFree(dB)); + static_cast(hipFree(dA)); + hipblasLtDestroy(handle); + return ok; + } + + /** + * Collects what the library writes to std::cerr while it is alive. It sees + * nothing sent to a log file or written straight to file descriptor 2. + */ + class CerrCapture + { + public: + CerrCapture() + : m_saved(std::cerr.rdbuf(m_buffer.rdbuf())) + { + } + + ~CerrCapture() + { + std::cerr.rdbuf(m_saved); + } + + CerrCapture(const CerrCapture&) = delete; + CerrCapture& operator=(const CerrCapture&) = delete; + + std::string str() const + { + return m_buffer.str(); + } + + private: + std::ostringstream m_buffer; + std::streambuf* m_saved; + }; + + /** + * The logger reads its environment once per process, so a run with a level + * or log file set sends tuning notices somewhere CerrCapture cannot see. + */ + bool loggingEnvIsClean() + { + return getenv("HIPBLASLT_LOG_LEVEL") == nullptr && getenv("HIPBLASLT_LOG_MASK") == nullptr + && getenv("HIPBLASLT_LOG_FILE") == nullptr; + } + struct Row { int index = -1; @@ -444,15 +656,20 @@ namespace protected: void SetUp() override { - const char* value = getenv("HIPBLASLT_TUNING_OVERRIDE_FILE"); - m_savedOverride = value ? std::optional(value) : std::nullopt; + for(const char* name : kTuningVariables) + { + const char* value = getenv(name); + m_savedEnv.emplace_back(name, + value ? std::optional(value) : std::nullopt); + } if(!gpuAvailable()) GTEST_SKIP() << "No GPU available"; - // Solutions are read with no file in play, so a user's own override - // cannot decide what the cases below take as the default. - unsetenv("HIPBLASLT_TUNING_OVERRIDE_FILE"); + // Solutions are read with no file in play, so a user's own tuning + // settings cannot decide what the cases below take as the default. + for(const char* name : kTuningVariables) + unsetenv(name); hipblaslt_tuning_reset_for_test(); m_path @@ -466,20 +683,34 @@ namespace void TearDown() override { std::remove(m_path.c_str()); - if(m_savedOverride) - setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_savedOverride->c_str(), 1); - else - unsetenv("HIPBLASLT_TUNING_OVERRIDE_FILE"); + for(const auto& [name, value] : m_savedEnv) + { + if(value) + setenv(name.c_str(), value->c_str(), 1); + else + unsetenv(name.c_str()); + } hipblaslt_tuning_reset_for_test(); } - /** Point the library at the file just written, starting clean. */ + /** Point the override variable at the file just written, starting clean. */ void useTuningFile() { setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); hipblaslt_tuning_reset_for_test(); } + /** Cache mode on a path, which may be empty, starting clean. */ + void useCache(const std::string& path) + { + setenv("HIPBLASLT_TUNING_MODE", "cache", 1); + if(path.empty()) + unsetenv("HIPBLASLT_TUNING_CACHE_PATH"); + else + setenv("HIPBLASLT_TUNING_CACHE_PATH", path.c_str(), 1); + hipblaslt_tuning_reset_for_test(); + } + /** * Skip unless the heuristic offers at least `count` solutions. A case * records a non-default one so that replaying it is distinguishable from @@ -490,10 +721,14 @@ namespace return m_identities.size() >= count; } - std::string m_path; - std::string m_stamp; - std::vector m_identities; - std::optional m_savedOverride; + static constexpr const char* kTuningVariables[] = {"HIPBLASLT_TUNING_OVERRIDE_FILE", + "HIPBLASLT_TUNING_MODE", + "HIPBLASLT_TUNING_CACHE_PATH"}; + + std::string m_path; + std::string m_stamp; + std::vector m_identities; + std::vector>> m_savedEnv; }; // A row whose recorded name still matches what its index resolves to is @@ -744,4 +979,296 @@ namespace ASSERT_TRUE(runGemm(&selected)); EXPECT_EQ(selected, recorded.index); } + + // Cache mode replays a named row through the heuristic, on both APIs. + TEST_F(TuningCache_pre_checkin, CacheModeReplaysANamedRow) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + const auto& recorded = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + useCache(m_path); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, recorded.index); + + const auto c = counters(); + EXPECT_EQ(c.loaded, 1u); + EXPECT_EQ(c.hits, 1u) << "the matmul with the heuristic's algo was counted again"; + + int extSelected = -1; + ASSERT_TRUE(extHeuristicIndex(&extSelected)); + EXPECT_EQ(extSelected, recorded.index); + } + + // A cache file that is not there yet when the first query looks for it is + // read once it appears. + TEST_F(TuningCache_pre_checkin, CacheModeReadsAFileCreatedAfterTheFirstQuery) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + useCache(m_path); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, m_identities[0].index); + + const auto& recorded = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, recorded.index); + EXPECT_EQ(counters().loaded, 1u); + } + + // A mode with nowhere to read from does nothing, and does not fall back to + // the override file. + TEST_F(TuningCache_pre_checkin, CacheModeWithoutPathIsInert) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); + useCache(""); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, m_identities[0].index); + EXPECT_EQ(counters().hits, 0u); + } + + // With a tuning mode set only the cache is consulted, never the override + // file. + TEST_F(TuningCache_pre_checkin, CacheModeIgnoresTheOverrideFile) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); + useCache(m_path + ".absent"); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, m_identities[0].index); + EXPECT_EQ(counters().loaded, 0u); + } + + TEST_F(TuningCache_pre_checkin, CachedEntryWhoseNameNoLongerMatchesIsRejected) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}}); + useCache(m_path); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + EXPECT_EQ(selected, m_identities[0].index); + + const auto c = counters(); + EXPECT_EQ(c.invalidated, 1u); + EXPECT_EQ(c.misses, 1u); + EXPECT_EQ(c.hits, 0u); + } + + // The heuristic lookup and the execution path both meet a stale row, and it + // is still one rejected entry. + TEST_F(TuningCache_pre_checkin, StaleEntryIsCountedAsOneInvalidation) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}}); + useCache(m_path); + + int selected = -1; + int launched = -1; + ASSERT_TRUE(runGemm(&selected)); + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); + EXPECT_EQ(counters().invalidated, 1u); + } + + // A caller that passes an algo gets that algo, whatever the cache holds. + TEST_F(TuningCache_pre_checkin, ExplicitAlgoLaunchesAsGivenInCacheMode) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + const auto& given = m_identities[0]; + const auto& cached = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}}); + useCache(m_path); + + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Index, given.index, &launched)); + EXPECT_EQ(launched, given.index) << "the cache replaced the caller's algo"; + EXPECT_EQ(counters().hits, 0u) << "a hit was counted for a launch the cache did not choose"; + } + + // A matmul given no algo leaves the choice to the library, so a usable + // entry is what it launches, and it computes the right result. + TEST_F(TuningCache_pre_checkin, NullAlgoLaunchesTheCachedKernel) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + const auto& cached = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}}); + useCache(m_path); + + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched, true)); + EXPECT_EQ(launched, cached.index) << "a matmul with no algo ignored the cached entry"; + EXPECT_EQ(counters().hits, 1u); + } + + // With no usable entry a null algo still runs, on default selection's pick, + // and the lookup counts as the fallback it is. + TEST_F(TuningCache_pre_checkin, NullAlgoWithoutAnEntryLaunchesDefaultSelection) + { + if(!haveSolutions(1)) + GTEST_SKIP() << "the heuristic offers no solution for this problem"; + + useCache(m_path); + + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched, true)); + EXPECT_EQ(launched, m_identities[0].index); + + const auto c = counters(); + EXPECT_EQ(c.hits, 0u); + EXPECT_EQ(c.misses, 1u); + } + + // A cache that served nothing must not read like a process that asked + // nothing. + TEST_F(TuningCache_pre_checkin, ShapesThatMissAnEmptyCacheAreStillCounted) + { + useCache(m_path); + + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + + const auto c = counters(); + EXPECT_EQ(c.shapes, 1u); + EXPECT_EQ(c.matched, 0u); + EXPECT_EQ(c.fellback, 1u); + } + + // The summary is read against loaded=N, so it counts problems rather than + // lookups. + TEST_F(TuningCache_pre_checkin, LookupTallyCountsShapesNotCalls) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + useCache(m_path); + + for(int i = 0; i < 4; i++) + { + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + } + + const auto c = counters(); + EXPECT_EQ(c.shapes, 1u); + EXPECT_EQ(c.matched, 1u); + EXPECT_EQ(c.fellback, 0u); + } + + // The first lookup says what mode is on and what loaded, without any log + // level set, and names an override file it is ignoring. + TEST_F(TuningCache_pre_checkin, StartupLineReportsWhatLoaded) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + if(!loggingEnvIsClean()) + GTEST_SKIP() << "external logging environment redirects tuning output"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + useCache(m_path); + + std::string loaded; + { + CerrCapture capture; + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + loaded = capture.str(); + } + EXPECT_NE(loaded.find("tuning-cache: mode=cache path=" + m_path + " load=ok loaded=1"), + std::string::npos) + << loaded; + + setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); + useCache(m_path + ".absent"); + + std::string missing; + { + CerrCapture capture; + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + missing = capture.str(); + } + EXPECT_NE(missing.find("load=not-found loaded=0"), std::string::npos) << missing; + EXPECT_NE(missing.find("ignoring HIPBLASLT_TUNING_OVERRIDE_FILE"), std::string::npos) + << missing; + } + + // Replay is the steady state, so it must be silent per call. + TEST_F(TuningCache_pre_checkin, RepeatedCacheHitsAddNoOutput) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + if(!loggingEnvIsClean()) + GTEST_SKIP() << "external logging environment redirects tuning output"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + useCache(m_path); + + { + CerrCapture capture; + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + } + + std::string later; + { + CerrCapture capture; + int selected = -1; + int launched = -1; + ASSERT_TRUE(runGemm(&selected)); + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); + later = capture.str(); + } + EXPECT_EQ(later, ""); + } + + // With no tuning mode set, the override file behaves as it always has and + // prints no tuning notices. + TEST_F(TuningCache_pre_checkin, OverrideFilePrintsNoTuningNotices) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + if(!loggingEnvIsClean()) + GTEST_SKIP() << "external logging environment redirects tuning output"; + + writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + useTuningFile(); + + std::string out; + int selected = -1; + { + CerrCapture capture; + ASSERT_TRUE(runGemm(&selected)); + out = capture.str(); + } + ASSERT_EQ(selected, m_identities[1].index) << "the override file was not consulted"; + EXPECT_EQ(out.find("tuning-cache:"), std::string::npos) << out; + } } // namespace diff --git a/projects/hipblaslt/clients/tests/src/tuning_store_test.cpp b/projects/hipblaslt/clients/tests/src/tuning_store_test.cpp index 35b80f04847f..6ef1fcee1a4d 100644 --- a/projects/hipblaslt/clients/tests/src/tuning_store_test.cpp +++ b/projects/hipblaslt/clients/tests/src/tuning_store_test.cpp @@ -12,6 +12,7 @@ #include #include +#include #include #include #include @@ -26,6 +27,17 @@ #include #endif +#ifdef WIN32 +static int setenv(const char* name, const char* value, int) +{ + return _putenv_s(name, value); +} +static int unsetenv(const char* name) +{ + return _putenv_s(name, ""); +} +#endif + using namespace TensileLite; namespace @@ -192,6 +204,29 @@ namespace class TuningStore : public ::testing::Test { + protected: + void SetUp() override + { + for(const char* name : {"HIPBLASLT_TUNING_MODE", "HIPBLASLT_TUNING_CACHE_PATH"}) + { + const char* value = std::getenv(name); + m_savedEnv.emplace_back(name, + value ? std::optional(value) : std::nullopt); + } + } + + void TearDown() override + { + for(const auto& [name, value] : m_savedEnv) + { + if(value) + setenv(name.c_str(), value->c_str(), 1); + else + unsetenv(name.c_str()); + } + } + + std::vector>> m_savedEnv; }; // Every column the writer emits reads back into the key and entry it came @@ -676,6 +711,41 @@ namespace } } + // A process in a secure execution context (set-user-ID and the like) must + // not let an inherited environment choose a file for it to read. + TEST_F(TuningStore, PrivilegedProcessIgnoresTheTuningEnvironment) + { + setenv("HIPBLASLT_TUNING_MODE", "cache", 1); + setenv("HIPBLASLT_TUNING_CACHE_PATH", "tuning.txt", 1); + + const auto ordinary = TuningModeConfig::fromEnvironment(false); + EXPECT_EQ(ordinary.mode, TuningMode::Cache); + EXPECT_EQ(ordinary.cachePath, "tuning.txt"); + EXPECT_TRUE(ordinary.reads()); + EXPECT_FALSE(ordinary.suppressedForSecurity); + + const auto privileged = TuningModeConfig::fromEnvironment(true); + EXPECT_EQ(privileged.mode, TuningMode::Off); + EXPECT_TRUE(privileged.cachePath.empty()); + EXPECT_FALSE(privileged.reads()); + EXPECT_TRUE(privileged.suppressedForSecurity); + } + + // An unknown mode is off, and a mode with no path reads nothing. + TEST_F(TuningStore, ModeNeedsAKnownValueAndAPath) + { + setenv("HIPBLASLT_TUNING_MODE", "online", 1); + setenv("HIPBLASLT_TUNING_CACHE_PATH", "tuning.txt", 1); + EXPECT_EQ(TuningModeConfig::fromEnvironment(false).mode, TuningMode::Off); + + setenv("HIPBLASLT_TUNING_MODE", "cache", 1); + unsetenv("HIPBLASLT_TUNING_CACHE_PATH"); + const auto noPath = TuningModeConfig::fromEnvironment(false); + EXPECT_EQ(noPath.mode, TuningMode::Cache); + EXPECT_FALSE(noPath.reads()); + EXPECT_FALSE(noPath.suppressedForSecurity); + } + // The file starts with the build stamp once, and every appended row reads // back. TEST_F(TuningStore, AppendedRowsReadBackFromTheFile) diff --git a/projects/hipblaslt/docs/how-to/how-to-use-hipblaslt-offline-tuning.rst b/projects/hipblaslt/docs/how-to/how-to-use-hipblaslt-offline-tuning.rst index a2d05719f870..6e27043dad56 100644 --- a/projects/hipblaslt/docs/how-to/how-to-use-hipblaslt-offline-tuning.rst +++ b/projects/hipblaslt/docs/how-to/how-to-use-hipblaslt-offline-tuning.rst @@ -115,3 +115,68 @@ guarantees is that the index resolves to the same compiled kernel as when the ro does not detect a rebuild that changes only that solution's defaults, whose row is used with the new defaults, and it cannot find a kernel that moved to a different index: that row is rejected and the shape falls back to heuristic selection. + +Replaying a tuning cache +======================== + +Instead of ``HIPBLASLT_TUNING_OVERRIDE_FILE``, a tuning file can be replayed as a runtime cache: + +.. code-block:: bash + + export HIPBLASLT_TUNING_MODE=cache + export HIPBLASLT_TUNING_CACHE_PATH=tuning.txt + +``cache`` replays what the file contains, validating each entry as described above. ``off``, the +default, changes nothing. Cache mode only reads the file, so any number of processes can share one, +for example the ranks of an MPI job. + +Which kernel a matmul launches depends on whether it passes an algorithm. A call to +``hipblasLtMatmul`` with an explicit ``algo`` always launches that algorithm; such a caller gets the +cached entry from ``hipblasLtMatmulAlgoGetHeuristic``, which returns it as its first result. A call +with ``algo=nullptr`` leaves the choice to the library and launches the cached entry when there is a +usable one, which the override file never did. + +Rows written by ``hipblaslt-bench`` match on transpose, shape and the principal datatypes only, +exactly as they do through ``HIPBLASLT_TUNING_OVERRIDE_FILE``: they do not distinguish leading +dimensions, batch strides, epilogue details or the device. Rows that carry a ``schema_version`` +column are keyed on the complete problem, and every column of such a row is required, so a row with +a missing or malformed value, including one cut short by an interrupted write, is ignored. + +``HIPBLASLT_TUNING_CACHE_PATH`` and ``HIPBLASLT_TUNING_OVERRIDE_FILE`` are mutually exclusive. With a +tuning mode set, only the cache is consulted and the override file is ignored. + +A process running in a secure execution context (set-user-ID, set-group-ID, or another +credential-changing exec such as file capabilities) ignores both tuning variables, so the cache stays +off in it. + +Reading the cache output +------------------------ + +``cache`` mode reports what it loaded without needing any logging variable: + +.. code-block:: none + + tuning-cache: mode=cache path=tuning.txt load=ok loaded=12 + tuning-cache: summary shapes=14 matched=12 fellback=2 invalidated=0 + +The first line is written when the file is first read. ``load`` is ``ok``, ``not-found``, +``read-error`` or ``no-path``, the last meaning ``HIPBLASLT_TUNING_CACHE_PATH`` is unset and the cache +does nothing. + +The summary, written as the process exits normally, is how you confirm a deployed cache is being +used. ``shapes`` counts distinct problems the process looked up, not calls, so it stays comparable +with ``loaded``. ``matched`` is how many of them a cache entry served and ``fellback`` how many fell +through to default selection; those are the shapes still worth tuning. ``invalidated`` counts entries +rejected because a rebuild moved the kernel they named. A process that is killed or aborts writes no +summary, and one that loaded nothing and looked nothing up leaves it out. + +Setting ``HIPBLASLT_LOG_LEVEL=4`` adds a cache hit, miss or invalidation line once per problem. See +:ref:`environment-variables` for where each notice is written. + +Limitations +----------- + +* Grouped GEMM is not served from the cache, and neither is a pointer-array batch called with + ``algo=nullptr``. +* ``alpha``, ``beta`` and whether ``C`` and ``D`` alias are not part of the lookup key, so an entry + can serve a caller whose values differ from those it was recorded with. diff --git a/projects/hipblaslt/docs/how-to/use-logging-heuristics.rst b/projects/hipblaslt/docs/how-to/use-logging-heuristics.rst index 19e147be6613..a6748fe15e5a 100644 --- a/projects/hipblaslt/docs/how-to/use-logging-heuristics.rst +++ b/projects/hipblaslt/docs/how-to/use-logging-heuristics.rst @@ -44,9 +44,15 @@ You can enable the hipBLASLt logging mechanism by setting the following environm "``64``","Profile" "``128``","Extended profile" + The levels are cumulative: level ``4`` also enables errors, trace and hints, and level ``5`` enables + everything above it. ``HIPBLASLT_LOG_MASK`` is only consulted when ``HIPBLASLT_LOG_LEVEL`` is unset, + and unlike a level it selects exactly the flags you list. + * ``HIPBLASLT_LOG_FILE=``: The ``file_name`` is a path to a logging file. The file name can contain ``%i``, which is replaced with the process ID, for example, ``_%i.log``. - If ``HIPBLASLT_LOG_FILE`` is not defined, the log messages are printed to stdout. + If ``HIPBLASLT_LOG_FILE`` is not defined, the log messages are printed to stderr. + This variable has no effect on its own: the log file is opened only when ``HIPBLASLT_LOG_LEVEL`` + or ``HIPBLASLT_LOG_MASK`` has enabled logging. * ``HIPBLASLT_ENABLE_MARKER=1``: Setting ``HIPBLASLT_ENABLE_MARKER`` to ``1`` enables marker trace for :doc:`ROCProfiler ` profiling. diff --git a/projects/hipblaslt/docs/reference/env-variables.rst b/projects/hipblaslt/docs/reference/env-variables.rst index 2a791d6af4b1..c261129ce8e1 100644 --- a/projects/hipblaslt/docs/reference/env-variables.rst +++ b/projects/hipblaslt/docs/reference/env-variables.rst @@ -26,6 +26,7 @@ For more information, see :doc:`Use logging and heuristics <../how-to/use-loggin * - | ``HIPBLASLT_LOG_LEVEL`` | Controls the verbosity level of hipBLASLt logging output. + | Levels are cumulative: each one also enables the levels below it. - | 0: Off (logging disabled, default) | 1: Error (only errors are logged) | 2: Trace (API calls with kernel launches log parameters) @@ -35,6 +36,7 @@ For more information, see :doc:`Use logging and heuristics <../how-to/use-loggin * - | ``HIPBLASLT_LOG_MASK`` | Controls logging output using bit mask flags (can be combined). + | Consulted only when ``HIPBLASLT_LOG_LEVEL`` is unset. - | 0: Off | 1: Error | 2: Trace @@ -47,8 +49,9 @@ For more information, see :doc:`Use logging and heuristics <../how-to/use-loggin * - | ``HIPBLASLT_LOG_FILE`` | Specifies path to logging file. Can contain ``%i`` for process ID replacement. + | Has no effect unless a level or mask has enabled logging. - | Path to log file (for example, ``logfile_%i.log``) - | If not defined: log messages printed to stdout + | If not defined: log messages printed to stderr * - | ``HIPBLASLT_ENABLE_MARKER`` | Enables marker trace for ROCProfiler profiling. @@ -84,6 +87,43 @@ For more information, see :doc:`Use hipBLASLt offline tuning <../how-to/how-to-u - | Integer value in bytes (default: 128 * 1024 * 1024) | Limits workspace size for solution selection +Runtime tuning cache +==================== + +The runtime tuning cache is opt-in and off by default. ``HIPBLASLT_TUNING_MODE`` and +``HIPBLASLT_TUNING_CACHE_PATH`` are read the first time they are needed in a process, so set them +before the first hipBLASLt call. A process running in a secure execution context (set-user-ID, +set-group-ID, or another credential-changing exec such as file capabilities) ignores both, and the +cache stays off. For more information, see +:doc:`Use hipBLASLt offline tuning <../how-to/how-to-use-hipblaslt-offline-tuning>`. + +``cache`` mode writes two notices without any logging variable: one line naming the mode, the file +and how many entries loaded, and a closing summary. Replaying the cache adds no output per call. +Where the notices go depends on logging: + +* No level or mask: stderr. ``HIPBLASLT_LOG_FILE`` alone does not open a log file. +* ``HIPBLASLT_LOG_LEVEL`` 1 to 3, or a mask without the info bit: the stream logging already opened, + including ``HIPBLASLT_LOG_FILE``. +* ``HIPBLASLT_LOG_LEVEL=4`` or higher, or ``HIPBLASLT_LOG_MASK`` including ``8``: the same stream, + formatted like every other log line, plus a cache hit, miss or invalidation line once per problem. + +.. list-table:: + :header-rows: 1 + :widths: 70,30 + + * - **Environment variable** + - **Value** + + * - | ``HIPBLASLT_TUNING_MODE`` + | Selects runtime tuning behavior. + - | ``off``: Disable runtime tuning (default) + | ``cache``: Replay valid entries from the cache file + + * - | ``HIPBLASLT_TUNING_CACHE_PATH`` + | Specifies the runtime cache file. + - | Path to a tuning file + | Required for ``cache`` mode + Origami with Stream-K configuration =================================== diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/TuningCacheStore.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/TuningCacheStore.cpp index 3fe9819d9ac4..927dd8ab8896 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/TuningCacheStore.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/TuningCacheStore.cpp @@ -3,6 +3,7 @@ #include "TuningCacheStore.hpp" #include "hipblaslt_type_strings.hpp" +#include "rocblaslt_secure_env.hpp" #include "tensile_type_map.hpp" #include @@ -15,6 +16,28 @@ namespace TensileLite { + TuningModeConfig TuningModeConfig::fromEnvironment(bool isPrivileged) + { + TuningModeConfig config; + + if(const char* env = rocblaslt_secure_getenv_impl("HIPBLASLT_TUNING_MODE", isPrivileged)) + { + if(std::string(env) == "cache") + config.mode = TuningMode::Cache; + } + + if(const char* path + = rocblaslt_secure_getenv_impl("HIPBLASLT_TUNING_CACHE_PATH", isPrivileged)) + config.cachePath = path; + + config.suppressedForSecurity + = rocblaslt_env_suppressed_for_security_impl("HIPBLASLT_TUNING_MODE", isPrivileged) + || rocblaslt_env_suppressed_for_security_impl("HIPBLASLT_TUNING_CACHE_PATH", + isPrivileged); + + return config; + } + namespace { const char* const kGitVersionHeader = "Git Version:"; diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp index 23284de6e057..32bad51786e8 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp @@ -26,9 +26,16 @@ * ************************************************************************ */ #include "UserDrivenTuningParser.hpp" +#include "rocblaslt_secure_env.hpp" +#include "utility.hpp" +#include #include #include +#include +#include +#include +#include #ifndef TO_STR2 #define TO_STR2(x) #x @@ -47,8 +54,269 @@ namespace TensileLite return stamp; } + void TuningModeSingleton::load() + { + m_config = TuningModeConfig::fromEnvironment(rocblaslt_process_is_privileged()); + + if(m_config.suppressedForSecurity) + log_error(__func__, + "Ignoring HIPBLASLT_TUNING_MODE and HIPBLASLT_TUNING_CACHE_PATH because the " + "process is running in a secure execution context (set-uid/set-gid or " + "another credential-changing exec, such as file capabilities); tuning " + "stays off."); + } + + TuningFileSelection selectTuningFile() + { + const auto& tuning = TuningModeSingleton::getInstance(); + + if(tuning.mode() == TuningMode::Off) + { + const OverrideSingleton& legacy = OverrideSingleton::getInstance(); + if(!legacy.env_mode) + return {}; + return {true, legacy.file_path}; + } + + if(tuning.cachePath().empty()) + { + // Without a path nothing is ever loaded, so this is the only place + // that can tell a misconfigured process why nothing happens. + announceTuningModeOnce(TuningLoadStatus::NoPath); + return {}; + } + + return {true, tuning.cachePath()}; + } + + namespace + { + /** + * Process-lifetime state behind the tuning diagnostics. In the source + * rather than the header so its containers can use + * std::hash. + */ + struct DiagnosticsState + { + static DiagnosticsState& instance() + { + static DiagnosticsState gInstance; + return gInstance; + } + + bool infoOn() const + { + return (get_logger_layer_mode() & rocblaslt_layer_mode_log_info) != 0; + } + + std::atomic announced{false}; + + mutable std::mutex mutex; + std::unordered_set hitKeys; + std::unordered_set missKeys; + std::unordered_set invalidKeys; + std::unordered_map lookups; + + // Indexes already counted as rejected, per key. A key can hold + // several entries, and each is its own rejected entry. + std::unordered_map> invalidEntries; + }; + + void summaryTally(uint64_t* shapes, uint64_t* matched, uint64_t* fellback) + { + DiagnosticsState& state = DiagnosticsState::instance(); + std::lock_guard lock(state.mutex); + + uint64_t served = 0; + for(const auto& lookup : state.lookups) + if(lookup.second) + served++; + + *shapes = state.lookups.size(); + *matched = served; + *fellback = state.lookups.size() - served; + } + + /** + * Writes the closing summary when the process unwinds its statics. + * + * The sink and layer mode are captured at registration rather than read + * back here, so nothing queries the logger while statics are being torn + * down. + */ + struct SummaryEmitter + { + SummaryEmitter(std::ostream* sink, uint32_t layerMode) + : m_sink(sink) + , m_layerMode(layerMode) + { + } + + ~SummaryEmitter() + { + // Runs during static destruction, where an exception terminates + // the process, and the sink throws on failure. + try + { + emit(); + } + catch(...) + { + } + } + + void emit() const + { + uint64_t shapes = 0, matched = 0, fellback = 0; + summaryTally(&shapes, &matched, &fellback); + + const auto& counters = TuningCounters::instance(); + const uint64_t invalidated = counters.invalidated.load(); + + // All zeroes would read as "the cache served nothing" when the + // cache was never reached at all. + if(shapes == 0 && invalidated == 0 && counters.entriesLoaded.load() == 0) + return; + + std::ostringstream msg; + msg << "tuning-cache: summary shapes=" << shapes << " matched=" << matched + << " fellback=" << fellback << " invalidated=" << invalidated; + + std::ostream* os = m_sink ? m_sink : &std::cerr; + std::lock_guard lock(log_mutex); + + // Laid out by hand rather than through log_info, which would + // query the logger during static destruction, with the spacing + // log_arguments would have used. + if(m_layerMode & rocblaslt_layer_mode_log_info) + *os << prefix(rocblaslt_layer_mode2string(rocblaslt_layer_mode_log_info), kFunc) + << " " << msg.str() << std::endl; + else + *os << msg.str() << std::endl; + } + + // Stands in for __func__, which here would name a destructor that + // runs at exit. + static constexpr const char* kFunc = "tuning_summary"; + + std::ostream* m_sink; + uint32_t m_layerMode; + }; + } // namespace + + void announceTuningModeOnce(TuningLoadStatus status) + { + const auto& tuning = TuningModeSingleton::getInstance(); + if(tuning.mode() == TuningMode::Off) + return; + + // A plain load first: selectTuningFile calls this on every heuristic + // call when the path is missing. + DiagnosticsState& state = DiagnosticsState::instance(); + if(state.announced.load(std::memory_order_relaxed) || state.announced.exchange(true)) + return; + + const char* statusName = "ok"; + switch(status) + { + case TuningLoadStatus::Ok: + statusName = "ok"; + break; + case TuningLoadStatus::NotFound: + statusName = "not-found"; + break; + case TuningLoadStatus::ReadError: + statusName = "read-error"; + break; + case TuningLoadStatus::NoPath: + statusName = "no-path"; + break; + } + + std::ostringstream msg; + msg << "mode=cache path=" << tuning.cachePath() << " load=" << statusName + << " loaded=" << TuningCounters::instance().entriesLoaded.load(); + + if(status == TuningLoadStatus::NoPath) + msg << "; set HIPBLASLT_TUNING_CACHE_PATH or the cache does nothing"; + + if(status == TuningLoadStatus::ReadError) + msg << "; the file is there but could not be read"; + + if(OverrideSingleton::getInstance().env_mode) + msg << "; ignoring HIPBLASLT_TUNING_OVERRIDE_FILE because a tuning mode is set"; + + log_tuning_lifecycle(__func__, msg.str()); + + // Registered only after the line above has gone through the logger, so + // LoggerSingleton is already constructed and outlives this emitter: + // statics are destroyed in reverse order of construction. + static SummaryEmitter emitter(get_logger_os(), get_logger_layer_mode()); + static_cast(emitter); + } + + void recordTuningLookup(const ProblemOverride& key, bool matched) + { + if(TuningModeSingleton::getInstance().mode() == TuningMode::Off) + return; + + DiagnosticsState& state = DiagnosticsState::instance(); + std::lock_guard lock(state.mutex); + + auto inserted = state.lookups.try_emplace(key, matched); + if(!inserted.second && matched) + inserted.first->second = true; + } + + bool recordTuningInvalidation(const ProblemOverride& key, int solutionIndex) + { + DiagnosticsState& state = DiagnosticsState::instance(); + std::lock_guard lock(state.mutex); + return state.invalidEntries[key].insert(solutionIndex).second; + } + + bool shouldLogTuningKeyEvent(TuningKeyEvent kind, const ProblemOverride& key) + { + DiagnosticsState& state = DiagnosticsState::instance(); + if(!state.infoOn()) + return false; + + std::lock_guard lock(state.mutex); + switch(kind) + { + case TuningKeyEvent::Hit: + return state.hitKeys.insert(key).second; + case TuningKeyEvent::Miss: + return state.missKeys.insert(key).second; + case TuningKeyEvent::Invalid: + return state.invalidKeys.insert(key).second; + } + return false; + } + + void resetTuningDiagnosticsForTest() + { + DiagnosticsState& state = DiagnosticsState::instance(); + state.announced = false; + + std::lock_guard lock(state.mutex); + state.hitKeys.clear(); + state.missKeys.clear(); + state.invalidKeys.clear(); + state.invalidEntries.clear(); + state.lookups.clear(); + } + + void tuningLookupTallyForTest(uint64_t* shapes, uint64_t* matched, uint64_t* fellback) + { + summaryTally(shapes, matched, fellback); + } + void getContractionProblemsFromFile(const std::string& path) { + if(path.empty()) + return; + OverrideMap& map = OverrideMap::getMap(); // Runs on every heuristic query. Once the file is loaded, the read lock @@ -60,29 +328,49 @@ namespace TensileLite if(map.isLoaded(path)) return; - // A file that does not open is looked for again on the next query. - std::ifstream file(path); + const auto& tuning = TuningModeSingleton::getInstance(); + const bool managed = tuning.mode() != TuningMode::Off && path == tuning.cachePath(); + + TuningLoadStatus status = TuningLoadStatus::Ok; + std::ifstream file(path); if(!file.is_open()) - return; + { + // A missing file is a status rather than an error, and is looked + // for again on the next query, since tune mode creates it. A file + // that is there and still will not open is a different problem, + // usually permissions, and reporting it as missing would send the + // user looking for a file they already have. + std::error_code ec; + status = std::filesystem::exists(path, ec) ? TuningLoadStatus::ReadError + : TuningLoadStatus::NotFound; + } + else + { + const auto loaded = loadTuningRows(file, map, currentBuildStamp()); + TuningCounters::instance().entriesLoaded += loaded.accepted; - const auto loaded = loadTuningRows(file, map, currentBuildStamp()); + if(loaded.skippedUnnamed > 0) + log_error(__func__, + "Ignored " + std::to_string(loaded.skippedUnnamed) + + " entries without a kernel name in " + path + + ": its Git Version line does not match this build. Re-run the " + "tuning with this build to use them."); + else if(currentBuildStamp().empty() || loaded.fileBuildStamp != currentBuildStamp()) + log_info(__func__, + path + + " has no Git Version line matching this build; each entry is " + "used only while its kernel_name still matches."); - if(loaded.skippedUnnamed > 0) - log_error(__func__, - "Ignored " + std::to_string(loaded.skippedUnnamed) - + " entries without a kernel name in " + path - + ": its Git Version line does not match this build. Re-run the " - "tuning with this build to use them."); - else if(currentBuildStamp().empty() || loaded.fileBuildStamp != currentBuildStamp()) - log_info(__func__, - path - + " has no Git Version line matching this build; each entry is used " - "only while its kernel_name still matches."); - - // Only a clean read counts as loaded. A read that stopped on an I/O - // error partway through would otherwise leave a partial map that is - // never completed. - if(!loaded.readError) - map.markLoaded(path); + // Only a clean read counts as loaded. A read that stopped on an I/O + // error partway through would otherwise leave a partial map that is + // never completed. + if(loaded.readError) + status = TuningLoadStatus::ReadError; + else + map.markLoaded(path); + } + + if(managed) + announceTuningModeOnce(status); } } // namespace TensileLite diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp index 770307f97850..3c0382a5bcb4 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp @@ -27,6 +27,37 @@ namespace TensileLite { + enum class TuningMode : uint32_t + { + Off = 0, + Cache = 1, + }; + + /** + * HIPBLASLT_TUNING_MODE and HIPBLASLT_TUNING_CACHE_PATH as the environment + * sets them. + */ + struct TuningModeConfig + { + TuningMode mode = TuningMode::Off; + std::string cachePath; + + // A process in a secure execution context ignores both variables and + // stays off: they choose a file for it to read, which an inherited + // environment must not be able to impose on it. Set when that is why + // the mode is off, so the caller can say so. + bool suppressedForSecurity = false; + + static TuningModeConfig fromEnvironment(bool isPrivileged); + + // There is no default cache location, so a mode without a path does + // nothing. + bool reads() const + { + return mode != TuningMode::Off && !cachePath.empty(); + } + }; + /** * The schema_version a row was written with. * diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp index a2d0fbea12d3..226ae3d059b3 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp @@ -32,6 +32,8 @@ #include "tensile_host.hpp" #include +#include +#include #include class OverrideSingleton @@ -85,6 +87,93 @@ class OverrideSingleton namespace TensileLite { + /** + * HIPBLASLT_TUNING_MODE and HIPBLASLT_TUNING_CACHE_PATH, read on first use. + * + * Read once rather than per call, so the hot path costs nothing and a + * process cannot change mode halfway through a run. Setting either variable + * after the first hipBLASLt call has no effect. + */ + class TuningModeSingleton + { + public: + static TuningModeSingleton& getInstance() + { + static TuningModeSingleton gInstance; + return gInstance; + } + + TuningModeSingleton(const TuningModeSingleton&) = delete; + TuningModeSingleton& operator=(const TuningModeSingleton&) = delete; + + TuningMode mode() const + { + return m_config.mode; + } + const std::string& cachePath() const + { + return m_config.cachePath; + } + bool reads() const + { + return m_config.reads(); + } + + /** Re-read the environment. Tests only, like OverrideSingleton::reloadForTest. */ + void reloadForTest() + { + load(); + } + + private: + TuningModeSingleton() + { + load(); + } + + void load(); + + TuningModeConfig m_config; + }; + + /** Running tallies behind the closing summary and the test hooks. */ + struct TuningCounters + { + static TuningCounters& instance() + { + static TuningCounters gInstance; + return gInstance; + } + + std::atomic entriesLoaded{0}; + std::atomic hits{0}; + std::atomic misses{0}; + std::atomic invalidated{0}; + + std::string summary() const + { + return "loaded=" + std::to_string(entriesLoaded.load()) + " hits=" + + std::to_string(hits.load()) + " misses=" + std::to_string(misses.load()) + + " invalidated=" + std::to_string(invalidated.load()); + } + }; + + /** + * Which tuning file this process consults, if any. + * + * HIPBLASLT_TUNING_CACHE_PATH and HIPBLASLT_TUNING_OVERRIDE_FILE are + * mutually exclusive rather than merged: with no tuning mode set the + * override file behaves as it always has, and with one set only the cache + * is consulted and the override is ignored, which the startup line says. + */ + struct TuningFileSelection + { + bool active = false; + std::string path; + }; + + TuningFileSelection selectTuningFile(); + /** The build rows are trusted against: the running library's own. */ const std::string& currentBuildStamp(); @@ -93,4 +182,56 @@ namespace TensileLite * once per process. */ void getContractionProblemsFromFile(const std::string& path); + + /** How the cache file was read, for the startup line. */ + enum class TuningLoadStatus : uint32_t + { + Ok = 0, + NotFound, + ReadError, + NoPath, + }; + + /** + * Announce the mode, path and load result once per process, and arrange for + * the closing summary. Does nothing in off mode. + */ + void announceTuningModeOnce(TuningLoadStatus status); + + /** + * Note that a lookup for this problem did or did not find a usable entry. + * + * Counted by distinct key rather than by call, because the summary is read + * against loaded=N and a hot loop over one uncached shape would otherwise + * report thousands of fallbacks for one missing row. A key that matches + * once counts as matched. Does nothing in off mode. + */ + void recordTuningLookup(const ProblemOverride& key, bool matched); + + /** + * True only the first time this key's entry at this index is rejected. The + * heuristic lookup and the execution path can both meet one stale row, and + * it is still one rejected entry. + */ + bool recordTuningInvalidation(const ProblemOverride& key, int solutionIndex); + + /** Cache events logged at most once per key. */ + enum class TuningKeyEvent : uint32_t + { + Hit = 0, + Miss, + Invalid, + }; + + /** + * Whether this key's event is worth logging: only with the info bit set, + * and once per key, since replay meets the same key on every call. + */ + bool shouldLogTuningKeyEvent(TuningKeyEvent kind, const ProblemOverride& key); + + /** Drop the announcement latch and every per-key set. Tests only. */ + void resetTuningDiagnosticsForTest(); + + /** The distinct-key tally behind the summary line, for tests. */ + void tuningLookupTallyForTest(uint64_t* shapes, uint64_t* matched, uint64_t* fellback); } // namespace TensileLite diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp index 11c8c9cb90be..5fd26ab3e251 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp @@ -181,6 +181,17 @@ rocblaslt_status isSolutionSupported(rocblaslt_handle handle, rocblaslt_matmul_algo* algo, size_t* workspaceSizeInBytes); +/** + * isSolutionSupported without leaving gemmData's Tensile problem modified. A + * cache probe may test a math mode or algorithm other than the one the call + * goes on to launch. + */ +rocblaslt_status isSolutionSupportedNoMutation(rocblaslt_handle handle, + const RocblasltContractionProblem& prob, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo, + size_t* workspaceSizeInBytes); + template rocblaslt_status isSolutionSupported(rocblaslt_handle handle, const rocblaslt::RocGemmType& gemmType, @@ -225,6 +236,25 @@ TensileLite::ProblemOverride TensileLite::ProblemOverride TensileDataGemm2ProblemOverride(std::shared_ptr); +/** + * The solution index of the first entry for this key that still resolves to its + * recorded kernel and supports this problem, or -1 when there is none. + * + * A pure probe: the caller counts the lookup, because only it knows whether the + * probe decided which kernel the call launches. + */ +int tuning_cache_find_valid_entry(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key, + const RocblasltContractionProblem& problem, + std::shared_ptr gemmData, + size_t max_workspace_bytes); + +/** + * The solution index this thread last launched through runContractionProblem, + * or -1, clearing it. + */ +int tuningLastLaunchedIndexForTest(); + TensileLite::ContractionProblemGemm* ExtractProblemGemm(std::shared_ptr); // Push the GemmPreference-supplied StreamK tile scheduling mode onto every diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp index 8055880a64cd..8de33dbe93ae 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/utility.hpp @@ -199,6 +199,46 @@ void log_info(const char* func, H head, Ts&&... xs) log_base(rocblaslt_layer_mode_log_info, func, head, std::forward(xs)...); } +/** + * A tuning notice that stays visible without a log level. + * + * With the info bit set it goes through log_info like any other diagnostic, so + * a level-4 run never prints it twice. Otherwise it is written bare to the + * stream logging already opened, so levels 1 to 3 still honour + * HIPBLASLT_LOG_FILE, or to stderr when logging is off. It never opens the log + * file itself, which would create and truncate it for a process that asked for + * tuning but not for logging. + * + * Callers check the tuning mode: reading it here would make this header, which + * most of the library includes, depend on Tensile. + * + * Sink failures are swallowed. open_log_stream arms the log file to throw, and + * these notices are written from destructors, where an escaping exception + * terminates the process, and from inside hipBLASLt calls, where it would turn + * a logging problem into a failed call. + */ +inline void log_tuning_lifecycle(const char* func, const std::string& body) +{ + try + { + const std::string full = "tuning-cache: " + body; + + // Outside the lock: log_base takes log_mutex itself. + if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info) + { + log_info(func, full); + return; + } + + std::lock_guard lock(log_mutex); + std::ostream* os = get_logger_os(); + *(os ? os : &std::cerr) << full << std::endl; + } + catch(...) + { + } +} + // if trace logging is turned on with // (handle->layer_mode & rocblaslt_layer_mode_log_api) == true // then diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp index dfd0cb283f18..5934817240de 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp @@ -176,6 +176,13 @@ inline bool // Preload problem/solution mappings namespace { + // With no tuning mode set, the file in play is HIPBLASLT_TUNING_OVERRIDE_FILE. + bool override_file_in_play() + { + return TensileLite::TuningModeSingleton::getInstance().mode() + == TensileLite::TuningMode::Off; + } + /** * Whether an entry's recorded name still matches what its index resolves to. * @@ -184,6 +191,7 @@ namespace * or refused on the file's build stamp when it was loaded. */ bool tuned_entry_identity_matches(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key, const TensileLite::TunedEntry& entry, const rocblaslt_matmul_heuristic_result& resolved) { @@ -210,16 +218,68 @@ namespace if(current == recorded) return true; - if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info) + auto& counters = TensileLite::TuningCounters::instance(); + if(TensileLite::recordTuningInvalidation(key, entry.solutionIndex)) + counters.invalidated++; + + // Replay meets a stale row on every call for its shape, so cache mode + // reports it once per key. + if(override_file_in_play()) + { + if(get_logger_layer_mode() & rocblaslt_layer_mode_log_info) + { + std::ostringstream msg; + msg << "Ignoring tuning file entry: index " << entry.solutionIndex + << " now resolves to '" << current << "', recorded '" << recorded << "'"; + log_info(__func__, msg.str()); + } + } + else if(TensileLite::shouldLogTuningKeyEvent(TensileLite::TuningKeyEvent::Invalid, key)) { std::ostringstream msg; - msg << "Ignoring tuning file entry: index " << entry.solutionIndex - << " now resolves to '" << current << "', recorded '" << recorded << "'"; + msg << "tuning-cache: cache-invalid index=" << entry.solutionIndex + << " now resolves to '" << current << "', recorded '" << recorded << "' [" + << counters.summary() << "]"; log_info(__func__, msg.str()); } return false; } + /** + * Count a lookup and report its result: in the override file's historical + * wording with no tuning mode set, and once per key in cache mode, where + * replay repeats the same lookup on every call. + */ + void record_lookup(const char* func, + const TensileLite::ProblemOverride& key, + bool matched, + int index) + { + auto& counters = TensileLite::TuningCounters::instance(); + if(matched) + counters.hits++; + else + counters.misses++; + TensileLite::recordTuningLookup(key, matched); + + if(override_file_in_play()) + { + if(matched) + log_info(func, "Find solution with index: " + std::to_string(index)); + else + log_info(func, "No valid solution index found in override file."); + } + else if(TensileLite::shouldLogTuningKeyEvent(matched ? TensileLite::TuningKeyEvent::Hit + : TensileLite::TuningKeyEvent::Miss, + key)) + { + if(matched) + log_info(func, "tuning-cache: cache-hit index=" + std::to_string(index)); + else + log_info(func, "tuning-cache: cache-miss, no valid entry for this problem"); + } + } + /** * Every entry for a key: current-schema rows first, then legacy rows, * which can only be matched on the fields the legacy format records. @@ -248,6 +308,14 @@ bool problem_override_from_file(rocblaslt_handle& handle, if(m_override.size() == 0) { + // Still a lookup that fell back, and counted as one, or a cache that + // loaded nothing would report that nothing was asked. The key is built + // only in cache mode, where it is tallied. + TensileLite::TuningCounters::instance().misses++; + if(!override_file_in_play()) + TensileLite::recordTuningLookup(RocblasltContractionProblem2ProblemOverride(problem), + false); + log_info(__func__, "No valid entries found in override file."); } else @@ -272,7 +340,7 @@ bool problem_override_from_file(rocblaslt_handle& handle, handle, solutionIndex, overrideResults, max_workspace_bytes) && !overrideResults.empty()) { - if(!tuned_entry_identity_matches(handle, entry, overrideResults[0])) + if(!tuned_entry_identity_matches(handle, prob_key, entry, overrideResults[0])) continue; size_t required_workspace_size = 0; @@ -318,16 +386,7 @@ bool problem_override_from_file(rocblaslt_handle& handle, } } - if(!success) - { - log_info(__func__, "No valid solution index found in override file."); - } - else - { - std::string mapping_result = "Find solution with index: "; - mapping_result += std::to_string(solutionIndex[0]); - log_info(__func__, mapping_result); - } + record_lookup(__func__, prob_key, success, solutionIndex[0]); } return success; @@ -345,7 +404,7 @@ bool problem_override_from_file_cpp( // gemmData as a TensileDataGemm. if(gemmType != rocblaslt::RocGemmType::ROCBLASLT_GEMM) { - log_info(__func__, "Grouped GEMM does not use the override file."); + log_info(__func__, "Grouped GEMM does not use the tuning file."); return false; } @@ -355,6 +414,11 @@ bool problem_override_from_file_cpp( if(m_override.size() == 0) { + // Counted for the same reason as the C API. The key is already on the + // gemm data, so there is nothing to build. + TensileLite::TuningCounters::instance().misses++; + TensileLite::recordTuningLookup(TensileDataGemm2ProblemOverride(gemmData), false); + log_info(__func__, "No valid entries found in override file."); } else @@ -378,7 +442,7 @@ bool problem_override_from_file_cpp( handle, solutionIndex, overrideResults, max_workspace_bytes) && !overrideResults.empty()) { - if(!tuned_entry_identity_matches(handle, entry, overrideResults[0])) + if(!tuned_entry_identity_matches(handle, prob_key, entry, overrideResults[0])) continue; size_t required_workspace_size = 0; @@ -429,19 +493,61 @@ bool problem_override_from_file_cpp( } } - if(!success) - { - log_info(__func__, "No valid solution index found in override file."); - } - else + record_lookup(__func__, prob_key, success, solutionIndex[0]); + } + + return success; +} + +int tuning_cache_find_valid_entry(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key, + const RocblasltContractionProblem& problem, + std::shared_ptr gemmData, + size_t max_workspace_bytes) +{ + std::vector resolved; + std::vector index(1); + + for(const auto& entry : tuned_entries_for(key)) + { + index[0] = entry.solutionIndex; + + // getSolutionsFromIndex appends, and everything below reads [0]. + resolved.clear(); + + if(rocblaslt_status_success + != getSolutionsFromIndex(handle, index, resolved, max_workspace_bytes) + || resolved.empty()) + continue; + if(!tuned_entry_identity_matches(handle, key, entry, resolved[0])) + continue; + + // The key leaves out values such as beta and C/D aliasing, so an entry + // whose identity holds can still fail this call's support predicates. + RocblasltContractionProblem supportProblem = problem; + size_t required = 0; + auto algo = resolved[0].algo; + if(rocblaslt_status_success + == isSolutionSupportedNoMutation(handle, supportProblem, gemmData, &algo, &required) + && required <= max_workspace_bytes) + return entry.solutionIndex; + + // The heuristic replay's XF32 fallback, as a pure probe: the call may go + // on to launch some other algorithm, which must not find gemmData left + // in FP32 mode. + if(problem.compute_type == rocblaslt_compute_f32_fast_xf32) { - std::string mapping_result = "Find solution with index: "; - mapping_result += std::to_string(solutionIndex[0]); - log_info(__func__, mapping_result); + supportProblem.compute_type = rocblaslt_compute_f32; + required = 0; + if(rocblaslt_status_success + == isSolutionSupportedNoMutation( + handle, supportProblem, gemmData, &algo, &required) + && required <= max_workspace_bytes) + return entry.solutionIndex; } } - return success; + return -1; } /****************************************************************************** @@ -2394,20 +2500,20 @@ rocblaslt_status } #endif - OverrideSingleton& override = OverrideSingleton::getInstance(); - bool override_success = false; + const auto tuningFile = TensileLite::selectTuningFile(); + bool override_success = false; // Set before the lookup: a hit for a single-algo request skips // getBestSolutions, and the dedup below still reads this count. *returnAlgoCount = 0; - if(override.env_mode) + if(tuningFile.active) { override_success = problem_override_from_file(handle, prob, matmul_desc, heuristicResultsArray, - override.file_path, + tuningFile.path, pref->max_workspace_bytes); if(override_success) requestedAlgoCount--; @@ -2678,14 +2784,14 @@ rocblaslt_status rocblaslt_status status = rocblaslt_status_success; try { - OverrideSingleton& override = OverrideSingleton::getInstance(); + const auto tuningFile = TensileLite::selectTuningFile(); bool override_success = false; std::vector override_result; - if(override.env_mode) + if(tuningFile.active) { override_success = problem_override_from_file_cpp( - handle, gemmType, gemmData, override_result, override.file_path, maxWorkspaceBytes); + handle, gemmType, gemmData, override_result, tuningFile.path, maxWorkspaceBytes); log_api(__func__, "OverrideAlgoCount", override_success ? 1 : 0); } @@ -2989,11 +3095,60 @@ extern "C" HIPBLASLT_EXPORT void hipblaslt_debug_reload() TensileLite::Debug::Instance().reloadDebugBitsForTest(); } -// Test support, like hipblaslt_debug_reload: HIPBLASLT_TUNING_OVERRIDE_FILE is -// read on first use and each file is loaded once per process. Not part of any -// supported interface. +// Test support, like hipblaslt_debug_reload: the tuning variables are read on +// first use, each file is loaded once per process, and the diagnostics latch +// for the process. Not part of any supported interface. extern "C" HIPBLASLT_EXPORT void hipblaslt_tuning_reset_for_test() { + TensileLite::TuningModeSingleton::getInstance().reloadForTest(); TensileLite::OverrideMap::getMap().resetForTest(); OverrideSingleton::getInstance().reloadForTest(); + TensileLite::resetTuningDiagnosticsForTest(); + + auto& counters = TensileLite::TuningCounters::instance(); + counters.entriesLoaded = 0; + counters.hits = 0; + counters.misses = 0; + counters.invalidated = 0; + + static_cast(tuningLastLaunchedIndexForTest()); +} + +// The solution the calling thread's last hipblasLtMatmul launched, or -1, +// cleared by reading. The counters say that a lookup matched; only this says +// which kernel ran. +extern "C" HIPBLASLT_EXPORT int hipblaslt_tuning_last_launch_for_test() +{ + return tuningLastLaunchedIndexForTest(); +} + +// The tallies tests assert on, so they need not scrape log output: +// +// 0 loaded, 1 hits, 2 misses, 3 invalidated +// The counters, which count lookups. +// 4 shapes, 5 matched, 6 fell back +// The distinct-shape tally behind the summary line, which is written +// during static destruction, after any capture a test could install. +// +// Fills up to count values in that order and returns how many there are, so a +// test expecting a different list notices instead of misreading it. New values +// are only ever appended. +extern "C" HIPBLASLT_EXPORT size_t hipblaslt_tuning_stats_for_test(uint64_t* values, size_t count) +{ + const auto& c = TensileLite::TuningCounters::instance(); + uint64_t shapes = 0, matched = 0, fellback = 0; + TensileLite::tuningLookupTallyForTest(&shapes, &matched, &fellback); + + const uint64_t all[] = {c.entriesLoaded.load(), + c.hits.load(), + c.misses.load(), + c.invalidated.load(), + shapes, + matched, + fellback}; + + constexpr size_t known = sizeof(all) / sizeof(all[0]); + for(size_t i = 0; values && i < count && i < known; i++) + values[i] = all[i]; + return known; } diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 53d8ce67ad5a..93a885de8f2c 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -3650,6 +3650,18 @@ static rocblaslt_status bindSynchronizers(rocblaslt_handle return rocblaslt_status_success; } +namespace +{ + thread_local int t_lastLaunchedIndex = -1; +} // namespace + +int tuningLastLaunchedIndexForTest() +{ + const int index = t_lastLaunchedIndex; + t_lastLaunchedIndex = -1; + return index; +} + /****************************************************************************** * runContractionProblem calls Tensile to run a contraction problem described * * by RocblasltContractionProblem * @@ -3686,6 +3698,11 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle if(prob.trans_b == HIPBLAS_OP_C) data->problem.setBOps({TensileLite::TensorOp::ComplexConjugate()}); + // Noted before algo is filled in below. An explicit algo launches as + // given, whatever the cache holds; the cache chooses the kernel only + // when the caller leaves that choice to the library. + const bool callerSuppliedAlgo = (algo != nullptr); + if(algo == nullptr) { int returnAlgoCount = 0; @@ -3705,6 +3722,53 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle int32_t coldIterations = ClientArguments.GetColdIterationsValue(); int* solutionIndex = (int*)algo->data; + + // Cache replay for a call that passed no algo. hipblasLtMatmul accepts + // algo == nullptr and then runs getBestSolutions above without ever + // entering hipblasLtMatmulAlgoGetHeuristic, so without this such a call + // would never be served from the cache. Only the index used for this + // launch changes; the caller's algo is const. + int launchIndex = -1; + try + { + const auto& tuning = TensileLite::TuningModeSingleton::getInstance(); + + if(tuning.reads() && !callerSuppliedAlgo && !prob.grouped_gemm + && prob.batchMode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY) + { + // A matmul-only caller never reaches the heuristic entry point, + // which is the only other place that loads the file. + TensileLite::getContractionProblemsFromFile(tuning.cachePath()); + + const TensileLite::ProblemOverride key + = RocblasltContractionProblem2ProblemOverride(prob); + const int cachedIndex = tuning_cache_find_valid_entry( + handle, key, prob, gemmData, prob.workspaceSize); + + TensileLite::recordTuningLookup(key, cachedIndex >= 0); + auto& counters = TensileLite::TuningCounters::instance(); + if(cachedIndex >= 0) + { + counters.hits++; + launchIndex = cachedIndex; + } + else + { + counters.misses++; + } + } + } + catch(...) + { + // The cache is an optimisation: a lookup that throws leaves the call + // on default selection rather than failing it. + static_cast(hipGetLastError()); + launchIndex = -1; + } + + if(launchIndex >= 0) + solutionIndex = &launchIndex; + data->algoIndex = *solutionIndex; data->inputs = GetTensileInputs(prob); @@ -3875,6 +3939,7 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle } isPreloaded = true; } + t_lastLaunchedIndex = solution->index; status = hip2RocStatus( adapter->launchKernels(kernels, prob.stream, nullptr, nullptr, isPreloaded)); if(rocblaslt::Debug::Instance().printLogAsMarker()) @@ -3924,7 +3989,7 @@ rocblaslt_status gemmCreate(RocblasltContractionProblem const& problem, // Building the key queries the device, and this runs on every gemm // object creation, so only when a tuning file will be consulted. - const bool cacheTuningKey = OverrideSingleton::getInstance().env_mode; + const bool cacheTuningKey = TensileLite::selectTuningFile().active; if(gemmData) { @@ -5475,6 +5540,38 @@ rocblaslt_status isSolutionSupported(rocblaslt_handle handle, return isSolutionSupported(handle, data->problem, prob, algo, tuning, workspaceSizeInBytes); } +rocblaslt_status isSolutionSupportedNoMutation(rocblaslt_handle handle, + const RocblasltContractionProblem& prob, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo, + size_t* workspaceSizeInBytes) +{ +#ifdef HIPBLASLT_USE_ROCROLLER + if(useRocRoller(handle, prob)) + { + RocblasltContractionProblem probe = prob; + return isRocRollerSolutionSupported(handle, probe, algo, workspaceSizeInBytes); + } +#endif + + std::shared_ptr data = std::static_pointer_cast(gemmData); + auto savedProblem = data->problem; + RocblasltContractionProblem probe = prob; + + try + { + const rocblaslt_status status + = isSolutionSupported(handle, probe, gemmData, algo, workspaceSizeInBytes); + data->problem = std::move(savedProblem); + return status; + } + catch(...) + { + data->problem = std::move(savedProblem); + throw; + } +} + template void setRestrictions(TensileLite::ContractionProblemGemm& tensile_prob, const T* alpha, From 031d4b0d271936beb9617d141e3a9f46182f66e5 Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 6 Oct 2026 12:37:52 -0700 Subject: [PATCH 3/5] fix(hipblaslt): profile the kernel selected by cache replay --- .../library/src/amd_detail/rocblaslt/src/tensile_host.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 93a885de8f2c..edb8d380c26a 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -3812,8 +3812,10 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle if(get_logger_layer_mode() & rocblaslt_layer_mode_log_extended_profile) { - std::string kernel_name = getKernelNameFromAlgoIndex(handle, *algo); - std::string Solution_name = getSolutionNameFromAlgoIndex(handle, *algo); + auto loggedAlgo = *algo; + *(int*)loggedAlgo.data = *solutionIndex; + std::string kernel_name = getKernelNameFromAlgoIndex(handle, loggedAlgo); + std::string Solution_name = getSolutionNameFromAlgoIndex(handle, loggedAlgo); logExtendedProfileFromTensileDataGemm(data->problem, data->inputs, From 5a8dc64ece2472b4c15f4497a89188c82427a392 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Wed, 7 Oct 2026 01:39:32 -0500 Subject: [PATCH 4/5] perf(hipblaslt): cut cache replay's cost on matmuls given no algo Replaying a cached entry made hipblasLtMatmul with no algo about 2.7 us slower than default selection on a 256^3 fp16 GEMM, and a cache file that was not there added 1.5 to 2.5 us to every call. Replay ran the heuristic before looking at the cache, then on every call checked each entry's kernel name against its index and probed support on a copy of the whole Tensile problem, and a missing file was opened and stat'ed each time. Replay now runs first, and default selection only when no entry serves the call. Name checks depend only on the loaded entries and the build, so they are made once per shape and kept until an entry is added, which OverrideMap::generation now reports. Support depends on the call and is still checked each time, but on the problem the launch goes on to use, putting back the parameters the check sets rather than copying every tensor. The lookup tally hears about a shape only when its answer improves, and a missing file is looked for again at most once a second; the heuristic still looks on every query. The extended profile names the kernel from algo again, which now always holds the index launched. Measured on gfx950 as host time spent in hipblasLtMatmul over 200k calls with no algo, 256^3 fp16 with f32 compute, against 7.0 us with the cache off: replaying a row went from +2.74 to +0.02 us, a file with no row for the shape from +0.23 to +0.18, and a missing file to +0.22. Co-authored-by: Cursor --- .../clients/tests/src/tuning_cache_test.cpp | 97 ++++++++++ .../rocblaslt/src/UserDrivenTuningParser.cpp | 31 +++ .../src/include/TuningCacheStore.hpp | 17 +- .../src/include/UserDrivenTuningParser.hpp | 12 ++ .../rocblaslt/src/include/tensile_host.hpp | 34 +++- .../rocblaslt/src/rocblaslt_auxiliary.cpp | 182 ++++++++++++++---- .../amd_detail/rocblaslt/src/tensile_host.cpp | 129 +++++++------ 7 files changed, 396 insertions(+), 106 deletions(-) diff --git a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp index cb2382c5542c..0a169fd75d36 100644 --- a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp +++ b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp @@ -18,6 +18,7 @@ #include #include +#include #include #include #include @@ -28,6 +29,7 @@ #include #include #include +#include #include #include @@ -179,6 +181,24 @@ namespace return found; } + /** The kernel an index resolves to in this build, or empty. */ + std::string kernelNameOf(int index) + { + std::string name; + hipblasLtHandle_t handle = nullptr; + if(hipblasLtCreate(&handle) != HIPBLAS_STATUS_SUCCESS) + return name; + + std::vector indexes{index}; + std::vector results; + if(hipblaslt_ext::getAlgosFromIndex(handle, indexes, results) == HIPBLAS_STATUS_SUCCESS + && !results.empty()) + name = hipblaslt_ext::getKernelNameFromAlgo(handle, results[0].algo); + + hipblasLtDestroy(handle); + return name; + } + /** The build this library reports, as hipblaslt-bench records it. */ std::string buildStamp() { @@ -1145,6 +1165,83 @@ namespace EXPECT_EQ(c.misses, 1u); } + // An entry whose name holds but whose solution cannot run the call falls + // back, and the support check leaves nothing behind for the launch that + // default selection then makes. + TEST_F(TuningCache_pre_checkin, NullAlgoFallsBackFromAnEntryThatCannotRunTheCall) + { + if(!haveSolutions(1)) + GTEST_SKIP() << "the heuristic offers no solution for this problem"; + + // An fp32 solution recorded for the fp16 problem: its index still + // names its kernel, but it cannot run fp16 data. + int fp32 = -1; + ASSERT_TRUE(extHeuristicIndex(&fp32, HIP_R_32F, HIPBLAS_COMPUTE_32F)); + const auto fp32Name = kernelNameOf(fp32); + ASSERT_FALSE(fp32Name.empty()); + writeTuningFile(m_path, m_stamp, {{fp32, fp32Name}}); + useCache(m_path); + + for(int call = 0; call < 2; call++) + { + SCOPED_TRACE(call); + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched, true)); + EXPECT_EQ(launched, m_identities[0].index); + } + + const auto c = counters(); + EXPECT_EQ(c.hits, 0u); + EXPECT_EQ(c.misses, 2u); + EXPECT_EQ(c.invalidated, 0u); + } + + // Replay keeps which entries still name their kernels per shape, so an + // entry loaded after a shape has fallen back must still be found. + TEST_F(TuningCache_pre_checkin, NullAlgoReplaysAnEntryLoadedAfterItsShapeFellBack) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + useCache(m_path); + + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); + EXPECT_EQ(launched, m_identities[0].index); + + // The heuristic looks for the file on every query, so this one loads it. + const auto& recorded = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + int selected = -1; + ASSERT_TRUE(runGemm(&selected)); + ASSERT_EQ(selected, recorded.index); + + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched, true)); + EXPECT_EQ(launched, recorded.index); + } + + // Replay looks for a missing file at most once a second rather than on + // every call, but it does look again. + TEST_F(TuningCache_pre_checkin, NullAlgoFindsAFileCreatedAfterItsFirstLook) + { + if(!haveSolutions(2)) + GTEST_SKIP() << "the heuristic offers one solution for this problem"; + + useCache(m_path); + + int launched = -1; + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); + EXPECT_EQ(launched, m_identities[0].index); + + const auto& recorded = m_identities[1]; + writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + std::this_thread::sleep_for(std::chrono::milliseconds(1100)); + + ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); + EXPECT_EQ(launched, recorded.index); + EXPECT_EQ(counters().loaded, 1u); + } + // A cache that served nothing must not read like a process that asked // nothing. TEST_F(TuningCache_pre_checkin, ShapesThatMissAnEmptyCacheAreStillCounted) diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp index 32bad51786e8..e7d7ad12559f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/UserDrivenTuningParser.cpp @@ -29,6 +29,7 @@ #include "rocblaslt_secure_env.hpp" #include "utility.hpp" +#include #include #include #include @@ -373,4 +374,34 @@ namespace TensileLite if(managed) announceTuningModeOnce(status); } + + namespace + { + using ReplayClock = std::chrono::steady_clock; + + // When replay may next look for a file that was not there. + std::atomic nextReplayLoad{0}; + } + + void loadTuningFileForReplay(const std::string& path) + { + if(path.empty() || OverrideMap::getMap().isLoaded(path)) + return; + + const auto now = ReplayClock::now().time_since_epoch().count(); + if(now < nextReplayLoad.load(std::memory_order_relaxed)) + return; + nextReplayLoad.store( + now + + std::chrono::duration_cast(std::chrono::seconds(1)) + .count(), + std::memory_order_relaxed); + + getContractionProblemsFromFile(path); + } + + void resetReplayLoadForTest() + { + nextReplayLoad.store(0, std::memory_order_relaxed); + } } // namespace TensileLite diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp index 3c0382a5bcb4..f90e2bba941f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/TuningCacheStore.hpp @@ -8,6 +8,7 @@ // here touches a device, the solution library or the logger, so the same code // builds into hipBLASLt and into a host-only test. +#include #include #include #include @@ -448,6 +449,17 @@ namespace TensileLite return addTo(m_legacy, key.legacyKey(), entry); } + /** + * Changes whenever an entry is added or the map is cleared, but not when + * a repeated row refreshes an entry's metadata. Lets a caller that + * derives something from a key's entries tell, without the lock, that + * it has to derive it again. + */ + uint64_t generation() const + { + return m_generation.load(std::memory_order_acquire); + } + /** * Whether a path has already been read. Tracked per path rather than * inferred from the map, so a file that yields no usable rows is still @@ -477,12 +489,13 @@ namespace TensileLite m_override.clear(); m_legacy.clear(); m_loaded.clear(); + m_generation.fetch_add(1, std::memory_order_release); } private: using Entries = std::multimap; - static bool addTo(Entries& entries, const ProblemOverride& key, const TunedEntry& entry) + bool addTo(Entries& entries, const ProblemOverride& key, const TunedEntry& entry) { auto range = entries.equal_range(key); for(auto it = range.first; it != range.second; ++it) @@ -493,6 +506,7 @@ namespace TensileLite } entries.emplace(key, entry); + m_generation.fetch_add(1, std::memory_order_release); return true; } @@ -501,6 +515,7 @@ namespace TensileLite std::set m_loaded; std::mutex m_guard; mutable std::shared_timed_mutex m_mutex; + std::atomic m_generation{0}; }; struct TuningRowsLoaded diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp index 226ae3d059b3..c62e1c83fcc9 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/UserDrivenTuningParser.hpp @@ -183,6 +183,18 @@ namespace TensileLite */ void getContractionProblemsFromFile(const std::string& path); + /** + * getContractionProblemsFromFile for cache replay, which runs on every + * matmul that passes no algo. A file that is not there is looked for again + * at most once a second rather than on every call, since each look is a + * failed open and a stat. The heuristic entry point still looks on every + * query. + */ + void loadTuningFileForReplay(const std::string& path); + + /** Let the next loadTuningFileForReplay look at once. Tests only. */ + void resetReplayLoadForTest(); + /** How the cache file was read, for the startup line. */ enum class TuningLoadStatus : uint32_t { diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp index 5fd26ab3e251..28711ec4250e 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/tensile_host.hpp @@ -192,6 +192,19 @@ rocblaslt_status isSolutionSupportedNoMutation(rocblaslt_handle rocblaslt_matmul_algo* algo, size_t* workspaceSizeInBytes); +/** + * isSolutionSupported, leaving gemmData's Tensile problem as + * updateTensileProblem(prob) makes it rather than as the check leaves it. It + * copies no problem, as isSolutionSupportedNoMutation does at an allocation per + * tensor, so it suits a caller that goes on to launch prob, as cache replay + * does. + */ +rocblaslt_status isSolutionSupportedInPlace(rocblaslt_handle handle, + const RocblasltContractionProblem& prob, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo, + size_t* workspaceSizeInBytes); + template rocblaslt_status isSolutionSupported(rocblaslt_handle handle, const rocblaslt::RocGemmType& gemmType, @@ -237,17 +250,20 @@ TensileLite::ProblemOverride TensileLite::ProblemOverride TensileDataGemm2ProblemOverride(std::shared_ptr); /** - * The solution index of the first entry for this key that still resolves to its - * recorded kernel and supports this problem, or -1 when there is none. + * Cache replay for a matmul that passed no algo: the solution index of the first + * entry for this key that still resolves to its recorded kernel and supports + * this problem within its workspace, or -1 when there is none. Fills *algo to + * launch that entry, and counts the lookup. * - * A pure probe: the caller counts the lookup, because only it knows whether the - * probe decided which kernel the call launches. + * Which entries still resolve to their recorded kernels is worked out once per + * key and kept until the loaded entries change; support depends on the call and + * is checked every time. */ -int tuning_cache_find_valid_entry(rocblaslt_handle handle, - const TensileLite::ProblemOverride& key, - const RocblasltContractionProblem& problem, - std::shared_ptr gemmData, - size_t max_workspace_bytes); +int tuning_cache_replay(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key, + const RocblasltContractionProblem& problem, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo); /** * The solution index this thread last launched through runContractionProblem, diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp index 5934817240de..ede21b97fe32 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_auxiliary.cpp @@ -53,9 +53,13 @@ #include +#include #include #include +#include +#include #include +#include #include #define TO_STR2(x) #x @@ -499,55 +503,147 @@ bool problem_override_from_file_cpp( return success; } -int tuning_cache_find_valid_entry(rocblaslt_handle handle, - const TensileLite::ProblemOverride& key, - const RocblasltContractionProblem& problem, - std::shared_ptr gemmData, - size_t max_workspace_bytes) +namespace { - std::vector resolved; - std::vector index(1); - - for(const auto& entry : tuned_entries_for(key)) - { - index[0] = entry.solutionIndex; - - // getSolutionsFromIndex appends, and everything below reads [0]. - resolved.clear(); - - if(rocblaslt_status_success - != getSolutionsFromIndex(handle, index, resolved, max_workspace_bytes) - || resolved.empty()) - continue; - if(!tuned_entry_identity_matches(handle, key, entry, resolved[0])) - continue; - - // The key leaves out values such as beta and C/D aliasing, so an entry - // whose identity holds can still fail this call's support predicates. - RocblasltContractionProblem supportProblem = problem; - size_t required = 0; - auto algo = resolved[0].algo; - if(rocblaslt_status_success - == isSolutionSupportedNoMutation(handle, supportProblem, gemmData, &algo, &required) - && required <= max_workspace_bytes) - return entry.solutionIndex; - - // The heuristic replay's XF32 fallback, as a pure probe: the call may go - // on to launch some other algorithm, which must not find gemmData left - // in FP32 mode. - if(problem.compute_type == rocblaslt_compute_f32_fast_xf32) + /** + * A key's entries that still name the kernel their index resolves to, in + * the order lookups try them. Neither resolving an index nor comparing + * names depends on the call, so replay does both once per key, again only + * after the loaded entries change. Support is still checked on every call: + * it also depends on values the key leaves out, such as beta and C/D + * aliasing. + */ + struct ReplayCandidates + { + uint64_t generation = 0; + std::vector indexes; + + // What the lookup tally has been told about this key: 0 nothing, 1 + // that it fell back, 2 that it matched. Replay asks on every call; the + // tally only needs to hear when the answer improves. + std::atomic tallied{0}; + }; + + class ReplayCandidatesByKey + { + public: + static ReplayCandidatesByKey& instance() + { + static ReplayCandidatesByKey gInstance; + return gInstance; + } + + std::shared_ptr find(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key) + { + // Read before the entries, so entries added meanwhile leave this + // result stale rather than missing them. + const uint64_t generation = TensileLite::OverrideMap::getMap().generation(); + { + std::shared_lock lock(m_mutex); + auto found = m_byKey.find(key); + if(found != m_byKey.end() && found->second->generation == generation) + return found->second; + } + + auto fresh = std::make_shared(); + fresh->generation = generation; + + std::vector resolved; + std::vector index(1); + for(const auto& entry : tuned_entries_for(key)) + { + index[0] = entry.solutionIndex; + + // getSolutionsFromIndex appends, and the check below reads [0]. + resolved.clear(); + if(rocblaslt_status_success == getSolutionsFromIndex(handle, index, resolved, 0) + && !resolved.empty() + && tuned_entry_identity_matches(handle, key, entry, resolved[0])) + fresh->indexes.push_back(entry.solutionIndex); + } + + std::unique_lock lock(m_mutex); + m_byKey[key] = fresh; + return fresh; + } + + void clear() { - supportProblem.compute_type = rocblaslt_compute_f32; - required = 0; + std::unique_lock lock(m_mutex); + m_byKey.clear(); + } + + private: + std::shared_mutex m_mutex; + std::unordered_map> m_byKey; + }; + + /** The first candidate that supports this call, with the algo to launch it, or -1. */ + int first_supported_candidate(rocblaslt_handle handle, + const ReplayCandidates& candidates, + const RocblasltContractionProblem& problem, + std::shared_ptr gemmData, + size_t max_workspace_bytes, + rocblaslt_matmul_algo* algo) + { + for(const int index : candidates.indexes) + { + // What getSolutionsFromIndex returns for this index. + *algo = rocblaslt_matmul_algo{}; + *reinterpret_cast(algo->data) = index; + algo->max_workspace_bytes = max_workspace_bytes; + + size_t required = 0; if(rocblaslt_status_success - == isSolutionSupportedNoMutation( - handle, supportProblem, gemmData, &algo, &required) + == isSolutionSupportedInPlace(handle, problem, gemmData, algo, &required) && required <= max_workspace_bytes) - return entry.solutionIndex; + return index; + + // The heuristic replay's XF32 fallback, as a pure probe: the call + // may go on to launch some other algorithm, which must not find + // gemmData left in FP32 mode. + if(problem.compute_type == rocblaslt_compute_f32_fast_xf32) + { + RocblasltContractionProblem supportProblem = problem; + supportProblem.compute_type = rocblaslt_compute_f32; + required = 0; + if(rocblaslt_status_success + == isSolutionSupportedNoMutation( + handle, supportProblem, gemmData, algo, &required) + && required <= max_workspace_bytes) + return index; + } } + return -1; + } +} // namespace + +int tuning_cache_replay(rocblaslt_handle handle, + const TensileLite::ProblemOverride& key, + const RocblasltContractionProblem& problem, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo) +{ + const auto candidates = ReplayCandidatesByKey::instance().find(handle, key); + const int index = first_supported_candidate( + handle, *candidates, problem, gemmData, problem.workspaceSize, algo); + + auto& counters = TensileLite::TuningCounters::instance(); + if(index >= 0) + counters.hits++; + else + counters.misses++; + + const int tally = index >= 0 ? 2 : 1; + int told = candidates->tallied.load(std::memory_order_relaxed); + while(told < tally && !candidates->tallied.compare_exchange_weak(told, tally)) + { } + if(told < tally) + TensileLite::recordTuningLookup(key, index >= 0); - return -1; + return index; } /****************************************************************************** @@ -3104,6 +3200,8 @@ extern "C" HIPBLASLT_EXPORT void hipblaslt_tuning_reset_for_test() TensileLite::OverrideMap::getMap().resetForTest(); OverrideSingleton::getInstance().reloadForTest(); TensileLite::resetTuningDiagnosticsForTest(); + TensileLite::resetReplayLoadForTest(); + ReplayCandidatesByKey::instance().clear(); auto& counters = TensileLite::TuningCounters::instance(); counters.entriesLoaded = 0; diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index edb8d380c26a..d5522fc25c7f 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -3698,10 +3698,41 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle if(prob.trans_b == HIPBLAS_OP_C) data->problem.setBOps({TensileLite::TensorOp::ComplexConjugate()}); - // Noted before algo is filled in below. An explicit algo launches as - // given, whatever the cache holds; the cache chooses the kernel only - // when the caller leaves that choice to the library. - const bool callerSuppliedAlgo = (algo != nullptr); + if(algo == nullptr) + { + // Cache replay for a call that passed no algo. hipblasLtMatmul + // accepts algo == nullptr and then runs getBestSolutions below + // without ever entering hipblasLtMatmulAlgoGetHeuristic, so without + // this such a call would never be served from the cache. An + // explicit algo launches as given, whatever the cache holds. + try + { + const auto& tuning = TensileLite::TuningModeSingleton::getInstance(); + + if(tuning.reads() && !prob.grouped_gemm + && prob.batchMode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY) + { + // A matmul-only caller never reaches the heuristic entry + // point, which is the only other place that loads the file. + TensileLite::loadTuningFileForReplay(tuning.cachePath()); + + if(tuning_cache_replay(handle, + RocblasltContractionProblem2ProblemOverride(prob), + prob, + gemmData, + &heuristicResult.algo) + >= 0) + algo = &heuristicResult.algo; + } + } + catch(...) + { + // The cache is an optimisation: a lookup that throws leaves the + // call on default selection rather than failing it. + static_cast(hipGetLastError()); + algo = nullptr; + } + } if(algo == nullptr) { @@ -3723,52 +3754,6 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle int* solutionIndex = (int*)algo->data; - // Cache replay for a call that passed no algo. hipblasLtMatmul accepts - // algo == nullptr and then runs getBestSolutions above without ever - // entering hipblasLtMatmulAlgoGetHeuristic, so without this such a call - // would never be served from the cache. Only the index used for this - // launch changes; the caller's algo is const. - int launchIndex = -1; - try - { - const auto& tuning = TensileLite::TuningModeSingleton::getInstance(); - - if(tuning.reads() && !callerSuppliedAlgo && !prob.grouped_gemm - && prob.batchMode != HIPBLASLT_BATCH_MODE_POINTER_ARRAY) - { - // A matmul-only caller never reaches the heuristic entry point, - // which is the only other place that loads the file. - TensileLite::getContractionProblemsFromFile(tuning.cachePath()); - - const TensileLite::ProblemOverride key - = RocblasltContractionProblem2ProblemOverride(prob); - const int cachedIndex = tuning_cache_find_valid_entry( - handle, key, prob, gemmData, prob.workspaceSize); - - TensileLite::recordTuningLookup(key, cachedIndex >= 0); - auto& counters = TensileLite::TuningCounters::instance(); - if(cachedIndex >= 0) - { - counters.hits++; - launchIndex = cachedIndex; - } - else - { - counters.misses++; - } - } - } - catch(...) - { - // The cache is an optimisation: a lookup that throws leaves the call - // on default selection rather than failing it. - static_cast(hipGetLastError()); - launchIndex = -1; - } - - if(launchIndex >= 0) - solutionIndex = &launchIndex; - data->algoIndex = *solutionIndex; data->inputs = GetTensileInputs(prob); @@ -3812,10 +3797,8 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle if(get_logger_layer_mode() & rocblaslt_layer_mode_log_extended_profile) { - auto loggedAlgo = *algo; - *(int*)loggedAlgo.data = *solutionIndex; - std::string kernel_name = getKernelNameFromAlgoIndex(handle, loggedAlgo); - std::string Solution_name = getSolutionNameFromAlgoIndex(handle, loggedAlgo); + std::string kernel_name = getKernelNameFromAlgoIndex(handle, *algo); + std::string Solution_name = getSolutionNameFromAlgoIndex(handle, *algo); logExtendedProfileFromTensileDataGemm(data->problem, data->inputs, @@ -5574,6 +5557,44 @@ rocblaslt_status isSolutionSupportedNoMutation(rocblaslt_handle } } +rocblaslt_status isSolutionSupportedInPlace(rocblaslt_handle handle, + const RocblasltContractionProblem& prob, + std::shared_ptr gemmData, + rocblaslt_matmul_algo* algo, + size_t* workspaceSizeInBytes) +{ +#ifdef HIPBLASLT_USE_ROCROLLER + if(useRocRoller(handle, prob)) + return isSolutionSupportedNoMutation(handle, prob, gemmData, algo, workspaceSizeInBytes); +#endif + + std::shared_ptr data = std::static_pointer_cast(gemmData); + updateTensileProblem(prob, data->problem); + + // All that the check below sets on the problem. + const auto params = data->problem.getParams(); + const auto workspace = data->problem.workspaceSize(); + auto restore = [&] { + data->problem.setParams() = params; + data->problem.setWorkspaceSize(workspace); + }; + + RocblasltContractionProblem inputs = prob; + rocblaslt::RocTuningV2* tuning = nullptr; + try + { + const rocblaslt_status status = isSolutionSupported( + handle, data->problem, inputs, algo, tuning, workspaceSizeInBytes); + restore(); + return status; + } + catch(...) + { + restore(); + throw; + } +} + template void setRestrictions(TensileLite::ContractionProblemGemm& tensile_prob, const T* alpha, From 7bf99f94b104e36ab4eda22df1e66f36d7689881 Mon Sep 17 00:00:00 2001 From: jinchen62 Date: Mon, 12 Oct 2026 01:10:50 -0500 Subject: [PATCH 5/5] test(hipblaslt): check tuning file writes in the cache-mode tests #13354 made the tuning file writers return a GoogleTest assertion result and checked it in every existing test. The cache-mode tests added here still ignored it, so a write that failed surfaced later as the wrong kernel. They now stop at the write, as the others do. Co-authored-by: Cursor --- .../clients/tests/src/tuning_cache_test.cpp | 38 +++++++++++-------- 1 file changed, 23 insertions(+), 15 deletions(-) diff --git a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp index 36589462cd4b..edcf645e8971 100644 --- a/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp +++ b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp @@ -1021,7 +1021,7 @@ namespace GTEST_SKIP() << "the heuristic offers one solution for this problem"; const auto& recorded = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}})); useCache(m_path); int selected = -1; @@ -1051,7 +1051,7 @@ namespace EXPECT_EQ(selected, m_identities[0].index); const auto& recorded = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}})); ASSERT_TRUE(runGemm(&selected)); EXPECT_EQ(selected, recorded.index); @@ -1065,7 +1065,8 @@ namespace if(!haveSolutions(2)) GTEST_SKIP() << "the heuristic offers one solution for this problem"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); useCache(""); @@ -1082,7 +1083,8 @@ namespace if(!haveSolutions(2)) GTEST_SKIP() << "the heuristic offers one solution for this problem"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); setenv("HIPBLASLT_TUNING_OVERRIDE_FILE", m_path.c_str(), 1); useCache(m_path + ".absent"); @@ -1097,7 +1099,8 @@ namespace if(!haveSolutions(2)) GTEST_SKIP() << "the heuristic offers one solution for this problem"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}})); useCache(m_path); int selected = -1; @@ -1117,7 +1120,8 @@ namespace if(!haveSolutions(2)) GTEST_SKIP() << "the heuristic offers one solution for this problem"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, std::string("NotARealKernel")}})); useCache(m_path); int selected = -1; @@ -1135,7 +1139,7 @@ namespace const auto& given = m_identities[0]; const auto& cached = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}})); useCache(m_path); int launched = -1; @@ -1152,7 +1156,7 @@ namespace GTEST_SKIP() << "the heuristic offers one solution for this problem"; const auto& cached = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{cached.index, cached.kernelName}})); useCache(m_path); int launched = -1; @@ -1193,7 +1197,7 @@ namespace ASSERT_TRUE(extHeuristicIndex(&fp32, HIP_R_32F, HIPBLAS_COMPUTE_32F)); const auto fp32Name = kernelNameOf(fp32); ASSERT_FALSE(fp32Name.empty()); - writeTuningFile(m_path, m_stamp, {{fp32, fp32Name}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{fp32, fp32Name}})); useCache(m_path); for(int call = 0; call < 2; call++) @@ -1225,7 +1229,7 @@ namespace // The heuristic looks for the file on every query, so this one loads it. const auto& recorded = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}})); int selected = -1; ASSERT_TRUE(runGemm(&selected)); ASSERT_EQ(selected, recorded.index); @@ -1248,7 +1252,7 @@ namespace EXPECT_EQ(launched, m_identities[0].index); const auto& recorded = m_identities[1]; - writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}}); + ASSERT_TRUE(writeTuningFile(m_path, m_stamp, {{recorded.index, recorded.kernelName}})); std::this_thread::sleep_for(std::chrono::milliseconds(1100)); ASSERT_TRUE(launchGemm(AlgoFrom::Null, -1, &launched)); @@ -1278,7 +1282,8 @@ namespace if(!haveSolutions(2)) GTEST_SKIP() << "the heuristic offers one solution for this problem"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); useCache(m_path); for(int i = 0; i < 4; i++) @@ -1302,7 +1307,8 @@ namespace if(!loggingEnvIsClean()) GTEST_SKIP() << "external logging environment redirects tuning output"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); useCache(m_path); std::string loaded; @@ -1339,7 +1345,8 @@ namespace if(!loggingEnvIsClean()) GTEST_SKIP() << "external logging environment redirects tuning output"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); useCache(m_path); { @@ -1369,7 +1376,8 @@ namespace if(!loggingEnvIsClean()) GTEST_SKIP() << "external logging environment redirects tuning output"; - writeTuningFile(m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}}); + ASSERT_TRUE(writeTuningFile( + m_path, m_stamp, {{m_identities[1].index, m_identities[1].kernelName}})); useTuningFile(); std::string out;