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