fix(cuda): set explicit cuBLAS/hipBLAS workspace to stop VRAM creep

Addresses upstream llama.cpp issue ggml-org#19979: on ROCm/HIP, VRAM creeps up
during long prompt processing when using non-f16 KV cache types (q4_0,
q5_1, q8_0, turbo2/3/4), eventually causing OOM. Does not occur with
f16/f16 KV cache.

Root cause: cublasCreate() is called but cublasSetWorkspace() is not,
so rocBLAS lazily allocates its own internal workspace on first GEMM
and grows it monotonically as it encounters new Tensile kernel shapes
across the varying batch sizes / KV cache lengths of prompt processing.
The growth is invisible to llama.cpp's memory accounting (hipMemGetInfo
snapshots it but the fit algorithm runs before the workspace exists).

Fix: after cublasCreate(), allocate a fixed 32 MiB buffer via cudaMalloc
and install it as the explicit workspace via cublasSetWorkspace (aliased
to hipblasSetWorkspace on HIP). rocBLAS then uses our workspace instead
of growing its own, and the workspace is freed deterministically in the
backend context destructor.

Added:
 - GGML_CUDA_CUBLAS_WORKSPACE_SIZE constant (32 MiB) in common.cuh
 - cublas_workspaces[] array in ggml_backend_cuda_context
 - cublasSetWorkspace -> hipblasSetWorkspace alias in vendors/hip.h
 - Workspace alloc + install in cublas_handle()
 - Workspace free in ~ggml_backend_cuda_context() (after cublasDestroy)

Upstream reference:
  ggml-org#19979

The 32 MiB size matches NVIDIA cuBLAS recommendations for typical
workloads. Bump if you see rocBLAS falling back to smaller kernels.

Affects all ROCm GPU archs, especially GFX11 (RDNA3/3.5) where Tensile
solution selection triggers more workspace growth.

Author: terrysimons

ggml/src/ggml-cuda/common.cuh

@@ -1354,13 +1354,26 @@ struct ggml_cuda_stream_context {
     }
 };
 
+// Fixed cuBLAS / hipBLAS workspace size. Setting this explicitly via
+// cublasSetWorkspace() on handle creation prevents the underlying BLAS
+// library (particularly rocBLAS on GFX11/RDNA3/3.5) from lazily growing
+// its own untracked workspace as it encounters new GEMM shapes during
+// long prompt processing. Without this, VRAM creeps up until OOM.
+//
+// 32 MiB is large enough for the Tensile kernels rocBLAS ships and matches
+// the NVIDIA cuBLAS recommended default for most workloads. Bump if you
+// see rocBLAS falling back to smaller kernels in traces.
+// Upstream bug: https://github.com/ggml-org/llama.cpp/issues/19979
+#define GGML_CUDA_CUBLAS_WORKSPACE_SIZE (size_t(32) * 1024 * 1024)
+
 struct ggml_backend_cuda_context {
     int device;
     std::string name;
     cudaEvent_t copy_event = nullptr;
 
     cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } };
     cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES] = {nullptr};
+    void * cublas_workspaces[GGML_CUDA_MAX_DEVICES] = {nullptr};
 
     int curr_stream_no = 0;
 
@@ -1426,6 +1439,16 @@ struct ggml_backend_cuda_context {
             ggml_cuda_set_device(device);
             CUBLAS_CHECK(cublasCreate(&cublas_handles[device]));
             CUBLAS_CHECK(cublasSetMathMode(cublas_handles[device], CUBLAS_TF32_TENSOR_OP_MATH));
+
+            // Allocate and install a fixed cuBLAS/hipBLAS workspace.
+            // Without this call, rocBLAS lazily allocates its own workspace
+            // and grows it as it encounters new GEMM shapes, producing
+            // untracked VRAM creep on GFX11/RDNA3/3.5. See the workspace
+            // size constant definition above and upstream issue #19979.
+            CUDA_CHECK(cudaMalloc(&cublas_workspaces[device], GGML_CUDA_CUBLAS_WORKSPACE_SIZE));
+            CUBLAS_CHECK(cublasSetWorkspace(cublas_handles[device],
+                                            cublas_workspaces[device],
+                                            GGML_CUDA_CUBLAS_WORKSPACE_SIZE));
         }
         return cublas_handles[device];
     }

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -577,6 +577,14 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() {
         }
         if (cublas_handles[i] != nullptr) {
             CUBLAS_CHECK(cublasDestroy(cublas_handles[i]));
+            cublas_handles[i] = nullptr;
+        }
+        // Free the fixed cuBLAS/hipBLAS workspace installed in cublas_handle()
+        // above. Must happen after cublasDestroy so the BLAS library doesn't
+        // touch the buffer during its own cleanup.
+        if (cublas_workspaces[i] != nullptr) {
+            CUDA_CHECK(cudaFree(cublas_workspaces[i]));
+            cublas_workspaces[i] = nullptr;
         }
     }
 }

ggml/src/ggml-cuda/vendors/hip.h

@@ -62,6 +62,12 @@
 #define cublasGemmStridedBatchedEx hipblasGemmStridedBatchedEx
 #define cublasHandle_t hipblasHandle_t
 #define cublasSetMathMode(handle, mode) CUBLAS_STATUS_SUCCESS
+// hipblasSetWorkspace is available in hipBLAS 2.0+ (ROCm 6.0+). We call this
+// on handle creation with a fixed-size buffer to prevent rocBLAS from lazily
+// growing its own untracked workspace across varying GEMM shapes during long
+// prompt processing. Upstream issue:
+//     https://github.com/ggml-org/llama.cpp/issues/19979
+#define cublasSetWorkspace hipblasSetWorkspace
 #define cublasSetStream hipblasSetStream
 #define cublasSgemm hipblasSgemm
 #define cublasStatus_t hipblasStatus_t