diff --git a/projects/hipblaslt/CHANGELOG.md b/projects/hipblaslt/CHANGELOG.md index d59440bf5ffc..0f62c7739d14 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 @@ -90,6 +91,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/clients/tests/src/tuning_cache_test.cpp b/projects/hipblaslt/clients/tests/src/tuning_cache_test.cpp index 2bd1c94b7d3b..edcf645e8971 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 @@ -17,13 +18,18 @@ #include #include +#include #include #include #include +#include #include +#include #include #include +#include #include +#include #include #include @@ -33,7 +39,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) @@ -174,6 +182,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() { @@ -388,6 +414,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; @@ -453,15 +685,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 @@ -475,20 +712,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 @@ -499,10 +750,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 @@ -758,4 +1013,381 @@ 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]; + ASSERT_TRUE(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]; + ASSERT_TRUE(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"; + + 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(""); + + 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"; + + 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"); + + 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"; + + ASSERT_TRUE(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"; + + ASSERT_TRUE(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]; + ASSERT_TRUE(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]; + ASSERT_TRUE(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); + } + + // 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()); + ASSERT_TRUE(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]; + ASSERT_TRUE(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]; + 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)); + 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) + { + 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"; + + 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++) + { + 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"; + + ASSERT_TRUE(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"; + + ASSERT_TRUE(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"; + + ASSERT_TRUE(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..e7d7ad12559f 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,17 @@ * ************************************************************************ */ #include "UserDrivenTuningParser.hpp" +#include "rocblaslt_secure_env.hpp" +#include "utility.hpp" +#include +#include #include #include +#include +#include +#include +#include #ifndef TO_STR2 #define TO_STR2(x) #x @@ -47,8 +55,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 +329,79 @@ 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()) + { + // 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; + + 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) + status = TuningLoadStatus::ReadError; + else + map.markLoaded(path); + } + + 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 loaded = loadTuningRows(file, map, currentBuildStamp()); + 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); - 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); + 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 770307f97850..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 @@ -27,6 +28,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. * @@ -417,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 @@ -446,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) @@ -462,6 +506,7 @@ namespace TensileLite } entries.emplace(key, entry); + m_generation.fetch_add(1, std::memory_order_release); return true; } @@ -470,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 a2d0fbea12d3..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 @@ -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,68 @@ namespace TensileLite * once per process. */ 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 + { + 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..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 @@ -181,6 +181,30 @@ 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); + +/** + * 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, @@ -225,6 +249,28 @@ TensileLite::ProblemOverride TensileLite::ProblemOverride TensileDataGemm2ProblemOverride(std::shared_ptr); +/** + * 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. + * + * 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_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, + * 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 94cd6b5fa0a7..1de91ca95f36 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 @@ -176,6 +180,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 +195,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 +222,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 +312,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 +344,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 +390,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 +408,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 +418,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 +446,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 +497,153 @@ bool problem_override_from_file_cpp( } } - if(!success) + record_lookup(__func__, prob_key, success, solutionIndex[0]); + } + + return success; +} + +namespace +{ + /** + * 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() { - log_info(__func__, "No valid solution index found in override file."); + static ReplayCandidatesByKey gInstance; + return gInstance; } - else + + 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() + { + 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) { - std::string mapping_result = "Find solution with index: "; - mapping_result += std::to_string(solutionIndex[0]); - log_info(__func__, mapping_result); + // 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 + == isSolutionSupportedInPlace(handle, problem, gemmData, algo, &required) + && required <= max_workspace_bytes) + 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 - return success; +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 index; } /****************************************************************************** @@ -2391,20 +2593,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--; @@ -2673,14 +2875,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, workspaceLimit); + handle, gemmType, gemmData, override_result, tuningFile.path, workspaceLimit); log_api(__func__, "OverrideAlgoCount", override_success ? 1 : 0); } @@ -2983,11 +3185,62 @@ 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(); + TensileLite::resetReplayLoadForTest(); + ReplayCandidatesByKey::instance().clear(); + + 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 3fe9def63b6b..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 @@ -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,42 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle if(prob.trans_b == HIPBLAS_OP_C) data->problem.setBOps({TensileLite::TensorOp::ComplexConjugate()}); + 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) { int returnAlgoCount = 0; @@ -3705,6 +3753,7 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle int32_t coldIterations = ClientArguments.GetColdIterationsValue(); int* solutionIndex = (int*)algo->data; + data->algoIndex = *solutionIndex; data->inputs = GetTensileInputs(prob); @@ -3875,6 +3924,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()) @@ -3891,10 +3941,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; @@ -3920,7 +3974,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) { @@ -5471,6 +5525,76 @@ 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; + } +} + +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,