tony-davis · Oct 10, 2026 · Oct 9, 2026
@@ -35,6 +35,7 @@ Full documentation for hipBLASLt is available at [rocm.docs.amd.com/projects/hip
 * Solution cache key now includes `HIPBLASLT_MATMUL_DESC_SM_COUNT_TARGET` and the StreamK tile scheduling mode, so the same problem can select a different kernel than before.
 * A tuning file is now trusted one entry at a time instead of all or nothing. `HIPBLASLT_TUNING_OVERRIDE_FILE` records solution indices, which are positions in one build's kernel library. Previously the C API ignored the whole file when its build-version line did not match the running build, while the C++ API applied it regardless and could run kernels it was never tuned on. `hipblaslt-bench` now records a `kernel_name` beside each `solution_index` in `HIPBLASLT_TUNING_FILE`, and on both APIs a row that records a name is checked at replay and dropped only if its index no longer names that kernel. A row without a name is used only when the file's `Git Version` line matches the running build; a build made outside a git checkout has no version, so it uses no such rows. Problems whose rows are dropped fall back to normal kernel selection.
 * Extension API `initialize` compares against the size set with `setMaxWorkspaceBytes` on the `Gemm` or `GroupedGemm` instance, which defaults to 0. A caller that passes a large enough buffer without setting it is now rejected with `HIPBLAS_STATUS_INVALID_VALUE` for solutions with a fixed workspace requirement, such as GSU. Stream-K and zero-workspace solutions are unaffected.
+* Workspace limits and sizes above `UINT32_MAX` (4 GiB - 1), such as `SIZE_MAX` for no limit, are clamped to `UINT32_MAX`, because kernels address the workspace with 32-bit offsets. This applies to `HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES` (which then reads back as `UINT32_MAX`), the `hipblasLtMatmul` workspace size, and the extension API `setMaxWorkspaceBytes`.
 
 ### Removed
 

@@ -2,7 +2,8 @@
 // SPDX-License-Identifier: MIT
 
 // The C++ heuristic query must apply a workspace limit from 2 GiB to UINT32_MAX as the C query
-// does. Kernels address the workspace with 32-bit offsets, so both APIs reject anything larger.
+// does. Kernels address the workspace with 32-bit offsets, so both APIs clamp anything larger to
+// UINT32_MAX. Callers pass SIZE_MAX to mean "no limit", so a larger value must not be an error.
 
 #include <gtest/gtest.h>
 #include <hip/hip_runtime.h>
@@ -40,7 +41,7 @@ namespace
         return out;
     }
 
-    // The buffers are never read: only the heuristic queries run.
+    // The buffers hold no data: the tests check statuses and heuristic results only.
     class HeuristicWorkspaceLimit : public ::testing::Test
     {
     protected:
@@ -173,19 +174,52 @@ namespace
         }
     }
 
-    // Every entry point that takes a workspace size or limit rejects it before selecting or
-    // launching a kernel, so no solution library is needed.
-    TEST_F(HeuristicWorkspaceLimit, smoke_RejectsWorkspaceAboveUint32Max)
+    // The preference stores a limit above UINT32_MAX as UINT32_MAX, so no solution library is
+    // needed.
+    TEST_F(HeuristicWorkspaceLimit, smoke_PreferenceClampsLimitAboveUint32Max)
     {
         for(uint64_t limit : {4 * kGiB, std::numeric_limits<uint64_t>::max()})
         {
             SCOPED_TRACE(limit);
-            EXPECT_EQ(setLimit(limit), HIPBLAS_STATUS_INVALID_VALUE);
+            hipblasLtMatmulPreference_t pref = nullptr;
+            ASSERT_EQ(hipblasLtMatmulPreferenceCreate(&pref), HIPBLAS_STATUS_SUCCESS);
+            EXPECT_EQ(hipblasLtMatmulPreferenceSetAttribute(
+                          pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &limit, sizeof(limit)),
+                      HIPBLAS_STATUS_SUCCESS);
+            uint64_t stored  = 0;
+            size_t   written = 0;
+            EXPECT_EQ(
+                hipblasLtMatmulPreferenceGetAttribute(pref,
+                                                      HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                                                      &stored,
+                                                      sizeof(stored),
+                                                      &written),
+                HIPBLAS_STATUS_SUCCESS);
+            EXPECT_EQ(stored, kUint32Max);
+            hipblasLtMatmulPreferenceDestroy(pref);
+        }
+    }
+
+    // Every entry point that takes a workspace size or limit treats a value above UINT32_MAX
+    // as UINT32_MAX.
+    TEST_F(HeuristicWorkspaceLimit, smoke_ClampsWorkspaceAboveUint32Max)
+    {
+        // Results at the cap are the reference. Solutions that fit the real buffer are run.
+        std::vector<hipblasLtMatmulHeuristicResult_t> capped, fitting;
+        if(cHeuristic(kUint32Max, capped) != HIPBLAS_STATUS_SUCCESS || capped.empty()
+           || cHeuristic(kWorkspace, fitting) != HIPBLAS_STATUS_SUCCESS || fitting.empty())
+            GTEST_SKIP() << "No solution for this problem in the loaded library";
 
-            std::vector<hipblasLtMatmulHeuristicResult_t> results;
-            EXPECT_EQ(cppHeuristic(limit, results), HIPBLAS_STATUS_INVALID_VALUE);
-            EXPECT_TRUE(results.empty());
+        for(uint64_t limit : {4 * kGiB, std::numeric_limits<uint64_t>::max()})
+        {
+            SCOPED_TRACE(limit);
+            std::vector<hipblasLtMatmulHeuristicResult_t> c, cpp;
+            ASSERT_EQ(cHeuristic(limit, c), HIPBLAS_STATUS_SUCCESS);
+            EXPECT_EQ(summary(c), summary(capped));
+            ASSERT_EQ(cppHeuristic(limit, cpp), HIPBLAS_STATUS_SUCCESS);
+            EXPECT_EQ(summary(cpp), summary(capped));
 
+            // The algo's own workspace limit keeps the kernel inside the real buffer.
             EXPECT_EQ(hipblasLtMatmul(handle,
                                       desc,
                                       &alpha,
@@ -198,11 +232,12 @@ namespace
                                       layoutD,
                                       d,
                                       layoutD,
-                                      nullptr,
+                                      &fitting.front().algo,
                                       ws,
                                       limit,
                                       nullptr),
-                      HIPBLAS_STATUS_INVALID_VALUE);
+                      HIPBLAS_STATUS_SUCCESS);
+            EXPECT_EQ(hipStreamSynchronize(nullptr), hipSuccess);
 
             hipblaslt_ext::Gemm gemm(handle,
                                      HIPBLAS_OP_N,
@@ -214,8 +249,8 @@ namespace
                                      HIPBLAS_COMPUTE_32F);
             ASSERT_EQ(setProblem(gemm), HIPBLAS_STATUS_SUCCESS);
             gemm.setMaxWorkspaceBytes(limit);
-            hipblasLtMatmulAlgo_t algo{};
-            EXPECT_EQ(gemm.initialize(algo, ws), HIPBLAS_STATUS_INVALID_VALUE);
+            hipblasLtMatmulAlgo_t algo = fitting.front().algo;
+            EXPECT_EQ(gemm.initialize(algo, ws), HIPBLAS_STATUS_SUCCESS);
         }
     }
 } // namespace
@@ -74,8 +74,8 @@ namespace hipblaslt_ext
          *  \brief This function sets the maximum workspace size.
          *
          *  @param[in]
-         *  workspaceBytes  Set the maximum workspace size in bytes. Above ``UINT32_MAX``
-         *  (4 GiB - 1), algoGetHeuristic returns HIPBLAS_STATUS_INVALID_VALUE.
+         *  workspaceBytes  Set the maximum workspace size in bytes. algoGetHeuristic clamps
+         *  values above ``UINT32_MAX`` (4 GiB - 1) to ``UINT32_MAX``.
          */
         HIPBLASLT_EXPORT void setMaxWorkspaceBytes(size_t workspaceBytes);
 
@@ -490,8 +490,8 @@ namespace hipblaslt_ext
          *  \brief This function sets the maximum workspace size.
          *
          *  @param[in]
-         *  workspaceBytes  Sets the maximum workspace size in bytes. Above ``UINT32_MAX``
-         *  (4 GiB - 1), initialize returns HIPBLAS_STATUS_INVALID_VALUE.
+         *  workspaceBytes  Sets the maximum workspace size in bytes. initialize clamps
+         *  values above ``UINT32_MAX`` (4 GiB - 1) to ``UINT32_MAX``.
          */
         HIPBLASLT_EXPORT void setMaxWorkspaceBytes(size_t workspaceBytes);
 

@@ -384,7 +384,7 @@ typedef enum {
  */
 typedef enum {
   HIPBLASLT_MATMUL_PREF_SEARCH_MODE = 0,          /**<Search mode. Data type: ``uint32_t``. */
-  HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES = 1,  /**<Maximum allowed workspace memory. Default is 0 (no workspace memory allowed). Values above ``UINT32_MAX`` (4 GiB - 1) are rejected with ``HIPBLAS_STATUS_INVALID_VALUE``. Data type: ``uint64_t``. */
+  HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES = 1,  /**<Maximum allowed workspace memory. Default is 0 (no workspace memory allowed). Values above ``UINT32_MAX`` (4 GiB - 1), such as ``SIZE_MAX`` for no limit, are clamped to ``UINT32_MAX``, which is the value read back. Data type: ``uint64_t``. */
   HIPBLASLT_MATMUL_PREF_SM_COUNT_TARGET = 2,      /**<Bias heuristic algorithm selection toward kernels that perform well at this targeted compute-unit count. ``0`` (default) means no constraint. Negative values are rejected with ``HIPBLAS_STATUS_INVALID_VALUE``. Data type: ``int32_t``. */
   HIPBLASLT_MATMUL_PREF_MAX = 3
 } hipblasLtMatmulPreferenceAttributes_t;
@@ -1320,7 +1320,8 @@ hipblasStatus_t
  * memory. Pointer must be 16B aligned (that is, the lowest 4 bits of the address must
  * be 0).
  *  @param[in]
- *  workspaceSizeInBytes    Size of the workspace.
+ *  workspaceSizeInBytes    Size of the workspace. Values above ``UINT32_MAX``
+ * (4 GiB - 1) are clamped to ``UINT32_MAX``.
  *  @param[in]
  *  stream                  The HIP stream where all GPU work is
  * submitted.
@@ -1333,8 +1334,7 @@ hipblasStatus_t
  * selected device doesn't support the configured operation. \retval
  * HIPBLAS_STATUS_INVALID_VALUE     If the parameters are unexpectedly NULL, in
  * conflict, or in an impossible configuration. For example, when
- * workspaceSizeInBytes is less than the workspace required by the configured algorithm,
- * or above ``UINT32_MAX`` (4 GiB - 1).
+ * workspaceSizeInBytes is less than the workspace required by the configured algorithm.
  *  \retval HIBLAS_STATUS_NOT_INITIALIZED    If the hipBLASLt handle has not been
  * initialized.
  */

@@ -34,23 +34,25 @@
 #include <limits>
 
 /*******************************************************************************
- * Validate Workspace Size
+ * Clamp Workspace Size
  * Kernels and helper kernels address the workspace with 32-bit byte offsets
- * and buffer sizes, so a larger workspace cannot be used.
+ * and buffer sizes, so a larger workspace cannot be used. A workspace size or
+ * limit is a ceiling, so larger values (e.g. SIZE_MAX for "no limit") are
+ * clamped.
  ******************************************************************************/
 constexpr size_t rocblaslt_max_workspace_bytes = std::numeric_limits<uint32_t>::max();
 
-inline rocblaslt_status validateWorkspaceSize(const char* func, size_t workspaceBytes)
+inline size_t clampWorkspaceSize(const char* func, size_t workspaceBytes)
 {
     if(workspaceBytes <= rocblaslt_max_workspace_bytes)
-        return rocblaslt_status_success;
-    log_error(func,
+        return workspaceBytes;
+    log_hints(func,
               "workspace size",
               workspaceBytes,
-              "exceeds the maximum of",
+              "clamped to the maximum of",
               rocblaslt_max_workspace_bytes,
               "bytes");
-    return rocblaslt_status_invalid_value;
+    return rocblaslt_max_workspace_bytes;
 }
 
 inline bool isValidOrderForDatatype(hipDataType datatype, hipblasLtOrder_t order)

@@ -2124,10 +2124,7 @@ rocblaslt_status
                     pref->search_mode);
             break;
         case ROCBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES:
-            if(auto status = validateWorkspaceSize(__func__, *(uint64_t*)data);
-               status != rocblaslt_status_success)
-                return status;
-            pref->max_workspace_bytes = *(uint64_t*)data;
+            pref->max_workspace_bytes = clampWorkspaceSize(__func__, *(uint64_t*)data);
             log_api(__func__,
                     "matmulPref",
                     pref,
@@ -2655,9 +2652,7 @@ rocblaslt_status
         log_error(__func__, "invalid requested count", requestedAlgoCount);
         return rocblaslt_status_invalid_value;
     }
-    if(auto status = validateWorkspaceSize(__func__, maxWorkspaceBytes);
-       status != rocblaslt_status_success)
-        return status;
+    const size_t workspaceLimit = clampWorkspaceSize(__func__, maxWorkspaceBytes);
     if(gemmType == rocblaslt::RocGemmType::ROCBLASLT_GROUPED_GEMM)
     {
         log_api(
@@ -2685,7 +2680,7 @@ rocblaslt_status
         if(override.env_mode)
         {
             override_success = problem_override_from_file_cpp(
-                handle, gemmType, gemmData, override_result, override.file_path, maxWorkspaceBytes);
+                handle, gemmType, gemmData, override_result, override.file_path, workspaceLimit);
 
             log_api(__func__, "OverrideAlgoCount", override_success ? 1 : 0);
         }
@@ -2695,7 +2690,7 @@ rocblaslt_status
                 = getBestSolutions(handle,
                                    gemmType,
                                    gemmData,
-                                   maxWorkspaceBytes,
+                                   workspaceLimit,
                                    override_success ? requestedAlgoCount - 1 : requestedAlgoCount,
                                    results);
 
@@ -2723,8 +2718,7 @@ rocblaslt_status
             std::vector<rocblaslt_matmul_heuristic_result> allSolutionsResults;
             size_t                                         workspaceSizeInBytes = 0;
             if(rocblaslt_status_success
-               == getAllSolutions(
-                   gemmData, handle, gemmType, allSolutionsResults, maxWorkspaceBytes))
+               == getAllSolutions(gemmData, handle, gemmType, allSolutionsResults, workspaceLimit))
             {
                 status = rocblaslt_status_success;
                 int oriReturnAlgoCount = results.size();

@@ -829,9 +829,7 @@ rocblaslt_status rocblaslt_matmul(rocblaslt_handle             handle,
         log_error(__func__, "invalid workspace pointer");
         return rocblaslt_status_invalid_pointer;
     }
-    if(auto status = validateWorkspaceSize(__func__, workspaceSizeInBytes);
-       status != rocblaslt_status_success)
-        return status;
+    workspaceSizeInBytes = clampWorkspaceSize(__func__, workspaceSizeInBytes);
 
     if(matC->type != matD->type)
     {
@@ -1617,15 +1615,12 @@ rocblaslt_status rocblaslt_makeArgument_cpp(rocblaslt_handle              handle
                                             hipStream_t                   stream,
                                             std::shared_ptr<void>         gemmData)
 {
-    if(auto status = validateWorkspaceSize(__func__, workspaceSizeInBytes);
-       status != rocblaslt_status_success)
-        return status;
     return makeArgument(handle,
                         gemmType,
                         algo,
                         tuning,
                         workspace,
-                        workspaceSizeInBytes,
+                        clampWorkspaceSize(__func__, workspaceSizeInBytes),
                         useUserArgs,
                         stream,
                         gemmData);