From 83f230ebbcde16303e6d432896b258e154976b03 Mon Sep 17 00:00:00 2001 From: Marcos Damasceno Date: Mon, 28 Sep 2026 05:42:06 -0500 Subject: [PATCH 1/2] perf(cuda): the CUDA graphs a context keeps are capped by their nodes, not their count A CUDA graph's executable holds device memory for each node of its graph, from a pool of the driver's that never shrinks. On an RTX 5080 (driver 610.43), executables of 1 to 2,048 kernel nodes took 1.9 to 3.0 KiB a node, at 16 B to 2 KiB of kernel parameters. Destroying them gave nothing back to cudaMemGetInfo, and instantiating them again took nothing more. Under a cap on the nodes held, executables of 20 to 4,000 nodes were instantiated and evicted 3,500 times over, and the pool stayed where the most held at once had put it: 14 MiB at 4,096 nodes, 44 at 16,384, 88 at 32,768. d0f8bae41 capped their count at 8. A -sm tensor split computes one graph per all-reduce step on each device, two a layer every token, and each is its own CUDA graph. Under a cap of 8 each was evicted before its turn came round again, then warmed up and captured anew every token. qwen3-0.6b Q4_0 (llama-bench tg128, -ts 1/1, RTX 5080 + 5070 Ti) decoded at 257.0 tok/s with -sm tensor against 558.0 uncapped. GGML_CUDA_GRAPH_NODES (default 32768, 0 for no cap) now bounds the nodes a context's executables hold: the least recently used are evicted before a new one is instantiated. GGML_CUDA_GRAPH_MAX still bounds their count, but only when set (default 0, no cap). The nodes of a captured graph are counted (cudaGraphGetNodes). The log gives the nodes held beside the count, and the device memory all the executables took; it logs at INFO when that memory grows or the nodes held reach a power of two, and at DEBUG otherwise. A -sm tensor context holds hundreds of graphs, each a new most. Decode graphs on the 5080: qwen3-0.6b's is 538 nodes, and Ternary Bonsai 2 27B's two are 1,192 and 1,096 (8 MiB in all), so the default holds 27 of the 27B's. test-cuda-graph-cap-nodes sets a cap of 100 nodes, with graphs of eight nodes (x halved eight times) that are built once and computed again. - Twelve computed in turn, five times round: each is captured once, then replayed from the third time round (captures 0 12 0 0 0). All twelve are held, 96 nodes. - Sixteen in turn: at most 96 nodes are held. It fails under a cap of 8 graphs (captures 0 12 0 4 4, 8 held), when the nodes are not counted (0 held), and when they are not capped (16 graphs, 128 nodes). test-cuda-graph-cap still passes under a cap of 4 graphs, and so do test-cuda-graph-key and test-cuda-graph-src-type. --- ggml/src/ggml-cuda/common.cuh | 56 +++++++--- ggml/src/ggml-cuda/ggml-cuda.cu | 40 +++++-- tests/CMakeLists.txt | 1 + tests/test-cuda-graph-cap.cpp | 178 +++++++++++++++++++++++++++----- 4 files changed, 221 insertions(+), 54 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index cb70a62408b0..f2713cae10dd 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1300,7 +1300,7 @@ struct ggml_cuda_graph { } cudaGraph_t graph = nullptr; cudaGraphExec_t instance = nullptr; - size_t num_nodes = 0; + size_t num_nodes = 0; // the captured graph's: its executable holds device memory for each (cuda_graph_make_room) std::vector nodes; bool disable_due_to_gpu_arch = false; bool warmup_complete = false; @@ -1746,38 +1746,60 @@ struct ggml_backend_cuda_context { return it->second.get(); } - // The most CUDA graph executables a context keeps (GGML_CUDA_GRAPH_MAX, 0 for no cap). Each holds device memory - // beside the buffers the scheduler sizes, and the graphs a context computes follow its traffic (one per graph shape, - // and per graph slot for the same shape), so without a count they would be bounded only by the 10 s eviction above: - // a new one evicts the least recently used. + // The CUDA graph executables a context keeps. Each holds device memory beside the buffers the scheduler sizes, a few + // KiB for each node of its graph (2.75 on an RTX 5080), from a pool of the driver's that never shrinks: destroying + // one returns nothing to the device, a later one reuses it, and the pool stays at the most nodes held at once. The + // graphs a context computes follow its traffic (one per graph shape, and per graph slot for the same shape), bounded + // otherwise only by the 10 s eviction above, so a new one evicts the least recently used others until the nodes held + // fit under GGML_CUDA_GRAPH_NODES (default 32768, 88 MiB on that card; 0 for no cap) and the executables under + // GGML_CUDA_GRAPH_MAX (default 0, no cap). The memory follows the nodes, not the count: a -sm tensor split computes a + // small graph for each all-reduce step on each device, hundreds a token, and a cap on the count below that evicts + // every one of them before its turn comes round again, to be captured anew every token. static size_t cuda_graphs_max() { static const size_t n = [] { const char * env = getenv("GGML_CUDA_GRAPH_MAX"); - return env != nullptr ? (size_t) std::max(0, atoi(env)) : (size_t) 8; + return env != nullptr ? (size_t) std::max(0, atoi(env)) : (size_t) 0; }(); return n; } - size_t cuda_graphs_held_max = 0; // the most executables held at once here - size_t cuda_graph_bytes_max = 0; // the most device memory one took to instantiate here (cudaMemGetInfo around it) + static size_t cuda_graph_nodes_max() { + static const size_t n = [] { + const char * env = getenv("GGML_CUDA_GRAPH_NODES"); + return env != nullptr ? (size_t) std::max(0, atoi(env)) : (size_t) 32768; + }(); + return n; + } + size_t cuda_graphs_held_max = 0; // the most executables held at once here + size_t cuda_graph_nodes_held_max = 0; // the most nodes they held at once here + size_t cuda_graph_bytes_max = 0; // the most device memory one took to instantiate here (cudaMemGetInfo around it) + size_t cuda_graph_bytes = 0; // and all of them + + struct cuda_graphs_held { + size_t graphs; + size_t nodes; + }; - // before `graph` gets an executable: evict the least recently used others until it fits under the cap; returns how - // many are held with it - size_t cuda_graph_make_room(const ggml_cuda_graph * graph) { - const size_t max = cuda_graphs_max(); + // before `graph` gets an executable: evict the least recently used others until it fits under the caps (alone, it is + // kept whatever its nodes); returns what is held with it + cuda_graphs_held cuda_graph_make_room(const ggml_cuda_graph * graph) { + const size_t max = cuda_graphs_max(); + const size_t max_nodes = cuda_graph_nodes_max(); for (;;) { - size_t held = 0; - auto lru = cuda_graphs.end(); + cuda_graphs_held held = { 1, graph->num_nodes }; + auto lru = cuda_graphs.end(); for (auto it = cuda_graphs.begin(); it != cuda_graphs.end(); ++it) { if (it->second->instance == nullptr || it->second.get() == graph) { continue; } - ++held; + held.graphs++; + held.nodes += it->second->num_nodes; if (lru == cuda_graphs.end() || it->second->last_used_time < lru->second->last_used_time) { lru = it; } } - if (max == 0 || held < max) { - return held + 1; + const bool fits = (max == 0 || held.graphs <= max) && (max_nodes == 0 || held.nodes <= max_nodes); + if (fits || lru == cuda_graphs.end()) { + return held; } cuda_graphs.erase(lru); } diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 139d13c9d66c..da94cb9ff7c9 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -2986,10 +2986,12 @@ static bool ggml_cuda_graph_update_required(ggml_backend_cuda_context * cuda_ctx return res; } -// create `graph`'s executable from its captured graph, first evicting what the cap needs (cuda_graphs_max); the device -// memory that takes is measured, and logged whenever it or the count held is the most on this context so far +// create `graph`'s executable from its captured graph, first evicting what the caps need (cuda_graph_make_room); the +// device memory that takes is measured (what the driver's pool grew by, if it did), and logged whenever what is held or +// the most one took is the most on this context so far: at INFO when that is the most memory, or the nodes held reach a +// power of two (a -sm tensor split's context holds hundreds of small graphs), at DEBUG otherwise static void ggml_cuda_graph_instantiate(ggml_backend_cuda_context * cuda_ctx, ggml_cuda_graph * graph) { - const size_t held = cuda_ctx->cuda_graph_make_room(graph); + const ggml_backend_cuda_context::cuda_graphs_held held = cuda_ctx->cuda_graph_make_room(graph); size_t free_before; size_t free_after; @@ -2998,13 +3000,29 @@ static void ggml_cuda_graph_instantiate(ggml_backend_cuda_context * cuda_ctx, gg CUDA_CHECK(cudaGraphInstantiate(&graph->instance, graph->graph, NULL, NULL, 0)); CUDA_CHECK(cudaMemGetInfo(&free_after, &total)); const size_t bytes = free_before > free_after ? free_before - free_after : 0; - - if (held > cuda_ctx->cuda_graphs_held_max || bytes > cuda_ctx->cuda_graph_bytes_max) { - cuda_ctx->cuda_graphs_held_max = std::max(cuda_ctx->cuda_graphs_held_max, held); - cuda_ctx->cuda_graph_bytes_max = std::max(cuda_ctx->cuda_graph_bytes_max, bytes); - GGML_LOG_INFO("%s: %s (context %p): %zu CUDA graphs held at most (cap %zu), the largest took %.2f MiB to instantiate\n", + cuda_ctx->cuda_graph_bytes += bytes; + + if (held.graphs > cuda_ctx->cuda_graphs_held_max || held.nodes > cuda_ctx->cuda_graph_nodes_held_max || + bytes > cuda_ctx->cuda_graph_bytes_max) { + const auto pow2_floor = [](size_t n) { + size_t p = 0; + for (size_t q = 1; q != 0 && q <= n; q <<= 1) { + p = q; + } + return p; + }; + const bool info = bytes > cuda_ctx->cuda_graph_bytes_max || + pow2_floor(held.nodes) > pow2_floor(cuda_ctx->cuda_graph_nodes_held_max); + cuda_ctx->cuda_graphs_held_max = std::max(cuda_ctx->cuda_graphs_held_max, held.graphs); + cuda_ctx->cuda_graph_nodes_held_max = std::max(cuda_ctx->cuda_graph_nodes_held_max, held.nodes); + cuda_ctx->cuda_graph_bytes_max = std::max(cuda_ctx->cuda_graph_bytes_max, bytes); + ggml_log_internal(info ? GGML_LOG_LEVEL_INFO : GGML_LOG_LEVEL_DEBUG, + "%s: %s (context %p): %zu CUDA graphs held at most, %zu nodes (caps %zu and %zu, 0 for none); " + "instantiating one took %.2f MiB at most, all of them %.2f MiB\n", __func__, cuda_ctx->name.c_str(), (void *) cuda_ctx, cuda_ctx->cuda_graphs_held_max, - ggml_backend_cuda_context::cuda_graphs_max(), cuda_ctx->cuda_graph_bytes_max / 1024.0 / 1024.0); + cuda_ctx->cuda_graph_nodes_held_max, ggml_backend_cuda_context::cuda_graphs_max(), + ggml_backend_cuda_context::cuda_graph_nodes_max(), cuda_ctx->cuda_graph_bytes_max / 1024.0 / 1024.0, + cuda_ctx->cuda_graph_bytes / 1024.0 / 1024.0); } } @@ -5509,6 +5527,7 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } CUDA_CHECK(cudaStreamEndCapture(cuda_ctx->stream(), &graph->graph)); + CUDA_CHECK(cudaGraphGetNodes(graph->graph, nullptr, &graph->num_nodes)); graph_evaluated_or_captured = true; // CUDA graph has been captured std::lock_guard lock(ggml_cuda_lock); @@ -6906,7 +6925,8 @@ ggml_backend_t ggml_backend_cuda_init(int device) { #ifdef USE_CUDA_GRAPH static std::once_flag cuda_graphs_max_logged; std::call_once(cuda_graphs_max_logged, [] { - GGML_LOG_INFO("%s: at most %zu CUDA graphs kept per context (GGML_CUDA_GRAPH_MAX; 0 = no cap)\n", __func__, + GGML_LOG_INFO("%s: CUDA graphs kept per context: at most %zu nodes (GGML_CUDA_GRAPH_NODES) and %zu graphs " + "(GGML_CUDA_GRAPH_MAX); 0 = no cap\n", __func__, ggml_backend_cuda_context::cuda_graph_nodes_max(), ggml_backend_cuda_context::cuda_graphs_max()); }); #endif // USE_CUDA_GRAPH diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 04ba9db4ff16..3c9d25514b0f 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -334,6 +334,7 @@ if (NOT GGML_BACKEND_DL) llama_build_and_test(test-cuda-graph-key.cpp) llama_build_and_test(test-cuda-graph-src-type.cpp) llama_build_and_test(test-cuda-graph-cap.cpp) + llama_test(test-cuda-graph-cap NAME test-cuda-graph-cap-nodes ARGS nodes) llama_build_and_test(test-pq2-mma-device-memory.cpp) llama_build_and_test(test-ptq1_0-element-map.cpp) llama_build_and_test(test-ptq1_0-cuda-dot.cpp) diff --git a/tests/test-cuda-graph-cap.cpp b/tests/test-cuda-graph-cap.cpp index 613654f77587..817f6decc5ee 100644 --- a/tests/test-cuda-graph-cap.cpp +++ b/tests/test-cuda-graph-cap.cpp @@ -1,15 +1,25 @@ // The CUDA backend keeps a CUDA graph per graph shape (test-cuda-graph-key), and each captured graph's executable holds -// device memory beside the buffers a scheduler sizes. The shapes a context computes follow its traffic, one per graph -// shape and per graph slot for the same shape, so the executables were bounded only by the 10 s eviction: under four -// concurrent requests a 27B hybrid's contexts held 84 MiB of them on a 5080. A context keeps at most -// GGML_CUDA_GRAPH_MAX of them, the least recently used evicted for a new one. +// device memory beside the buffers a scheduler sizes, a few KiB for each node of the graph. The shapes a context +// computes follow its traffic, one per graph shape and per graph slot for the same shape, so the executables were +// bounded only by the 10 s eviction: under four concurrent requests a 27B hybrid's contexts held 84 MiB of them on a +// 5080. A context keeps at most GGML_CUDA_GRAPH_NODES nodes in them and at most GGML_CUDA_GRAPH_MAX of them, the least +// recently used evicted for a new one. // -// Here the cap is 4, and y = w x is computed for ten row counts of x, three times each: every shape is captured on its -// second compute and replayed on its third, every output matching the CPU backend. The backend's log reports the most -// executables it held at once, which must be the cap: never over it, and reached, since ten shapes were captured. The -// first shape, evicted by then, is computed again and must be captured again and give the right values. A control -// first computes one shape three times: a device that does not capture it does not use CUDA graphs and is skipped. -// Without the eviction all ten are held and it fails. +// With no argument the cap is 4 graphs, and y = w x is computed for ten row counts of x, three times each: every shape +// is captured on its second compute and replayed on its third, every output matching the CPU backend. The backend's log +// reports the most executables it held at once, which must be the cap: never over it, and reached, since ten shapes +// were captured. The first shape, evicted by then, is computed again and must be captured again and give the right +// values. Without the eviction all ten are held and it fails. +// +// With `nodes` the cap is 100 nodes and none on the count, and each graph halves x eight times, eight nodes, built once +// and computed again and again, as a -sm tensor split computes its graph for each all-reduce step every token. Twelve +// row counts in turn, one compute each, five times round: all twelve are held (96 nodes, more graphs than the old cap +// of 8), so each is captured once, the second time round, and replayed from the third. Then sixteen in turn (128 +// nodes): the nodes held reach the cap and never pass it. Under a cap of 8 graphs the twelve are evicted before their +// turn comes round again and captured anew every time; without the cap on the nodes all sixteen are held. +// +// A control first computes one shape three times: a device that does not capture it does not use CUDA graphs and is +// skipped. #include "ggml.h" #include "ggml-alloc.h" @@ -24,11 +34,14 @@ #include #include -static constexpr int cap = 4; +static constexpr int cap = 4; // graphs, with no argument +static constexpr int cap_nodes = 100; // nodes, with `nodes` +static constexpr int n_halvings = 8; // the nodes of each graph with `nodes` struct graph_log { - int captured = 0; - size_t held_max = 0; // "N CUDA graphs held at most", the most the backend reported + int captured = 0; + size_t held_max = 0; // "N CUDA graphs held at most", the most the backend reported + size_t nodes_max = 0; // "..., K nodes" }; static void log_callback(ggml_log_level level, const char * text, void * user_data) { @@ -37,7 +50,7 @@ static void log_callback(ggml_log_level level, const char * text, void * user_da log->captured++; } if (const char * held = strstr(text, "CUDA graphs held at most")) { - // the count is the number right before it: "...: CUDA0: 4 CUDA graphs held at most (cap 4), ..." + // the count is the number right before it: "...: CUDA0 (context 0x...): 4 CUDA graphs held at most, 12 nodes ..." const char * p = held; while (p > text && p[-1] == ' ') { --p; @@ -46,6 +59,10 @@ static void log_callback(ggml_log_level level, const char * text, void * user_da --p; } log->held_max = std::max(log->held_max, (size_t) strtoull(p, nullptr, 10)); + const char * nodes = "CUDA graphs held at most, "; + if (strncmp(held, nodes, strlen(nodes)) == 0) { + log->nodes_max = std::max(log->nodes_max, (size_t) strtoull(held + strlen(nodes), nullptr, 10)); + } } if (level != GGML_LOG_LEVEL_DEBUG) { fputs(text, stderr); @@ -160,14 +177,97 @@ static bool compute_shapes(ggml_backend_dev_t dev, ggml_backend_t cpu, graph_log return ok; } -int main(void) { - // the cap is read once, when the first CUDA backend is made - const std::string cap_env = std::to_string(cap); +// halves x n_halvings times on one backend for each row count, each graph built once in its own context and buffer and +// computed again, as a -sm tensor split keeps its graphs; `rounds` times round in turn, one compute each, every output +// checked (halving is exact). Returns the captures each time round; sets ok to false if an output is wrong. +static std::vector compute_rounds(ggml_backend_dev_t dev, graph_log & log, const std::vector & rows, + int rounds, bool & ok) { + const int64_t k = 256; + std::mt19937 rng(42); + log = graph_log(); + ggml_backend_t backend = ggml_backend_dev_init(dev, nullptr); + + struct shape { + ggml_context * ctx; + ggml_backend_buffer_t buf; + ggml_tensor * x; + ggml_tensor * y; + ggml_cgraph * gf; + }; + std::vector shapes; + for (const int64_t m : rows) { + ggml_init_params params = { ggml_tensor_overhead() * (1 + n_halvings) + ggml_graph_overhead(), nullptr, true }; + shape s; + s.ctx = ggml_init(params); + s.x = ggml_new_tensor_2d(s.ctx, GGML_TYPE_F32, k, m); + s.y = s.x; + for (int i = 0; i < n_halvings; ++i) { + s.y = ggml_scale(s.ctx, s.y, 0.5f); + } + s.gf = ggml_new_graph(s.ctx); + ggml_build_forward_expand(s.gf, s.y); + s.buf = ggml_backend_alloc_ctx_tensors(s.ctx, backend); + shapes.push_back(s); + } + + std::vector captured(rounds, 0); + for (int r = 0; r < rounds; ++r) { + const int captured_before = log.captured; + for (const shape & s : shapes) { + const std::vector xv = random_values(rng, ggml_nelements(s.x)); + ggml_backend_tensor_set(s.x, xv.data(), 0, xv.size() * sizeof(float)); + GGML_ASSERT(ggml_backend_graph_compute(backend, s.gf) == GGML_STATUS_SUCCESS); + + std::vector yv(xv.size()); + ggml_backend_tensor_get(s.y, yv.data(), 0, yv.size() * sizeof(float)); + for (size_t i = 0; i < yv.size(); ++i) { + if (yv[i] != xv[i] / (1 << n_halvings)) { + printf(" %lld rows, time %d round: element %zu is %g, not %g: FAILED\n", (long long) s.x->ne[1], + r + 1, i, yv[i], xv[i] / (1 << n_halvings)); + ok = false; + break; + } + } + } + captured[r] = log.captured - captured_before; + } + + for (const shape & s : shapes) { + ggml_backend_buffer_free(s.buf); + ggml_free(s.ctx); + } + ggml_backend_free(backend); + return captured; +} + +static void set_env(const char * name, const char * value) { #ifdef _WIN32 - _putenv_s("GGML_CUDA_GRAPH_MAX", cap_env.c_str()); + _putenv_s(name, value != nullptr ? value : ""); #else - setenv("GGML_CUDA_GRAPH_MAX", cap_env.c_str(), 1); + if (value != nullptr) { + setenv(name, value, 1); + } else { + unsetenv(name); + } #endif +} + +static std::string join(const std::vector & v) { + std::string s; + for (const int i : v) { + s += (s.empty() ? "" : " ") + std::to_string(i); + } + return s; +} + +int main(int argc, char ** argv) { + const bool nodes = argc > 1 && strcmp(argv[1], "nodes") == 0; + + // the caps are read once, when the first CUDA backend is made + const std::string cap_env = std::to_string(cap); + const std::string cap_nodes_env = std::to_string(cap_nodes); + set_env("GGML_CUDA_GRAPH_MAX", nodes ? nullptr : cap_env.c_str()); + set_env("GGML_CUDA_GRAPH_NODES", nodes ? cap_nodes_env.c_str() : nullptr); ggml_backend_load_all(); @@ -197,14 +297,38 @@ int main(void) { continue; } - // ten shapes, then the first again: evicted by the fifth, it is captured anew - printf(" ten shapes, then the first again, under a cap of %d:\n", cap); - ok &= compute_shapes(dev, cpu, log, { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 1 }, 3, supported); - const bool all_captured = log.captured == 11; - const bool at_cap = log.held_max == (size_t) cap; - ok &= all_captured && at_cap; - printf(" captures %d (11 expected), held at most %zu (the cap, %d, expected): %s\n", log.captured, log.held_max, - cap, all_captured && at_cap ? "ok" : "FAILED"); + if (!nodes) { + // ten shapes, then the first again: evicted by the fifth, it is captured anew + printf(" ten shapes, then the first again, under a cap of %d:\n", cap); + ok &= compute_shapes(dev, cpu, log, { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 1 }, 3, supported); + const bool all_captured = log.captured == 11; + const bool at_cap = log.held_max == (size_t) cap; + ok &= all_captured && at_cap; + printf(" captures %d (11 expected), held at most %zu (the cap, %d, expected): %s\n", log.captured, + log.held_max, cap, all_captured && at_cap ? "ok" : "FAILED"); + } else { + // twelve shapes in turn fit under the cap: each captured once, then replayed every time round + printf(" twelve shapes in turn, five times round, under a cap of %d nodes:\n", cap_nodes); + bool values = true; + const std::vector each = compute_rounds(dev, log, { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12 }, 5, values); + const bool replayed = each == std::vector{ 0, 12, 0, 0, 0 }; + const bool all_held = log.held_max == 12 && log.nodes_max == 12 * n_halvings; + ok &= values && replayed && all_held; + printf(" captures each time round %s (0 12 0 0 0 expected), held at most %zu graphs of %zu nodes (12 of " + "%d expected), outputs %s: %s\n", join(each).c_str(), log.held_max, log.nodes_max, + 12 * n_halvings, values ? "exact" : "WRONG", values && replayed && all_held ? "ok" : "FAILED"); + + // sixteen do not: the nodes held reach the cap and never pass it + printf(" sixteen shapes in turn, four times round, under a cap of %d nodes:\n", cap_nodes); + values = true; + const std::vector over = compute_rounds(dev, log, + { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16 }, 4, values); + const bool under_cap = log.nodes_max <= (size_t) cap_nodes && log.nodes_max > (size_t) (cap_nodes - n_halvings); + ok &= values && under_cap; + printf(" captures each time round %s, held at most %zu graphs of %zu nodes (at most %d, more than %d " + "expected), outputs %s: %s\n", join(over).c_str(), log.held_max, log.nodes_max, cap_nodes, + cap_nodes - n_halvings, values ? "exact" : "WRONG", values && under_cap ? "ok" : "FAILED"); + } n_tested++; } ggml_backend_free(cpu); From f021f83431cadca80cd86664dd11bf05d9fdffce Mon Sep 17 00:00:00 2001 From: Marcos Damasceno Date: Mon, 28 Sep 2026 05:55:14 -0500 Subject: [PATCH 2/2] docs(torad): the row for the CUDA graphs capped by their nodes (83f230ebb) --- TORAD.md | 1 + 1 file changed, 1 insertion(+) diff --git a/TORAD.md b/TORAD.md index 70ff7de97216..a9aff1aee714 100644 --- a/TORAD.md +++ b/TORAD.md @@ -106,6 +106,7 @@ which pins a commit of this branch as a submodule. | `78e0fadf7` | `topk_radix` also takes a row of up to 1,024 columns at any k: with CUB's top-k available such a row had no path but CUB one row after another, four launches a row (a DSA indexer's ubatch under 4,096 cells: 2,048 launches a layer). RTX 5080: 3.3-3.9 us against 8.2-12.7 (1 row) and 129-171 (16 rows). | `GGML_CUDA_TOPK_RADIX_LEGACY` | | `640a12c41` | A SwiGLU limit (GLM-5.3's routed experts, shared experts and dense FFN: the gate clamped to `[-inf, 10]`, the up projection to `[-10, 10]`, then the GLU) fuses into the quantized mat-vec's GLU epilogue: `ggml_cuda_can_fuse` takes `{MUL_MAT(_ID), CLAMP, MUL_MAT(_ID), CLAMP, GLU}` when the clamps are exactly a limit and the GLU a SwiGLU, and `mul_mat_vec_q` applies it as one float, where the two nodes had kept the gate/up fusion from matching (five launches for one). In place (GLM-5.3-Flash's routed FFN, 8 layers of 32 experts, RTX 5080): 807 against 861 us at 1 token; a verify's 3 tokens are unchanged, their experts unfused. | `GGML_CUDA_GLU_LIMIT_FUSE_LEGACY` | | `b8d44d4b8` | Routed experts at 1-8 tokens stream through `mmvq-moe.cu`: one block an SM, a producer warp listing the distinct experts once (every token/slot pair that routes to one reads it once) and streaming tiles of their rows through a ring of shared-memory slots by `cp.async.bulk`, two teams of eight warps on the dot products; an IQ3_XXS warp frees its slot before the math, and the last tiles go by tickets on the stream's tile counter. It takes IQ3_XXS, and Q8_0 past 1 token: in place (GLM-5.3-Flash's routed FFN on 32 experts, CUDA graphs and PDL, RTX 5080) IQ3_XXS 815 against 818 us for 8 layers at 1 token and 1,823 against 1,953 at 3, Q8_0 2,265 against 2,382 for 4 layers at 3; IQ4_XS measured slower and keeps `mul_mat_vec_q`. Its steady state runs at DRAM's ceiling (~913 GB/s); a launch's first q8_1 vector, loaded from global memory behind the stream, costs 3-6 us. | `GGML_CUDA_MMVQ_MOE_LEGACY` | +| `83f230ebb` | The CUDA graphs a context keeps are capped by their nodes: `GGML_CUDA_GRAPH_NODES` (default 32768) bounds the nodes their executables hold, evicting the least recently used, and `GGML_CUDA_GRAPH_MAX` bounds their count only when set. An executable holds a few KiB of device memory a node (1.9-3.0 on an RTX 5080), from a driver pool that never shrinks and stays at the most nodes held at once: 14, 44 and 88 MiB at 4,096, 16,384 and 32,768 nodes, under 3,500 evictions of mixed sizes. d0f8bae41's cap of 8 graphs broke `-sm tensor`: it computes a graph per all-reduce step on each device, two a layer every token, and each was evicted before its turn came round again, captured anew every token (qwen3-0.6b Q4_0 on a 5080 + 5070 Ti, 257.0 tok/s against 558.0 uncapped). `test-cuda-graph-cap-nodes`: twelve graphs of 8 nodes computed in turn under a cap of 100 nodes, each captured once and then replayed (a cap of 8 graphs captures 4 anew every time round). | `GGML_CUDA_GRAPH_NODES=0 GGML_CUDA_GRAPH_MAX=8` | Every switch in the last column is read once per process and parses as an integer: a `*_LEGACY` switch set to `0` is the same as unset (the change stays on), and `=0` turns off `GGML_CUDA_LORA_RANK1_FUSE` and