From 4c697cb29ba13f4a55c59563e94fbac20d43d1eb Mon Sep 17 00:00:00 2001 From: Caio Salvador Rohwedder Date: Sat, 17 Oct 2026 00:17:14 -0500 Subject: [PATCH 1/3] Adapt launches to the CUs a probe measures With HIPBLASLT_ADAPTIVE_SM_COUNT set, a co-tenant that takes whole CUs otherwise leaves SK5 grids sized for the full device. Every few C API launches on a stream now carry the occupancy probe; its reading feeds the stream's estimator, whose hint becomes the implicit sm_count_target for the heuristic query and for launches. The top query result reports enough workspace for a regrid under a later hint. A per-handle controller owns the probe buffer, the per-stream estimators and the published hints. Each launch holds its adaptive state in an RAII object whose end records the probe's completion event, and TensileLite marks the kernel invocations that carry the probe. The probe events are created with the controller; if creating one fails, the handle runs without adaptive mode and logs why. --- .../clients/tests/src/CMakeLists.txt | 1 + .../tests/src/adaptive_sm_count_gtest.cpp | 485 ++++++++++++++++++ .../rocblaslt/include/rocblaslt-types.h | 14 + .../amd_detail/rocblaslt/src/CMakeLists.txt | 1 + .../src/amd_detail/rocblaslt/src/handle.cpp | 3 + .../amd_detail/rocblaslt/src/include/handle.h | 41 +- .../rocblaslt_adaptive_sm_controller.hpp | 140 +++++ .../include/rocblaslt_adaptive_sm_count.hpp | 4 + .../src/rocblaslt_adaptive_sm_controller.cpp | 274 ++++++++++ .../rocblaslt/src/rocblaslt_auxiliary.cpp | 55 ++ .../rocblaslt/src/rocblaslt_mat.cpp | 16 + .../amd_detail/rocblaslt/src/tensile_host.cpp | 21 +- .../tensilelite/include/Tensile/Tensile.hpp | 3 + .../tensilelite/src/ContractionSolution.cpp | 1 + .../tests/PersistentArgumentLayout_test.cpp | 3 + 15 files changed, 1055 insertions(+), 7 deletions(-) create mode 100644 projects/hipblaslt/clients/tests/src/adaptive_sm_count_gtest.cpp create mode 100644 projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_controller.hpp create mode 100644 projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp diff --git a/projects/hipblaslt/clients/tests/src/CMakeLists.txt b/projects/hipblaslt/clients/tests/src/CMakeLists.txt index b537a6a15b73..faf594880a27 100755 --- a/projects/hipblaslt/clients/tests/src/CMakeLists.txt +++ b/projects/hipblaslt/clients/tests/src/CMakeLists.txt @@ -10,6 +10,7 @@ target_sources(hipblaslt-test ${CMAKE_CURRENT_SOURCE_DIR}/hip_solution_adapter_gtest.cpp ${CMAKE_CURRENT_SOURCE_DIR}/secure_env_gtest.cpp ${CMAKE_CURRENT_SOURCE_DIR}/arch_revision_gtest.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/adaptive_sm_count_gtest.cpp ${CMAKE_CURRENT_SOURCE_DIR}/efficiency_monitor_gtest.cpp ${CMAKE_CURRENT_SOURCE_DIR}/matmul_gtest.cpp ${CMAKE_CURRENT_SOURCE_DIR}/math_gtest.cpp diff --git a/projects/hipblaslt/clients/tests/src/adaptive_sm_count_gtest.cpp b/projects/hipblaslt/clients/tests/src/adaptive_sm_count_gtest.cpp new file mode 100644 index 000000000000..e7bb4ea90514 --- /dev/null +++ b/projects/hipblaslt/clients/tests/src/adaptive_sm_count_gtest.cpp @@ -0,0 +1,485 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Adaptive sm_count_target through the C API: the heuristic-query workspace, +// launches under a forced hint, probe readings, graph capture and the stream +// lock. The mode is read from the environment once per process, so +// the cases for a mode the environment does not select skip; run with +// HIPBLASLT_ADAPTIVE_SM_COUNT=1 or HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE= (and +// HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE=1) to cover them. The cotenant cases +// are Linux only; the exact one is manual (HIPBLASLT_TEST_COTENANT_PROBE=1, +// idle GPU). Host-only cases are in +// clients/tests/unit/adaptive_sm_count_gtest.cpp. + +#include +#include +#include +#include + +#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_controller.hpp" +#include "streamk_test_util.hpp" +#ifndef _WIN32 +#include "cotenant.hpp" +#endif + +#include +#include +#include +#include +#include +#include +#include +#include + +extern "C" bool hipblaslt_adaptive_sm_state_for_test(void* handle, + hipStream_t stream, + void* state, + size_t size); +extern "C" bool hipblaslt_adaptive_sm_launch_lock_for_test( + void* handle, hipStream_t stream, bool tiny, bool* probe, bool* heldBetween, bool* heldAfter); + +namespace rocblaslt::adaptive_sm +{ + inline void PrintTo(const StreamStateForTest& s, std::ostream* os) + { + *os << "{published=" << s.published << " epoch=" << s.epoch << " pending=" << s.pending + << " missed=" << s.missed << " sinceTick=" << s.sinceTick << " due=" << s.due + << " probesRead=" << s.probesRead << " lastCount=" << s.lastCount << "}"; + } +} // namespace rocblaslt::adaptive_sm + +namespace +{ + using namespace rocblaslt::adaptive_sm; + using namespace streamk_test; + + // BF16 NN, large enough to probe (not tiny) and to rank Stream-K solutions. + constexpr int64_t kM = 512, kN = 8192, kK = 8192; + + class AdaptiveSmCount_pre_checkin : public ::testing::Test + { + protected: + int64_t m = kM, n = kN, k = kK; + Resources r; + hipStream_t stream = nullptr; + const Config cfg = Config::fromEnv(); + float alpha = 1.0f, beta = 0.0f; + + void SetUp() override + { + if(!gpuAvailable()) + GTEST_SKIP() << "No GPU available"; + createProblem(r, m, n, k); + if(IsSkipped() || HasFatalFailure()) + return; + setStreamK(HIPBLASLT_STREAMK_TILE_SCHEDULING_AUTO); + ASSERT_EQ(hipMalloc(&r.dA, m * k * sizeof(uint16_t)), hipSuccess); + ASSERT_EQ(hipMalloc(&r.dB, k * n * sizeof(uint16_t)), hipSuccess); + r.dD.assign(1, nullptr); + ASSERT_EQ(hipMalloc(&r.dD[0], m * n * sizeof(uint16_t)), hipSuccess); + r.dWs.assign(1, nullptr); + ASSERT_EQ(hipMalloc(&r.dWs[0], kWsBudgetBytes), hipSuccess); + r.streams.assign(1, nullptr); + ASSERT_EQ(hipStreamCreate(&r.streams[0]), hipSuccess); + stream = r.streams[0]; + } + + void TearDown() override + { + if(stream) + static_cast(hipStreamSynchronize(stream)); + } + + void setStreamK(int32_t sk) + { + ASSERT_EQ( + hipblasLtMatmulDescSetAttribute( + r.desc, HIPBLASLT_MATMUL_DESC_STREAMK_TILE_SCHEDULING_EXT, &sk, sizeof(sk)), + HIPBLAS_STATUS_SUCCESS); + } + + // Up to `count` heuristic results under workspace limit `limit`. + std::vector query(int count = 1, + uint64_t limit = kWsBudgetBytes) + { + std::vector out(count); + hipblasLtMatmulPreference_t pref = nullptr; + if(hipblasLtMatmulPreferenceCreate(&pref) != HIPBLAS_STATUS_SUCCESS) + return {}; + int returned = 0; + if(hipblasLtMatmulPreferenceSetAttribute( + pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &limit, sizeof(limit)) + != HIPBLAS_STATUS_SUCCESS + || hipblasLtMatmulAlgoGetHeuristic(r.handle, + r.desc, + r.layA, + r.layB, + r.layD, + r.layD, + pref, + count, + out.data(), + &returned) + != HIPBLAS_STATUS_SUCCESS) + returned = 0; + hipblasLtMatmulPreferenceDestroy(pref); + out.resize(returned); + return out; + } + + // Workspace `algo` needs under the desc's current settings. + size_t required(hipblasLtMatmulAlgo_t algo) + { + size_t bytes = 0; + EXPECT_EQ( + hipblaslt_ext::matmulIsAlgoSupported( + r.handle, r.desc, &alpha, r.layA, r.layB, &beta, r.layD, r.layD, algo, bytes), + HIPBLAS_STATUS_SUCCESS); + return bytes; + } + + hipblasStatus_t run(const hipblasLtMatmulAlgo_t& algo, size_t workspace) + { + return hipblasLtMatmul(r.handle, + r.desc, + &alpha, + r.dA, + r.layA, + r.dB, + r.layB, + &beta, + r.dD[0], + r.layD, + r.dD[0], + r.layD, + &algo, + r.dWs[0], + workspace, + stream); + } + + StreamStateForTest state() + { + StreamStateForTest st{}; + EXPECT_TRUE(hipblaslt_adaptive_sm_state_for_test(r.handle, stream, &st, sizeof(st))); + return st; + } + + // Under FORCE every launch takes the forced hint. Probes run only with + // FORCE_PROBE, and their readings never move the hint. + void checkForcedHint(int32_t sk) + { + if(!cfg.force) + GTEST_SKIP() << "set HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE"; + setStreamK(sk); + const auto results = query(); + if(results.empty()) + GTEST_SKIP() << "No solution for this problem in the loaded library"; + const auto& result = results[0]; + + for(uint32_t i = 0; i < cfg.probeEvery + 1; ++i) + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + const StreamStateForTest st = state(); + EXPECT_EQ(st.published, 0u) << "nothing is published under FORCE"; + if(cfg.forceProbe) + EXPECT_GT(st.epoch, 0u); + else + EXPECT_EQ(st.epoch, 0u); + } + }; + + // Adaptive mode reports min(limit, max(plain, c_adaptiveWorkspace)) for the + // top result; with it off the result reports the plain size. + TEST_F(AdaptiveSmCount_pre_checkin, QueryReportsTheAdaptiveWorkspace) + { + for(uint64_t limit : {kWsBudgetBytes / 4, kWsBudgetBytes * 2}) + { + SCOPED_TRACE("limit " + std::to_string(limit)); + const auto results = query(1, limit); + if(results.empty()) + GTEST_SKIP() << "No solution for this problem in the loaded library"; + const auto& result = results[0]; + const size_t plain = required(result.algo); + ASSERT_LE(plain, limit); + const size_t expected + = cfg.enabled ? std::min(limit, std::max(plain, c_adaptiveWorkspace)) + : plain; + EXPECT_EQ(result.workspaceSize, expected); + } + } + + TEST_F(AdaptiveSmCount_pre_checkin, ForcedHintForQueryAndLaunchAuto) + { + checkForcedHint(HIPBLASLT_STREAMK_TILE_SCHEDULING_AUTO); + } + + TEST_F(AdaptiveSmCount_pre_checkin, ForcedHintForQueryAndLaunchOn) + { + checkForcedHint(HIPBLASLT_STREAMK_TILE_SCHEDULING_ON); + } + + // A probe recorded on the stream completes and is read by the next launch. + TEST_F(AdaptiveSmCount_pre_checkin, ProbeProducesAReading) + { + if(!cfg.force || !cfg.forceProbe) + GTEST_SKIP() << "needs FORCE and FORCE_PROBE, so the solution and grid stay fixed"; + const auto results = query(); + if(results.empty()) + GTEST_SKIP() << "No solution for this problem in the loaded library"; + const auto& result = results[0]; + for(uint32_t i = 0; i < cfg.probeEvery; ++i) + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + const StreamStateForTest before = state(); + if(!before.pending) + GTEST_SKIP() << "no probe recorded: the launched solution does not probe"; + + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + const StreamStateForTest after = state(); + EXPECT_GT(after.probesRead, before.probesRead); + int device = 0, cus = 0; + ASSERT_EQ(hipGetDevice(&device), hipSuccess); + ASSERT_EQ(hipDeviceGetAttribute(&cus, hipDeviceAttributeMultiprocessorCount, device), + hipSuccess); + EXPECT_GT(after.lastCount, 0u); + EXPECT_LE(after.lastCount, uint32_t(cus)); + EXPECT_EQ(after.published, 0u) << "the reading is discarded"; + } + +#ifndef _WIN32 + // The cotenant holds this fraction of the CUs. + constexpr int kHeldCuDivisor = 4; + // One cotenant workgroup per CU, reserving all of its LDS, so no GEMM + // or probe-only workgroup fits beside it. + constexpr int kWholeCuOccupancy = 1; + // Opt-in for the manual exact-count test. + constexpr const char* kEnvCotenantProbe = "HIPBLASLT_TEST_COTENANT_PROBE"; + // Readings per phase. + constexpr int kReadings = 3; + + // Probes beside a cotenant holding `held` whole CUs. Runs under FORCE with + // FORCE_PROBE, so the solution and grid stay fixed; with FORCE at most the + // free CU count, the probe-only padding workgroups reach the free CUs the + // StreamK grid does not. + class AdaptiveSmCountCotenant_pre_checkin : public AdaptiveSmCount_pre_checkin + { + protected: + hipblasLtMatmulHeuristicResult_t result{}; + int cus = 0; + int held = 0; + + // Skips when the mode, library or device does not fit. + void prepare() + { + if(!cfg.force || !cfg.forceProbe) + GTEST_SKIP() << "needs FORCE and FORCE_PROBE, so the solution and grid stay fixed"; + const auto results = query(); + if(results.empty()) + GTEST_SKIP() << "No solution for this problem in the loaded library"; + result = results[0]; + + int device = 0, ldsPerCu = 0, optin = 0; + ASSERT_EQ(hipGetDevice(&device), hipSuccess); + ASSERT_EQ(hipDeviceGetAttribute(&cus, hipDeviceAttributeMultiprocessorCount, device), + hipSuccess); + ASSERT_EQ(hipDeviceGetAttribute( + &ldsPerCu, hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, device), + hipSuccess); + ASSERT_EQ( + hipDeviceGetAttribute(&optin, hipDeviceAttributeSharedMemPerBlockOptin, device), + hipSuccess); + try + { + hipblaslt_cotenant::lds_reservation(ldsPerCu, optin, kWholeCuOccupancy); + } + catch(const std::runtime_error& e) + { + GTEST_SKIP() << "a workgroup cannot reserve a whole CU's LDS: " << e.what(); + } + held = cus / kHeldCuDivisor; + if(held == 0) + GTEST_SKIP() << "fewer than " << kHeldCuDivisor << " CUs"; + if(cfg.force > uint32_t(cus - held)) + GTEST_SKIP() << "FORCE must not exceed the " << cus - held + << " CUs the cotenant leaves free"; + } + + // Launches until a probe is pending, waits for it, then reads it + // with one more launch. Launch failures are fatal; `*probed` is false + // when no launch probes. + void nextReading(bool* probed, uint32_t* count) + { + *probed = false; + const uint32_t read = state().probesRead; + for(uint32_t i = 0; i <= cfg.probeEvery && !state().pending; ++i) + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + if(!state().pending) + return; + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + const StreamStateForTest st = state(); + ASSERT_EQ(st.probesRead, read + 1); + *probed = true; + *count = st.lastCount; + } + + // Starts the cotenant and drops any probe launched before it was + // resident. Skips if other work keeps its CUs from becoming free. + void startCotenant(std::optional* cotenant) + { + try + { + cotenant->emplace(held, kWholeCuOccupancy, stream); + } + catch(const std::runtime_error& e) + { + GTEST_SKIP() << "cotenant did not become resident: " << e.what(); + } + if(state().pending) + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + } + + // Takes kReadings readings beside the cotenant, checking each with + // `check`. + template + void readBesideCotenant(Check check) + { + // Destroyed on every exit, so the cotenant stops on failure too. + std::optional cotenant; + ASSERT_NO_FATAL_FAILURE(startCotenant(&cotenant)); + if(IsSkipped()) + return; + for(int i = 0; i < kReadings; ++i) + { + SCOPED_TRACE("reading " + std::to_string(i)); + bool probed = false; + uint32_t count = 0; + ASSERT_NO_FATAL_FAILURE(nextReading(&probed, &count)); + if(!probed && i == 0) + GTEST_SKIP() << "no probe reading: the launched solution does not probe"; + ASSERT_TRUE(probed) << "no probe reading under the cotenant"; + check(count); + } + } + }; + + // Other work can only take CUs, so the reading never exceeds what the + // cotenant leaves free. + TEST_F(AdaptiveSmCountCotenant_pre_checkin, ProbeCountsAtMostTheCusACotenantLeaves) + { + ASSERT_NO_FATAL_FAILURE(prepare()); + if(IsSkipped()) + return; + readBesideCotenant([&](uint32_t count) { + EXPECT_GT(count, 0u); + EXPECT_LE(count, uint32_t(cus - held)); + }); + } + + // Manual: other processes that take CUs change the count, so this runs + // only when HIPBLASLT_TEST_COTENANT_PROBE=1, on an otherwise idle GPU. + TEST_F(AdaptiveSmCountCotenant_pre_checkin, ProbeCountsExactlyTheCusACotenantLeaves) + { + const char* optIn = std::getenv(kEnvCotenantProbe); + if(!optIn || std::string(optIn) != "1") + GTEST_SKIP() << "manual: set " << kEnvCotenantProbe + << "=1 on an otherwise idle GPU; other processes that take CUs " + "change the count"; + ASSERT_NO_FATAL_FAILURE(prepare()); + if(IsSkipped()) + return; + bool probed = false; + uint32_t idle = 0; + ASSERT_NO_FATAL_FAILURE(nextReading(&probed, &idle)); + if(!probed) + GTEST_SKIP() << "no probe reading: the launched solution does not probe"; + if(idle != uint32_t(cus)) + GTEST_SKIP() << "idle reading " << idle << " of " << cus + << " CUs: other work occupies CUs"; + readBesideCotenant([&](uint32_t count) { EXPECT_EQ(count, idle - held); }); + } +#endif // _WIN32 + + // A launch on a capturing stream neither takes a probe nor counts toward + // the next tick: the stream's state after the capture is the state before. + TEST_F(AdaptiveSmCount_pre_checkin, CaptureLeavesTheStateAlone) + { + if(!cfg.enabled || !cfg.tracksLaunches()) + GTEST_SKIP() << "launches leave the state alone in this mode"; + if(cfg.probeEvery < 2) + GTEST_SKIP() << "needs a launch between ticks"; + const auto results = query(); + if(results.empty()) + GTEST_SKIP() << "No solution for this problem in the loaded library"; + const auto& result = results[0]; + + // A probe falls due, completes and is read, leaving the stream one + // launch past a tick. + for(uint32_t i = 0; i < cfg.probeEvery; ++i) + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + const bool probed = state().pending; + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + const StreamStateForTest before = state(); + ASSERT_NE(before.sinceTick, 0u); + if(probed) + ASSERT_GT(before.probesRead, 0u); + + ASSERT_EQ(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal), hipSuccess); + for(uint32_t i = 0; i < 3 * cfg.probeEvery; ++i) + EXPECT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + hipGraph_t graph = nullptr; + ASSERT_EQ(hipStreamEndCapture(stream, &graph), hipSuccess); + hipGraphExec_t exec = nullptr; + const hipError_t inst = hipGraphInstantiate(&exec, graph, nullptr, nullptr, 0); + if(inst == hipSuccess) + { + EXPECT_EQ(hipGraphLaunch(exec, stream), hipSuccess); + EXPECT_EQ(hipStreamSynchronize(stream), hipSuccess); + static_cast(hipGraphExecDestroy(exec)); + } + static_cast(hipGraphDestroy(graph)); + ASSERT_EQ(inst, hipSuccess); + + EXPECT_EQ(state(), before); + + // The next live launch carries on from there. + ASSERT_EQ(run(result.algo, result.workspaceSize), HIPBLAS_STATUS_SUCCESS); + EXPECT_EQ(state().sinceTick, (before.sinceTick + 1) % cfg.probeEvery); + } + + // A probe launch holds its stream's adaptive lock from its first phase to + // its second, so another thread cannot enqueue a later probe in between. + TEST_F(AdaptiveSmCount_pre_checkin, ProbeLaunchHoldsTheStreamLockBetweenItsPhases) + { + if(!cfg.enabled || !cfg.tracksLaunches()) + GTEST_SKIP() << "launches probe only in adaptive mode or under FORCE_PROBE"; + bool probe = false, heldBetween = false, heldAfter = false; + for(uint32_t i = 0; i < cfg.probeEvery && !probe; ++i) + ASSERT_TRUE(hipblaslt_adaptive_sm_launch_lock_for_test( + r.handle, stream, /*tiny=*/false, &probe, &heldBetween, &heldAfter)); + ASSERT_TRUE(probe); + EXPECT_TRUE(heldBetween); + EXPECT_FALSE(heldAfter); + } + + // A launch that does not probe holds the lock only inside its first phase. + TEST_F(AdaptiveSmCount_pre_checkin, NonProbeLaunchReleasesTheStreamLock) + { + if(!cfg.enabled || !cfg.tracksLaunches()) + GTEST_SKIP() << "launches leave the state alone in this mode"; + for(uint32_t i = 0; i < cfg.probeEvery; ++i) + { + SCOPED_TRACE("launch " + std::to_string(i)); + bool probe = true, heldBetween = true, heldAfter = true; + ASSERT_TRUE(hipblaslt_adaptive_sm_launch_lock_for_test( + r.handle, stream, /*tiny=*/true, &probe, &heldBetween, &heldAfter)); + EXPECT_FALSE(probe); + EXPECT_FALSE(heldBetween); + EXPECT_FALSE(heldAfter); + } + } +} // namespace diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h index 596b8beb4c9e..f73d6cec12fb 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/include/rocblaslt-types.h @@ -482,6 +482,14 @@ typedef struct _rocblaslt_matrix_transform_desc } #endif +namespace rocblaslt +{ + namespace adaptive_sm + { + class Launch; + } +} + /******************************************************************** * RocblasltContractionProblem captures the arguments for a GEMM-like * * contraction problem, to be passed to runContractionProblem. * @@ -632,6 +640,12 @@ struct RocblasltContractionProblem // 1 = on. Forwarded into ContractionProblemParameters::setUniformSummationOrder // by tensile_host.cpp. int32_t uniform_summation_order = 0; + // The C API launch's adaptive state (null: none), set post-construction. + // runContractionProblem forwards its probe target into + // ContractionProblemParameters::setOccupancyProbe and sets its `probed`. + rocblaslt::adaptive_sm::Launch* adaptive_launch = nullptr; + // sm_count_target came from the adaptive estimator rather than the user. + bool adaptive_sm_count = false; #if HIPBLASLT_HAS_GEMM_A2A_FUSION // The four below are set post-construction, not through the constructor. diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/CMakeLists.txt b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/CMakeLists.txt index ea914d2c986d..01fd45431f41 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/CMakeLists.txt +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/CMakeLists.txt @@ -7,6 +7,7 @@ target_sources(hipblaslt PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/Debug.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/handle.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/rocblaslt_adaptive_sm_controller.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/rocblaslt_auxiliary.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/rocblaslt_mat.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/status.cpp" diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/handle.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/handle.cpp index 5f15cab97b36..c45dccb08afb 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/handle.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/handle.cpp @@ -139,6 +139,9 @@ _rocblaslt_handle::_rocblaslt_handle() check_numerics_stop_on_first = (s == "1" || s == "on" || s == "true"); } + adaptive_sm = rocblaslt::adaptive_sm::Controller::create( + c_syncSkStreamSlots, uint32_t(properties.multiProcessorCount)); + // Allocate the device flag. STOP_ON_FIRST uses hipHostMalloc(MAPPED) so // scan_D can poll without sync; on any failure we fall back to hipMalloc // (the cross-call short-circuit then only trips on drains). If hipMalloc diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h index 8e71da706d40..32f5040c156b 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/handle.h @@ -28,12 +28,14 @@ #pragma once #include "rocblaslt.h" +#include "rocblaslt_adaptive_sm_controller.hpp" #include //#include "rocblaslt_ostream.hpp" #include #include #include #include +#include #include struct _rocblaslt_attribute @@ -212,6 +214,19 @@ struct _rocblaslt_handle if(problemIndex >= c_syncSkSlotsPerStream) return rocblaslt_status_internal_error; + const int i = streamSlot(stream); + if(i < 0) + return rocblaslt_status_internal_error; + + *out = static_cast(StreamKFlags) + + (i * c_syncSkSlotsPerStream + problemIndex) * c_syncSkSlotBytes; + return rocblaslt_status_success; + } + + // Index of the per-stream block owned by `stream`, claiming a free one on + // first use; -1 once all c_syncSkStreamSlots are taken. + int streamSlot(hipStream_t stream) + { // hipStreamPerThread is a sentinel that resolves to a different stream // for every host thread, so every thread would present the same key and // share one block. Key those by thread instead. The legacy null stream @@ -244,19 +259,19 @@ struct _rocblaslt_handle : expected; } if(owner == key) - { - *out = static_cast(StreamKFlags) - + (i * c_syncSkSlotsPerStream + problemIndex) * c_syncSkSlotBytes; - return rocblaslt_status_success; - } + return static_cast(i); } - return rocblaslt_status_internal_error; + return -1; } // Value-initialised so every block starts unowned. std::atomic m_skSlotOwner[c_syncSkStreamSlots] = {}; + // Adaptive sm_count_target state, per stream slot; null when the mode is + // off or its setup failed. + std::unique_ptr adaptive_sm; + #if HIPBLASLT_HAS_GEMM_A2A_FUSION // Device communicator, registered once by hipblasLtSetDeviceComm. Every rank // registers on its own handle; the W handles form one communicator, of which @@ -470,6 +485,20 @@ inline int32_t effective_sm_count_target(const _rocblaslt_handle* han return 0; } +// True when the hint for a matmul comes from the adaptive estimator: adaptive +// mode on, StreamK tile scheduling AUTO or ON, and no desc or handle hint. +inline bool adaptive_sm_count_applies(const _rocblaslt_handle* handle, + const _rocblaslt_matmul_desc* desc) +{ + return handle && desc + && rocblaslt::adaptive_sm::applies( + handle->adaptive_sm != nullptr, + desc->streamk_tile_scheduling_ext == HIPBLASLT_STREAMK_TILE_SCHEDULING_AUTO + || desc->streamk_tile_scheduling_ext == HIPBLASLT_STREAMK_TILE_SCHEDULING_ON, + desc->sm_count_target, + handle->sm_count_target); +} + // Resolve the effective uniform-summation-order request for a matmul launch // using first-on-wins precedence (C++ preference true, then desc==1, then // handle==1). 0 on desc/pref means inherit. There is no per-GEMM opt-out diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_controller.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_controller.hpp new file mode 100644 index 000000000000..71320d6e5d3b --- /dev/null +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_controller.hpp @@ -0,0 +1,140 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +// The HIP side of adaptive sm_count_target (see +// rocblaslt_adaptive_sm_count.hpp): a handle's probe slot arrays, one +// estimator per stream slot and the hints they publish. + +#include "rocblaslt_adaptive_sm_count.hpp" + +#include + +#include +#include +#include +#include +#include +#include + +// Adaptive sm_count_target settings, read from the environment once. +const rocblaslt::adaptive_sm::Config& rocblaslt_adaptive_sm_config(); + +// Writes one adaptive sm_count_target log line to the library log stream. +void rocblaslt_adaptive_sm_log(const std::string& line); + +namespace rocblaslt +{ + namespace adaptive_sm + { + class Launch; + + // Whether `stream` is capturing into a graph. A launch on a capturing + // stream is not adaptive: its graph replays it with frozen kernargs. + bool streamCapturing(hipStream_t stream); + + // Snapshot of one stream slot's state, for tests. + struct StreamStateForTest + { + uint32_t published; // published hint, 0: none + uint32_t epoch; // last probe epoch + bool pending; // a recorded probe is unread + uint32_t missed; + uint32_t sinceTick; + bool due; + uint32_t probesRead; + uint32_t lastCount; + + bool operator==(const StreamStateForTest& o) const + { + return published == o.published && epoch == o.epoch && pending == o.pending + && missed == o.missed && sinceTick == o.sinceTick && due == o.due + && probesRead == o.probesRead && lastCount == o.lastCount; + } + }; + + // A handle's adaptive state. Each stream slot has a probe slot array + // in host-mapped memory, written by the GPU with plain stores and read + // by the host without a sync, plus an estimator, a completion event + // and the hint it published. HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE is + // applied here only. + class Controller + { + public: + // For `streams` stream slots on the current device, which has `nCu` + // CUs. Null when adaptive mode is off or setup fails. + static std::unique_ptr create(size_t streams, uint32_t nCu); + + ~Controller(); + + // First phase of a C API launch on `stream`, in stream slot `slot` + // (-1: none): its hint, and the probe target when it should probe. + // A tiny launch never probes; a launch finding the slot's lock held + // leaves the slot's state alone. Not for a capturing stream. + Launch begin(hipStream_t stream, int slot, bool tiny); + + // Hint for the heuristic query, which has no stream: the latest + // published on any stream slot. + uint32_t queryHint() const; + + StreamStateForTest stateForTest(int slot); + + // Whether another thread holds `slot`'s lock. + bool lockedForTest(int slot); + + private: + friend class Launch; + struct Stream; + + Controller(const Config& cfg, size_t streams, uint32_t nCu); + + // Publishes `st`'s estimator hint if it changed. Called once per + // tracked launch, by begin(). + void publish(Stream& st, hipStream_t stream); + + const Config& m_cfg; + const uint32_t m_nCu; + std::unique_ptr m_streams; + uint32_t* m_probeHost = nullptr; + void* m_probeDev = nullptr; + std::atomic m_queryHint{0}; + }; + + // One C API launch, from Controller::begin() to its destruction, which + // records the completion event of a probe the launch ran. A probe + // launch holds its stream slot's lock throughout, so probes enqueue + // in epoch order; other launches on the slot meanwhile leave its + // state alone. + class Launch + { + public: + Launch(Launch&&) = default; + Launch& operator=(Launch&&) = delete; + ~Launch(); + + uint32_t hint = 0; // sm_count_target, 0: none + void* probeAddr = nullptr; // null: no probe + uint32_t epoch = 0; + uint32_t minGrid = 0; // workgroups a probe covers + bool probed = false; // set by the launch: a kernel carried the probe + + bool holdsLockForTest() const + { + return m_lock.owns_lock(); + } + + private: + friend class Controller; + + Launch() = default; + + // Second phase of a probe launch: records its probe, if it probed. + void finish(); + + hipStream_t m_stream = nullptr; + Controller::Stream* m_state = nullptr; + std::unique_lock m_lock; // held exactly when probing + }; + } // namespace adaptive_sm +} // namespace rocblaslt diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp index 373f6f0d101f..2c62ca11960b 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp @@ -166,6 +166,10 @@ namespace rocblaslt constexpr uint32_t c_probeSlots = 512; // Misses in a row before the hint drops to 0. constexpr uint32_t c_resetMisses = 2; + // Least workspace reported for the top heuristic result in adaptive + // mode, capped by the caller's limit: origami's static-grid + // partial-tile limit, enough for StreamK hybrid re-grids under any hint. + constexpr size_t c_adaptiveWorkspace = size_t(128) << 20; // The adaptive hint applies in StreamK AUTO or ON mode when neither the // desc nor the handle sets an explicit sm_count_target. diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp new file mode 100644 index 000000000000..30444f8ced2d --- /dev/null +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_adaptive_sm_controller.cpp @@ -0,0 +1,274 @@ +// Copyright Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "rocblaslt_adaptive_sm_controller.hpp" +#include "check_numerics_matrix.hpp" +#include "utility.hpp" + +#include + +#include +#include +#include +#include + +// Gated by HIPBLASLT_ADAPTIVE_SM_COUNT_LOG rather than the log mask, written +// where the library log goes, timestamped like CHECK_NUMERICS lines. +void rocblaslt_adaptive_sm_log(const std::string& line) +{ + std::lock_guard lk(log_mutex); + std::ostream* sink = get_logger_os(); + if(!sink) + sink = &std::cerr; + *sink << hipblaslt_check_numerics_ts() << "[hipBLASLt ADAPTIVE_SM] " << line << std::endl; +} + +const rocblaslt::adaptive_sm::Config& rocblaslt_adaptive_sm_config() +{ + static const rocblaslt::adaptive_sm::Config cfg = [] { + std::vector warnings; + const auto c = rocblaslt::adaptive_sm::Config::fromEnv(&warnings); + for(auto const& w : warnings) + rocblaslt_adaptive_sm_log(w); + if(c.log) + rocblaslt_adaptive_sm_log(c.describe()); + return c; + }(); + return cfg; +} + +namespace rocblaslt +{ + namespace adaptive_sm + { + bool streamCapturing(hipStream_t stream) + { + hipStreamCaptureStatus cap = hipStreamCaptureStatusNone; + return hipStreamIsCapturing(stream, &cap) == hipSuccess + && cap != hipStreamCaptureStatusNone; + } + + namespace + { + // Logs, under _LOG, a probe event call `what` that failed with + // `err` on `stream`, and clears the error: the probe gives no + // reading. + void eventFailed(const char* what, hipError_t err, hipStream_t stream) + { + if(rocblaslt_adaptive_sm_config().log) + { + std::ostringstream msg; + msg << "probe event " << what + << " failed stream=" << static_cast(stream) << " (" + << hipGetErrorName(err) << ")"; + rocblaslt_adaptive_sm_log(msg.str()); + } + static_cast(hipGetLastError()); + } + } // namespace + + struct Controller::Stream + { + ~Stream() + { + if(event) + static_cast(hipEventDestroy(event)); + } + + size_t slot = 0; // index in the handle's stream slots + std::atomic hint{0}; // published, 0: none; read without the lock + std::mutex lock; // guards the members below + hipEvent_t event = nullptr; // the pending probe's completion + Estimator estimator; + }; + + Launch::~Launch() + { + if(m_lock.owns_lock()) + finish(); + } + + std::unique_ptr Controller::create(size_t streams, uint32_t nCu) + { + const Config& cfg = rocblaslt_adaptive_sm_config(); + if(!cfg.enabled) + return nullptr; + + // Warns once per process, for the first device it exceeds. + static std::atomic warnedForce{false}; + if(cfg.force > nCu && !warnedForce.exchange(true)) + rocblaslt_adaptive_sm_log(std::string(kEnvForce) + "=" + std::to_string(cfg.force) + + " exceeds the device's " + std::to_string(nCu) + + " CUs"); + + std::unique_ptr c(new Controller(cfg, streams, nCu)); + const size_t bytes = streams * c_probeSlots * sizeof(uint32_t); + void* host = nullptr; + hipError_t err + = hipHostMalloc(&host, bytes, hipHostMallocCoherent | hipHostMallocMapped); + if(err == hipSuccess) + { + c->m_probeHost = static_cast(host); + std::memset(host, 0, bytes); + err = hipHostGetDevicePointer(&c->m_probeDev, host, 0); + } + if(err != hipSuccess) + { + rocblaslt_adaptive_sm_log("disabled: probe allocation failed (hipError=" + + std::to_string(static_cast(err)) + ")"); + return nullptr; + } + for(size_t i = 0; i < streams; ++i) + { + Stream& st = c->m_streams[i]; + st.slot = i; + st.estimator = Estimator(cfg, c->m_probeHost + i * c_probeSlots, nCu); + err = hipEventCreateWithFlags(&st.event, hipEventDisableTiming); + if(err != hipSuccess) + { + st.event = nullptr; + static_cast(hipGetLastError()); + rocblaslt_adaptive_sm_log("disabled: probe event creation failed (hipError=" + + std::to_string(static_cast(err)) + ")"); + return nullptr; + } + } + return c; + } + + Controller::Controller(const Config& cfg, size_t streams, uint32_t nCu) + : m_cfg(cfg) + , m_nCu(nCu) + , m_streams(std::make_unique(streams)) + { + } + + Controller::~Controller() + { + if(m_probeHost) + static_cast(hipHostFree(m_probeHost)); + } + + Launch Controller::begin(hipStream_t stream, int slot, bool tiny) + { + Launch launch; + launch.minGrid = m_nCu; + Stream* st = slot < 0 ? nullptr : &m_streams[slot]; + // Under force the state is kept only to probe under forceProbe. + if(st && m_cfg.tracksLaunches()) + { + // Another thread on this stream holds its state: launch untracked. + std::unique_lock lock(st->lock, std::try_to_lock); + if(lock.owns_lock()) + { + Estimator& est = st->estimator; + ProbeState state = ProbeState::Pending; + if(est.pending()) + { + // hipErrorNotReady does not set the last error; other failures do. + const hipError_t q = hipEventQuery(st->event); + if(q == hipSuccess) + state = ProbeState::Done; + else if(q != hipErrorNotReady) + { + eventFailed("query", q, stream); + state = ProbeState::Failed; + } + } + const uint32_t read = est.probesRead(); + launch.epoch = est.beforeLaunch(!tiny, state); + if(m_cfg.log >= 2 && est.probesRead() != read) + { + std::ostringstream msg; + msg << "reading stream=" << static_cast(stream) + << " slot=" << slot << " count=" << est.lastCount() + << (m_cfg.force ? " (discarded)" : ""); + rocblaslt_adaptive_sm_log(msg.str()); + } + publish(*st, stream); + if(launch.epoch) + { + launch.probeAddr = static_cast(m_probeDev) + + slot * c_probeSlots * sizeof(uint32_t); + launch.m_stream = stream; + launch.m_state = st; + launch.m_lock = std::move(lock); + } + } + } + if(m_cfg.force) + launch.hint = m_cfg.force; + else if(st) + launch.hint = st->hint.load(std::memory_order_relaxed); + return launch; + } + + void Launch::finish() + { + bool recorded = false; + // A capture begun since begin() would take the record into its + // graph: the probe then stays due. + if(probed && !streamCapturing(m_stream)) + { + const hipError_t err = hipEventRecord(m_state->event, m_stream); + recorded = err == hipSuccess; + if(!recorded) + eventFailed("record", err, m_stream); + } + m_state->estimator.afterProbe(recorded); + } + + void Controller::publish(Stream& st, hipStream_t stream) + { + const uint32_t hint = st.estimator.hint(); + const uint32_t old = st.hint.load(std::memory_order_relaxed); + // Readings under force move no published hint. + if(m_cfg.force || hint == old) + return; + st.hint.store(hint, std::memory_order_relaxed); + m_queryHint.store(hint, std::memory_order_relaxed); + if(!m_cfg.log) + return; + std::ostringstream msg; + msg << "hint stream=" << static_cast(stream) << " slot=" << st.slot + << " hint=" << old << "->" << hint; + const auto& w = st.estimator.window(); + if(w.size == 0) + msg << " missed=" << st.estimator.missed(); + else + msg << " readings="; + for(uint32_t i = 0; i < w.size; ++i) + msg << (i ? "," : "") << w.values[i]; + rocblaslt_adaptive_sm_log(msg.str()); + } + + uint32_t Controller::queryHint() const + { + return m_cfg.force ? m_cfg.force : m_queryHint.load(std::memory_order_relaxed); + } + + StreamStateForTest Controller::stateForTest(int slot) + { + Stream& st = m_streams[slot]; + const Estimator& est = st.estimator; + std::lock_guard guard(st.lock); + return {st.hint.load(std::memory_order_relaxed), + est.epoch(), + est.pending(), + est.missed(), + est.sinceTick(), + est.due(), + est.probesRead(), + est.lastCount()}; + } + + bool Controller::lockedForTest(int slot) + { + std::mutex& lock = m_streams[slot].lock; + if(!lock.try_lock()) + return true; + lock.unlock(); + return false; + } + } // namespace adaptive_sm +} // namespace rocblaslt 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 e4e13f6227a4..fda191586280 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 @@ -935,6 +935,12 @@ RocblasltContractionProblem construct_rocblaslt_problem(rocblaslt_handle // Encoding is a solution predicate; set it on every search path. problem.int4EncodingA = matmul_descr->int4_encoding_a_ext; + if(adaptive_sm_count_applies(handle, matmul_descr)) + { + problem.sm_count_target = int32_t(handle->adaptive_sm->queryHint()); + problem.adaptive_sm_count = true; + } + if(scaleAlphaVec) { // Fill owned storage with "1" for the compute type, and repoint alpha into it. @@ -2908,6 +2914,19 @@ rocblaslt_status pref->max_workspace_bytes); } + // The top Tensile result reports enough workspace for a StreamK hybrid + // re-grid under a later adaptive hint. A tuning-file result, the other + // results and AlgoCheck report the plain size. + if(prob.adaptive_sm_count && !override_success && *returnAlgoCount > 0 + && heuristicResultsArray[0].state == rocblaslt_status_success + && *(int*)heuristicResultsArray[0].algo.data >= 0) + { + using namespace rocblaslt::adaptive_sm; + auto& r0 = heuristicResultsArray[0]; + r0.workspaceSize = std::min(pref->max_workspace_bytes, + std::max(r0.workspaceSize, c_adaptiveWorkspace)); + } + if(override_success) { @@ -3614,3 +3633,39 @@ extern "C" HIPBLASLT_EXPORT size_t hipblaslt_tuning_stats_for_test(uint64_t* val values[i] = all[i]; return known; } + +// Test support: the adaptive sm_count_target state of `stream` into `state`, +// a rocblaslt::adaptive_sm::StreamStateForTest of `size` bytes. Returns false +// when adaptive mode is off, the stream has no slot or the size differs. +extern "C" HIPBLASLT_EXPORT bool + hipblaslt_adaptive_sm_state_for_test(void* handle, hipStream_t stream, void* state, size_t size) +{ + auto* h = static_cast(handle); + const int slot = h && h->adaptive_sm ? h->streamSlot(stream) : -1; + if(slot < 0 || size != sizeof(rocblaslt::adaptive_sm::StreamStateForTest)) + return false; + *static_cast(state) + = h->adaptive_sm->stateForTest(slot); + return true; +} + +// Test support: an adaptive launch on `stream` with no kernel. `*probe`: the +// launch planned a probe; `*heldBetween`, `*heldAfter`: the slot's lock is +// held while the launch lives and after it ends. Returns false when adaptive +// mode is off or the stream has no slot. +extern "C" HIPBLASLT_EXPORT bool hipblaslt_adaptive_sm_launch_lock_for_test( + void* handle, hipStream_t stream, bool tiny, bool* probe, bool* heldBetween, bool* heldAfter) +{ + auto* h = static_cast(handle); + const int slot = h && h->adaptive_sm ? h->streamSlot(stream) : -1; + if(slot < 0) + return false; + { + const auto launch = h->adaptive_sm->begin(stream, slot, tiny); + *probe = launch.epoch != 0; + *heldBetween = launch.holdsLockForTest(); + } + // This thread no longer holds the lock, so try_lock is well defined. + *heldAfter = h->adaptive_sm->lockedForTest(slot); + return true; +} diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp index 95293b1397c6..6adca16668e7 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/rocblaslt_mat.cpp @@ -32,6 +32,7 @@ #include "rocblaslt_mat_utils.hpp" #include "tensile_host.hpp" #include +#include #include @@ -267,7 +268,22 @@ rocblaslt_status rocblaslt_matmul_impl(const rocblaslt_handle handle, return rocblaslt_status_invalid_value; #endif + // Ends right after the launch, recording the completion of a probe it ran. + std::optional adaptive; + if(adaptive_sm_count_applies(handle, matmul_descr) + && !rocblaslt::adaptive_sm::streamCapturing(stream)) + { + const bool tiny = rocblaslt::adaptive_sm::tinyProblem( + m, n, k, matD->batch_count, rocblaslt_adaptive_sm_config()); + problem.adaptive_launch = &adaptive.emplace( + handle->adaptive_sm->begin(stream, handle->streamSlot(stream), tiny)); + problem.sm_count_target = int32_t(adaptive->hint); + problem.adaptive_sm_count = true; + } + rocblaslt_status st = runContractionProblem(handle, algo, problem, gemmData); + adaptive.reset(); + problem.adaptive_launch = nullptr; if(st == rocblaslt_status_success) { 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 e19e25a10138..7984d1d2320d 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 @@ -2030,6 +2030,16 @@ namespace return needsPadFreeDim; } + // The adaptive launch's occupancy probe target; none outside one. + void setOccupancyProbe(TensileLite::ContractionProblemParameters& params, + const RocblasltContractionProblem& prob) + { + if(const auto* launch = prob.adaptive_launch) + params.setOccupancyProbe(launch->probeAddr, launch->epoch, launch->minGrid); + else + params.setOccupancyProbe(nullptr, 0, 0); + } + /**************************************************************** * Construct a Tensile Problem from a RocblasltContractionProblem * ****************************************************************/ @@ -2355,6 +2365,7 @@ namespace // the effective CU budget. Non-StreamK=5 solutions ignore it. tensileProblem.setParams().setStreamKTileSchedulingMode(prob.streamk_tile_scheduling_ext); tensileProblem.setParams().setSmCountTarget(prob.sm_count_target); + setOccupancyProbe(tensileProblem.setParams(), prob); tensileProblem.setParams().setUniformSummationOrder(prob.uniform_summation_order != 0); @@ -2647,6 +2658,7 @@ namespace // companion block in ConstructTensileProblem for details. tensileProblem.setParams().setStreamKTileSchedulingMode(prob.streamk_tile_scheduling_ext); tensileProblem.setParams().setSmCountTarget(prob.sm_count_target); + setOccupancyProbe(tensileProblem.setParams(), prob); tensileProblem.setParams().setUniformSummationOrder(prob.uniform_summation_order != 0); @@ -3644,7 +3656,8 @@ TensileLite::ProblemOverride po.swizzleA = problem.swizzleA; po.swizzleB = problem.swizzleB; po.streamkTileScheduling = problem.streamk_tile_scheduling_ext; - po.smCountTarget = problem.sm_count_target; + // An adaptive hint is transient; the requested target is 0 then. + po.smCountTarget = problem.adaptive_sm_count ? 0 : problem.sm_count_target; po.uniformSummationOrder = problem.uniform_summation_order != 0; const auto& device = getDeviceIdentity(); @@ -6055,6 +6068,12 @@ rocblaslt_status runContractionProblem(rocblaslt_handle handle t_lastLaunchedIndex = solution->index; status = hip2RocStatus( adapter->launchKernels(kernels, prob.stream, nullptr, nullptr, isPreloaded)); + if(prob.adaptive_launch) + prob.adaptive_launch->probed + = status == rocblaslt_status_success + && std::any_of(kernels.begin(), kernels.end(), [](auto const& k) { + return k.occupancyProbe; + }); if(rocblaslt::Debug::Instance().printLogAsMarker()) rocblaslt::Debug::Instance().logMarkerStop(); } diff --git a/projects/hipblaslt/tensilelite/include/Tensile/Tensile.hpp b/projects/hipblaslt/tensilelite/include/Tensile/Tensile.hpp index 61b6ef3fe0d7..424a746b1099 100644 --- a/projects/hipblaslt/tensilelite/include/Tensile/Tensile.hpp +++ b/projects/hipblaslt/tensilelite/include/Tensile/Tensile.hpp @@ -134,6 +134,9 @@ namespace TensileLite size_t sharedMemBytes = 0; KernelArguments args; + + // Carries a non-null occupancy-probe address. + bool occupancyProbe = false; }; /** diff --git a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp index 1d6911b3bdd4..3c34c6f10f46 100644 --- a/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp +++ b/projects/hipblaslt/tensilelite/src/ContractionSolution.cpp @@ -2919,6 +2919,7 @@ namespace TensileLite if(needsPad && !canPad) probeAddr = nullptr; + rv.occupancyProbe = probeAddr != nullptr; rv.args.appendAligned("ProbeAddr", probeAddr); rv.args.append("ProbeEpoch", params.occupancyProbeEpoch()); rv.args.append("ProbeGrid", static_cast(realGrid)); diff --git a/projects/hipblaslt/tensilelite/tests/PersistentArgumentLayout_test.cpp b/projects/hipblaslt/tensilelite/tests/PersistentArgumentLayout_test.cpp index 35a8610de457..7a413c67733e 100644 --- a/projects/hipblaslt/tensilelite/tests/PersistentArgumentLayout_test.cpp +++ b/projects/hipblaslt/tensilelite/tests/PersistentArgumentLayout_test.cpp @@ -973,6 +973,7 @@ TEST(PersistentArgumentLayout, OccupancyProbeAbsentWithoutSupportOrForGroupedGem EXPECT_FALSE(hasArgument(call.args, "ProbeAddr")); EXPECT_FALSE(hasArgument(call.args, "ProbeEpoch")); EXPECT_FALSE(hasArgument(call.args, "ProbeGrid")); + EXPECT_FALSE(call.occupancyProbe); EXPECT_EQ(call.numWorkGroups.x, 7u); } } @@ -1020,6 +1021,7 @@ TEST(PersistentArgumentLayout, OccupancyProbeTailArgs) EXPECT_EQ(value(args, "numWorkGroups"), 7u); EXPECT_EQ(value(args, "ProbeGrid"), 7u); EXPECT_EQ(value(args, "ProbeAddr"), c.addr); + EXPECT_EQ(call.occupancyProbe, c.addr != nullptr); EXPECT_EQ(value(args, "ProbeEpoch"), 41 + c.minGrid); auto addrOffset = offset(args, "ProbeAddr"); EXPECT_EQ(addrOffset % 8, 0u); @@ -1052,6 +1054,7 @@ TEST(PersistentArgumentLayout, OccupancyProbeDisabledForMultiDimGridBelowMinGrid EXPECT_LT(call.numWorkGroups.x * call.numWorkGroups.y, 256u); EXPECT_EQ(call.numWorkItems.x, call.workGroupSize.x * call.numWorkGroups.x); EXPECT_EQ(value(call.args, "ProbeAddr"), nullptr); + EXPECT_FALSE(call.occupancyProbe); EXPECT_EQ(value(call.args, "ProbeGrid"), call.numWorkGroups.x * call.numWorkGroups.y * call.numWorkGroups.z); } From daab76b6bc2b35da7f2203766459976c1972b911 Mon Sep 17 00:00:00 2001 From: Caio Salvador Rohwedder Date: Sat, 17 Oct 2026 00:17:14 -0500 Subject: [PATCH 2/3] Refuse benchmark timing with an unforced adaptive sm_count_target A varying adaptive hint can change the timed kernel mid-run, so the timing clients exit unless HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE fixes it. --- projects/hipblaslt/clients/bench/README.md | 6 ++ .../hipblaslt/clients/bench/src/client.cpp | 3 + .../clients/bench/src/client_api_overhead.cpp | 5 ++ .../clients/bench/src/client_sequence.cpp | 2 + .../clients/common/include/utility.hpp | 5 ++ .../hipblaslt/clients/common/src/utility.cpp | 11 +++ .../tests/unit/adaptive_sm_count_gtest.cpp | 71 ++++++++++++++++++- 7 files changed, 100 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/clients/bench/README.md b/projects/hipblaslt/clients/bench/README.md index 018540b75617..1dfbbbb6d7e9 100644 --- a/projects/hipblaslt/clients/bench/README.md +++ b/projects/hipblaslt/clients/bench/README.md @@ -109,6 +109,12 @@ For debugging, `TENSILE_PERSISTENT_HYBRID_FORCE_MODE` overrides the policy: `-1` The legacy environment alias `TENSILE_STREAMK5_FORCE_MODE` remains supported; the preferred environment name takes precedence when present. +`hipblaslt-bench`, `hipblaslt-sequence`, and `hipblaslt-api-overhead` exit with an error when +`HIPBLASLT_ADAPTIVE_SM_COUNT` is enabled without `HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE`, because the +adaptive hint can change the timed solution during a run. Set `HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE=` +to time the solution adaptive mode picks at that CU count, and add +`HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE=1` to include the probe cost. + # demo Run fp32 GEMM with validation ``` diff --git a/projects/hipblaslt/clients/bench/src/client.cpp b/projects/hipblaslt/clients/bench/src/client.cpp index 31cdbf3d25c7..630fb0960dbc 100644 --- a/projects/hipblaslt/clients/bench/src/client.cpp +++ b/projects/hipblaslt/clients/bench/src/client.cpp @@ -821,6 +821,9 @@ try return 0; } + if(!hipblaslt_timing_allowed_with_adaptive_sm_count()) + return EXIT_FAILURE; + if(api_method_str.compare("c") == 0) { api_method = 0; diff --git a/projects/hipblaslt/clients/bench/src/client_api_overhead.cpp b/projects/hipblaslt/clients/bench/src/client_api_overhead.cpp index 0922602ec26d..07777250f77c 100644 --- a/projects/hipblaslt/clients/bench/src/client_api_overhead.cpp +++ b/projects/hipblaslt/clients/bench/src/client_api_overhead.cpp @@ -33,6 +33,9 @@ #include "helper.h" +// From the clients' utility.hpp, whose CHECK_* macros clash with helper.h's. +bool hipblaslt_timing_allowed_with_adaptive_sm_count(); + using clock_point = std::chrono::time_point; using double_micro = std::chrono::duration; using HHSRunner = Runner; @@ -201,6 +204,8 @@ int main(int argc, char** argv) printUsage(argv[0]); return err; } + if(!hipblaslt_timing_allowed_with_adaptive_sm_count()) + return EXIT_FAILURE; int returned_sol; // new: test get heuristic with requested-solution = 100, -1 diff --git a/projects/hipblaslt/clients/bench/src/client_sequence.cpp b/projects/hipblaslt/clients/bench/src/client_sequence.cpp index 0faff25e4926..40e4a3def3b6 100644 --- a/projects/hipblaslt/clients/bench/src/client_sequence.cpp +++ b/projects/hipblaslt/clients/bench/src/client_sequence.cpp @@ -381,6 +381,8 @@ int main(int argc, char** argv) hipblaslt_cerr << "Usage: " << argv[0] << " yaml-config.yaml" << std::endl; exit(1); } + if(!hipblaslt_timing_allowed_with_adaptive_sm_count()) + exit(1); auto inputFile = llvm::MemoryBuffer::getFile(argv[1]); LayerConfigIO rv; llvm::yaml::Input yin((*inputFile)->getMemBufferRef()); diff --git a/projects/hipblaslt/clients/common/include/utility.hpp b/projects/hipblaslt/clients/common/include/utility.hpp index c57ce77fb984..53aad2702664 100644 --- a/projects/hipblaslt/clients/common/include/utility.hpp +++ b/projects/hipblaslt/clients/common/include/utility.hpp @@ -398,6 +398,11 @@ int32_t hipblaslt_get_arch(); int32_t hipblaslt_get_arch_major(); void hipblaslt_print_version(); +/*! \brief Timing clients call this first: returns false, after printing why, when + * HIPBLASLT_ADAPTIVE_SM_COUNT may change the timed solution (adaptive mode + * without HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE). */ +bool hipblaslt_timing_allowed_with_adaptive_sm_count(); + /* ==================================================================== */ /*! \brief write a matrix to file. */ void hipblasltDispatchValuesToFile(hipblasOperation_t transA, hipDataType TiA, diff --git a/projects/hipblaslt/clients/common/src/utility.cpp b/projects/hipblaslt/clients/common/src/utility.cpp index aa31bcf6e3ee..69c146bfc05d 100644 --- a/projects/hipblaslt/clients/common/src/utility.cpp +++ b/projects/hipblaslt/clients/common/src/utility.cpp @@ -30,6 +30,7 @@ #include #endif +#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp" #include "d_vector.hpp" #include "utility.hpp" #include @@ -476,3 +477,13 @@ void hipblasltDispatchValuesToFile(hipblasOperation_t transA, hipDataType T, "condition to write to file" << std::endl; } + +bool hipblaslt_timing_allowed_with_adaptive_sm_count() +{ + using namespace rocblaslt::adaptive_sm; + if(!Config::fromEnv().hintVaries()) + return true; + hipblaslt_cerr << "error: " << kEnvEnable << " changes the timed solution during timing; set " + << kEnvForce << "= to time a fixed hint" << std::endl; + return false; +} diff --git a/projects/hipblaslt/clients/tests/unit/adaptive_sm_count_gtest.cpp b/projects/hipblaslt/clients/tests/unit/adaptive_sm_count_gtest.cpp index 3795061e5255..e976e043d9a7 100644 --- a/projects/hipblaslt/clients/tests/unit/adaptive_sm_count_gtest.cpp +++ b/projects/hipblaslt/clients/tests/unit/adaptive_sm_count_gtest.cpp @@ -2,17 +2,20 @@ // SPDX-License-Identifier: MIT // Host-only tests for the adaptive sm_count_target estimator and helpers in -// rocblaslt_adaptive_sm_count.hpp, reached by relative #include. Launch -// scenarios drive Estimator::beforeLaunch/afterProbe over a hand-filled probe -// slot array. +// rocblaslt_adaptive_sm_count.hpp, reached by relative #include, and for the +// benchmark timing guard built on them. Launch scenarios drive +// Estimator::beforeLaunch/afterProbe over a hand-filled probe slot array. #include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp" +#include "utility.hpp" #include #include #include +#include #include +#include #include #include #include @@ -635,4 +638,66 @@ namespace EXPECT_EQ(w[0].rfind(std::string(name) + "=", 0), 0u) << w[0]; } } + + // hipblaslt_timing_allowed_with_adaptive_sm_count reads the process + // environment on each call; the variables are restored afterwards. + class HostUnitAdaptiveSmCountTiming : public ::testing::Test + { + protected: + void SetUp() override + { + for(const char* name : kVariables) + { + const char* value = std::getenv(name); + m_saved.emplace_back(name, + value ? std::optional(value) : std::nullopt); + set(name, nullptr); + } + } + + void TearDown() override + { + for(const auto& [name, value] : m_saved) + set(name, value ? value->c_str() : nullptr); + } + + static void set(const char* name, const char* value) + { +#if defined(_WIN32) + _putenv_s(name, value ? value : ""); +#else + if(value) + setenv(name, value, 1); + else + unsetenv(name); +#endif + } + + static constexpr const char* kVariables[] = {kEnvEnable, kEnvForce}; + + std::vector>> m_saved; + }; + + TEST_F(HostUnitAdaptiveSmCountTiming, smoke_RefusedOnlyWhenTheHintVaries) + { + struct Case + { + const char* label; + const char* enable; + const char* force; + bool allowed; + }; + const std::vector cases = { + {"off", nullptr, nullptr, true}, + {"adaptive", "1", nullptr, false}, + {"adaptive, forced", "1", "128", true}, + }; + for(const Case& c : cases) + { + SCOPED_TRACE(c.label); + set(kEnvEnable, c.enable); + set(kEnvForce, c.force); + EXPECT_EQ(hipblaslt_timing_allowed_with_adaptive_sm_count(), c.allowed); + } + } } // namespace From da2a1c17bbb8bd681dcf59a5d2658f6b2134caef Mon Sep 17 00:00:00 2001 From: Caio Salvador Rohwedder Date: Sat, 17 Oct 2026 00:17:14 -0500 Subject: [PATCH 3/3] Document the adaptive sm_count_target mode --- projects/hipblaslt/CHANGELOG.md | 1 + .../docs/how-to/how-to-use-streamk.rst | 91 +++++++++++++++++++ .../docs/reference/env-variables.rst | 54 +++++++++++ 3 files changed, 146 insertions(+) diff --git a/projects/hipblaslt/CHANGELOG.md b/projects/hipblaslt/CHANGELOG.md index f655cc8b77df..321e9c151b06 100644 --- a/projects/hipblaslt/CHANGELOG.md +++ b/projects/hipblaslt/CHANGELOG.md @@ -25,6 +25,7 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip * `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=off|cache|tune` and `HIPBLASLT_TUNING_CACHE_PATH=` let the library benchmark candidate kernels for an unseen problem at its first `hipblasLtMatmul`, record the winner, and replay it afterwards. `tune` measures candidates the way `hipblaslt-bench` does (warm-up and timed launches over rotating buffers, with an instruction-cache flush), serialises tuning within a process, and bounds each shape's search to five minutes by default (`HIPBLASLT_TUNING_BUDGET_MS_PER_SHAPE`). A search the budget stops records its best candidate as incomplete, which is never slower than the untuned kernel because that kernel is measured first, and tune mode revisits it when a later run can finish or widen the search. `cache` only replays, validating each entry as `HIPBLASLT_TUNING_OVERRIDE_FILE` does, and also serves `hipblasLtMatmul` calls that pass no `algo`. An explicit `algo` always launches as given. Both modes report what they loaded, each tuned shape's start and result, and a closing summary without needing a log level. With a tuning mode set, `HIPBLASLT_TUNING_OVERRIDE_FILE` is ignored, and a process in a secure execution context ignores every tuning variable. * `TENSILELITE_EMIT_OCCUPANCY_PROBE` CMake option and `invoke build --emit-occupancy-probe` (off by default) to emit a CU-occupancy probe into gfx950 Stream-K Hybrid kernels, needed by `HIPBLASLT_ADAPTIVE_SM_COUNT`. +* Adaptive `sm_count_target` mode, off by default. With `HIPBLASLT_ADAPTIVE_SM_COUNT=1`, `hipblasLtMatmul` periodically measures the CUs available to a stream with the occupancy probe and uses the count as an implicit `sm_count_target`; `HIPBLASLT_ADAPTIVE_SM_COUNT_*` variables tune or force it. In this mode the top heuristic result reports `workspaceSize = min(MAX_WORKSPACE_BYTES, max(required, 128 MiB))`, and `hipblaslt-bench`, `hipblaslt-sequence`, and `hipblaslt-api-overhead` refuse to time unless `HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE` is set. ### Changed diff --git a/projects/hipblaslt/docs/how-to/how-to-use-streamk.rst b/projects/hipblaslt/docs/how-to/how-to-use-streamk.rst index a119514d4f81..6820a94d188c 100644 --- a/projects/hipblaslt/docs/how-to/how-to-use-streamk.rst +++ b/projects/hipblaslt/docs/how-to/how-to-use-streamk.rst @@ -120,3 +120,94 @@ Follow these guidelines to optimize how Stream-K uses resources: .. code-block:: bash export TENSILE_PERSISTENT_MAX_CUS=32 + +.. _adaptive-sm-count: + +Adapting to the available compute units +======================================= + +When other work occupies part of the device, hipBLASLt can measure how many compute units (CUs) +are available to its GEMMs and use that count as an implicit ``sm_count_target``. +Every few launches, a GEMM kernel records which CUs its workgroups ran on. hipBLASLt reads the +result asynchronously and updates a per-stream CU estimate, the adaptive hint. +The hint affects solution selection and the Stream-K Hybrid grid size and assignment mode. + +Requirements +------------ + +Adaptive mode applies to a ``hipblasLtMatmul`` call when all of the following are true: + +* The call uses the C API (``hipblasLtMatmul``). The hipBLASLtExt API, grouped GEMM, rocRoller, and + launches during stream capture are not adapted. +* ``HIPBLASLT_MATMUL_DESC_STREAMK_TILE_SCHEDULING_EXT`` is set to ``HIPBLASLT_STREAMK_TILE_SCHEDULING_AUTO`` + or ``HIPBLASLT_STREAMK_TILE_SCHEDULING_ON``. Under ``ON`` the dynamic grid ignores the hint, + but solution selection still uses it. +* Neither ``HIPBLASLT_MATMUL_DESC_SM_COUNT_TARGET`` nor ``hipblasLtSetSmCountTarget`` sets a positive value. +* ``HIPBLASLT_ADAPTIVE_SM_COUNT=1`` is set. + +.. code-block:: bash + + export HIPBLASLT_ADAPTIVE_SM_COUNT=1 + +These conditions do not depend on the device or the device library. Whenever they hold, the top +heuristic result reports a larger workspace, launches are tracked for probing, and the timing +clients refuse to run (see the behavior notes and timing sections below). + +The hint itself only moves when the launched kernels carry the occupancy probe, which requires both: + +* A device library built with the occupancy probe, using ``inv build --emit-occupancy-probe`` + or the CMake option ``-DTENSILELITE_EMIT_OCCUPANCY_PROBE=ON``. +* A gfx950 device. The probe is only emitted for gfx950 Stream-K Hybrid kernels. + +Otherwise the hint stays ``0`` (no hint), unless ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE`` sets it. + +The environment variables that control adaptive mode are listed in the following table. +They are read once per process. Boolean variables accept ``1``, ``on``, or ``true`` and ``0``, ``off``, or ``false`` +(case-insensitive). An empty variable keeps its default. An invalid or out-of-range value keeps the default +and logs a warning. + +.. csv-table:: + :header: "Environment variable","Description" + :widths: 40, 100 + + "``HIPBLASLT_ADAPTIVE_SM_COUNT``","Enable adaptive mode. Default: ``0``." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_PROBE_EVERY``","Launches per stream between probes. Range: ``1`` or greater. Default: ``8``." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_TINY_MFLOPS``","GEMMs with fewer than this many MFLOPs (``2*M*N*K*batch / 10^6``) never probe. Default: ``256``." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_PROBE_WINDOW``","Number of latest probe readings that decide the hint. The hint changes only when all of them are below it or all are above it, and then takes the smallest reading. No hint counts as the device's CU count. A single reading of the device's CU count clears the hint at once, and a new hint then needs a full set of readings below that count. Range: ``1`` to ``16``. Default: ``3``." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_LOG``","``0``: no logging (default). ``1``: log the configuration, hint changes, and failed probe events. ``2``: also log every reading. Lines start with a timestamp and ``[hipBLASLt ADAPTIVE_SM]`` and are written to ``HIPBLASLT_LOG_FILE`` if set, otherwise to stderr." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``","Use this fixed CU count as the hint for every query and launch instead of measuring it. A positive value enables adaptive mode on its own. Range: ``0`` to ``65535``. Default: ``0`` (measure)." + "``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE``","With ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``, still run the probes and discard their readings, so timing includes the probe cost. Ignored with a warning without ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``. Default: ``0``." + +Behavior notes +-------------- + +* **Workspace**: In adaptive mode, the top heuristic result reports + ``workspaceSize = min(MAX_WORKSPACE_BYTES, max(required, 128 MiB))`` so a regridded launch + has room to run. Set ``HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES`` to your workspace budget. +* **Precedence**: An explicit ``sm_count_target`` on the descriptor or handle always wins over the adaptive hint. +* **Probing cost**: GEMMs below ``HIPBLASLT_ADAPTIVE_SM_COUNT_TINY_MFLOPS`` never probe. + Each probe records an event on the stream. While one thread launches a probe, other threads' + launches on the same stream run without updating the estimate. +* **Missed probes**: A probe comes due every ``HIPBLASLT_ADAPTIVE_SM_COUNT_PROBE_EVERY`` launches. + When the next one comes due and the previous one was still not taken, that counts as a missed probe, + for example when only tiny GEMMs ran or the selected solutions do not carry the probe, such as a + solution other than a Stream-K Hybrid kernel chosen at the hint. Two missed probes in a row clear + the hint until new readings arrive. +* **Graph capture**: Launches on a capturing stream are not adapted: they are not probed and run + without a hint, including ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``. +* **Detection scope**: The probe detects co-running work that takes whole CUs, for example a CU-masked stream. + It does not detect work that shares CUs with the GEMM. + +Timing with adaptive mode +------------------------- + +Measured performance depends on the hint, which can change during a run. For this reason +``hipblaslt-bench``, ``hipblaslt-sequence``, and ``hipblaslt-api-overhead`` exit with an error when +``HIPBLASLT_ADAPTIVE_SM_COUNT`` is enabled without ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``. +To time the solution adaptive mode picks at a given hint, set ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE`` +to that CU count. To include the probe cost, also set ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE=1``. + +.. code-block:: bash + + HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE=192 HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE=1 \ + hipblaslt-bench -m 4096 -n 4096 -k 4096 --precision f16_r --hybrid_assignment_policy Auto diff --git a/projects/hipblaslt/docs/reference/env-variables.rst b/projects/hipblaslt/docs/reference/env-variables.rst index a961542c7be7..434ac5c0383c 100644 --- a/projects/hipblaslt/docs/reference/env-variables.rst +++ b/projects/hipblaslt/docs/reference/env-variables.rst @@ -243,6 +243,60 @@ even if its value is invalid; it does not fall back to the legacy value. | 1: Force dynamic work-queue assignment | Malformed or out-of-range values are ignored. +Adaptive compute-unit target +============================ + +The adaptive ``sm_count_target`` environment variables for hipBLASLt are collected in the following table. +They are read once per process. Boolean values accept ``1``, ``on``, ``true``, ``0``, ``off``, and ``false`` +(case-insensitive). Invalid or out-of-range values keep the default and log a warning. +For requirements and behavior, see :ref:`Adapting to the available compute units `. + +.. list-table:: + :header-rows: 1 + :widths: 70,30 + + * - **Environment variable** + - **Value** + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT`` + | Measures the compute units available to GEMMs and uses the count as an implicit ``sm_count_target``. + | Applies to C API calls with Stream-K tile scheduling ``AUTO`` or ``ON`` and no explicit ``sm_count_target``. + | The hint only moves when the launched kernels carry the occupancy probe (a probe-built device library and gfx950 Stream-K Hybrid kernels); otherwise it stays 0. + - | 0: Off (default) + | 1: On + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_PROBE_EVERY`` + | Launches per stream between probes. + - | Integer, 1 or greater + | Default: 8 + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_TINY_MFLOPS`` + | GEMMs below this many MFLOPs (``2*M*N*K*batch / 10^6``) never probe. + - | Non-negative integer + | Default: 256 + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_PROBE_WINDOW`` + | Number of latest probe readings that decide the hint. The hint changes only when all of them are below it or all are above it, and then takes the smallest reading. No hint counts as the device's CU count. A single reading of the device's CU count clears the hint at once, and a new hint then needs a full set of readings below that count. + - | Integer from 1 to 16 + | Default: 3 + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_LOG`` + | Logs adaptive mode activity to ``HIPBLASLT_LOG_FILE``, or stderr if unset. + - | 0: Off (default) + | 1: Configuration, hint changes, and failed probe events + | 2: Also every probe reading + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE`` + | Uses a fixed CU count as the hint instead of measuring it. A positive value enables adaptive mode on its own. + - | Integer from 0 to 65535 + | Default: 0 (measure) + + * - | ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE_PROBE`` + | With ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``, still runs the probes and discards their readings, so timing includes the probe cost. + | Ignored with a warning without ``HIPBLASLT_ADAPTIVE_SM_COUNT_FORCE``. + - | 0: Off (default) + | 1: On + .. _env-type_overrides: Type overrides