Commit 2e325dcd5 for llama.cpp
commit 2e325dcd57cafe31a9deb660133ec71cafebb844
Author: Tarek Dakhran <tarek@liquid.ai>
Date: Wed Nov 11 13:00:45 2026 +0100
backend : re-instantiate CUDA graphs when the cluster dims of a kernel change (#31284)
* CUDA: re-instantiate graphs when the cluster dims of a kernel change
* CUDA: keep the cluster dims with the graph instance, check them only on Hopper+
Assisted-by: Claude Opus 5.5
* CUDA: use CUDA_CHECK and a helper struct for the graph cluster dims
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
* CUDA: add a TODO for the graph cluster dims test
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 17661c19f..aa51b72bf 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -3239,7 +3239,54 @@ static bool ggml_cuda_graph_update_required(ggml_backend_cuda_context * cuda_ctx
return res;
}
-static void ggml_cuda_graph_update_executable(cudaGraphExec_t & instance, cudaGraph_t cuda_graph) {
+#ifdef USE_CUDA_GRAPH_CLUSTER_CHECK
+static std::vector<ggml_cuda_graph_cluster> ggml_cuda_graph_get_clusters(cudaGraph_t cuda_graph, int device) {
+ std::vector<ggml_cuda_graph_cluster> clusters;
+ static const bool runtime_affected = [] {
+ int version = 0;
+ CUDA_CHECK(cudaRuntimeGetVersion(&version));
+ // TODO: upper bound once fixed in the CUDA runtime
+ return version >= 11080;
+ }();
+ if (!runtime_affected || ggml_cuda_info().devices[device].cc < GGML_CUDA_CC_HOPPER) {
+ return clusters;
+ }
+ size_t n_nodes = 0;
+ CUDA_CHECK(cudaGraphGetNodes(cuda_graph, nullptr, &n_nodes));
+ std::vector<cudaGraphNode_t> nodes(n_nodes);
+ CUDA_CHECK(cudaGraphGetNodes(cuda_graph, nodes.data(), &n_nodes));
+ for (size_t i = 0; i < n_nodes; ++i) {
+ cudaGraphNodeType type;
+ CUDA_CHECK(cudaGraphNodeGetType(nodes[i], &type));
+ if (type != cudaGraphNodeTypeKernel) {
+ continue;
+ }
+ cudaLaunchAttributeValue val = {};
+ CUDA_CHECK(cudaGraphKernelNodeGetAttribute(nodes[i], cudaLaunchAttributeClusterDimension, &val));
+ if (val.clusterDim.x * val.clusterDim.y * val.clusterDim.z > 1) {
+ clusters.push_back({ i, val.clusterDim.x, val.clusterDim.y, val.clusterDim.z });
+ }
+ }
+ return clusters;
+}
+#endif // USE_CUDA_GRAPH_CLUSTER_CHECK
+
+static void ggml_cuda_graph_update_executable(ggml_cuda_graph_instance & instance, cudaGraph_t cuda_graph, int device) {
+#ifdef USE_CUDA_GRAPH_CLUSTER_CHECK
+ std::vector<ggml_cuda_graph_cluster> clusters = ggml_cuda_graph_get_clusters(cuda_graph, device);
+ if (instance.clusters != clusters) {
+ if (ggml_cuda_graph_debug_enabled()) {
+ GGML_LOG_DEBUG("%s: cluster dims changed, re-instantiating CUDA graph\n", __func__);
+ }
+ CUDA_CHECK(cudaGraphExecDestroy(instance));
+ CUDA_CHECK(cudaGraphInstantiate(&instance.handle, cuda_graph, nullptr, nullptr, 0));
+ instance.clusters = std::move(clusters);
+ return;
+ }
+#else
+ GGML_UNUSED(device);
+#endif // USE_CUDA_GRAPH_CLUSTER_CHECK
+
#if CUDART_VERSION >= 12000
cudaGraphExecUpdateResultInfo result_info;
cudaError_t stat = cudaGraphExecUpdate(instance, cuda_graph, &result_info);
@@ -3258,8 +3305,8 @@ static void ggml_cuda_graph_update_executable(cudaGraphExec_t & instance, cudaGr
// so instead clear error and re-instantiate
(void)cudaGetLastError();
CUDA_CHECK(cudaGraphExecDestroy(instance));
- instance = nullptr;
- CUDA_CHECK(cudaGraphInstantiate(&instance, cuda_graph, nullptr, nullptr, 0));
+ instance.handle = nullptr;
+ CUDA_CHECK(cudaGraphInstantiate(&instance.handle, cuda_graph, nullptr, nullptr, 0));
} else {
GGML_ASSERT(stat == cudaSuccess);
}
@@ -3280,11 +3327,14 @@ static void ggml_cuda_graph_launch_chunk(ggml_backend_cuda_context * cuda_ctx, g
CUDA_CHECK(cudaStreamEndCapture(cuda_ctx->stream(), &cuda_graph));
if (idx < graph->instances.size()) {
- ggml_cuda_graph_update_executable(graph->instances[idx], cuda_graph);
+ ggml_cuda_graph_update_executable(graph->instances[idx], cuda_graph, cuda_ctx->device);
} else {
GGML_ASSERT(idx == graph->instances.size());
- graph->instances.emplace_back(nullptr);
- CUDA_CHECK(cudaGraphInstantiate(&graph->instances[idx], cuda_graph, nullptr, nullptr, 0));
+ graph->instances.emplace_back();
+ CUDA_CHECK(cudaGraphInstantiate(&graph->instances[idx].handle, cuda_graph, nullptr, nullptr, 0));
+#ifdef USE_CUDA_GRAPH_CLUSTER_CHECK
+ graph->instances[idx].clusters = ggml_cuda_graph_get_clusters(cuda_graph, cuda_ctx->device);
+#endif // USE_CUDA_GRAPH_CLUSTER_CHECK
}
CUDA_CHECK(cudaGraphDestroy(cuda_graph));
diff --git a/ggml/src/ggml-cuda/graph.cuh b/ggml/src/ggml-cuda/graph.cuh
index 3481de072..c26373700 100644
--- a/ggml/src/ggml-cuda/graph.cuh
+++ b/ggml/src/ggml-cuda/graph.cuh
@@ -9,6 +9,36 @@
#include <utility>
#include <vector>
+#ifdef USE_CUDA_GRAPH
+// cudaGraphExecUpdate keeps stale kernel cluster dims (cuBLAS sm90), re-instantiate when they change
+// TODO: add a test, fails on H100 (driver 590.48.01, CUDA runtime 12.9.79, cuBLAS 12.9.1.4) for F16 MUL_MAT with
+// ne01 1200 <-> 1040 (k=768, n=1024) between graph updates, see https://github.com/ggml-org/llama.cpp/pull/31284
+#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) && CUDART_VERSION >= 11080
+#define USE_CUDA_GRAPH_CLUSTER_CHECK
+#endif
+
+#ifdef USE_CUDA_GRAPH_CLUSTER_CHECK
+// cluster dims of a clustered kernel node
+struct ggml_cuda_graph_cluster {
+ size_t node;
+ uint32_t x, y, z;
+
+ bool operator==(const ggml_cuda_graph_cluster & other) const {
+ return node == other.node && x == other.x && y == other.y && z == other.z;
+ }
+};
+#endif
+
+struct ggml_cuda_graph_instance {
+ cudaGraphExec_t handle = nullptr;
+#ifdef USE_CUDA_GRAPH_CLUSTER_CHECK
+ std::vector<ggml_cuda_graph_cluster> clusters;
+#endif
+
+ operator cudaGraphExec_t() const { return handle; }
+};
+#endif // USE_CUDA_GRAPH
+
struct ggml_cuda_graph {
#ifdef USE_CUDA_GRAPH
~ggml_cuda_graph() {
@@ -17,7 +47,7 @@ struct ggml_cuda_graph {
}
}
// the ggml graph is captured as a series of CUDA graphs, so the GPU can run one chunk while the next is enqueued
- std::vector<cudaGraphExec_t> instances;
+ std::vector<ggml_cuda_graph_instance> instances;
size_t num_nodes = 0;
std::vector<cudaGraphNode_t> nodes;
bool disable_due_to_gpu_arch = false;