From c5351c0faa0f8f2520c428d54d3d2c7a8c503e44 Mon Sep 17 00:00:00 2001 From: aendk Date: Tue, 4 Aug 2026 11:14:51 +0200 Subject: [PATCH 01/11] Adds initial scheduler stress tests --- tests/CMakeLists.txt | 1 + tests/test-backend-sched.cpp | 362 +++++++++++++++++++++++++++++++++++ 2 files changed, 363 insertions(+) create mode 100644 tests/test-backend-sched.cpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 08c6f5a47963..de3c879ac4af 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -277,6 +277,7 @@ if (NOT LLAMA_SANITIZE_ADDRESS AND NOT GGML_SCHED_NO_REALLOC) llama_build_and_test(test-opt.cpp) endif() llama_build_and_test(test-backend-ops.cpp) +llama_build_and_test(test-backend-sched.cpp) llama_build_and_test(test-model-load-cancel.cpp LABEL "model") llama_build_and_test(test-autorelease.cpp LABEL "model") diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp new file mode 100644 index 000000000000..00b81d4eaf7d --- /dev/null +++ b/tests/test-backend-sched.cpp @@ -0,0 +1,362 @@ +// Stress test for ggml_backend_sched: graphs built only from increments, placed on alternating backends. +// Since every node adds one, the expected result of a counter is just the number of times it was incremented, +// which makes the correct output independent of how the scheduler splits and copies the graph. + +#include "ggml.h" +#include "ggml-alloc.h" +#include "ggml-backend.h" + +#include +#include +#include +#include + +static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t ne) { + ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + + const size_t graph_size = n_nodes + 16; + + ggml_backend_sched_t sched = ggml_backend_sched_new(backends, nullptr, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); + + // the inputs are allocated separately so that they can be written before the graph is computed + ggml_init_params params_static = { + /*.mem_size =*/ 2*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_static = ggml_init(params_static); + + ggml_tensor * x = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, ne); + ggml_set_name(x, "x"); + ggml_set_input(x); + + ggml_tensor * one = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, ne); + ggml_set_name(one, "one"); + ggml_set_input(one); + + ggml_backend_buffer_t buf_static = ggml_backend_alloc_ctx_tensors(ctx_static, backend_cpu); + + std::vector data(ne); + + std::fill(data.begin(), data.end(), 0.0f); + ggml_backend_tensor_set(x, data.data(), 0, ggml_nbytes(x)); + + std::fill(data.begin(), data.end(), 1.0f); + ggml_backend_tensor_set(one, data.data(), 0, ggml_nbytes(one)); + + ggml_init_params params_compute = { + /*.mem_size =*/ (n_nodes + 2)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); + + std::vector nodes; + + ggml_tensor * out = x; + for (int i = 0; i < n_nodes; i++) { + out = ggml_add(ctx_compute, out, one); + nodes.push_back(out); + } + ggml_set_output(out); + ggml_build_forward_expand(gf, out); + + // the assignments are only kept until the next reset, and ggml_backend_sched_graph_compute + // resets the scheduler unless the graph has already been allocated + ggml_backend_sched_reset(sched); + for (int i = 0; i < n_nodes; i++) { + ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[i % 2]); + } + + bool ok = true; + + if (!ggml_backend_sched_alloc_graph(sched, gf)) { + printf("\n failed to allocate the graph\n"); + ok = false; + } + + if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { + printf("\n failed to compute the graph\n"); + ok = false; + } + + if (ok) { + const int n_splits = ggml_backend_sched_get_n_splits(sched); + if (n_splits != n_nodes) { + printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_nodes); + ok = false; + } + } + + if (ok) { + ggml_backend_tensor_get(out, data.data(), 0, ggml_nbytes(out)); + for (int64_t i = 0; i < ne; i++) { + if (data[i] != float(n_nodes)) { + printf("\n out[%" PRId64 "] = %f, expected %d\n", i, data[i], n_nodes); + ok = false; + break; + } + } + } + + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + ggml_backend_buffer_free(buf_static); + ggml_free(ctx_static); + + return ok; +} + +// Same idea as test_ping_pong, but with 8 independent counters and a graph shape that covers +// three cases the plain chain does not: +// - splits with no inputs at all, from lanes restarted out of constants that already live on the split's backend +// - lanes with different histories, so a misrouted copy shows up as a wrong count in a single lane +// - one split producing two values followed by two splits that each consume one of them and do not depend on each other +static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_rounds, int64_t ne) { + + const int GPU = 0; + const int CPU = 1; + + const int n_lanes = 8; + + ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + + // the graph needs 2*n_lanes nodes per round plus one restart for half of the lanes, but the hash set of + // the scheduler is sized from the same value and GGML_SCHED_DEBUG inserts every copy tensor into it as + // well, so leave room for those + const size_t n_nodes = 2*n_lanes*n_rounds + n_lanes/2; + const size_t graph_size = 3*n_nodes + 64; + + ggml_backend_sched_t sched = ggml_backend_sched_new(backends, nullptr, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); + + ggml_init_params params_static = { + /*.mem_size =*/ 2*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + + // the constants are kept on both backends so that a lane can be restarted without any cross-backend copy + ggml_context * ctx_cpu = ggml_init(params_static); + + ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_set_name(zero_cpu, "zero_cpu"); + + ggml_tensor * one_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_set_name(one_cpu, "one_cpu"); + + ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, backend_cpu); + + ggml_context * ctx_gpu = ggml_init(params_static); + + ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_set_name(zero_gpu, "zero_gpu"); + + ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_set_name(one_gpu, "one_gpu"); + + ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, backend_gpu); + + std::vector data(ne); + + std::fill(data.begin(), data.end(), 0.0f); + ggml_backend_tensor_set(zero_cpu, data.data(), 0, ggml_nbytes(zero_cpu)); + ggml_backend_tensor_set(zero_gpu, data.data(), 0, ggml_nbytes(zero_gpu)); + + std::fill(data.begin(), data.end(), 1.0f); + ggml_backend_tensor_set(one_cpu, data.data(), 0, ggml_nbytes(one_cpu)); + ggml_backend_tensor_set(one_gpu, data.data(), 0, ggml_nbytes(one_gpu)); + + ggml_init_params params_compute = { + /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); + + ggml_tensor * lane[n_lanes]; + int expected[n_lanes]; + for (int l = 0; l < n_lanes; l++) { + lane[l] = zero_cpu; + expected[l] = 0; + } + + std::vector nodes; + std::vector node_backend; + + // expanding each node as it is created keeps the graph order equal to the creation order, + // which is what determines where the scheduler puts the split boundaries + auto append = [&](int l, ggml_tensor * node, int b) { + lane[l] = node; + nodes.push_back(node); + node_backend.push_back(b); + ggml_build_forward_expand(gf, node); + }; + + auto increment = [&](int l, int b) { + append(l, ggml_add(ctx_compute, lane[l], b == GPU ? one_gpu : one_cpu), b); + expected[l]++; + }; + + // both sources already live on backend b, so the resulting split has no inputs to copy + auto restart = [&](int l, int b) { + append(l, ggml_add(ctx_compute, b == GPU ? zero_gpu : zero_cpu, b == GPU ? one_gpu : one_cpu), b); + expected[l] = 1; + }; + + for (int round = 0; round < n_rounds; round++) { + // halfway through, the upper half of the lanes starts over from scratch + if (round == n_rounds/2) { + // start on the backend the previous node did not use, so that no restart is absorbed into the split before it + int b = node_backend.empty() || node_backend.back() == CPU ? GPU : CPU; + for (int l = n_lanes/2; l < n_lanes; l++) { + restart(l, b); + b = b == CPU ? GPU : CPU; + } + } + + for (int p = 0; p < n_lanes/2; p++) { + const int la = 2*p; + const int lb = 2*p + 1; + + const int prod = p % 2 == 0 ? CPU : GPU; + + // one split producing two values ... + increment(la, prod); + increment(lb, prod); + + // ... and two splits that each consume one of them, with no dependency on each other + increment(la, prod == CPU ? GPU : CPU); + increment(lb, prod); + } + } + + for (int l = 0; l < n_lanes; l++) { + ggml_set_output(lane[l]); + } + + // the scheduler starts a new split on every backend change, so the assignments below fix the split count + int n_splits_expected = nodes.empty() ? 0 : 1; + for (size_t i = 1; i < node_backend.size(); i++) { + n_splits_expected += node_backend[i] != node_backend[i - 1]; + } + + ggml_backend_sched_reset(sched); + for (size_t i = 0; i < nodes.size(); i++) { + ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[node_backend[i]]); + } + + bool ok = true; + + if (!ggml_backend_sched_alloc_graph(sched, gf)) { + printf("\n failed to allocate the graph\n"); + ok = false; + } + + if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { + printf("\n failed to compute the graph\n"); + ok = false; + } + + if (ok) { + const int n_splits = ggml_backend_sched_get_n_splits(sched); + if (n_splits != n_splits_expected) { + printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_splits_expected); + ok = false; + } + } + + if (ok) { + nvtx3::scoped_range sc_4{nvtx3::event_attributes{nvtx3::rgb{250, 128, 114}, "readback and verify"}}; // salmon + + for (int l = 0; l < n_lanes && ok; l++) { + ggml_backend_tensor_get(lane[l], data.data(), 0, ggml_nbytes(lane[l])); + for (int64_t i = 0; i < ne; i++) { + if (data[i] != float(expected[l])) { + printf("\n lane %d [%" PRId64 "] = %f, expected %d\n", l, i, data[i], expected[l]); + ok = false; + break; + } + } + } + } + + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + ggml_backend_buffer_free(buf_gpu); + ggml_free(ctx_gpu); + ggml_backend_buffer_free(buf_cpu); + ggml_free(ctx_cpu); + + return ok; +} + +int main() { + ggml_backend_load_all(); + + ggml_backend_dev_t dev_gpu = nullptr; + for (size_t i = 0; i < ggml_backend_dev_count(); i++) { + ggml_backend_dev_t dev = ggml_backend_dev_get(i); + const enum ggml_backend_dev_type type = ggml_backend_dev_type(dev); + if (type == GGML_BACKEND_DEVICE_TYPE_GPU || type == GGML_BACKEND_DEVICE_TYPE_IGPU) { + dev_gpu = dev; + break; + } + } + + if (dev_gpu == nullptr) { + printf("no GPU device found, skipping\n"); + return 0; + } + + ggml_backend_t backend_gpu = ggml_backend_dev_init(dev_gpu, nullptr); + ggml_backend_t backend_cpu = ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr); + GGML_ASSERT(backend_gpu != nullptr); + GGML_ASSERT(backend_cpu != nullptr); + + printf("GPU: %s (%s)\n", ggml_backend_name(backend_gpu), ggml_backend_dev_description(dev_gpu)); + printf("CPU: %s\n\n", ggml_backend_name(backend_cpu)); + + int n_ok = 0; + int n_test = 0; + + for (int n_nodes : { 1, 2, 3, 8, 33, 128, 1024 }) { + for (int ne : { 1, 4096 }) { + printf(" n_nodes = %4d, ne = %4d: ", n_nodes, ne); + fflush(stdout); + + const bool ok = test_ping_pong(backend_gpu, backend_cpu, n_nodes, ne); + + printf("%s\n", ok ? "OK" : "FAIL"); + + n_ok += ok; + n_test++; + } + } + + printf("\n"); + + for (int n_rounds : { 1, 2, 3, 8, 33, 64 }) { + for (int ne : { 1, 4096 }) { + printf(" n_rounds = %4d, ne = %4d: ", n_rounds, ne); + fflush(stdout); + + const bool ok = test_multi_lane(backend_gpu, backend_cpu, n_rounds, ne); + + printf("%s\n", ok ? "OK" : "FAIL"); + + n_ok += ok; + n_test++; + } + } + + ggml_backend_free(backend_gpu); + ggml_backend_free(backend_cpu); + + printf("\n%d/%d tests passed\n", n_ok, n_test); + + return n_ok == n_test ? 0 : 1; +} From 19f50ffe8111134ff39dbf67fbe048f529237c92 Mon Sep 17 00:00:00 2001 From: aendk Date: Wed, 5 Aug 2026 10:36:42 +0200 Subject: [PATCH 02/11] Adds new sleep op to simplify scheduler testing --- ggml/include/ggml.h | 8 +++++++ ggml/src/ggml-cpu/ggml-cpu.c | 5 ++++ ggml/src/ggml-cpu/ops.cpp | 22 ++++++++++++++++++ ggml/src/ggml-cpu/ops.h | 1 + ggml/src/ggml-cuda/ggml-cuda.cu | 10 ++++++++ ggml/src/ggml-cuda/sleep.cu | 41 +++++++++++++++++++++++++++++++++ ggml/src/ggml-cuda/sleep.cuh | 3 +++ ggml/src/ggml.c | 28 ++++++++++++++++++++-- tests/CMakeLists.txt | 1 + tests/test-backend-ops.cpp | 31 +++++++++++++++++++++++++ 10 files changed, 148 insertions(+), 2 deletions(-) create mode 100644 ggml/src/ggml-cuda/sleep.cu create mode 100644 ggml/src/ggml-cuda/sleep.cuh diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index c2ccd9725396..009d77ee1170 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -590,6 +590,8 @@ extern "C" { GGML_OP_GLU, + GGML_OP_SLEEP, + GGML_OP_COUNT, }; @@ -2722,6 +2724,12 @@ extern "C" { int n_tasks, void * userdata); + // Busy-wait operator for scheduling and backend testing. + GGML_API struct ggml_tensor * ggml_sleep( + struct ggml_context * ctx, + struct ggml_tensor * a, + int32_t us); + // loss function GGML_API struct ggml_tensor * ggml_cross_entropy_loss( diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 87ac0a702efc..80330c5d4330 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -1997,6 +1997,10 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_fill(params, tensor); } break; + case GGML_OP_SLEEP: + { + ggml_compute_forward_sleep(params, tensor); + } break; case GGML_OP_FLASH_ATTN_EXT: { ggml_compute_forward_flash_attn_ext(params, tensor); @@ -2265,6 +2269,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_REPEAT: case GGML_OP_REPEAT_BACK: case GGML_OP_LEAKY_RELU: + case GGML_OP_SLEEP: { n_tasks = 1; } break; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 001e1ae85a89..9b00b4a56f4d 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -10,6 +10,7 @@ #include #include +#include #include // ggml_compute_forward_dup @@ -2288,6 +2289,27 @@ void ggml_compute_forward_fill(const ggml_compute_params * params, ggml_tensor * } } +// ggml_compute_forward_sleep + +void ggml_compute_forward_sleep(const ggml_compute_params * params, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + + if (params->ith != 0) { + return; + } + + GGML_ASSERT(src0->type == dst->type); + GGML_ASSERT(ggml_are_same_shape(src0, dst)); + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + const auto t_end = std::chrono::steady_clock::now() + std::chrono::microseconds(ggml_get_op_params_i32(dst, 0)); + + memcpy(dst->data, src0->data, ggml_nbytes(dst)); + + while (std::chrono::steady_clock::now() < t_end) {} +} + // ggml_compute_tri static void ggml_compute_forward_tri_f32(const ggml_compute_params * params, ggml_tensor * dst) { diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h index 4c1642a67603..bdd438e90867 100644 --- a/ggml/src/ggml-cpu/ops.h +++ b/ggml/src/ggml-cpu/ops.h @@ -87,6 +87,7 @@ void ggml_compute_forward_top_k(const struct ggml_compute_params * params, struc void ggml_compute_forward_leaky_relu(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_tri(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_fill(const struct ggml_compute_params * params, struct ggml_tensor * dst); +void ggml_compute_forward_sleep(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_flash_attn_ext(const struct ggml_compute_params * params, struct ggml_tensor * dst); void ggml_compute_forward_flash_attn_back( const struct ggml_compute_params * params, diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 598f3228c401..9e90c2a7427f 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -42,6 +42,7 @@ #include "ggml-cuda/rope.cuh" #include "ggml-cuda/roll.cuh" #include "ggml-cuda/scale.cuh" +#include "ggml-cuda/sleep.cuh" #include "ggml-cuda/snake.cuh" #include "ggml-cuda/softcap.cuh" #include "ggml-cuda/softmax.cuh" @@ -2392,6 +2393,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_FILL: ggml_cuda_op_fill(ctx, dst); break; + case GGML_OP_SLEEP: + ggml_cuda_op_sleep(ctx, dst); + break; case GGML_OP_LIGHTNING_INDEXER: ggml_cuda_lightning_indexer(ctx, dst); break; @@ -5297,6 +5301,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g return true; case GGML_OP_LIGHTNING_INDEXER: return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op); + case GGML_OP_SLEEP: +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + return false; +#else + return op->type == op->src[0]->type && ggml_is_contiguous(op->src[0]); +#endif default: return false; diff --git a/ggml/src/ggml-cuda/sleep.cu b/ggml/src/ggml-cuda/sleep.cu new file mode 100644 index 000000000000..00baa07c110f --- /dev/null +++ b/ggml/src/ggml-cuda/sleep.cu @@ -0,0 +1,41 @@ +#include "sleep.cuh" + +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + +// %globaltimer is a nanosecond wall clock, unlike clock64() it is unaffected by the SM clock and by frequency scaling +static __device__ __forceinline__ uint64_t globaltimer_ns() { + uint64_t t; + asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(t)); + return t; +} + +// a single thread is enough, the following memcpy on the same stream cannot start before this kernel retires +static __global__ void sleep_ns(const uint64_t ns) { + const uint64_t t0 = globaltimer_ns(); + + while (globaltimer_ns() - t0 < ns) {} +} + +void ggml_cuda_op_sleep(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + + GGML_ASSERT(src0->type == dst->type); + GGML_ASSERT(ggml_are_same_shape(src0, dst)); + GGML_ASSERT(ggml_is_contiguous(src0)); + GGML_ASSERT(ggml_is_contiguous(dst)); + + cudaStream_t stream = ctx.stream(); + + sleep_ns<<<1, 1, 0, stream>>>(1000*(uint64_t) ggml_get_op_params_i32(dst, 0)); + CUDA_CHECK(cudaMemcpyAsync(dst->data, src0->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, stream)); +} + +#else + +void ggml_cuda_op_sleep(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + GGML_UNUSED(ctx); + GGML_UNUSED(dst); + GGML_ABORT("GGML_OP_SLEEP requires the %%globaltimer register, which is only available on CUDA"); +} + +#endif diff --git a/ggml/src/ggml-cuda/sleep.cuh b/ggml/src/ggml-cuda/sleep.cuh new file mode 100644 index 000000000000..d62c3178565c --- /dev/null +++ b/ggml/src/ggml-cuda/sleep.cuh @@ -0,0 +1,3 @@ +#include "common.cuh" + +void ggml_cuda_op_sleep(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index d0d369c41737..683fefc7b7ef 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1098,9 +1098,11 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "OPT_STEP_SGD", "GLU", + + "SLEEP", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1213,9 +1215,11 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "sgd(x)", "glu(x)", + + "sleep(x)", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -6135,6 +6139,26 @@ struct ggml_tensor * ggml_custom_inplace( return result; } + +// ggml_sleep + +struct ggml_tensor * ggml_sleep( + struct ggml_context * ctx, + struct ggml_tensor * a, + int32_t us) { + GGML_ASSERT(us >= 0); + GGML_ASSERT(ggml_is_contiguous(a)); + + struct ggml_tensor * result = ggml_dup_tensor(ctx, a); + + ggml_set_op_params_i32(result, 0, us); + + result->op = GGML_OP_SLEEP; + result->src[0] = a; + + return result; +} + // ggml_cross_entropy_loss struct ggml_tensor * ggml_cross_entropy_loss( diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index de3c879ac4af..9105f39d96d9 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -277,6 +277,7 @@ if (NOT LLAMA_SANITIZE_ADDRESS AND NOT GGML_SCHED_NO_REALLOC) llama_build_and_test(test-opt.cpp) endif() llama_build_and_test(test-backend-ops.cpp) +llama_build_and_test(test-sleep.cpp) llama_build_and_test(test-backend-sched.cpp) llama_build_and_test(test-model-load-cancel.cpp LABEL "model") diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 3349a64b174d..795fe46c3ae6 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -7451,6 +7451,33 @@ struct test_fill : public test_case { } }; +// GGML_OP_SLEEP (for scheduler testing only) +struct test_sleep : public test_case { + const ggml_type type; + const std::array ne; + const int32_t us; + + std::string vars() override { return VARS_TO_STR3(type, ne, us); } + + // the op is a plain copy, any deviation at all is a bug + double max_nmse_err() override { return 0.0; } + + test_sleep(int32_t us = 100, ggml_type type = GGML_TYPE_F32, + std::array ne = { 10, 10, 4, 3 }) + : type(type), ne(ne), us(us) {} + + ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * a = ggml_new_tensor_4d(ctx, type, ne[0], ne[1], ne[2], ne[3]); + ggml_set_name(a, "a"); + + ggml_tensor * out = ggml_sleep(ctx, a, us); + + ggml_set_name(out, "out"); + + return out; + } +}; + // GGML_OP_SOLVE_TRI struct test_solve_tri : public test_case { const ggml_type type; @@ -9759,6 +9786,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_fill(-152.0f, GGML_TYPE_F32, { 800, 600, 4, 4 })); test_cases.emplace_back(new test_fill(3.5f, GGML_TYPE_F32, { 2048, 512, 2, 2 })); + test_cases.emplace_back(new test_sleep(0)); + test_cases.emplace_back(new test_sleep(10, GGML_TYPE_F16, { 128, 4, 2, 2 })); + test_cases.emplace_back(new test_sleep(10, GGML_TYPE_F32, { 128, 4, 2, 2 })); + test_cases.emplace_back(new test_diag()); test_cases.emplace_back(new test_diag(GGML_TYPE_F32, { 79, 1, 19, 13 })); test_cases.emplace_back(new test_diag(GGML_TYPE_F32, { 256, 1, 8, 16 })); From 31c5a7f1dbe91894717bdddceeee380a030606b0 Mon Sep 17 00:00:00 2001 From: aendk Date: Wed, 5 Aug 2026 15:57:43 +0200 Subject: [PATCH 03/11] Adds testcase addressing https://github.com/ggml-org/llama.cpp/issues/23321 --- tests/test-backend-sched.cpp | 305 ++++++++++++++++++++++++++++++++--- 1 file changed, 279 insertions(+), 26 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index 00b81d4eaf7d..a421b76321e5 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -9,14 +9,36 @@ #include #include #include +#include #include -static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t ne) { + +static ggml_backend_buffer_type_t sched_cpu_buft(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, bool use_device_host_buft) { + if (use_device_host_buft) { + ggml_backend_buffer_type_t host = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend_gpu)); + if (host != nullptr) { + return host; + } + } + return ggml_backend_get_default_buffer_type(backend_cpu); +} + +static ggml_backend_sched_t create_test_scheduler(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, size_t graph_size, bool use_device_host_buft) { + ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + ggml_backend_buffer_type_t bufts[2] = { + ggml_backend_get_default_buffer_type(backend_gpu), + sched_cpu_buft(backend_gpu, backend_cpu, use_device_host_buft), + }; + return ggml_backend_sched_new(backends, bufts, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); +} + +static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t ne, bool use_device_host_buft) { + ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; const size_t graph_size = n_nodes + 16; - ggml_backend_sched_t sched = ggml_backend_sched_new(backends, nullptr, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); + ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); // the inputs are allocated separately so that they can be written before the graph is computed ggml_init_params params_static = { @@ -114,7 +136,7 @@ static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cp // - splits with no inputs at all, from lanes restarted out of constants that already live on the split's backend // - lanes with different histories, so a misrouted copy shows up as a wrong count in a single lane // - one split producing two values followed by two splits that each consume one of them and do not depend on each other -static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_rounds, int64_t ne) { +static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_rounds, int64_t ne, bool use_device_host_buft) { const int GPU = 0; const int CPU = 1; @@ -129,7 +151,7 @@ static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_c const size_t n_nodes = 2*n_lanes*n_rounds + n_lanes/2; const size_t graph_size = 3*n_nodes + 64; - ggml_backend_sched_t sched = ggml_backend_sched_new(backends, nullptr, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); + ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); ggml_init_params params_static = { /*.mem_size =*/ 2*ggml_tensor_overhead(), @@ -270,8 +292,6 @@ static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_c } if (ok) { - nvtx3::scoped_range sc_4{nvtx3::event_attributes{nvtx3::rgb{250, 128, 114}, "readback and verify"}}; // salmon - for (int l = 0; l < n_lanes && ok; l++) { ggml_backend_tensor_get(lane[l], data.data(), 0, ggml_nbytes(lane[l])); for (int64_t i = 0; i < ne; i++) { @@ -294,6 +314,206 @@ static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_c return ok; } +// Created in response to https://github.com/ggml-org/llama.cpp/issues/23321 +// Currently, input copies to splits synchronize split execution. +// CPU split N+1 could overwrite output of split N if this was not copied to backend in time. +// Mechanism here: GPU split with long sleep to clog CUDA stream / vk command queue, +// followed by CPU splits without input, provoking a race condition in the current setup. +// GPU: 100ms sleep to clog stream/queue +// CPU: produce {55} +// GPU: request copy of {55} in CUDA stream, ADD {44}; Output should be {99} +// CPU: produce {66} +// GPU: increment both {99} and {66} -> {100} and {67} +// Correct result is thus {100}, incorrect output is {111} when {55} was overwritten by {66} before copy started. +// Note: currently only reproducible on async H2D copy (pinned memory in CUDA). +static bool test_sleep_race(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int64_t ne, int32_t sleep_us, bool use_device_host_buft) { + + const int GPU = 0; + const int CPU = 1; + + ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + + const size_t graph_size = 64; + + ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); + + ggml_init_params params_static = { + /*.mem_size =*/ 4*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + + ggml_context * ctx_cpu = ggml_init(params_static); + + ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_set_name(zero_cpu, "zero_cpu"); + + ggml_tensor * val55_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_set_name(val55_cpu, "val55_cpu"); + + ggml_tensor * val66_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_set_name(val66_cpu, "val66_cpu"); + + ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, backend_cpu); + + ggml_context * ctx_gpu = ggml_init(params_static); + + ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_set_name(zero_gpu, "zero_gpu"); + + ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_set_name(one_gpu, "one_gpu"); + + ggml_tensor * val44_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_set_name(val44_gpu, "val44_gpu"); + + ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, backend_gpu); + + std::vector data(ne); + + std::fill(data.begin(), data.end(), 0.0f); + ggml_backend_tensor_set(zero_cpu, data.data(), 0, ggml_nbytes(zero_cpu)); + ggml_backend_tensor_set(zero_gpu, data.data(), 0, ggml_nbytes(zero_gpu)); + + std::fill(data.begin(), data.end(), 1.0f); + ggml_backend_tensor_set(one_gpu, data.data(), 0, ggml_nbytes(one_gpu)); + + std::fill(data.begin(), data.end(), 44.0f); + ggml_backend_tensor_set(val44_gpu, data.data(), 0, ggml_nbytes(val44_gpu)); + + std::fill(data.begin(), data.end(), 55.0f); + ggml_backend_tensor_set(val55_cpu, data.data(), 0, ggml_nbytes(val55_cpu)); + + std::fill(data.begin(), data.end(), 66.0f); + ggml_backend_tensor_set(val66_cpu, data.data(), 0, ggml_nbytes(val66_cpu)); + + ggml_init_params params_compute = { + /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); + + std::vector nodes; + std::vector node_backend; + + auto append = [&](ggml_tensor * node, int b) { + nodes.push_back(node); + node_backend.push_back(b); + ggml_build_forward_expand(gf, node); + }; + + // GPU: add then sleep - occupies the GPU while later splits run + ggml_tensor * delayed = ggml_add(ctx_compute, zero_gpu, one_gpu); + append(delayed, GPU); + delayed = ggml_sleep(ctx_compute, delayed, sleep_us); + append(delayed, GPU); + + // CPU: 0 + 55 -> 55 + ggml_tensor * v55 = ggml_add(ctx_compute, zero_cpu, val55_cpu); + append(v55, CPU); + + // GPU: 55 + 44 -> 99 + ggml_tensor * v99 = ggml_add(ctx_compute, v55, val44_gpu); + append(v99, GPU); + + // CPU: 0 + 66 -> 66 + ggml_tensor * v66 = ggml_add(ctx_compute, zero_cpu, val66_cpu); + append(v66, CPU); + + // GPU: increment both of the previous split outputs + ggml_tensor * out99 = ggml_add(ctx_compute, v99, one_gpu); + append(out99, GPU); + ggml_tensor * out66 = ggml_add(ctx_compute, v66, one_gpu); + append(out66, GPU); + + ggml_set_output(delayed); + ggml_set_output(out99); + ggml_set_output(out66); + + const int n_splits_expected = 5; + + ggml_backend_sched_reset(sched); + for (size_t i = 0; i < nodes.size(); i++) { + ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[node_backend[i]]); + } + + bool ok = true; + + if (!ggml_backend_supports_op(backend_gpu, delayed)) { + printf("\n GGML_OP_SLEEP not supported on GPU, skipping\n"); + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + ggml_backend_buffer_free(buf_gpu); + ggml_free(ctx_gpu); + ggml_backend_buffer_free(buf_cpu); + ggml_free(ctx_cpu); + return true; + } + + if (!ggml_backend_sched_alloc_graph(sched, gf)) { + printf("\n failed to allocate the graph\n"); + ok = false; + } + + // checks if the the allocator reuses the same memory for output tensors between splits. + if (ok && v55->data == v66->data) { + printf("\n NOTE: v55 and v66 alias after alloc (v55=%p v66=%p)\n", v55->data, v66->data); + } else { + printf("\n NOTE: v55 and v66 do not alias after alloc (v55=%p v66=%p)\n", v55->data, v66->data); + } + + if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { + printf("\n failed to compute the graph\n"); + ok = false; + } + + if (ok) { + const int n_splits = ggml_backend_sched_get_n_splits(sched); + if (n_splits != n_splits_expected) { + printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_splits_expected); + ok = false; + } + } + + if (ok) { + const struct { + ggml_tensor * tensor; + float expected; + const char * name; + } checks[] = { + { delayed, 1.0f, "delayed" }, + { out99, 100.0f, "out99" }, + { out66, 67.0f, "out66" }, + }; + + for (const auto & c : checks) { + ggml_backend_tensor_get(c.tensor, data.data(), 0, ggml_nbytes(c.tensor)); + for (int64_t i = 0; i < ne; i++) { + if (data[i] != c.expected) { + printf("\n %s[%" PRId64 "] = %f, expected %f\n", c.name, i, data[i], c.expected); + ok = false; + break; + } + } + if (!ok) { + break; + } + } + } + + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + ggml_backend_buffer_free(buf_gpu); + ggml_free(ctx_gpu); + ggml_backend_buffer_free(buf_cpu); + ggml_free(ctx_cpu); + + return ok; +} + int main() { ggml_backend_load_all(); @@ -318,45 +538,78 @@ int main() { GGML_ASSERT(backend_cpu != nullptr); printf("GPU: %s (%s)\n", ggml_backend_name(backend_gpu), ggml_backend_dev_description(dev_gpu)); - printf("CPU: %s\n\n", ggml_backend_name(backend_cpu)); + printf("CPU: %s\n", ggml_backend_name(backend_cpu)); + + const bool have_device_host_buft = ggml_backend_dev_host_buffer_type(dev_gpu) != nullptr; + if (!have_device_host_buft) { + printf("GPU has no host buffer type; device_host cases will be skipped\n"); + } + printf("\n"); int n_ok = 0; int n_test = 0; - for (int n_nodes : { 1, 2, 3, 8, 33, 128, 1024 }) { - for (int ne : { 1, 4096 }) { - printf(" n_nodes = %4d, ne = %4d: ", n_nodes, ne); - fflush(stdout); + for (bool use_device_host_buft : { false, true }) { + if (use_device_host_buft && !have_device_host_buft) { + continue; + } + + const char * buft = use_device_host_buft ? "device_host" : "pageable"; + printf("=== CPU sched buft: %s ===\n\n", buft); - const bool ok = test_ping_pong(backend_gpu, backend_cpu, n_nodes, ne); + for (int n_nodes : { 1, 2, 3, 8, 33, 128, 1024 }) { + for (int ne : { 1, 4096 }) { + printf(" n_nodes = %4d, ne = %4d: ", n_nodes, ne); + fflush(stdout); - printf("%s\n", ok ? "OK" : "FAIL"); + const bool ok = test_ping_pong(backend_gpu, backend_cpu, n_nodes, ne, use_device_host_buft); - n_ok += ok; - n_test++; + printf("%s\n", ok ? "OK" : "FAIL"); + + n_ok += ok; + n_test++; + } } - } - printf("\n"); + printf("\n"); - for (int n_rounds : { 1, 2, 3, 8, 33, 64 }) { - for (int ne : { 1, 4096 }) { - printf(" n_rounds = %4d, ne = %4d: ", n_rounds, ne); - fflush(stdout); + for (int n_rounds : { 1, 2, 3, 8, 33, 64 }) { + for (int ne : { 1, 4096 }) { + printf(" n_rounds = %4d, ne = %4d: ", n_rounds, ne); + fflush(stdout); - const bool ok = test_multi_lane(backend_gpu, backend_cpu, n_rounds, ne); + const bool ok = test_multi_lane(backend_gpu, backend_cpu, n_rounds, ne, use_device_host_buft); - printf("%s\n", ok ? "OK" : "FAIL"); + printf("%s\n", ok ? "OK" : "FAIL"); - n_ok += ok; - n_test++; + n_ok += ok; + n_test++; + } } + + printf("\n"); + + for (int32_t sleep_us : { 100000, 0 }) { + for (int ne : { 1, 4096 }) { + printf(" sleep_race sleep_us = %6d, ne = %4d: ", sleep_us, ne); + fflush(stdout); + + const bool ok = test_sleep_race(backend_gpu, backend_cpu, ne, sleep_us, use_device_host_buft); + + printf("%s\n", ok ? "OK" : "FAIL"); + + n_ok += ok; + n_test++; + } + } + + printf("\n"); } ggml_backend_free(backend_gpu); ggml_backend_free(backend_cpu); - printf("\n%d/%d tests passed\n", n_ok, n_test); + printf("%d/%d tests passed\n", n_ok, n_test); return n_ok == n_test ? 0 : 1; } From 9ebe3673bdd483e0a1eb2f10bb7c2d64ea682cc6 Mon Sep 17 00:00:00 2001 From: aendk Date: Wed, 5 Aug 2026 17:13:45 +0200 Subject: [PATCH 04/11] Adds experimental vulkan GGML_OP_SLEEP --- ggml/src/ggml-vulkan/CMakeLists.txt | 6 ++ ggml/src/ggml-vulkan/ggml-vulkan.cpp | 68 +++++++++++++++++++ .../ggml-vulkan/vulkan-shaders/CMakeLists.txt | 4 ++ .../feature-tests/shader_clock.comp | 7 ++ .../src/ggml-vulkan/vulkan-shaders/sleep.comp | 36 ++++++++++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 3 + 6 files changed, 124 insertions(+) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/feature-tests/shader_clock.comp create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/sleep.comp diff --git a/ggml/src/ggml-vulkan/CMakeLists.txt b/ggml/src/ggml-vulkan/CMakeLists.txt index 1dc6a145de1e..130a1ca2fceb 100644 --- a/ggml/src/ggml-vulkan/CMakeLists.txt +++ b/ggml/src/ggml-vulkan/CMakeLists.txt @@ -109,6 +109,12 @@ if (Vulkan_FOUND) "GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT" ) + test_shader_extension_support( + "GL_EXT_shader_realtime_clock" + "${CMAKE_CURRENT_SOURCE_DIR}/vulkan-shaders/feature-tests/shader_clock.comp" + "GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT" + ) + target_link_libraries(ggml-vulkan PRIVATE Vulkan::Vulkan) target_include_directories(ggml-vulkan PRIVATE ${CMAKE_CURRENT_BINARY_DIR}) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 585e10d45ac7..44be39397851 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -818,6 +818,7 @@ struct vk_device_struct { bool shader_int64; bool buffer_device_address; bool vulkan_memory_model; + bool shader_clock; bool add_rms_fusion; uint32_t partials_binding_alignment; @@ -1013,6 +1014,8 @@ struct vk_device_struct { vk_pipeline pipeline_fill_f32; vk_pipeline pipeline_fill_f16; + vk_pipeline pipeline_sleep; + vk_pipeline pipeline_geglu[2]; vk_pipeline pipeline_reglu[2]; vk_pipeline pipeline_swiglu[2]; @@ -1394,6 +1397,11 @@ struct vk_op_push_constants { float param4; }; +struct vk_op_sleep_push_constants { + uint32_t ne; + uint32_t ticks; +}; + struct vk_op_fwht_push_constants { uint32_t n_rows; uint32_t src_offset; @@ -5675,6 +5683,12 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_fill_f32, "fill_f32", fill_f32_len, fill_f32_data, "main", 1, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_fill_f16, "fill_f16", fill_f16_len, fill_f16_data, "main", 1, sizeof(vk_op_push_constants), {512, 1, 1}, {}, 1); +#if defined(GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT) + if (device->shader_clock) { + ggml_vk_create_pipeline(device, device->pipeline_sleep, "sleep", sleep_len, sleep_data, "main", 2, sizeof(vk_op_sleep_push_constants), {512, 1, 1}, {}, 1); + } +#endif + #define CREATE_GLU(name) \ ggml_vk_create_pipeline(device, device->pipeline_ ## name [0], #name "_f32", name ## _f32_len, name ## _f32_data, "main", 3, sizeof(vk_op_glu_push_constants), {512, 1, 1}, {}, 1, true); \ ggml_vk_create_pipeline(device, device->pipeline_ ## name [1], #name "_f16", name ## _f16_len, name ## _f16_data, "main", 3, sizeof(vk_op_glu_push_constants), {512, 1, 1}, {}, 1, true); @@ -6203,6 +6217,7 @@ static vk_device ggml_vk_get_device(size_t idx) { bool dot2_f16_support = false; bool ocp_microscaling_extension = false; bool shader_float8_extension = false; + bool shader_clock_support = false; for (const auto& properties : ext_props) { if (strcmp("VK_KHR_maintenance4", properties.extensionName) == 0) { @@ -6252,6 +6267,11 @@ static vk_device ggml_vk_get_device(size_t idx) { #if defined(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT) } else if (strcmp(VK_EXT_SHADER_FLOAT8_EXTENSION_NAME, properties.extensionName) == 0) { shader_float8_extension = true; +#endif +#if defined(GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT) + } else if (strcmp("VK_KHR_shader_clock", properties.extensionName) == 0 && + !getenv("GGML_VK_DISABLE_SHADER_CLOCK")) { + shader_clock_support = true; #endif } else if (strcmp("VK_VALVE_shader_mixed_float_dot_product", properties.extensionName) == 0 && !getenv("GGML_VK_DISABLE_DOT2")) { @@ -6595,6 +6615,14 @@ static vk_device ggml_vk_get_device(size_t idx) { device_extensions.push_back("VK_KHR_shader_integer_dot_product"); } + VkPhysicalDeviceShaderClockFeaturesKHR shader_clock_features {}; + shader_clock_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_CLOCK_FEATURES_KHR; + if (shader_clock_support) { + last_struct->pNext = (VkBaseOutStructure *)&shader_clock_features; + last_struct = (VkBaseOutStructure *)&shader_clock_features; + device_extensions.push_back("VK_KHR_shader_clock"); + } + VkPhysicalDeviceShaderMixedFloatDotProductFeaturesVALVE dot2_features {}; dot2_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_MIXED_FLOAT_DOT_PRODUCT_FEATURES_VALVE; if (dot2_f16_support) { @@ -6678,6 +6706,7 @@ static vk_device ggml_vk_get_device(size_t idx) { device->shader_int64 = device_features2.features.shaderInt64; device->buffer_device_address = vk12_features.bufferDeviceAddress; device->vulkan_memory_model = vk12_features.vulkanMemoryModel; + device->shader_clock = shader_clock_support && shader_clock_features.shaderDeviceClock; if (device->subgroup_size_control) { device->subgroup_min_size = subgroup_size_control_props.minSubgroupSize; @@ -11779,6 +11808,8 @@ static vk_pipeline ggml_vk_op_get_pipeline(ggml_backend_vk_context * ctx, const return ctx->device->pipeline_arange_f32; } return nullptr; + case GGML_OP_SLEEP: + return ctx->device->pipeline_sleep; case GGML_OP_FILL: if (dst->type == GGML_TYPE_F32) { return ctx->device->pipeline_fill_f32; @@ -12980,6 +13011,32 @@ static void ggml_vk_fill(ggml_backend_vk_context * ctx, vk_context& subctx, ggml ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { dst_buf }, pc, elements); } +static void ggml_vk_sleep(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) { + VK_LOG_DEBUG("ggml_vk_sleep(dst=" << dst << ", us=" << ggml_get_op_params_i32(dst, 0) << ")"); + + // Vulkan does not specify the period of the shader realtime clock, assume it matches the timestamp counter + const float ns_per_tick = ctx->device->properties.limits.timestampPeriod; + const uint64_t ns = 1000 * (uint64_t) ggml_get_op_params_i32(dst, 0); + const uint64_t ticks = std::min(ns_per_tick > 0.0f ? (uint64_t)(ns / ns_per_tick) : ns, std::numeric_limits::max()); + + vk_op_sleep_push_constants pc = { + (uint32_t)(ggml_nbytes(dst) / sizeof(uint32_t)), + (uint32_t) ticks, + }; + + vk_pipeline pipeline = ggml_vk_op_get_pipeline(ctx, src0, nullptr, nullptr, dst, GGML_OP_SLEEP); + GGML_ASSERT(pipeline != nullptr); + + ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1); + vk_subbuffer src_buf = ggml_vk_tensor_subbuffer(ctx, src0); + vk_subbuffer dst_buf = ggml_vk_tensor_subbuffer(ctx, dst); + + // dispatch a single workgroup, all of its invocations spin concurrently so the delay does not accumulate + std::array elements = { pipeline->wg_denoms[0], 1, 1 }; + + ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { src_buf, dst_buf }, pc, elements); +} + static void ggml_vk_sin(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) { ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_SIN, vk_op_unary_push_constants_init(src0, dst)); } @@ -15386,6 +15443,10 @@ static bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgr case GGML_OP_FILL: ggml_vk_fill(ctx, compute_ctx, node); + break; + case GGML_OP_SLEEP: + ggml_vk_sleep(ctx, compute_ctx, src0, node); + break; case GGML_OP_SCALE: ggml_vk_scale(ctx, compute_ctx, src0, node); @@ -18349,6 +18410,11 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm return op->type == GGML_TYPE_F32; case GGML_OP_FILL: return op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16; + case GGML_OP_SLEEP: + // the shader copies 4-byte units, the delay comes from the device clock + return device->shader_clock && op->type == op->src[0]->type && + ggml_is_contiguous(op->src[0]) && ggml_is_contiguous(op) && + (ggml_nbytes(op) % sizeof(uint32_t)) == 0; case GGML_OP_SCALE: return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; case GGML_OP_PAD: @@ -19114,6 +19180,8 @@ static void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * } else if (tensor->op == GGML_OP_FILL) { const float value = ggml_get_op_params_f32(tensor, 0); tensor_clone = ggml_fill(ggml_ctx, src_clone[0], value); + } else if (tensor->op == GGML_OP_SLEEP) { + tensor_clone = ggml_sleep(ggml_ctx, src_clone[0], ggml_get_op_params_i32(tensor, 0)); } else if (tensor->op == GGML_OP_SQR) { tensor_clone = ggml_sqr(ggml_ctx, src_clone[0]); } else if (tensor->op == GGML_OP_SQRT) { diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/CMakeLists.txt b/ggml/src/ggml-vulkan/vulkan-shaders/CMakeLists.txt index cbe7a68bf372..0879994525d7 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/CMakeLists.txt +++ b/ggml/src/ggml-vulkan/vulkan-shaders/CMakeLists.txt @@ -31,6 +31,10 @@ if (GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT) add_compile_definitions(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT) message(STATUS "Enabling E4M3 glslc support") endif() +if (GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT) + add_compile_definitions(GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT) + message(STATUS "Enabling shader clock glslc support") +endif() if (GGML_VULKAN_SHADER_DEBUG_INFO) add_compile_definitions(GGML_VULKAN_SHADER_DEBUG_INFO) message(STATUS "Enabling shader debug info") diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/feature-tests/shader_clock.comp b/ggml/src/ggml-vulkan/vulkan-shaders/feature-tests/shader_clock.comp new file mode 100644 index 000000000000..0a9b0a4bd286 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/feature-tests/shader_clock.comp @@ -0,0 +1,7 @@ +#version 460 + +#extension GL_EXT_shader_realtime_clock : require + +void main() +{ +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/sleep.comp b/ggml/src/ggml-vulkan/vulkan-shaders/sleep.comp new file mode 100644 index 000000000000..e018f93d1236 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/sleep.comp @@ -0,0 +1,36 @@ +#version 450 + +#extension GL_EXT_shader_realtime_clock : require + +layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in; + +layout (binding = 0) readonly buffer A {uint data_a[];}; +layout (binding = 1) writeonly buffer D {uint data_d[];}; + +layout (push_constant) uniform parameter { + uint ne; + uint ticks; +} p; + +// bounds the spin in case the clock does not advance, turning a hang into a too-short sleep +const uint MAX_ITER = 1u << 27; + +void main() { + // the dispatch is a single workgroup, so all invocations spin concurrently and the delay does not accumulate + const uvec2 t0 = clockRealtime2x32EXT(); + + for (uint iter = 0; iter < MAX_ITER; ++iter) { + const uvec2 t = clockRealtime2x32EXT(); + + const uint lo = t.x - t0.x; + const uint hi = t.y - t0.y - (t.x < t0.x ? 1u : 0u); + + if (hi != 0u || lo >= p.ticks) { + break; + } + } + + for (uint i = gl_LocalInvocationID.x; i < p.ne; i += gl_WorkGroupSize.x) { + data_d[i] = data_a[i]; + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 6c9f76af1c9c..099212d6357b 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -960,6 +960,9 @@ void process_shaders() { string_to_spv("arange_f32", "arange.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}}); string_to_spv("fill_f32", "fill.comp", {{"D_TYPE", "float"}, {"FLOAT_TYPE", "float"}}); string_to_spv("fill_f16", "fill.comp", {{"D_TYPE", "float16_t"}, {"FLOAT_TYPE", "float"}}); +#if defined(GGML_VULKAN_SHADER_CLOCK_GLSLC_SUPPORT) + string_to_spv("sleep", "sleep.comp", {}); +#endif string_to_spv("step_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"OP", "op_step"}}); string_to_spv("step_f32", "unary.comp", {{"A_TYPE", "float"}, {"D_TYPE", "float"}, {"OP", "op_step"}}); string_to_spv("round_f16", "unary.comp", {{"A_TYPE", "float16_t"}, {"D_TYPE", "float16_t"}, {"OP", "op_round"}}); From fb335f758e0a1d179b9e1fb9f433a99f4652a656 Mon Sep 17 00:00:00 2001 From: aendk Date: Wed, 12 Aug 2026 14:15:02 +0200 Subject: [PATCH 05/11] Improves structure and naming of previous tests. Adds new tests, including one which detects the nkvo-bug correctly, and validates the proposed hotfix. --- tests/test-backend-sched.cpp | 985 ++++++++++++++++++++++------------- 1 file changed, 637 insertions(+), 348 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index a421b76321e5..ff5fbb09d3e9 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -1,6 +1,16 @@ -// Stress test for ggml_backend_sched: graphs built only from increments, placed on alternating backends. -// Since every node adds one, the expected result of a counter is just the number of times it was incremented, -// which makes the correct output independent of how the scheduler splits and copies the graph. +// The tests below define the scheduler and backend behavior for every combination of asynchronous backends +// On a machine with CPU and CUDA backends only, that means CPU->CPU, CPU->CUDA, CUDA->CPU and CUDA->CUDA + +// Formalized scheduling behavior: +// - a single inference pass can run on several backends. The subset of nodes running on a single backend is a split +// - synchronous backends: +// - explicit synchronization (ggml_backend_synchronize()) is required between each operation +// - asynchronous backends: +// - between splits, activations are copied asynchronously (output of split N -> input of split N+1) +// - Several scheduling patterns must be supported by async backends. The scheduler may: +// - not explicitly synchronize between CPU->backend memcpy and graph execution on backend +// - dispatch several parallel memcpys to the same backend at once +// - schedule splits without inputs at any point in the graph #include "ggml.h" #include "ggml-alloc.h" @@ -8,10 +18,73 @@ #include #include +#include #include #include #include +#include + +static int n_ok = 0; +static int n_test = 0; + +// pretty-print helpers start +static int case_len = 0; +static char case_label[256]; + +static void case_end(bool ok) { + if (case_len == 0) { + case_len = printf(" %s", case_label); + } + + printf("%*s%s\n", std::max(1, 82 - case_len), "", ok ? "OK" : "FAIL"); + + n_ok += ok; + n_test++; +} + +static void vnote(const char * fmt, va_list args) { + if (case_len != 0) { + printf("\n"); + } + printf(" "); + vprintf(fmt, args); + printf("\n"); + + case_len = 0; +} + +GGML_ATTRIBUTE_FORMAT(1, 2) +static void note(const char * fmt, ...) { + va_list args; + va_start(args, fmt); + vnote(fmt, args); + va_end(args); +} + +GGML_ATTRIBUTE_FORMAT(1, 2) +static bool fail(const char * fmt, ...) { + va_list args; + va_start(args, fmt); + vnote(fmt, args); + va_end(args); + + return false; +} + +GGML_ATTRIBUTE_FORMAT(1, 2) +static void case_begin(const char * fmt, ...) { + va_list args; + va_start(args, fmt); + vsnprintf(case_label, sizeof(case_label), fmt, args); + va_end(args); + + case_len = printf(" %s", case_label); + + fflush(stdout); +} +// pretty-print helpers end + static ggml_backend_buffer_type_t sched_cpu_buft(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, bool use_device_host_buft) { if (use_device_host_buft) { @@ -23,22 +96,158 @@ static ggml_backend_buffer_type_t sched_cpu_buft(ggml_backend_t backend_gpu, ggm return ggml_backend_get_default_buffer_type(backend_cpu); } -static ggml_backend_sched_t create_test_scheduler(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, size_t graph_size, bool use_device_host_buft) { - ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; - ggml_backend_buffer_type_t bufts[2] = { - ggml_backend_get_default_buffer_type(backend_gpu), - sched_cpu_buft(backend_gpu, backend_cpu, use_device_host_buft), +static ggml_backend_sched_t create_test_scheduler(std::vector backends, size_t graph_size, + bool use_device_host_buft, bool parallel = false) { + + std::vector bufts(backends.size()); + for (size_t b = 0; b < backends.size(); b++) { + bufts[b] = ggml_backend_get_default_buffer_type(backends[b]); + } + // CPU backend is always last + bufts.back() = sched_cpu_buft(backends[0], backends.back(), use_device_host_buft); + + return ggml_backend_sched_new(backends.data(), bufts.data(), (int) backends.size(), graph_size, parallel, /*op_offload =*/ false); +} + +// the nodes of the graph together with the backend each of them is assigned to +struct sched_graph { + ggml_cgraph * gf; + std::vector nodes; + std::vector backend_id; + + sched_graph(ggml_context * ctx, size_t graph_size) : gf(ggml_new_graph_custom(ctx, graph_size, false)) {} + + ggml_tensor * add(ggml_tensor * node, int backend) { + nodes.push_back(node); + backend_id.push_back(backend); + ggml_build_forward_expand(gf, node); + + return node; + } +}; + +struct sched_check { + ggml_tensor * tensor; + float expected; + std::string name; +}; + +// executes the graph and evaluates results +static bool run_and_check(ggml_backend_sched_t sched, const sched_graph & g, + const std::vector & backends, const std::vector & checks, int64_t ne) { + + int n_splits_expected = g.nodes.empty() ? 0 : 1; + for (size_t i = 1; i < g.backend_id.size(); i++) { + n_splits_expected += g.backend_id[i] != g.backend_id[i - 1]; + } + + ggml_backend_sched_reset(sched); + for (size_t i = 0; i < g.nodes.size(); i++) { + ggml_backend_sched_set_tensor_backend(sched, g.nodes[i], backends[g.backend_id[i]]); + } + + if (!ggml_backend_sched_alloc_graph(sched, g.gf)) { + return fail("failed to allocate the graph"); + } + + if (ggml_backend_sched_graph_compute(sched, g.gf) != GGML_STATUS_SUCCESS) { + return fail("failed to compute the graph"); + } + + const int n_splits = ggml_backend_sched_get_n_splits(sched); + if (n_splits != n_splits_expected) { + return fail("n_splits = %d, expected %d - the backend assignments were not respected", n_splits, n_splits_expected); + } + + std::vector data(ne); + for (const sched_check & c : checks) { + ggml_backend_tensor_get(c.tensor, data.data(), 0, ggml_nbytes(c.tensor)); + for (int64_t i = 0; i < ne; i++) { + if (data[i] != c.expected) { + return fail("%s[%" PRId64 "] = %f, expected %f", c.name.c_str(), i, data[i], c.expected); + } + } + } + + return true; +} + +static bool backend_supports(ggml_backend_t backend, ggml_tensor * (*build)(ggml_context *, ggml_tensor *)) { + ggml_init_params params = { + /*.mem_size =*/ 4*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, }; - return ggml_backend_sched_new(backends, bufts, 2, graph_size, /*parallel =*/ false, /*op_offload =*/ false); + ggml_context * ctx = ggml_init(params); + + ggml_tensor * a = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 4); + + const bool ok = ggml_backend_supports_op(backend, build(ctx, a)); + + ggml_free(ctx); + + return ok; } -static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t ne, bool use_device_host_buft) { +//helper struct for inputless splits, and constants like 0 and 1 +struct backend_consts { + std::vector ctxs; + std::vector bufs; + std::vector zero; + std::vector one; + + backend_consts(const std::vector & backends, int64_t ne) { + std::vector data(ne); + + for (size_t b = 0; b < backends.size(); b++) { + ggml_init_params params = { + /*.mem_size =*/ 2*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx = ggml_init(params); + + ggml_tensor * z = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, ne); + ggml_format_name(z, "zero_%s", ggml_backend_name(backends[b])); + + ggml_tensor * o = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, ne); + ggml_format_name(o, "one_%s", ggml_backend_name(backends[b])); + + ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backends[b]); + + std::fill(data.begin(), data.end(), 0.0f); + ggml_backend_tensor_set(z, data.data(), 0, ggml_nbytes(z)); + + std::fill(data.begin(), data.end(), 1.0f); + ggml_backend_tensor_set(o, data.data(), 0, ggml_nbytes(o)); + + ctxs.push_back(ctx); + bufs.push_back(buf); + zero.push_back(z); + one.push_back(o); + } + } + + ~backend_consts() { + for (size_t b = 0; b < ctxs.size(); b++) { + ggml_backend_buffer_free(bufs[b]); + ggml_free(ctxs[b]); + } + } +}; + +// stress test: a tensor is incremented and sent back and forth between a backend and CPU in a ping-pong pattern +static bool stress_test_linked_list_cpu_device(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t tensor_len, bool use_device_host_buft) { + nvtx3::scoped_range sc_4{nvtx3::event_attributes{nvtx3::rgb{0, 0, 255}, // blue + ("stress_test_linked_list_cpu_device " + std::string(ggml_backend_name(backend_gpu)) + " + " + ggml_backend_name(backend_cpu) + + " n_nodes=" + std::to_string(n_nodes) + " ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; - ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + const std::vector backends = { backend_gpu, backend_cpu }; - const size_t graph_size = n_nodes + 16; + const size_t graph_size = n_nodes + 2; // see sched->hash_set FIXME, 2+ needed to account for leafs - ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); // the inputs are allocated separately so that they can be written before the graph is computed ggml_init_params params_static = { @@ -48,17 +257,17 @@ static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cp }; ggml_context * ctx_static = ggml_init(params_static); - ggml_tensor * x = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, ne); + ggml_tensor * x = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, tensor_len); ggml_set_name(x, "x"); ggml_set_input(x); - ggml_tensor * one = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, ne); + ggml_tensor * one = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, tensor_len); ggml_set_name(one, "one"); ggml_set_input(one); ggml_backend_buffer_t buf_static = ggml_backend_alloc_ctx_tensors(ctx_static, backend_cpu); - std::vector data(ne); + std::vector data(tensor_len); std::fill(data.begin(), data.end(), 0.0f); ggml_backend_tensor_set(x, data.data(), 0, ggml_nbytes(x)); @@ -73,55 +282,15 @@ static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cp }; ggml_context * ctx_compute = ggml_init(params_compute); - ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); - - std::vector nodes; + sched_graph g(ctx_compute, graph_size); ggml_tensor * out = x; for (int i = 0; i < n_nodes; i++) { - out = ggml_add(ctx_compute, out, one); - nodes.push_back(out); + out = g.add(ggml_add(ctx_compute, out, one), i % 2); } ggml_set_output(out); - ggml_build_forward_expand(gf, out); - - // the assignments are only kept until the next reset, and ggml_backend_sched_graph_compute - // resets the scheduler unless the graph has already been allocated - ggml_backend_sched_reset(sched); - for (int i = 0; i < n_nodes; i++) { - ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[i % 2]); - } - - bool ok = true; - - if (!ggml_backend_sched_alloc_graph(sched, gf)) { - printf("\n failed to allocate the graph\n"); - ok = false; - } - - if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { - printf("\n failed to compute the graph\n"); - ok = false; - } - if (ok) { - const int n_splits = ggml_backend_sched_get_n_splits(sched); - if (n_splits != n_nodes) { - printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_nodes); - ok = false; - } - } - - if (ok) { - ggml_backend_tensor_get(out, data.data(), 0, ggml_nbytes(out)); - for (int64_t i = 0; i < ne; i++) { - if (data[i] != float(n_nodes)) { - printf("\n out[%" PRId64 "] = %f, expected %d\n", i, data[i], n_nodes); - ok = false; - break; - } - } - } + const bool ok = run_and_check(sched, g, backends, {{ out, float(n_nodes), "out" }}, tensor_len); ggml_backend_sched_free(sched); ggml_free(ctx_compute); @@ -131,64 +300,29 @@ static bool test_ping_pong(ggml_backend_t backend_gpu, ggml_backend_t backend_cp return ok; } -// Same idea as test_ping_pong, but with 8 independent counters and a graph shape that covers -// three cases the plain chain does not: -// - splits with no inputs at all, from lanes restarted out of constants that already live on the split's backend +// Same idea as stress_test_linked_list_cpu_device, but with the compute graph being a direct acyclic graph (DAG). +// n_lanes independent counters run through rounds of four shapes, so that one graph covers: // - lanes with different histories, so a misrouted copy shows up as a wrong count in a single lane // - one split producing two values followed by two splits that each consume one of them and do not depend on each other -static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_rounds, int64_t ne, bool use_device_host_buft) { - - const int GPU = 0; - const int CPU = 1; - - const int n_lanes = 8; - - ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; - - // the graph needs 2*n_lanes nodes per round plus one restart for half of the lanes, but the hash set of - // the scheduler is sized from the same value and GGML_SCHED_DEBUG inserts every copy tensor into it as - // well, so leave room for those - const size_t n_nodes = 2*n_lanes*n_rounds + n_lanes/2; - const size_t graph_size = 3*n_nodes + 64; - - ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); - - ggml_init_params params_static = { - /*.mem_size =*/ 2*ggml_tensor_overhead(), - /*.mem_buffer =*/ nullptr, - /*.no_alloc =*/ true, - }; - - // the constants are kept on both backends so that a lane can be restarted without any cross-backend copy - ggml_context * ctx_cpu = ggml_init(params_static); - - ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); - ggml_set_name(zero_cpu, "zero_cpu"); - - ggml_tensor * one_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); - ggml_set_name(one_cpu, "one_cpu"); +// - splits with no inputs at all, from lanes re-seeded out of constants that already live on the split's backend +// - splits taking one activation per lane from the previous split, one of them from an older split instead +static bool stress_test_dag(const std::vector & backends, int n_lanes, int n_rounds, + int64_t tensor_len, bool use_device_host_buft) { + nvtx3::scoped_range sc_5{nvtx3::event_attributes{nvtx3::rgb{255, 165, 0}, // orange + ("stress_test_dag n_lanes=" + std::to_string(n_lanes) + + " n_rounds=" + std::to_string(n_rounds) + " ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; - ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, backend_cpu); - - ggml_context * ctx_gpu = ggml_init(params_static); - - ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); - ggml_set_name(zero_gpu, "zero_gpu"); + const int n_backends = (int) backends.size(); - ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); - ggml_set_name(one_gpu, "one_gpu"); + GGML_ASSERT(n_lanes % 2 == 0); - ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, backend_gpu); + // sized for the worst case + const size_t graph_size = n_lanes*(2*n_rounds + 1) + 2*n_backends; - std::vector data(ne); - - std::fill(data.begin(), data.end(), 0.0f); - ggml_backend_tensor_set(zero_cpu, data.data(), 0, ggml_nbytes(zero_cpu)); - ggml_backend_tensor_set(zero_gpu, data.data(), 0, ggml_nbytes(zero_gpu)); + backend_consts consts(backends, tensor_len); - std::fill(data.begin(), data.end(), 1.0f); - ggml_backend_tensor_set(one_cpu, data.data(), 0, ggml_nbytes(one_cpu)); - ggml_backend_tensor_set(one_gpu, data.data(), 0, ggml_nbytes(one_gpu)); + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); ggml_init_params params_compute = { /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), @@ -197,119 +331,103 @@ static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_c }; ggml_context * ctx_compute = ggml_init(params_compute); - ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); - - ggml_tensor * lane[n_lanes]; - int expected[n_lanes]; - for (int l = 0; l < n_lanes; l++) { - lane[l] = zero_cpu; - expected[l] = 0; - } + sched_graph g(ctx_compute, graph_size); - std::vector nodes; - std::vector node_backend; + std::vector lane(n_lanes); + std::vector value(n_lanes); + std::vector checks; - // expanding each node as it is created keeps the graph order equal to the creation order, - // which is what determines where the scheduler puts the split boundaries - auto append = [&](int l, ggml_tensor * node, int b) { - lane[l] = node; - nodes.push_back(node); - node_backend.push_back(b); - ggml_build_forward_expand(gf, node); - }; - - auto increment = [&](int l, int b) { - append(l, ggml_add(ctx_compute, lane[l], b == GPU ? one_gpu : one_cpu), b); - expected[l]++; + // every lane starts at a different value, so that a misrouted copy changes a sum instead of going unnoticed + auto seed = [&](int b, const char * prefix) { + for (int l = 0; l < n_lanes; l++) { + lane[l] = g.add(ggml_add(ctx_compute, l == 0 ? consts.zero[b] : lane[l - 1], consts.one[b]), b); + value[l] = l + 1; + ggml_format_name(lane[l], "%s%d", prefix, l); + } }; - // both sources already live on backend b, so the resulting split has no inputs to copy - auto restart = [&](int l, int b) { - append(l, ggml_add(ctx_compute, b == GPU ? zero_gpu : zero_cpu, b == GPU ? one_gpu : one_cpu), b); - expected[l] = 1; + auto check_lanes = [&]() { + for (int l = 0; l < n_lanes; l++) { + ggml_set_output(lane[l]); + checks.push_back({ lane[l], float(value[l]), ggml_get_name(lane[l]) }); + } }; - for (int round = 0; round < n_rounds; round++) { - // halfway through, the upper half of the lanes starts over from scratch - if (round == n_rounds/2) { - // start on the backend the previous node did not use, so that no restart is absorbed into the split before it - int b = node_backend.empty() || node_backend.back() == CPU ? GPU : CPU; - for (int l = n_lanes/2; l < n_lanes; l++) { - restart(l, b); - b = b == CPU ? GPU : CPU; - } - } + seed(0, "seed"); - for (int p = 0; p < n_lanes/2; p++) { - const int la = 2*p; - const int lb = 2*p + 1; + std::vector older = lane; + std::vector value_older = value; - const int prod = p % 2 == 0 ? CPU : GPU; + for (int round = 0; round < n_rounds; round++) { + // a backend the previous node did not use, so that the round is not absorbed into the split before it + const int b = (g.backend_id.back() + 1) % n_backends; + + const std::vector prev = lane; + const std::vector value_prev = value; + + switch (round % 5) { + case 0: + for (int p = 0; p < n_lanes/2; p++) { + const int la = 2*p; + const int lb = 2*p + 1; + const int prod = (b + p) % n_backends; + const int cons = (prod + 1) % n_backends; + + // one split producing two values ... + for (int l : { la, lb }) { + lane[l] = g.add(ggml_add(ctx_compute, lane[l], consts.one[prod]), prod); + value[l]++; + ggml_format_name(lane[l], "r%d_prod%d", round, l); + } + + // ... and two splits that each consume one of them, with no dependency on each other + lane[la] = g.add(ggml_add(ctx_compute, lane[la], consts.one[cons]), cons); + value[la]++; + ggml_format_name(lane[la], "r%d_cons%d", round, la); + + lane[lb] = g.add(ggml_add(ctx_compute, lane[lb], consts.one[prod]), prod); + value[lb]++; + ggml_format_name(lane[lb], "r%d_cons%d", round, lb); + } + break; + case 1: + check_lanes(); + seed(b, "restart"); + break; + case 4: + lane[0] = g.add(ggml_add(ctx_compute, prev[0], consts.one[b]), b); + value[0] = value_prev[0] + 1; + ggml_format_name(lane[0], "r%d_single", round); + break; + default: + for (int l = 0; l < n_lanes; l++) { + const int partner = (l + 1 + round) % n_lanes; + const bool reach_back = round >= 2 && l == round % n_lanes; - // one split producing two values ... - increment(la, prod); - increment(lb, prod); + ggml_tensor * a = reach_back ? older[l] : prev[l]; + const int64_t va = reach_back ? value_older[l] : value_prev[l]; - // ... and two splits that each consume one of them, with no dependency on each other - increment(la, prod == CPU ? GPU : CPU); - increment(lb, prod); + lane[l] = g.add(ggml_add(ctx_compute, a, prev[partner]), b); + value[l] = va + value_prev[partner]; + ggml_format_name(lane[l], "r%d_lane%d", round, l); + } + break; } - } - for (int l = 0; l < n_lanes; l++) { - ggml_set_output(lane[l]); + older = prev; + value_older = value_prev; } - // the scheduler starts a new split on every backend change, so the assignments below fix the split count - int n_splits_expected = nodes.empty() ? 0 : 1; - for (size_t i = 1; i < node_backend.size(); i++) { - n_splits_expected += node_backend[i] != node_backend[i - 1]; - } + check_lanes(); - ggml_backend_sched_reset(sched); - for (size_t i = 0; i < nodes.size(); i++) { - ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[node_backend[i]]); + for (const sched_check & c : checks) { + GGML_ASSERT(c.expected < float(1 << 24)); // the counts have to stay exact in f32 } - bool ok = true; - - if (!ggml_backend_sched_alloc_graph(sched, gf)) { - printf("\n failed to allocate the graph\n"); - ok = false; - } - - if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { - printf("\n failed to compute the graph\n"); - ok = false; - } - - if (ok) { - const int n_splits = ggml_backend_sched_get_n_splits(sched); - if (n_splits != n_splits_expected) { - printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_splits_expected); - ok = false; - } - } - - if (ok) { - for (int l = 0; l < n_lanes && ok; l++) { - ggml_backend_tensor_get(lane[l], data.data(), 0, ggml_nbytes(lane[l])); - for (int64_t i = 0; i < ne; i++) { - if (data[i] != float(expected[l])) { - printf("\n lane %d [%" PRId64 "] = %f, expected %d\n", l, i, data[i], expected[l]); - ok = false; - break; - } - } - } - } + const bool ok = run_and_check(sched, g, backends, checks, tensor_len); ggml_backend_sched_free(sched); ggml_free(ctx_compute); - ggml_backend_buffer_free(buf_gpu); - ggml_free(ctx_gpu); - ggml_backend_buffer_free(buf_cpu); - ggml_free(ctx_cpu); return ok; } @@ -317,25 +435,29 @@ static bool test_multi_lane(ggml_backend_t backend_gpu, ggml_backend_t backend_c // Created in response to https://github.com/ggml-org/llama.cpp/issues/23321 // Currently, input copies to splits synchronize split execution. // CPU split N+1 could overwrite output of split N if this was not copied to backend in time. -// Mechanism here: GPU split with long sleep to clog CUDA stream / vk command queue, +// Mechanism to test: GPU split with long sleep to clog CUDA stream / vk command queue, // followed by CPU splits without input, provoking a race condition in the current setup. // GPU: 100ms sleep to clog stream/queue // CPU: produce {55} // GPU: request copy of {55} in CUDA stream, ADD {44}; Output should be {99} // CPU: produce {66} // GPU: increment both {99} and {66} -> {100} and {67} -// Correct result is thus {100}, incorrect output is {111} when {55} was overwritten by {66} before copy started. +// Correct result is thus {100}, incorrect output is {111} when {55} was overwritten by {66} before the copy started. // Note: currently only reproducible on async H2D copy (pinned memory in CUDA). -static bool test_sleep_race(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int64_t ne, int32_t sleep_us, bool use_device_host_buft) { +static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int64_t tensor_len, int32_t sleep_us, bool use_device_host_buft) { + nvtx3::scoped_range sc_6{nvtx3::event_attributes{nvtx3::rgb{255, 0, 0}, // red + ("test_inputless_splits_scheduling " + std::string(ggml_backend_name(backend_gpu)) + " + " + ggml_backend_name(backend_cpu) + + " sleep_us=" + std::to_string(sleep_us) + " ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const int GPU = 0; const int CPU = 1; - ggml_backend_t backends[2] = { backend_gpu, backend_cpu }; + const std::vector backends = { backend_gpu, backend_cpu }; const size_t graph_size = 64; - ggml_backend_sched_t sched = create_test_scheduler(backend_gpu, backend_cpu, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); ggml_init_params params_static = { /*.mem_size =*/ 4*ggml_tensor_overhead(), @@ -345,31 +467,31 @@ static bool test_sleep_race(ggml_backend_t backend_gpu, ggml_backend_t backend_c ggml_context * ctx_cpu = ggml_init(params_static); - ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, tensor_len); ggml_set_name(zero_cpu, "zero_cpu"); - ggml_tensor * val55_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_tensor * val55_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, tensor_len); ggml_set_name(val55_cpu, "val55_cpu"); - ggml_tensor * val66_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, ne); + ggml_tensor * val66_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, tensor_len); ggml_set_name(val66_cpu, "val66_cpu"); ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, backend_cpu); ggml_context * ctx_gpu = ggml_init(params_static); - ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, tensor_len); ggml_set_name(zero_gpu, "zero_gpu"); - ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, tensor_len); ggml_set_name(one_gpu, "one_gpu"); - ggml_tensor * val44_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, ne); + ggml_tensor * val44_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, tensor_len); ggml_set_name(val44_gpu, "val44_gpu"); ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, backend_gpu); - std::vector data(ne); + std::vector data(tensor_len); std::fill(data.begin(), data.end(), 0.0f); ggml_backend_tensor_set(zero_cpu, data.data(), 0, ggml_nbytes(zero_cpu)); @@ -394,111 +516,175 @@ static bool test_sleep_race(ggml_backend_t backend_gpu, ggml_backend_t backend_c }; ggml_context * ctx_compute = ggml_init(params_compute); - ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); - - std::vector nodes; - std::vector node_backend; - - auto append = [&](ggml_tensor * node, int b) { - nodes.push_back(node); - node_backend.push_back(b); - ggml_build_forward_expand(gf, node); - }; + sched_graph g(ctx_compute, graph_size); // GPU: add then sleep - occupies the GPU while later splits run - ggml_tensor * delayed = ggml_add(ctx_compute, zero_gpu, one_gpu); - append(delayed, GPU); - delayed = ggml_sleep(ctx_compute, delayed, sleep_us); - append(delayed, GPU); + ggml_tensor * delayed = g.add(ggml_add(ctx_compute, zero_gpu, one_gpu), GPU); + delayed = g.add(ggml_sleep(ctx_compute, delayed, sleep_us), GPU); // CPU: 0 + 55 -> 55 - ggml_tensor * v55 = ggml_add(ctx_compute, zero_cpu, val55_cpu); - append(v55, CPU); + ggml_tensor * v55 = g.add(ggml_add(ctx_compute, zero_cpu, val55_cpu), CPU); // GPU: 55 + 44 -> 99 - ggml_tensor * v99 = ggml_add(ctx_compute, v55, val44_gpu); - append(v99, GPU); + ggml_tensor * v99 = g.add(ggml_add(ctx_compute, v55, val44_gpu), GPU); // CPU: 0 + 66 -> 66 - ggml_tensor * v66 = ggml_add(ctx_compute, zero_cpu, val66_cpu); - append(v66, CPU); + ggml_tensor * v66 = g.add(ggml_add(ctx_compute, zero_cpu, val66_cpu), CPU); // GPU: increment both of the previous split outputs - ggml_tensor * out99 = ggml_add(ctx_compute, v99, one_gpu); - append(out99, GPU); - ggml_tensor * out66 = ggml_add(ctx_compute, v66, one_gpu); - append(out66, GPU); + ggml_tensor * out99 = g.add(ggml_add(ctx_compute, v99, one_gpu), GPU); + ggml_tensor * out66 = g.add(ggml_add(ctx_compute, v66, one_gpu), GPU); ggml_set_output(delayed); ggml_set_output(out99); ggml_set_output(out66); - const int n_splits_expected = 5; + const bool ok = run_and_check(sched, g, backends, { + { delayed, 1.0f, "delayed" }, + { out99, 100.0f, "out99" }, + { out66, 67.0f, "out66" }, + }, tensor_len); - ggml_backend_sched_reset(sched); - for (size_t i = 0; i < nodes.size(); i++) { - ggml_backend_sched_set_tensor_backend(sched, nodes[i], backends[node_backend[i]]); - } + // the race needs the allocator to reuse the memory of the first CPU split for the second one + note("v55 and v66 %s (v55=%p v66=%p)", v55->data == v66->data ? "alias" : "do not alias", v55->data, v66->data); + + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + ggml_backend_buffer_free(buf_gpu); + ggml_free(ctx_gpu); + ggml_backend_buffer_free(buf_cpu); + ggml_free(ctx_cpu); - bool ok = true; - - if (!ggml_backend_supports_op(backend_gpu, delayed)) { - printf("\n GGML_OP_SLEEP not supported on GPU, skipping\n"); - ggml_backend_sched_free(sched); - ggml_free(ctx_compute); - ggml_backend_buffer_free(buf_gpu); - ggml_free(ctx_gpu); - ggml_backend_buffer_free(buf_cpu); - ggml_free(ctx_cpu); - return true; + return ok; +} + +// Test that all async backends transmit their activations to the following async backend correctly. No user inputs tested. +static bool test_chain_all_backends(const std::vector & backends, int64_t tensor_len, bool use_device_host_buft) { + nvtx3::scoped_range sc_7{nvtx3::event_attributes{nvtx3::rgb{0, 255, 0}, // green + ("test_chain_all_backends ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; + + const int n_backends = (int) backends.size(); + + std::vector seq; + + auto push = [&](int b) { + if (seq.empty() || seq.back() != b) { + seq.push_back(b); + } + }; + + for (int i = 0; i < n_backends; i++) { + for (int j = 0; j < n_backends; j++) { + if (i != j) { + push(i); + push(j); + } + } } - if (!ggml_backend_sched_alloc_graph(sched, gf)) { - printf("\n failed to allocate the graph\n"); - ok = false; + const size_t n_nodes = seq.size(); + const size_t graph_size = n_nodes + n_backends + 1; // see sched->hash_set FIXME, one "one" per backend plus the "zero" that starts the chain + + backend_consts consts(backends, tensor_len); + + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + + ggml_init_params params_compute = { + /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + sched_graph g(ctx_compute, graph_size); + + ggml_tensor * out = nullptr; + for (size_t k = 0; k < n_nodes; k++) { + const int b = seq[k]; + + out = g.add(ggml_add(ctx_compute, k == 0 ? consts.zero[b] : out, consts.one[b]), b); } + ggml_set_output(out); + + const bool ok = run_and_check(sched, g, backends, {{ out, float(n_nodes), "out" }}, tensor_len); - // checks if the the allocator reuses the same memory for output tensors between splits. - if (ok && v55->data == v66->data) { - printf("\n NOTE: v55 and v66 alias after alloc (v55=%p v66=%p)\n", v55->data, v66->data); - } else { - printf("\n NOTE: v55 and v66 do not alias after alloc (v55=%p v66=%p)\n", v55->data, v66->data); + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); + + return ok; +} + +// Tests data transfer between all combinations of backend pairs +// Always tests between two backends only with a single activation and 4 parallel user inputs. +static bool test_pair_user_inputs(const std::vector & backends, int b_send, int b_recv, int64_t tensor_len, + int n_inputs, bool inputs_on_sender, bool parallel, bool use_device_host_buft) { + + nvtx3::scoped_range sc_8{nvtx3::event_attributes{nvtx3::rgb{255, 215, 0}, // gold + ("test_pair_user_inputs " + std::string(ggml_backend_name(backends[b_send])) + " -> " + ggml_backend_name(backends[b_recv]) + + (inputs_on_sender ? " inputs on sender" : " inputs on receiver") + + " parallel=" + std::to_string(parallel) + " ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; + + const size_t graph_size = 64; + + backend_consts consts(backends, tensor_len); + + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft, parallel); + + // placing the inputs on the sender makes the receiving split copy all of them, placing them on the + // receiver leaves the activation as its only input + const int b_inputs = inputs_on_sender ? b_send : b_recv; + + ggml_init_params params_inputs = { + /*.mem_size =*/ n_inputs*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_inputs = ggml_init(params_inputs); + + std::vector inputs; + for (int k = 0; k < n_inputs; k++) { + ggml_tensor * in = ggml_new_tensor_1d(ctx_inputs, GGML_TYPE_F32, tensor_len); + ggml_format_name(in, "in%d", k); + ggml_set_input(in); + inputs.push_back(in); } - if (ok && ggml_backend_sched_graph_compute(sched, gf) != GGML_STATUS_SUCCESS) { - printf("\n failed to compute the graph\n"); - ok = false; + ggml_backend_buffer_t buf_inputs = ggml_backend_alloc_ctx_tensors(ctx_inputs, backends[b_inputs]); + + std::vector data(tensor_len); + std::fill(data.begin(), data.end(), 1.0f); + for (ggml_tensor * in : inputs) { + ggml_backend_tensor_set(in, data.data(), 0, ggml_nbytes(in)); } - if (ok) { - const int n_splits = ggml_backend_sched_get_n_splits(sched); - if (n_splits != n_splits_expected) { - printf("\n n_splits = %d, expected %d - the backend assignments were not respected\n", n_splits, n_splits_expected); - ok = false; - } + ggml_init_params params_compute = { + /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + sched_graph g(ctx_compute, graph_size); + + ggml_tensor * out = g.add(ggml_add(ctx_compute, consts.zero[b_send], consts.one[b_send]), b_send); + + for (ggml_tensor * in : inputs) { + out = g.add(ggml_add(ctx_compute, out, in), b_recv); } + ggml_set_output(out); - if (ok) { - const struct { - ggml_tensor * tensor; - float expected; - const char * name; - } checks[] = { - { delayed, 1.0f, "delayed" }, - { out99, 100.0f, "out99" }, - { out66, 67.0f, "out66" }, - }; - - for (const auto & c : checks) { - ggml_backend_tensor_get(c.tensor, data.data(), 0, ggml_nbytes(c.tensor)); - for (int64_t i = 0; i < ne; i++) { - if (data[i] != c.expected) { - printf("\n %s[%" PRId64 "] = %f, expected %f\n", c.name, i, data[i], c.expected); - ok = false; - break; - } - } - if (!ok) { + bool ok = run_and_check(sched, g, backends, {{ out, float(1 + n_inputs), "out" }}, tensor_len); + + for (ggml_tensor * in : inputs) { + if (!ok) { + break; + } + ggml_backend_tensor_get(in, data.data(), 0, ggml_nbytes(in)); + for (int64_t i = 0; i < tensor_len; i++) { + if (data[i] != 1.0f) { + ok = fail("%s[%" PRId64 "] = %f after compute, expected 1.000000 - the input was modified", ggml_get_name(in), i, data[i]); break; } } @@ -506,108 +692,211 @@ static bool test_sleep_race(ggml_backend_t backend_gpu, ggml_backend_t backend_c ggml_backend_sched_free(sched); ggml_free(ctx_compute); - ggml_backend_buffer_free(buf_gpu); - ggml_free(ctx_gpu); - ggml_backend_buffer_free(buf_cpu); - ggml_free(ctx_cpu); + ggml_backend_buffer_free(buf_inputs); + ggml_free(ctx_inputs); return ok; } -int main() { - ggml_backend_load_all(); +// Tests Y-shaped scheduling: two parallel lanes merging into 1. Lane A and B, merging into a single split. +// The lanes join in a final split on b_send that receives one activation from each of the two b_recv splits. +// TODO improve function signature + hoist backend_consts out of it? +static bool test_y_shaped_graph(const std::vector & backends, int backend_a, int backend_b, int64_t tensor_len, + bool use_device_host_buft) { + + nvtx3::scoped_range sc_9{nvtx3::event_attributes{nvtx3::rgb{238, 130, 238}, // violet + ("test_y_shaped_graph " + std::string(ggml_backend_name(backends[backend_a])) + " -> " + ggml_backend_name(backends[backend_b]) + + " ne=" + std::to_string(tensor_len) + + (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; + + const size_t graph_size = 64; + + backend_consts consts(backends, tensor_len); + + ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + + ggml_init_params params_compute = { + /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + + sched_graph g(ctx_compute, graph_size); + + ggml_tensor * seed_a = g.add(ggml_add(ctx_compute, consts.zero[backend_a], consts.one[backend_a]), backend_a); + ggml_set_name(seed_a, "seed_a"); + + ggml_tensor * lane_a = g.add(ggml_add(ctx_compute, seed_a, consts.one[backend_b]), backend_b); + ggml_set_name(lane_a, "lane_a"); + + ggml_tensor * seed_b = g.add(ggml_add(ctx_compute, consts.zero[backend_a], consts.one[backend_a]), backend_a); + ggml_set_name(seed_b, "seed_b"); + + ggml_tensor * lane_b = g.add(ggml_add(ctx_compute, seed_b, consts.one[backend_b]), backend_b); + ggml_set_name(lane_b, "lane_b"); + + // neither lane ends on b_send, so both activations are inputs of the last split + ggml_tensor * out = g.add(ggml_add(ctx_compute, lane_a, lane_b), backend_a); + ggml_set_name(out, "out"); + ggml_set_output(out); + + const bool ok = run_and_check(sched, g, backends, {{ out, 4.0f, "out" }}, tensor_len); + + ggml_backend_sched_free(sched); + ggml_free(ctx_compute); - ggml_backend_dev_t dev_gpu = nullptr; - for (size_t i = 0; i < ggml_backend_dev_count(); i++) { + return ok; +} + +static bool initialize_gpu_backends(std::vector & backends, bool & have_device_host_buft, bool & have_sleep) { + + // cf. GGML_SCHED_MAX_BACKENDS + const size_t max_backends = 16; + + for (size_t i = 0; i < ggml_backend_dev_count() && backends.size() + 1 < max_backends; i++) { ggml_backend_dev_t dev = ggml_backend_dev_get(i); const enum ggml_backend_dev_type type = ggml_backend_dev_type(dev); - if (type == GGML_BACKEND_DEVICE_TYPE_GPU || type == GGML_BACKEND_DEVICE_TYPE_IGPU) { - dev_gpu = dev; - break; + //GPU or IGPU only, for now + if (type != GGML_BACKEND_DEVICE_TYPE_GPU && type != GGML_BACKEND_DEVICE_TYPE_IGPU) { + continue; + } + + ggml_backend_t backend = ggml_backend_dev_init(dev, nullptr); + if (backend == nullptr) { + printf("failed to initialize %s, skipping it\n", ggml_backend_dev_name(dev)); + continue; } + + backends.push_back(backend); } - if (dev_gpu == nullptr) { + if (backends.empty()) { printf("no GPU device found, skipping\n"); return 0; } - ggml_backend_t backend_gpu = ggml_backend_dev_init(dev_gpu, nullptr); - ggml_backend_t backend_cpu = ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr); - GGML_ASSERT(backend_gpu != nullptr); - GGML_ASSERT(backend_cpu != nullptr); - - printf("GPU: %s (%s)\n", ggml_backend_name(backend_gpu), ggml_backend_dev_description(dev_gpu)); - printf("CPU: %s\n", ggml_backend_name(backend_cpu)); + // ggml_backend_sched_new requires the CPU backend to be the last one + backends.push_back(ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr)); + GGML_ASSERT(backends.back() != nullptr); + // todo rework this, only selects a single GPU + // create struct incorporating have_device_host_buft and have_sleep flags + ggml_backend_t backend_gpu = backends[0]; + ggml_backend_t backend_cpu = backends.back(); + + for (ggml_backend_t backend : backends) { + ggml_backend_dev_t dev = ggml_backend_get_device(backend); + printf("backend: %-10s (%s)\n", ggml_backend_name(backend), ggml_backend_dev_description(dev)); + } - const bool have_device_host_buft = ggml_backend_dev_host_buffer_type(dev_gpu) != nullptr; + have_device_host_buft = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend_gpu)) != nullptr; if (!have_device_host_buft) { - printf("GPU has no host buffer type; device_host cases will be skipped\n"); + printf("GPU has no host buffer type, the device_host cases are skipped\n"); + } + + have_sleep = backend_supports(backend_gpu, [](ggml_context * ctx, ggml_tensor * a) { return ggml_sleep(ctx, a, 0); }); + if (!have_sleep) { + printf("GPU does not support GGML_OP_SLEEP, some tests are skipped\n"); } printf("\n"); + return 1; +} + + +int main() { + ggml_backend_load_all(); - int n_ok = 0; - int n_test = 0; + std::vector backends; // TODO merge into single struct + bool have_device_host_buft = false; + bool have_sleep = false; + bool gpu_initialized = initialize_gpu_backends(backends, have_device_host_buft, have_sleep); + ggml_backend_t backend_gpu = backends[0]; + ggml_backend_t backend_cpu = backends.back(); + + if (!gpu_initialized) { + return 0; + } for (bool use_device_host_buft : { false, true }) { if (use_device_host_buft && !have_device_host_buft) { continue; } - const char * buft = use_device_host_buft ? "device_host" : "pageable"; - printf("=== CPU sched buft: %s ===\n\n", buft); - - for (int n_nodes : { 1, 2, 3, 8, 33, 128, 1024 }) { - for (int ne : { 1, 4096 }) { - printf(" n_nodes = %4d, ne = %4d: ", n_nodes, ne); - fflush(stdout); + printf("=== CPU sched buft: %s ===\n\n", use_device_host_buft ? "device_host" : "pageable"); - const bool ok = test_ping_pong(backend_gpu, backend_cpu, n_nodes, ne, use_device_host_buft); + for (int n_nodes : { 2, 5, 128, 1024 }) { + for (int tensor_len : { 2, 4096 }) { + case_begin("test_linked_list n_nodes = %4d, tensor_len = %4d", n_nodes, tensor_len); + case_end(stress_test_linked_list_cpu_device(backend_gpu, backend_cpu, n_nodes, tensor_len, use_device_host_buft)); + } + } - printf("%s\n", ok ? "OK" : "FAIL"); + printf("\n"); - n_ok += ok; - n_test++; + for (int n_lanes : { 8, 16 }) { + for (int n_rounds : { 5, 13, 64 }) { + for (int tensor_len : { 2, 4096 }) { + case_begin("test_dag n_lanes = %2d, n_rounds = %2d, tensor_len = %4d", n_lanes, n_rounds, tensor_len); + case_end(stress_test_dag(backends, n_lanes, n_rounds, tensor_len, use_device_host_buft)); + } } } printf("\n"); - for (int n_rounds : { 1, 2, 3, 8, 33, 64 }) { - for (int ne : { 1, 4096 }) { - printf(" n_rounds = %4d, ne = %4d: ", n_rounds, ne); - fflush(stdout); - - const bool ok = test_multi_lane(backend_gpu, backend_cpu, n_rounds, ne, use_device_host_buft); + if (have_sleep) { + for (int32_t sleep_us : { 100000}) { + for (int tensor_len : { 2, 4096 }) { + case_begin("test_inputless_splits_scheduling sleep_us = %6d, tensor_len = %4d", sleep_us, tensor_len); + case_end(test_inputless_splits_scheduling(backend_gpu, backend_cpu, tensor_len, sleep_us, use_device_host_buft)); + } + } - printf("%s\n", ok ? "OK" : "FAIL"); + printf("\n"); + } - n_ok += ok; - n_test++; - } + for (int tensor_len : { 2, 4096 }) { + case_begin("test_chain_all_backends tensor_len = %4d", tensor_len); + case_end(test_chain_all_backends(backends, tensor_len, use_device_host_buft)); } printf("\n"); - for (int32_t sleep_us : { 100000, 0 }) { - for (int ne : { 1, 4096 }) { - printf(" sleep_race sleep_us = %6d, ne = %4d: ", sleep_us, ne); - fflush(stdout); - - const bool ok = test_sleep_race(backend_gpu, backend_cpu, ne, sleep_us, use_device_host_buft); + // every ordered pair of backends is the sender and the receiver of a copy + for (size_t b_send = 0; b_send < backends.size(); b_send++) { + for (size_t b_recv = 0; b_recv < backends.size(); b_recv++) { + if (b_send == b_recv) { + continue; + } - printf("%s\n", ok ? "OK" : "FAIL"); + const char * name_send = ggml_backend_name(backends[b_send]); + const char * name_recv = ggml_backend_name(backends[b_recv]); + + for (bool inputs_on_sender : { false, true }) { + for (bool parallel : { false, true }) { + for (int tensor_len : { 1, 4096 }) { + // todo aendk fix up. + case_begin("test_pair_user_inputs %-8s -> %-8s inputs on %-8s parallel = %d, tensor_len = %4d", + name_send, name_recv, inputs_on_sender ? "sender" : "receiver", parallel, tensor_len); + case_end(test_pair_user_inputs(backends, (int) b_send, (int) b_recv, tensor_len, + /*n_inputs =*/ 4, inputs_on_sender, parallel, use_device_host_buft)); + } + } + } - n_ok += ok; - n_test++; + for (int tensor_len : { 1, 4096 }) { + case_begin("test_y_shaped %-8s -> %-8s tensor_len = %4d", name_send, name_recv, tensor_len); + case_end(test_y_shaped_graph(backends, (int) b_send, (int) b_recv, tensor_len, use_device_host_buft)); + } } } printf("\n"); } - ggml_backend_free(backend_gpu); - ggml_backend_free(backend_cpu); + for (ggml_backend_t backend : backends) { + ggml_backend_free(backend); + } printf("%d/%d tests passed\n", n_ok, n_test); From cd53112a529fd3bcfd7e669724bf7eca3f99d1e7 Mon Sep 17 00:00:00 2001 From: aendk Date: Mon, 17 Aug 2026 11:48:19 +0200 Subject: [PATCH 06/11] Cleanup: Removes CUDA specific comments and nvtx-markers --- tests/test-backend-sched.cpp | 39 +++++++----------------------------- 1 file changed, 7 insertions(+), 32 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index ff5fbb09d3e9..4ad0251b98dd 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -1,12 +1,11 @@ -// The tests below define the scheduler and backend behavior for every combination of asynchronous backends -// On a machine with CPU and CUDA backends only, that means CPU->CPU, CPU->CUDA, CUDA->CPU and CUDA->CUDA +// Test suite for scheduler and backend behavior -// Formalized scheduling behavior: +// Towards formalized scheduling behavior: // - a single inference pass can run on several backends. The subset of nodes running on a single backend is a split // - synchronous backends: // - explicit synchronization (ggml_backend_synchronize()) is required between each operation // - asynchronous backends: -// - between splits, activations are copied asynchronously (output of split N -> input of split N+1) +// - Activations from one split can be copied asynchronously to the next (output of split N -> input of split N+1) // - Several scheduling patterns must be supported by async backends. The scheduler may: // - not explicitly synchronize between CPU->backend memcpy and graph execution on backend // - dispatch several parallel memcpys to the same backend at once @@ -23,8 +22,8 @@ #include #include -#include +// global counters for pass/fail static int n_ok = 0; static int n_test = 0; @@ -238,10 +237,6 @@ struct backend_consts { // stress test: a tensor is incremented and sent back and forth between a backend and CPU in a ping-pong pattern static bool stress_test_linked_list_cpu_device(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t tensor_len, bool use_device_host_buft) { - nvtx3::scoped_range sc_4{nvtx3::event_attributes{nvtx3::rgb{0, 0, 255}, // blue - ("stress_test_linked_list_cpu_device " + std::string(ggml_backend_name(backend_gpu)) + " + " + ggml_backend_name(backend_cpu) + - " n_nodes=" + std::to_string(n_nodes) + " ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const std::vector backends = { backend_gpu, backend_cpu }; @@ -308,10 +303,6 @@ static bool stress_test_linked_list_cpu_device(ggml_backend_t backend_gpu, ggml_ // - splits taking one activation per lane from the previous split, one of them from an older split instead static bool stress_test_dag(const std::vector & backends, int n_lanes, int n_rounds, int64_t tensor_len, bool use_device_host_buft) { - nvtx3::scoped_range sc_5{nvtx3::event_attributes{nvtx3::rgb{255, 165, 0}, // orange - ("stress_test_dag n_lanes=" + std::to_string(n_lanes) + - " n_rounds=" + std::to_string(n_rounds) + " ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const int n_backends = (int) backends.size(); @@ -443,12 +434,8 @@ static bool stress_test_dag(const std::vector & backends, int n_ // CPU: produce {66} // GPU: increment both {99} and {66} -> {100} and {67} // Correct result is thus {100}, incorrect output is {111} when {55} was overwritten by {66} before the copy started. -// Note: currently only reproducible on async H2D copy (pinned memory in CUDA). +// Note: currently only reproducible on async H2D copy (=pinned memory). static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int64_t tensor_len, int32_t sleep_us, bool use_device_host_buft) { - nvtx3::scoped_range sc_6{nvtx3::event_attributes{nvtx3::rgb{255, 0, 0}, // red - ("test_inputless_splits_scheduling " + std::string(ggml_backend_name(backend_gpu)) + " + " + ggml_backend_name(backend_cpu) + - " sleep_us=" + std::to_string(sleep_us) + " ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const int GPU = 0; const int CPU = 1; @@ -560,9 +547,6 @@ static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_ba // Test that all async backends transmit their activations to the following async backend correctly. No user inputs tested. static bool test_chain_all_backends(const std::vector & backends, int64_t tensor_len, bool use_device_host_buft) { - nvtx3::scoped_range sc_7{nvtx3::event_attributes{nvtx3::rgb{0, 255, 0}, // green - ("test_chain_all_backends ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const int n_backends = (int) backends.size(); @@ -620,11 +604,6 @@ static bool test_chain_all_backends(const std::vector & backends static bool test_pair_user_inputs(const std::vector & backends, int b_send, int b_recv, int64_t tensor_len, int n_inputs, bool inputs_on_sender, bool parallel, bool use_device_host_buft) { - nvtx3::scoped_range sc_8{nvtx3::event_attributes{nvtx3::rgb{255, 215, 0}, // gold - ("test_pair_user_inputs " + std::string(ggml_backend_name(backends[b_send])) + " -> " + ggml_backend_name(backends[b_recv]) + - (inputs_on_sender ? " inputs on sender" : " inputs on receiver") + - " parallel=" + std::to_string(parallel) + " ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const size_t graph_size = 64; @@ -704,10 +683,6 @@ static bool test_pair_user_inputs(const std::vector & backends, static bool test_y_shaped_graph(const std::vector & backends, int backend_a, int backend_b, int64_t tensor_len, bool use_device_host_buft) { - nvtx3::scoped_range sc_9{nvtx3::event_attributes{nvtx3::rgb{238, 130, 238}, // violet - ("test_y_shaped_graph " + std::string(ggml_backend_name(backends[backend_a])) + " -> " + ggml_backend_name(backends[backend_b]) + - " ne=" + std::to_string(tensor_len) + - (use_device_host_buft ? " device_host" : " pageable")).c_str()}}; const size_t graph_size = 64; @@ -845,8 +820,8 @@ int main() { printf("\n"); if (have_sleep) { - for (int32_t sleep_us : { 100000}) { - for (int tensor_len : { 2, 4096 }) { + for (int32_t sleep_us : {50, 60, 70, 80, 400, 500, 600, 700, 1000, 10000}) { + for (int tensor_len : { 2, 2048, 4096, 8192, 1600}) { case_begin("test_inputless_splits_scheduling sleep_us = %6d, tensor_len = %4d", sleep_us, tensor_len); case_end(test_inputless_splits_scheduling(backend_gpu, backend_cpu, tensor_len, sleep_us, use_device_host_buft)); } From d58702be72ddf87eb1ccd0fd28ff3c82cc0394f0 Mon Sep 17 00:00:00 2001 From: aendk Date: Mon, 17 Aug 2026 14:32:10 +0200 Subject: [PATCH 07/11] Fix: forgot to add sleep specific test case --- tests/test-sleep.cpp | 163 +++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 163 insertions(+) create mode 100644 tests/test-sleep.cpp diff --git a/tests/test-sleep.cpp b/tests/test-sleep.cpp new file mode 100644 index 000000000000..3bcf53d21fe6 --- /dev/null +++ b/tests/test-sleep.cpp @@ -0,0 +1,163 @@ +// Tests for GGML_OP_SLEEP, on every backend that supports it. Two properties are checked: +// - the op is a pass-through: the output must be a bit-for-bit copy of the input +// - the op does not return early: a chain of n sleeps must take at least n*us +// Only a lower bound is asserted for the duration. A busy-wait cannot finish early, so that bound is +// deterministic, whereas an upper bound would be flaky on a loaded machine. + +#include "ggml.h" +#include "ggml-alloc.h" +#include "ggml-backend.h" + +#include +#include +#include +#include +#include + +enum test_result { + TEST_OK, + TEST_FAIL, + TEST_SKIP, +}; + +static test_result test_sleep(ggml_backend_t backend, ggml_type type, int64_t ne, int32_t us, int n_nodes) { + const size_t graph_size = n_nodes + 16; + + ggml_init_params params_static = { + /*.mem_size =*/ ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_static = ggml_init(params_static); + + ggml_tensor * x = ggml_new_tensor_1d(ctx_static, type, ne); + ggml_set_name(x, "x"); + ggml_set_input(x); + + ggml_init_params params_compute = { + /*.mem_size =*/ n_nodes*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context * ctx_compute = ggml_init(params_compute); + ggml_cgraph * gf = ggml_new_graph_custom(ctx_compute, graph_size, false); + + ggml_tensor * out = x; + for (int i = 0; i < n_nodes; i++) { + out = ggml_sleep(ctx_compute, out, us); + } + ggml_set_name(out, "out"); + ggml_build_forward_expand(gf, out); + + // the backends log to stdout while computing, so the result line is only printed once it is complete + char desc[128]; + snprintf(desc, sizeof(desc), " %-6s %-4s ne=%-9" PRId64 " us=%-7d n_nodes=%d ", + ggml_backend_name(backend), ggml_type_name(type), ne, us, n_nodes); + + test_result result = TEST_OK; + + if (!ggml_backend_supports_op(backend, out)) { + printf("%snot supported\n", desc); + result = TEST_SKIP; + } else { + ggml_backend_buffer_t buf_static = ggml_backend_alloc_ctx_tensors(ctx_static, backend); + + // random bytes rather than random floats, so that bit patterns a value-wise copy could + // mangle, such as NaNs and denormals, are covered as well + std::vector data_in(ggml_nbytes(x)); + std::mt19937 rng(1234); + for (size_t i = 0; i < data_in.size(); i++) { + data_in[i] = rng() & 0xFF; + } + ggml_backend_tensor_set(x, data_in.data(), 0, data_in.size()); + + ggml_gallocr_t galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend)); + ggml_gallocr_alloc_graph(galloc, gf); + + // warm up, so that one-time costs such as lazy CUDA module loading are not timed + ggml_backend_graph_compute(backend, gf); + + const int64_t t_start_us = ggml_time_us(); + const ggml_status status = ggml_backend_graph_compute(backend, gf); + const int64_t t_us = ggml_time_us() - t_start_us; + + std::vector data_out(ggml_nbytes(out)); + ggml_backend_tensor_get(out, data_out.data(), 0, data_out.size()); + + const int64_t t_min_us = (int64_t) n_nodes * us; + + const bool ok_status = status == GGML_STATUS_SUCCESS; + const bool ok_data = memcmp(data_in.data(), data_out.data(), data_in.size()) == 0; + const bool ok_time = t_us >= t_min_us; + + printf("%selapsed=%7" PRId64 " us (>= %6" PRId64 ") data: %-4s time: %s\n", + desc, t_us, t_min_us, ok_data ? "OK" : "FAIL", ok_time ? "OK" : "FAIL"); + + if (!ok_status) { + printf(" compute failed: %s\n", ggml_status_to_string(status)); + } + + result = ok_status && ok_data && ok_time ? TEST_OK : TEST_FAIL; + + ggml_gallocr_free(galloc); + ggml_backend_buffer_free(buf_static); + } + + ggml_free(ctx_compute); + ggml_free(ctx_static); + + return result; +} + +int main() { + ggml_time_init(); + ggml_backend_load_all(); + + int n_ok = 0; + int n_fail = 0; + int n_skip = 0; + + for (size_t i = 0; i < ggml_backend_dev_count(); i++) { + ggml_backend_dev_t dev = ggml_backend_dev_get(i); + + const enum ggml_backend_dev_type type = ggml_backend_dev_type(dev); + if (type != GGML_BACKEND_DEVICE_TYPE_CPU && type != GGML_BACKEND_DEVICE_TYPE_GPU) { + continue; + } + + ggml_backend_t backend = ggml_backend_dev_init(dev, nullptr); + if (backend == nullptr) { + printf(" failed to initialize %s\n", ggml_backend_dev_name(dev)); + n_fail++; + continue; + } + + // a zero duration must still produce a valid copy, longer ones are checked against the clock, + // and the last case verifies that a chain of sleeps accumulates instead of collapsing into one + const struct { + ggml_type type; + int64_t ne; + int32_t us; + int n_nodes; + } cases[] = { + { GGML_TYPE_F32, 1, 0, 1 }, + { GGML_TYPE_F32, 4096, 10000, 1 }, + { GGML_TYPE_F16, 1024, 10000, 1 }, + { GGML_TYPE_F32, 1 << 20, 5000, 4 }, + }; + + for (const auto & c : cases) { + switch (test_sleep(backend, c.type, c.ne, c.us, c.n_nodes)) { + case TEST_OK: n_ok++; break; + case TEST_FAIL: n_fail++; break; + case TEST_SKIP: n_skip++; break; + } + } + + ggml_backend_free(backend); + } + + printf("%d passed, %d failed, %d skipped\n", n_ok, n_fail, n_skip); + + return n_fail > 0; +} From f159c80c0edfc8589bf57d72dfbb343dd632a2cd Mon Sep 17 00:00:00 2001 From: aendk Date: Mon, 17 Aug 2026 15:03:57 +0200 Subject: [PATCH 08/11] Fix oversights identified by CI --- ggml/include/ggml-rpc.h | 2 +- tests/test-backend-sched.cpp | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h index 276aea00ea1b..9e9baa642a62 100644 --- a/ggml/include/ggml-rpc.h +++ b/ggml/include/ggml-rpc.h @@ -11,7 +11,7 @@ extern "C" { #define RPC_PROTO_PATCH_VERSION 0 #ifdef __cplusplus -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); #endif #define GGML_RPC_MAX_SERVERS 16 diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index 4ad0251b98dd..c816afc11973 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -757,7 +757,6 @@ static bool initialize_gpu_backends(std::vector & backends, bool // todo rework this, only selects a single GPU // create struct incorporating have_device_host_buft and have_sleep flags ggml_backend_t backend_gpu = backends[0]; - ggml_backend_t backend_cpu = backends.back(); for (ggml_backend_t backend : backends) { ggml_backend_dev_t dev = ggml_backend_get_device(backend); From 615206b97c0b5f266550a378c7bf5aeab04eb609 Mon Sep 17 00:00:00 2001 From: aendk Date: Tue, 18 Aug 2026 16:16:56 +0200 Subject: [PATCH 09/11] Drops default prints so that it behaves like other ctests. Adds -v flag to toggle this behavior --- tests/test-backend-sched.cpp | 92 +++++++++++++++++++++++++++--------- 1 file changed, 70 insertions(+), 22 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index c816afc11973..20b34f2addb7 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -27,19 +28,41 @@ static int n_ok = 0; static int n_test = 0; +// default silent on success; --verbose / -v prints progress +static bool print_log = false; + +GGML_ATTRIBUTE_FORMAT(1, 2) +static void log_maybe(const char * fmt, ...) { + if (!print_log) { + return; + } + + va_list args; + va_start(args, fmt); + vprintf(fmt, args); + va_end(args); +} + // pretty-print helpers start +static constexpr int status_column = 82; // for ok/fail column alignment static int case_len = 0; -static char case_label[256]; +static char case_label[256]; // max length of a case label static void case_end(bool ok) { + n_ok += ok; + n_test++; + + if (ok && !print_log) { + case_len = 0; + return; + } + if (case_len == 0) { case_len = printf(" %s", case_label); } - printf("%*s%s\n", std::max(1, 82 - case_len), "", ok ? "OK" : "FAIL"); - - n_ok += ok; - n_test++; + printf("%*s%s\n", std::max(1, status_column - case_len), "", ok ? "OK" : "FAIL"); + case_len = 0; } static void vnote(const char * fmt, va_list args) { @@ -55,6 +78,10 @@ static void vnote(const char * fmt, va_list args) { GGML_ATTRIBUTE_FORMAT(1, 2) static void note(const char * fmt, ...) { + if (!print_log) { + return; + } + va_list args; va_start(args, fmt); vnote(fmt, args); @@ -63,6 +90,10 @@ static void note(const char * fmt, ...) { GGML_ATTRIBUTE_FORMAT(1, 2) static bool fail(const char * fmt, ...) { + if (case_len == 0) { + case_len = printf(" %s", case_label); + } + va_list args; va_start(args, fmt); vnote(fmt, args); @@ -78,9 +109,12 @@ static void case_begin(const char * fmt, ...) { vsnprintf(case_label, sizeof(case_label), fmt, args); va_end(args); - case_len = printf(" %s", case_label); - - fflush(stdout); + if (print_log) { + case_len = printf(" %s", case_label); + fflush(stdout); + } else { + case_len = 0; + } } // pretty-print helpers end @@ -604,7 +638,6 @@ static bool test_chain_all_backends(const std::vector & backends static bool test_pair_user_inputs(const std::vector & backends, int b_send, int b_recv, int64_t tensor_len, int n_inputs, bool inputs_on_sender, bool parallel, bool use_device_host_buft) { - const size_t graph_size = 64; backend_consts consts(backends, tensor_len); @@ -760,24 +793,37 @@ static bool initialize_gpu_backends(std::vector & backends, bool for (ggml_backend_t backend : backends) { ggml_backend_dev_t dev = ggml_backend_get_device(backend); - printf("backend: %-10s (%s)\n", ggml_backend_name(backend), ggml_backend_dev_description(dev)); + log_maybe("backend: %-10s (%s)\n", ggml_backend_name(backend), ggml_backend_dev_description(dev)); } have_device_host_buft = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend_gpu)) != nullptr; if (!have_device_host_buft) { - printf("GPU has no host buffer type, the device_host cases are skipped\n"); + log_maybe("GPU has no host buffer type, the device_host cases are skipped\n"); } have_sleep = backend_supports(backend_gpu, [](ggml_context * ctx, ggml_tensor * a) { return ggml_sleep(ctx, a, 0); }); if (!have_sleep) { - printf("GPU does not support GGML_OP_SLEEP, some tests are skipped\n"); + log_maybe("GPU does not support GGML_OP_SLEEP, some tests are skipped\n"); } - printf("\n"); + log_maybe("\n"); return 1; } -int main() { +int main(int argc, char ** argv) { + for (int i = 1; i < argc; i++) { + if (strcmp(argv[i], "-v") == 0 || strcmp(argv[i], "--verbose") == 0) { + print_log = true; + } else if (strcmp(argv[i], "-h") == 0 || strcmp(argv[i], "--help") == 0) { + printf("Usage: %s [-v|--verbose]\n", argv[0]); + return 0; + } else { + fprintf(stderr, "unknown argument: %s\n", argv[i]); + fprintf(stderr, "Usage: %s [-v|--verbose]\n", argv[0]); + return 1; + } + } + ggml_backend_load_all(); std::vector backends; // TODO merge into single struct @@ -796,8 +842,7 @@ int main() { continue; } - printf("=== CPU sched buft: %s ===\n\n", use_device_host_buft ? "device_host" : "pageable"); - + log_maybe("=== CPU sched buft: %s ===\n\n", use_device_host_buft ? "device_host" : "pageable"); for (int n_nodes : { 2, 5, 128, 1024 }) { for (int tensor_len : { 2, 4096 }) { case_begin("test_linked_list n_nodes = %4d, tensor_len = %4d", n_nodes, tensor_len); @@ -805,7 +850,8 @@ int main() { } } - printf("\n"); + log_maybe("\n"); + for (int n_lanes : { 8, 16 }) { for (int n_rounds : { 5, 13, 64 }) { @@ -816,7 +862,7 @@ int main() { } } - printf("\n"); + log_maybe("\n"); if (have_sleep) { for (int32_t sleep_us : {50, 60, 70, 80, 400, 500, 600, 700, 1000, 10000}) { @@ -826,7 +872,7 @@ int main() { } } - printf("\n"); + log_maybe("\n"); } for (int tensor_len : { 2, 4096 }) { @@ -834,7 +880,7 @@ int main() { case_end(test_chain_all_backends(backends, tensor_len, use_device_host_buft)); } - printf("\n"); + log_maybe("\n"); // every ordered pair of backends is the sender and the receiver of a copy for (size_t b_send = 0; b_send < backends.size(); b_send++) { @@ -865,14 +911,16 @@ int main() { } } - printf("\n"); + log_maybe("\n"); } for (ggml_backend_t backend : backends) { ggml_backend_free(backend); } - printf("%d/%d tests passed\n", n_ok, n_test); + if (print_log || n_ok != n_test) { + printf("%d/%d tests passed\n", n_ok, n_test); + } return n_ok == n_test ? 0 : 1; } From 1116bbe479bb26cad8e0ba64cbbc86b6b69a57a1 Mon Sep 17 00:00:00 2001 From: aendk Date: Tue, 18 Aug 2026 16:46:04 +0200 Subject: [PATCH 10/11] Fix: Early exit in CPU-only settings (cause for graviton CI fails) --- tests/test-backend-sched.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index 20b34f2addb7..eacf5736fc26 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -830,13 +830,15 @@ int main(int argc, char ** argv) { bool have_device_host_buft = false; bool have_sleep = false; bool gpu_initialized = initialize_gpu_backends(backends, have_device_host_buft, have_sleep); - ggml_backend_t backend_gpu = backends[0]; - ggml_backend_t backend_cpu = backends.back(); - if (!gpu_initialized) { + // nothing to test with synchronous CPU backend only + if (!gpu_initialized || backends.size() < 2) { return 0; } + ggml_backend_t backend_gpu = backends[0]; + ggml_backend_t backend_cpu = backends.back(); + for (bool use_device_host_buft : { false, true }) { if (use_device_host_buft && !have_device_host_buft) { continue; From e7287598404ee04603729a93f10e7eef261f9c72 Mon Sep 17 00:00:00 2001 From: aendk Date: Thu, 10 Sep 2026 18:49:41 +0200 Subject: [PATCH 11/11] Improves stress_test_linked_list by removing overlap with test_chain_all_backends and test_pair_user_inputs. The latter ones test either activations or user inputs in isolation and just once, whereas the stress test tests all combinations multiple rounds --- tests/test-backend-sched.cpp | 352 ++++++++++++++++++++++------------- 1 file changed, 222 insertions(+), 130 deletions(-) diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp index eacf5736fc26..ff8a25a85a6c 100644 --- a/tests/test-backend-sched.cpp +++ b/tests/test-backend-sched.cpp @@ -129,17 +129,26 @@ static ggml_backend_buffer_type_t sched_cpu_buft(ggml_backend_t backend_gpu, ggm return ggml_backend_get_default_buffer_type(backend_cpu); } -static ggml_backend_sched_t create_test_scheduler(std::vector backends, size_t graph_size, +// a backend together with the capabilities that decide which tests can run on it +struct sched_backend_caps { + ggml_backend_t backend = nullptr; + bool have_device_host_buft = false; + bool have_sleep = false; +}; + +static ggml_backend_sched_t create_test_scheduler(const std::vector & backends_w_caps, size_t graph_size, bool use_device_host_buft, bool parallel = false) { - std::vector bufts(backends.size()); - for (size_t b = 0; b < backends.size(); b++) { - bufts[b] = ggml_backend_get_default_buffer_type(backends[b]); + std::vector backend_handles(backends_w_caps.size()); + std::vector bufts(backends_w_caps.size()); + for (size_t b = 0; b < backends_w_caps.size(); b++) { + backend_handles[b] = backends_w_caps[b].backend; + bufts[b] = ggml_backend_get_default_buffer_type(backends_w_caps[b].backend); } - // CPU backend is always last - bufts.back() = sched_cpu_buft(backends[0], backends.back(), use_device_host_buft); + // sets either pinned or paged memory + bufts.back() = sched_cpu_buft(backends_w_caps[0].backend, backends_w_caps.back().backend, use_device_host_buft); - return ggml_backend_sched_new(backends.data(), bufts.data(), (int) backends.size(), graph_size, parallel, /*op_offload =*/ false); + return ggml_backend_sched_new(backend_handles.data(), bufts.data(), (int) backend_handles.size(), graph_size, parallel, /*op_offload =*/ false); } // the nodes of the graph together with the backend each of them is assigned to @@ -167,7 +176,7 @@ struct sched_check { // executes the graph and evaluates results static bool run_and_check(ggml_backend_sched_t sched, const sched_graph & g, - const std::vector & backends, const std::vector & checks, int64_t ne) { + const std::vector & backends_w_caps, const std::vector & checks, int64_t ne) { int n_splits_expected = g.nodes.empty() ? 0 : 1; for (size_t i = 1; i < g.backend_id.size(); i++) { @@ -176,7 +185,7 @@ static bool run_and_check(ggml_backend_sched_t sched, const sched_graph & g, ggml_backend_sched_reset(sched); for (size_t i = 0; i < g.nodes.size(); i++) { - ggml_backend_sched_set_tensor_backend(sched, g.nodes[i], backends[g.backend_id[i]]); + ggml_backend_sched_set_tensor_backend(sched, g.nodes[i], backends_w_caps[g.backend_id[i]].backend); } if (!ggml_backend_sched_alloc_graph(sched, g.gf)) { @@ -229,10 +238,10 @@ struct backend_consts { std::vector zero; std::vector one; - backend_consts(const std::vector & backends, int64_t ne) { + backend_consts(const std::vector & backends_w_caps, int64_t ne) { std::vector data(ne); - for (size_t b = 0; b < backends.size(); b++) { + for (size_t b = 0; b < backends_w_caps.size(); b++) { ggml_init_params params = { /*.mem_size =*/ 2*ggml_tensor_overhead(), /*.mem_buffer =*/ nullptr, @@ -241,12 +250,12 @@ struct backend_consts { ggml_context * ctx = ggml_init(params); ggml_tensor * z = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, ne); - ggml_format_name(z, "zero_%s", ggml_backend_name(backends[b])); + ggml_format_name(z, "zero_%s", ggml_backend_name(backends_w_caps[b].backend)); ggml_tensor * o = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, ne); - ggml_format_name(o, "one_%s", ggml_backend_name(backends[b])); + ggml_format_name(o, "one_%s", ggml_backend_name(backends_w_caps[b].backend)); - ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backends[b]); + ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backends_w_caps[b].backend); std::fill(data.begin(), data.end(), 0.0f); ggml_backend_tensor_set(z, data.data(), 0, ggml_nbytes(z)); @@ -269,40 +278,82 @@ struct backend_consts { } }; -// stress test: a tensor is incremented and sent back and forth between a backend and CPU in a ping-pong pattern -static bool stress_test_linked_list_cpu_device(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int n_nodes, int64_t tensor_len, bool use_device_host_buft) { +// Creates a vector which holds a chain of backend ids. There, every backend is sender to and receiver from every other backend, to test all data transfers. +static std::vector create_all_to_all_chain(int n_backends) { + std::vector chain; + for (int i = 0; i < n_backends; i++) { + for (int j = i + 1; j < n_backends; j++) { + chain.push_back(i); + chain.push_back(j); + } + } + return chain; +} - const std::vector backends = { backend_gpu, backend_cpu }; +// stress test: a tensor is incremented and sent back and forth between the backends in a ping-pong pattern. +// One lap tests data transfers from any backend to any other backend, with and without user inputs. +static bool stress_test_linked_list(const std::vector & backends_w_caps, int n_laps, + int64_t tensor_len, bool use_device_host_buft) { + const int n_backends = (int) backends_w_caps.size(); - const size_t graph_size = n_nodes + 2; // see sched->hash_set FIXME, 2+ needed to account for leafs + const std::vector chain = create_all_to_all_chain(n_backends); + const int chain_len = (int) chain.size(); - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + // one node seeds the chain, two passes over the chain per lap follow. the extra node closes the + // wrap-around pair of every pass + const int n_nodes = 2*chain_len*n_laps + 1; - // the inputs are allocated separately so that they can be written before the graph is computed - ggml_init_params params_static = { - /*.mem_size =*/ 2*ggml_tensor_overhead(), - /*.mem_buffer =*/ nullptr, - /*.no_alloc =*/ true, - }; - ggml_context * ctx_static = ggml_init(params_static); + // see sched->hash_set FIXME, the set has to hold the nodes and the leafs: two constants per backend + // plus one user input per node of a second pass + const size_t graph_size = n_nodes + 2*n_backends + chain_len*n_laps; + + backend_consts consts(backends_w_caps, tensor_len); - ggml_tensor * x = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, tensor_len); - ggml_set_name(x, "x"); - ggml_set_input(x); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft); - ggml_tensor * one = ggml_new_tensor_1d(ctx_static, GGML_TYPE_F32, tensor_len); - ggml_set_name(one, "one"); - ggml_set_input(one); + // the user input of a second pass node is placed on the backend that sends it the activation, so that the + // receiving split has to copy both of them + std::vector> inputs_of(n_backends); + std::vector input_value(n_nodes, 0); - ggml_backend_buffer_t buf_static = ggml_backend_alloc_ctx_tensors(ctx_static, backend_cpu); + for (int i = 1; i < n_nodes; i++) { + const int pass = (i - 1)/chain_len; + if (pass % 2 == 0) { + continue; + } + + inputs_of[chain[(i - 1) % chain_len]].push_back(i); + input_value[i] = (int) (chain_len*(pass/2) + (i - 1) % chain_len + 1); + } + + // the inputs are allocated separately so that they can be written before the graph is computed + std::vector ctxs_input(n_backends); + std::vector bufs_input(n_backends); + std::vector input(n_nodes, nullptr); std::vector data(tensor_len); - std::fill(data.begin(), data.end(), 0.0f); - ggml_backend_tensor_set(x, data.data(), 0, ggml_nbytes(x)); + for (int b = 0; b < n_backends; b++) { + ggml_init_params params_input = { + /*.mem_size =*/ inputs_of[b].size()*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ctxs_input[b] = ggml_init(params_input); + + for (int i : inputs_of[b]) { + input[i] = ggml_new_tensor_1d(ctxs_input[b], GGML_TYPE_F32, tensor_len); + ggml_format_name(input[i], "input%d", i); + ggml_set_input(input[i]); + } - std::fill(data.begin(), data.end(), 1.0f); - ggml_backend_tensor_set(one, data.data(), 0, ggml_nbytes(one)); + bufs_input[b] = ggml_backend_alloc_ctx_tensors(ctxs_input[b], backends_w_caps[b].backend); + + for (int i : inputs_of[b]) { + std::fill(data.begin(), data.end(), float(input_value[i])); + ggml_backend_tensor_set(input[i], data.data(), 0, ggml_nbytes(input[i])); + } + } ggml_init_params params_compute = { /*.mem_size =*/ (n_nodes + 2)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), @@ -313,41 +364,63 @@ static bool stress_test_linked_list_cpu_device(ggml_backend_t backend_gpu, ggml_ sched_graph g(ctx_compute, graph_size); - ggml_tensor * out = x; + int64_t sum = 0; + ggml_tensor * out = nullptr; + for (int i = 0; i < n_nodes; i++) { - out = g.add(ggml_add(ctx_compute, out, one), i % 2); + const int b = chain[i % chain_len]; + + if (i == 0) { + out = ggml_add(ctx_compute, consts.zero[b], consts.one[b]); + sum += 1; + } else if (input[i] == nullptr) { + out = ggml_add(ctx_compute, out, consts.one[b]); + sum += 1; + } else { + out = ggml_add(ctx_compute, out, input[i]); + sum += input_value[i]; + } + + g.add(out, b); } ggml_set_output(out); - const bool ok = run_and_check(sched, g, backends, {{ out, float(n_nodes), "out" }}, tensor_len); + bool ok = sum < (1 << 24); // the counter has to stay exact in f32 + if (!ok) { + fail("expected result %" PRId64 " does not fit into the f32 mantissa, lower n_laps", sum); + } else { + ok = run_and_check(sched, g, backends_w_caps, {{ out, float(sum), "out" }}, tensor_len); + } ggml_backend_sched_free(sched); ggml_free(ctx_compute); - ggml_backend_buffer_free(buf_static); - ggml_free(ctx_static); + for (int b = 0; b < n_backends; b++) { + ggml_backend_buffer_free(bufs_input[b]); + ggml_free(ctxs_input[b]); + } return ok; } -// Same idea as stress_test_linked_list_cpu_device, but with the compute graph being a direct acyclic graph (DAG). +// Same idea as stress_test_linked_list, but with the compute graph being a direct acyclic graph (DAG). // n_lanes independent counters run through rounds of four shapes, so that one graph covers: // - lanes with different histories, so a misrouted copy shows up as a wrong count in a single lane // - one split producing two values followed by two splits that each consume one of them and do not depend on each other // - splits with no inputs at all, from lanes re-seeded out of constants that already live on the split's backend // - splits taking one activation per lane from the previous split, one of them from an older split instead -static bool stress_test_dag(const std::vector & backends, int n_lanes, int n_rounds, +static bool stress_test_dag(const std::vector & backends_w_caps, int n_lanes, int n_rounds, int64_t tensor_len, bool use_device_host_buft) { - const int n_backends = (int) backends.size(); + const int n_backends = (int) backends_w_caps.size(); GGML_ASSERT(n_lanes % 2 == 0); // sized for the worst case const size_t graph_size = n_lanes*(2*n_rounds + 1) + 2*n_backends; - backend_consts consts(backends, tensor_len); + backend_consts consts(backends_w_caps, tensor_len); - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft); ggml_init_params params_compute = { /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), @@ -449,7 +522,7 @@ static bool stress_test_dag(const std::vector & backends, int n_ GGML_ASSERT(c.expected < float(1 << 24)); // the counts have to stay exact in f32 } - const bool ok = run_and_check(sched, g, backends, checks, tensor_len); + const bool ok = run_and_check(sched, g, backends_w_caps, checks, tensor_len); ggml_backend_sched_free(sched); ggml_free(ctx_compute); @@ -469,16 +542,16 @@ static bool stress_test_dag(const std::vector & backends, int n_ // GPU: increment both {99} and {66} -> {100} and {67} // Correct result is thus {100}, incorrect output is {111} when {55} was overwritten by {66} before the copy started. // Note: currently only reproducible on async H2D copy (=pinned memory). -static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_backend_t backend_cpu, int64_t tensor_len, int32_t sleep_us, bool use_device_host_buft) { +static bool test_inputless_splits_scheduling(const sched_backend_caps & gpu, const sched_backend_caps & cpu, int64_t tensor_len, int32_t sleep_us, bool use_device_host_buft) { const int GPU = 0; const int CPU = 1; - const std::vector backends = { backend_gpu, backend_cpu }; + const std::vector backends_w_caps = { gpu, cpu }; const size_t graph_size = 64; - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft); ggml_init_params params_static = { /*.mem_size =*/ 4*ggml_tensor_overhead(), @@ -497,7 +570,7 @@ static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_ba ggml_tensor * val66_cpu = ggml_new_tensor_1d(ctx_cpu, GGML_TYPE_F32, tensor_len); ggml_set_name(val66_cpu, "val66_cpu"); - ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, backend_cpu); + ggml_backend_buffer_t buf_cpu = ggml_backend_alloc_ctx_tensors(ctx_cpu, cpu.backend); ggml_context * ctx_gpu = ggml_init(params_static); @@ -510,7 +583,7 @@ static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_ba ggml_tensor * val44_gpu = ggml_new_tensor_1d(ctx_gpu, GGML_TYPE_F32, tensor_len); ggml_set_name(val44_gpu, "val44_gpu"); - ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, backend_gpu); + ggml_backend_buffer_t buf_gpu = ggml_backend_alloc_ctx_tensors(ctx_gpu, gpu.backend); std::vector data(tensor_len); @@ -560,7 +633,7 @@ static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_ba ggml_set_output(out99); ggml_set_output(out66); - const bool ok = run_and_check(sched, g, backends, { + const bool ok = run_and_check(sched, g, backends_w_caps, { { delayed, 1.0f, "delayed" }, { out99, 100.0f, "out99" }, { out66, 67.0f, "out66" }, @@ -580,9 +653,9 @@ static bool test_inputless_splits_scheduling(ggml_backend_t backend_gpu, ggml_ba } // Test that all async backends transmit their activations to the following async backend correctly. No user inputs tested. -static bool test_chain_all_backends(const std::vector & backends, int64_t tensor_len, bool use_device_host_buft) { +static bool test_chain_all_backends(const std::vector & backends_w_caps, int64_t tensor_len, bool use_device_host_buft) { - const int n_backends = (int) backends.size(); + const int n_backends = (int) backends_w_caps.size(); std::vector seq; @@ -604,9 +677,9 @@ static bool test_chain_all_backends(const std::vector & backends const size_t n_nodes = seq.size(); const size_t graph_size = n_nodes + n_backends + 1; // see sched->hash_set FIXME, one "one" per backend plus the "zero" that starts the chain - backend_consts consts(backends, tensor_len); + backend_consts consts(backends_w_caps, tensor_len); - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft); ggml_init_params params_compute = { /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), @@ -625,7 +698,7 @@ static bool test_chain_all_backends(const std::vector & backends } ggml_set_output(out); - const bool ok = run_and_check(sched, g, backends, {{ out, float(n_nodes), "out" }}, tensor_len); + const bool ok = run_and_check(sched, g, backends_w_caps, {{ out, float(n_nodes), "out" }}, tensor_len); ggml_backend_sched_free(sched); ggml_free(ctx_compute); @@ -635,14 +708,14 @@ static bool test_chain_all_backends(const std::vector & backends // Tests data transfer between all combinations of backend pairs // Always tests between two backends only with a single activation and 4 parallel user inputs. -static bool test_pair_user_inputs(const std::vector & backends, int b_send, int b_recv, int64_t tensor_len, +static bool test_pair_user_inputs(const std::vector & backends_w_caps, int b_send, int b_recv, int64_t tensor_len, int n_inputs, bool inputs_on_sender, bool parallel, bool use_device_host_buft) { const size_t graph_size = 64; - backend_consts consts(backends, tensor_len); + backend_consts consts(backends_w_caps, tensor_len); - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft, parallel); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft, parallel); // placing the inputs on the sender makes the receiving split copy all of them, placing them on the // receiver leaves the activation as its only input @@ -663,7 +736,7 @@ static bool test_pair_user_inputs(const std::vector & backends, inputs.push_back(in); } - ggml_backend_buffer_t buf_inputs = ggml_backend_alloc_ctx_tensors(ctx_inputs, backends[b_inputs]); + ggml_backend_buffer_t buf_inputs = ggml_backend_alloc_ctx_tensors(ctx_inputs, backends_w_caps[b_inputs].backend); std::vector data(tensor_len); std::fill(data.begin(), data.end(), 1.0f); @@ -687,7 +760,7 @@ static bool test_pair_user_inputs(const std::vector & backends, } ggml_set_output(out); - bool ok = run_and_check(sched, g, backends, {{ out, float(1 + n_inputs), "out" }}, tensor_len); + bool ok = run_and_check(sched, g, backends_w_caps, {{ out, float(1 + n_inputs), "out" }}, tensor_len); for (ggml_tensor * in : inputs) { if (!ok) { @@ -713,15 +786,15 @@ static bool test_pair_user_inputs(const std::vector & backends, // Tests Y-shaped scheduling: two parallel lanes merging into 1. Lane A and B, merging into a single split. // The lanes join in a final split on b_send that receives one activation from each of the two b_recv splits. // TODO improve function signature + hoist backend_consts out of it? -static bool test_y_shaped_graph(const std::vector & backends, int backend_a, int backend_b, int64_t tensor_len, +static bool test_y_shaped_graph(const std::vector & backends_w_caps, int backend_a, int backend_b, int64_t tensor_len, bool use_device_host_buft) { const size_t graph_size = 64; - backend_consts consts(backends, tensor_len); + backend_consts consts(backends_w_caps, tensor_len); - ggml_backend_sched_t sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + ggml_backend_sched_t sched = create_test_scheduler(backends_w_caps, graph_size, use_device_host_buft); ggml_init_params params_compute = { /*.mem_size =*/ (graph_size + 8)*ggml_tensor_overhead() + ggml_graph_overhead_custom(graph_size, false), @@ -749,7 +822,7 @@ static bool test_y_shaped_graph(const std::vector & backends, in ggml_set_name(out, "out"); ggml_set_output(out); - const bool ok = run_and_check(sched, g, backends, {{ out, 4.0f, "out" }}, tensor_len); + const bool ok = run_and_check(sched, g, backends_w_caps, {{ out, 4.0f, "out" }}, tensor_len); ggml_backend_sched_free(sched); ggml_free(ctx_compute); @@ -757,16 +830,29 @@ static bool test_y_shaped_graph(const std::vector & backends, in return ok; } -static bool initialize_gpu_backends(std::vector & backends, bool & have_device_host_buft, bool & have_sleep) { +static const char * dev_type_name(enum ggml_backend_dev_type type) { + switch (type) { + case GGML_BACKEND_DEVICE_TYPE_CPU: return "CPU"; + case GGML_BACKEND_DEVICE_TYPE_GPU: return "GPU"; + case GGML_BACKEND_DEVICE_TYPE_IGPU: return "IGPU"; + case GGML_BACKEND_DEVICE_TYPE_ACCEL: return "ACCEL"; + case GGML_BACKEND_DEVICE_TYPE_META: return "META"; + } + return "UNKNOWN"; +} + +static bool initialize_backends(std::vector & backends_w_caps) { // cf. GGML_SCHED_MAX_BACKENDS const size_t max_backends = 16; - for (size_t i = 0; i < ggml_backend_dev_count() && backends.size() + 1 < max_backends; i++) { + for (size_t i = 0; i < ggml_backend_dev_count() && backends_w_caps.size() + 1 < max_backends; i++) { ggml_backend_dev_t dev = ggml_backend_dev_get(i); + log_maybe("device %2zu: %-10s %-5s (%s)\n", i, ggml_backend_dev_name(dev), + dev_type_name(ggml_backend_dev_type(dev)), ggml_backend_dev_description(dev)); + const enum ggml_backend_dev_type type = ggml_backend_dev_type(dev); - //GPU or IGPU only, for now - if (type != GGML_BACKEND_DEVICE_TYPE_GPU && type != GGML_BACKEND_DEVICE_TYPE_IGPU) { + if (type == GGML_BACKEND_DEVICE_TYPE_CPU) { continue; } @@ -776,34 +862,26 @@ static bool initialize_gpu_backends(std::vector & backends, bool continue; } - backends.push_back(backend); + backends_w_caps.push_back({ backend }); } - if (backends.empty()) { - printf("no GPU device found, skipping\n"); + if (backends_w_caps.empty()) { + printf("no non-CPU backend found, skipping\n"); return 0; } // ggml_backend_sched_new requires the CPU backend to be the last one - backends.push_back(ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr)); - GGML_ASSERT(backends.back() != nullptr); - // todo rework this, only selects a single GPU - // create struct incorporating have_device_host_buft and have_sleep flags - ggml_backend_t backend_gpu = backends[0]; + backends_w_caps.push_back({ ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr) }); + GGML_ASSERT(backends_w_caps.back().backend != nullptr); - for (ggml_backend_t backend : backends) { - ggml_backend_dev_t dev = ggml_backend_get_device(backend); - log_maybe("backend: %-10s (%s)\n", ggml_backend_name(backend), ggml_backend_dev_description(dev)); - } + for (sched_backend_caps & sched_back_caps : backends_w_caps) { + ggml_backend_dev_t dev = ggml_backend_get_device(sched_back_caps.backend); - have_device_host_buft = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend_gpu)) != nullptr; - if (!have_device_host_buft) { - log_maybe("GPU has no host buffer type, the device_host cases are skipped\n"); - } + sched_back_caps.have_device_host_buft = ggml_backend_dev_host_buffer_type(dev) != nullptr; + sched_back_caps.have_sleep = backend_supports(sched_back_caps.backend, [](ggml_context * ctx, ggml_tensor * a) { return ggml_sleep(ctx, a, 0); }); - have_sleep = backend_supports(backend_gpu, [](ggml_context * ctx, ggml_tensor * a) { return ggml_sleep(ctx, a, 0); }); - if (!have_sleep) { - log_maybe("GPU does not support GGML_OP_SLEEP, some tests are skipped\n"); + log_maybe("backend: %-10s host_buft = %d, sleep = %d (%s)\n", ggml_backend_name(sched_back_caps.backend), + sched_back_caps.have_device_host_buft, sched_back_caps.have_sleep, ggml_backend_dev_description(dev)); } log_maybe("\n"); return 1; @@ -826,73 +904,87 @@ int main(int argc, char ** argv) { ggml_backend_load_all(); - std::vector backends; // TODO merge into single struct - bool have_device_host_buft = false; - bool have_sleep = false; - bool gpu_initialized = initialize_gpu_backends(backends, have_device_host_buft, have_sleep); + std::vector backends_w_caps; + bool non_cpu_backends_initialized = initialize_backends(backends_w_caps); // nothing to test with synchronous CPU backend only - if (!gpu_initialized || backends.size() < 2) { + if (!non_cpu_backends_initialized || backends_w_caps.size() < 2) { return 0; } - ggml_backend_t backend_gpu = backends[0]; - ggml_backend_t backend_cpu = backends.back(); + const sched_backend_caps & cpu = backends_w_caps.back(); + const size_t n_non_cpu_backends = backends_w_caps.size() - 1; - for (bool use_device_host_buft : { false, true }) { - if (use_device_host_buft && !have_device_host_buft) { - continue; - } + // Test cases for single backends against CPU backend + for (size_t b_id = 0; b_id < n_non_cpu_backends; b_id++) { + const sched_backend_caps & backend_w_caps = backends_w_caps[b_id]; + const char * name_backend = ggml_backend_name(backend_w_caps.backend); - log_maybe("=== CPU sched buft: %s ===\n\n", use_device_host_buft ? "device_host" : "pageable"); - for (int n_nodes : { 2, 5, 128, 1024 }) { - for (int tensor_len : { 2, 4096 }) { - case_begin("test_linked_list n_nodes = %4d, tensor_len = %4d", n_nodes, tensor_len); - case_end(stress_test_linked_list_cpu_device(backend_gpu, backend_cpu, n_nodes, tensor_len, use_device_host_buft)); + for (bool use_device_host_buft : { false, true }) { + if (use_device_host_buft && !backend_w_caps.have_device_host_buft) { + continue; + } + + log_maybe("=== Backend: %s, CPU sched buft: %s ===\n\n", name_backend, use_device_host_buft ? "device_host" : "pageable"); + + if (backend_w_caps.have_sleep) { + for (int32_t sleep_us : {50, 60, 70, 80, 400, 500, 600, 700, 1000, 10000}) { + for (int tensor_len : { 2, 2048, 4096, 8192, 1600}) { + case_begin("test_inputless_splits_scheduling %-8s sleep_us = %6d, tensor_len = %4d", name_backend, sleep_us, tensor_len); + case_end(test_inputless_splits_scheduling(backend_w_caps, cpu, tensor_len, sleep_us, use_device_host_buft)); + } + } + + log_maybe("\n"); } } + } - log_maybe("\n"); + // Test cases for all present backend subsets + for (bool use_device_host_buft : { false, true }) { + // create_test_scheduler takes the host buffer type from the first backend + if (use_device_host_buft && !backends_w_caps[0].have_device_host_buft) { + continue; + } + log_maybe("=== all backends, CPU sched buft: %s ===\n\n", use_device_host_buft ? "device_host" : "pageable"); for (int n_lanes : { 8, 16 }) { for (int n_rounds : { 5, 13, 64 }) { for (int tensor_len : { 2, 4096 }) { - case_begin("test_dag n_lanes = %2d, n_rounds = %2d, tensor_len = %4d", n_lanes, n_rounds, tensor_len); - case_end(stress_test_dag(backends, n_lanes, n_rounds, tensor_len, use_device_host_buft)); + case_begin("stress_test_dag n_lanes = %2d, n_rounds = %2d, tensor_len = %4d", n_lanes, n_rounds, tensor_len); + case_end(stress_test_dag(backends_w_caps, n_lanes, n_rounds, tensor_len, use_device_host_buft)); } } } log_maybe("\n"); - if (have_sleep) { - for (int32_t sleep_us : {50, 60, 70, 80, 400, 500, 600, 700, 1000, 10000}) { - for (int tensor_len : { 2, 2048, 4096, 8192, 1600}) { - case_begin("test_inputless_splits_scheduling sleep_us = %6d, tensor_len = %4d", sleep_us, tensor_len); - case_end(test_inputless_splits_scheduling(backend_gpu, backend_cpu, tensor_len, sleep_us, use_device_host_buft)); - } - } - - log_maybe("\n"); - } - for (int tensor_len : { 2, 4096 }) { case_begin("test_chain_all_backends tensor_len = %4d", tensor_len); - case_end(test_chain_all_backends(backends, tensor_len, use_device_host_buft)); + case_end(test_chain_all_backends(backends_w_caps, tensor_len, use_device_host_buft)); + } + + log_maybe("\n"); + + for (int n_laps : { 1, 4, 64 }) { + for (int tensor_len : { 2, 4096 }) { + case_begin("stress_test_linked_list n_laps = %3d, tensor_len = %4d", n_laps, tensor_len); + case_end(stress_test_linked_list(backends_w_caps, n_laps, tensor_len, use_device_host_buft)); + } } log_maybe("\n"); // every ordered pair of backends is the sender and the receiver of a copy - for (size_t b_send = 0; b_send < backends.size(); b_send++) { - for (size_t b_recv = 0; b_recv < backends.size(); b_recv++) { + for (size_t b_send = 0; b_send < backends_w_caps.size(); b_send++) { + for (size_t b_recv = 0; b_recv < backends_w_caps.size(); b_recv++) { if (b_send == b_recv) { continue; } - const char * name_send = ggml_backend_name(backends[b_send]); - const char * name_recv = ggml_backend_name(backends[b_recv]); + const char * name_send = ggml_backend_name(backends_w_caps[b_send].backend); + const char * name_recv = ggml_backend_name(backends_w_caps[b_recv].backend); for (bool inputs_on_sender : { false, true }) { for (bool parallel : { false, true }) { @@ -900,7 +992,7 @@ int main(int argc, char ** argv) { // todo aendk fix up. case_begin("test_pair_user_inputs %-8s -> %-8s inputs on %-8s parallel = %d, tensor_len = %4d", name_send, name_recv, inputs_on_sender ? "sender" : "receiver", parallel, tensor_len); - case_end(test_pair_user_inputs(backends, (int) b_send, (int) b_recv, tensor_len, + case_end(test_pair_user_inputs(backends_w_caps, (int) b_send, (int) b_recv, tensor_len, /*n_inputs =*/ 4, inputs_on_sender, parallel, use_device_host_buft)); } } @@ -908,7 +1000,7 @@ int main(int argc, char ** argv) { for (int tensor_len : { 1, 4096 }) { case_begin("test_y_shaped %-8s -> %-8s tensor_len = %4d", name_send, name_recv, tensor_len); - case_end(test_y_shaped_graph(backends, (int) b_send, (int) b_recv, tensor_len, use_device_host_buft)); + case_end(test_y_shaped_graph(backends_w_caps, (int) b_send, (int) b_recv, tensor_len, use_device_host_buft)); } } } @@ -916,8 +1008,8 @@ int main(int argc, char ** argv) { log_maybe("\n"); } - for (ggml_backend_t backend : backends) { - ggml_backend_free(backend); + for (const sched_backend_caps & caps : backends_w_caps) { + ggml_backend_free(caps.backend); } if (print_log || n_ok != n_test) {