caio96 · Oct 28, 2026 · Oct 17, 2026 · Oct 17, 2026 · Oct 17, 2026
@@ -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=<file>` 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
 

@@ -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=<cus>`
+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
 ```

@@ -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;

@@ -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<std::chrono::steady_clock>;
 using double_micro = std::chrono::duration<double, std::micro>;
 using HHSRunner    = Runner<hipblasLtHalf, hipblasLtHalf, hipblasLtHalf, float, float>;
@@ -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

@@ -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());

@@ -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,

@@ -30,6 +30,7 @@
 #include <libloaderapi.h>
 #endif
 
+#include "../../../library/src/amd_detail/rocblaslt/src/include/rocblaslt_adaptive_sm_count.hpp"
 #include "d_vector.hpp"
 #include "utility.hpp"
 #include <atomic>
@@ -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 << "=<cus> to time a fixed hint" << std::endl;
+    return false;
+}
@@ -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