diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h index ecf0a3a6dc25..00be60f0fbc7 100644 --- a/ggml/include/ggml-rpc.h +++ b/ggml/include/ggml-rpc.h @@ -8,10 +8,10 @@ extern "C" { #define RPC_PROTO_MAJOR_VERSION 8 #define RPC_PROTO_MINOR_VERSION 0 -#define RPC_PROTO_PATCH_VERSION 1 +#define RPC_PROTO_PATCH_VERSION 2 #ifdef __cplusplus -static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); #endif #define GGML_RPC_MAX_SERVERS 16 diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index d62a631043a3..c4ddceb56839 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -602,6 +602,8 @@ extern "C" { GGML_OP_GLU, + GGML_OP_SLEEP, + GGML_OP_COUNT, }; @@ -2846,6 +2848,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 53d39d2491ed..4e4ecac4c19d 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -2031,6 +2031,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); @@ -2304,6 +2308,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 5735f96f8e8b..d7fc37bde8eb 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 @@ -2352,6 +2353,18 @@ 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) { + if (params->ith != 0) { + return; + } + + const auto t_end = std::chrono::steady_clock::now() + std::chrono::microseconds(ggml_get_op_params_i32(dst, 0)); + + 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 00a2ea7dc2de..6283fdf1bc58 100644 --- a/ggml/src/ggml-cpu/ops.h +++ b/ggml/src/ggml-cpu/ops.h @@ -86,6 +86,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 b795951b2f7b..81f14e688911 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -45,6 +45,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" @@ -2607,6 +2608,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; @@ -6109,6 +6113,12 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g return ggml_cuda_lightning_indexer_supported(dev_ctx->device, op); case GGML_OP_TOP_K_MASK: return op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_F32; + case GGML_OP_SLEEP: +#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) + return false; +#else + return true; +#endif // defined(GGML_USE_HIP) || defined(GGML_USE_MUSA) 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..e6a40e5f213a --- /dev/null +++ b/ggml/src/ggml-cuda/sleep.cu @@ -0,0 +1,32 @@ +#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 next kernel on the same stream cannot start before this one 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) { + sleep_ns<<<1, 1, 0, ctx.stream()>>>(1000*(uint64_t) ggml_get_op_params_i32(dst, 0)); + CUDA_CHECK(cudaGetLastError()); +} + +#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 // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) 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-vulkan/ggml-vulkan-debug.cpp b/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp index 97747e866df3..c6795e2566ec 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan-debug.cpp @@ -993,6 +993,8 @@ void ggml_vk_check_results_0(ggml_backend_vk_context * ctx, ggml_cgraph * 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/ggml-vulkan-push-constants.h b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h index 219c477c6170..3b6c13be945e 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h +++ b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h @@ -151,6 +151,10 @@ struct vk_op_push_constants { float param4; }; +struct vk_op_sleep_push_constants { + uint32_t ticks; +}; + struct vk_op_fwht_push_constants { uint32_t n_rows; uint32_t src_offset; diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-types.h b/ggml/src/ggml-vulkan/ggml-vulkan-types.h index 5cf46c89085f..0e24494c96d3 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-types.h +++ b/ggml/src/ggml-vulkan/ggml-vulkan-types.h @@ -710,6 +710,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; @@ -903,6 +904,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]; diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 33936acae8cd..3c425921a906 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -3470,6 +3470,10 @@ 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 (device->shader_clock) { + ggml_vk_create_pipeline(device, device->pipeline_sleep, "sleep", sleep_len, sleep_data, "main", 1, sizeof(vk_op_sleep_push_constants), {1, 1, 1}, {}, 1); + } + #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); @@ -4055,6 +4059,7 @@ 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) { @@ -4105,6 +4110,9 @@ vk_device ggml_vk_get_device(size_t idx) { } else if (strcmp(VK_EXT_SHADER_FLOAT8_EXTENSION_NAME, properties.extensionName) == 0) { shader_float8_extension = true; #endif + } else if (strcmp("VK_KHR_shader_clock", properties.extensionName) == 0 && + !getenv("GGML_VK_DISABLE_SHADER_CLOCK")) { + shader_clock_support = true; } else if (strcmp("VK_VALVE_shader_mixed_float_dot_product", properties.extensionName) == 0 && !getenv("GGML_VK_DISABLE_DOT2")) { dot2_f16_support = true; @@ -4443,6 +4451,14 @@ 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) { @@ -4526,6 +4542,7 @@ 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; @@ -9474,6 +9491,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; @@ -10689,6 +10708,31 @@ void ggml_vk_fill(ggml_backend_vk_context * ctx, vk_context& subctx, ggml_tensor ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { dst_buf }, pc, elements); } +static void ggml_vk_sleep(ggml_backend_vk_context * ctx, vk_context& subctx, 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 this assumes it matches the timestamp counter. + // A driver where this does not hold could sleep too short or too long by that ratio + 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) ticks, + }; + + vk_pipeline pipeline = ggml_vk_op_get_pipeline(ctx, nullptr, nullptr, nullptr, dst, GGML_OP_SLEEP); + GGML_ASSERT(pipeline != nullptr); + + ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1); + vk_subbuffer dst_buf = ggml_vk_tensor_subbuffer(ctx, dst); + + // a single invocation spins, the buffer is bound only to keep the spin from being optimized out + std::array elements = { 1, 1, 1 }; + + ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { dst_buf }, pc, elements); +} + 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)); } @@ -12551,6 +12595,10 @@ bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, in case GGML_OP_FILL: ggml_vk_fill(ctx, compute_ctx, node); + break; + case GGML_OP_SLEEP: + ggml_vk_sleep(ctx, compute_ctx, node); + break; case GGML_OP_SCALE: if (ctx->fused_hc_post_gate) { @@ -15912,6 +15960,8 @@ 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: + return device->shader_clock; case GGML_OP_SCALE: return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32; case GGML_OP_PAD: 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..07f6cafc9715 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/sleep.comp @@ -0,0 +1,37 @@ +#version 450 + +#extension GL_EXT_shader_realtime_clock : require + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +// never written, it only anchors the spin below so the compiler cannot drop it +layout (binding = 0) writeonly buffer D {uint data_d[];}; + +layout (push_constant) uniform parameter { + 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() { + const uvec2 t0 = clockRealtime2x32EXT(); + + uvec2 t = t0; + + for (uint iter = 0; iter < MAX_ITER; ++iter) { + 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; + } + } + + // the store never runs, but the compiler cannot prove it and has to keep the spin + if (t.x == 0xFFFFFFFFu && t.y == 0xFFFFFFFFu) { + data_d[0] = t.x; + } +} 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 f545015464f1..becab254f75a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -1039,6 +1039,7 @@ 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"}}); + string_to_spv("sleep", "sleep.comp", {}); 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"}}); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index ed8554b14173..02ebdc7d05ab 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1108,9 +1108,11 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "OPT_STEP_SGD", "GLU", + + "SLEEP", }; -static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1224,9 +1226,11 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "sgd(x)", "glu(x)", + + "sleep(x)", }; -static_assert(GGML_OP_COUNT == 102, "GGML_OP_COUNT != 102"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -6292,6 +6296,27 @@ struct ggml_tensor * ggml_custom_inplace( return result; } + +// ggml_sleep + +// the result aliases a, the op only delays and never touches the data +struct ggml_tensor * ggml_sleep( + struct ggml_context * ctx, + struct ggml_tensor * a, + int32_t us) { + GGML_ASSERT(us >= 0); + + struct ggml_tensor * result = ggml_view_tensor(ctx, a); + ggml_format_name(result, "%s (sleep)", a->name); + + 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 2e768cb38550..f891c23799cc 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -335,6 +335,8 @@ if (GGML_METAL AND NOT GGML_BACKEND_DL) endif() llama_build(test-backend-ops.cpp LINK ggml) +llama_build_and_test(test-sleep.cpp LINK ggml) +llama_build_and_test(test-backend-sched.cpp LINK ggml) llama_build_and_test(test-tiled-mulmat.cpp LABEL "main behavior backend" LINK ggml) llama_build_and_test(test-model-load-cancel.cpp LABEL "model behavior") diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 9a4ddaa29a35..13c656977667 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9035,6 +9035,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 only delays, the output aliases the input and must match it exactly + 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; @@ -12025,6 +12052,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_fill(-152.0f, GGML_TYPE_F16, { 800, 600, 4, 4 })); test_cases.emplace_back(new test_fill(3.5f, GGML_TYPE_F16, { 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 })); diff --git a/tests/test-backend-sched.cpp b/tests/test-backend-sched.cpp new file mode 100644 index 000000000000..eb042c206027 --- /dev/null +++ b/tests/test-backend-sched.cpp @@ -0,0 +1,1055 @@ +// Test suite for scheduler and backend 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: +// - 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 +// - schedule splits without inputs at any point in the graph + +#include "ggml.h" +#include "ggml-alloc.h" +#include "ggml-backend.h" +#include "ggml-cpp.h" + +#include +#include +#include +#include +#include +#include +#include + + + +enum test_result { + TEST_FAIL = 0, + TEST_OK = 1, + TEST_SKIP = 2, // in case a backend does not support a required OP (e.g. OP_SLEEP) +}; + +// global counters for pass/fail/skip +static int n_ok = 0; +static int n_skip = 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]; // max length of a case label + +static void case_end(test_result res) { + n_ok += res == TEST_OK; + n_skip += res == TEST_SKIP; + n_test++; + + if (res != TEST_FAIL && !print_log) { + case_len = 0; + return; + } + + if (case_len == 0) { + case_len = printf(" %s", case_label); + } + + const char * status = "FAIL"; + if (res == TEST_OK) { + status = "OK"; + } else if (res == TEST_SKIP) { + status = "SKIP"; + } + + printf("%*s%s\n", std::max(1, status_column - case_len), "", status); + case_len = 0; +} + +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, ...) { + if (!print_log) { + return; + } + + va_list args; + va_start(args, fmt); + vnote(fmt, args); + va_end(args); +} + +GGML_ATTRIBUTE_FORMAT(1, 2) +static test_result fail(const char * fmt, ...) { + if (case_len == 0) { + case_len = printf(" %s", case_label); + } + + va_list args; + va_start(args, fmt); + vnote(fmt, args); + va_end(args); + + return TEST_FAIL; +} + +GGML_ATTRIBUTE_FORMAT(1, 2) +static test_result skip(const char * fmt, ...) { + if (print_log) { + va_list args; + va_start(args, fmt); + vnote(fmt, args); + va_end(args); + } + + return TEST_SKIP; +} + +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); + + if (print_log) { + case_len = printf(" %s", case_label); + fflush(stdout); + } else { + case_len = 0; + } +} +// pretty-print helpers end + + +// NB: the last backend passed has to be the CPU backend (cf. ggml_backend_sched_new) +static ggml_backend_sched_ptr 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]); + } + + // pinned memory for the CPU backend (always required to be the last backend) enables async copies to other devices + // the host buffer type comes from the first backend, (cf. llama_context::llama_context()) + if (use_device_host_buft) { + ggml_backend_buffer_type_t host = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backends[0])); + if (host != nullptr) { + bufts.back() = host; + } + } + + return ggml_backend_sched_ptr(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 helper_graph { + ggml_cgraph * gf; + std::vector nodes; + std::vector backend_id; + + helper_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 helper_check { + ggml_tensor * tensor; // Tensor for which to check the data. + float expected; // The expected value that the tensor data should have after graph execution. +}; + +// executes the graph, then checks the number of splits and the results +static test_result run_and_check(ggml_backend_sched_t sched, const helper_graph & g, + const std::vector & backends, const std::vector & checks, int64_t ne) { + + // ggml_backend_sched_set_tensor_backend does not validate the op, an unsupported one would abort inside the backend + for (size_t i = 0; i < g.nodes.size(); i++) { + ggml_backend_t backend = backends[g.backend_id[i]]; + if (!ggml_backend_supports_op(backend, g.nodes[i])) { + return skip("%s does not support %s", ggml_backend_name(backend), ggml_op_name(g.nodes[i]->op)); + } + } + + // to later check if any further graph logic alters the number of splits + 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, each split consisting of just a single node", n_splits, n_splits_expected); + } + + std::vector data(ne); + for (const helper_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", ggml_get_name(c.tensor), i, data[i], c.expected); + } + } + } + + return TEST_OK; +} + +// helper struct for inputless splits, and constants like 0 and 1 +struct backend_consts { + // the buffers are declared after the contexts, so that they are released first + 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]); + GGML_ASSERT(buf != nullptr); + + 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.emplace_back(ctx); + bufs.emplace_back(buf); + zero.push_back(z); + one.push_back(o); + } + } +}; + +// stress test: a tensor is incremented and sent back and forth between the backends in a ping-pong pattern. +// One pass pairs every backend a with every other backend b via nested loops. +// The counter is brought to a, then bounced to b and back for every partner in turn. +// So one pass copies an activation across every ordered pair of backends. +// Every lap runs two passes: one where the nodes add the constant 1, one where they add a user input +// from the CPU backend, so that every transition is tested with and without an input next to the activation. +// The input values come from a fixed seed. +static test_result stress_test_linked_list(const std::vector & backends, int n_laps, + int64_t tensor_len, bool use_device_host_buft) { + const int n_backends = (int) backends.size(); + + // per pass and backend a: one node to bring the counter to a, then two per partner for the ping-pong + const int n_nodes_pass = n_backends*(2*n_backends - 1); + const int n_nodes = 2*n_laps*n_nodes_pass; + const int n_inputs = n_laps*n_nodes_pass; // one per node of the pass that takes user inputs + + // initial pseudo-random value draw so that the expected result is known before graph execution + std::mt19937 rng(42); + std::uniform_int_distribution draw_value(1, 10); + + std::vector input_value(n_inputs); // random values store + int64_t sum_inputs = 0; + + for (int i = 0; i < n_inputs; i++) { + input_value[i] = draw_value(rng); + sum_inputs += input_value[i]; + } + + // see sched->hash_set FIXME, the set has to hold the nodes and the leafs: two constants per backend + // plus the user inputs + const size_t graph_size = n_nodes + 2*n_backends + n_inputs; + + backend_consts consts(backends, tensor_len); + + ggml_backend_sched_ptr sched = create_test_scheduler(backends, graph_size, use_device_host_buft); + + // the user inputs all live on the CPU backend, with the same buffer type the scheduler uses for it + ggml_init_params params_input = { + /*.mem_size =*/ n_inputs*ggml_tensor_overhead(), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context_ptr ctx_input(ggml_init(params_input)); + ggml_backend_buffer_ptr buf_input; + + std::vector input(n_inputs); + for (int i = 0; i < n_inputs; i++) { + input[i] = ggml_new_tensor_1d(ctx_input.get(), GGML_TYPE_F32, tensor_len); + ggml_format_name(input[i], "input%d", i); + ggml_set_input(input[i]); + } + + // the inputs follow the buffer type the scheduler uses for the CPU backend (cf. create_test_scheduler) + ggml_backend_buffer_type_t buft = ggml_backend_get_default_buffer_type(backends.back()); + if (use_device_host_buft) { + ggml_backend_buffer_type_t host = ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backends[0])); + if (host != nullptr) { + buft = host; + } + } + + buf_input.reset(ggml_backend_alloc_ctx_tensors_from_buft(ctx_input.get(), buft)); + GGML_ASSERT(buf_input != nullptr); + + // write inputs before graph execution + std::vector data(tensor_len); + for (int i = 0; i < n_inputs; i++) { + 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), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_context_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), graph_size); + + int64_t sum = 0; + ggml_tensor * out = nullptr; + + int i = 0; // next user input to consume + + // increments the counter on backend b, which makes the scheduler copy it there from the previous backend + auto step = [&](int b, bool with_input) { + ggml_tensor * src = out == nullptr ? consts.zero[b] : out; + + if (with_input) { + out = g.add(ggml_add(ctx_compute.get(), src, input[i]), b); + i++; + } else { + out = g.add(ggml_add(ctx_compute.get(), src, consts.one[b]), b); + sum += 1; + } + }; + + for (int lap = 0; lap < n_laps; lap++) { + for (bool with_input : { false, true }) { // test each transition with and without input + for (int a = 0; a < n_backends; a++) { + step(a, with_input); + + for (int b = 0; b < n_backends; b++) { + if (b == a) { + continue; + } + + step(b, with_input); + step(a, with_input); + } + } + } + } + ggml_set_name(out, "out"); + ggml_set_output(out); + + // every node adds either the constant 1 or its drawn user input + const int64_t expected = sum + sum_inputs; + + note("%d of %d nodes take a user input, they contribute %" PRId64 " of %" PRId64, n_inputs, n_nodes, sum_inputs, expected); + + test_result ok; + if (expected >= (1 << 24)) { // the counter has to stay exact in f32 + ok = fail("expected result %" PRId64 " does not fit into the f32 mantissa, lower n_laps", expected); + } else { + ok = run_and_check(sched.get(), g, backends, {{ out, float(expected) }}, tensor_len); + } + + return ok; +} + +// Same idea as stress_test_linked_list, but with the compute graph being a directed 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 test_result stress_test_dag(const std::vector & backends, int n_lanes, int n_rounds, + int64_t tensor_len, bool use_device_host_buft) { + + const int n_backends = (int) backends.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); + + ggml_backend_sched_ptr 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_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), graph_size); + + std::vector lane(n_lanes); + std::vector value(n_lanes); + std::vector checks; + + // 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.get(), 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); + } + }; + + auto check_lanes = [&]() { + for (int l = 0; l < n_lanes; l++) { + ggml_set_output(lane[l]); + checks.push_back({ lane[l], float(value[l]) }); + } + }; + + seed(0, "seed"); + + std::vector older = lane; + std::vector value_older = value; + + 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.get(), 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.get(), 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.get(), 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 2: + case 3: + 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; + + ggml_tensor * a = reach_back ? older[l] : prev[l]; + const int64_t va = reach_back ? value_older[l] : value_prev[l]; + + lane[l] = g.add(ggml_add(ctx_compute.get(), a, prev[partner]), b); + value[l] = va + value_prev[partner]; + ggml_format_name(lane[l], "r%d_lane%d", round, l); + } + break; + case 4: + lane[0] = g.add(ggml_add(ctx_compute.get(), prev[0], consts.one[b]), b); + value[0] = value_prev[0] + 1; + ggml_format_name(lane[0], "r%d_single", round); + break; + default: + GGML_ABORT("unexpected round %d", round); + } + + older = prev; + value_older = value_prev; + } + + check_lanes(); + + test_result ok = TEST_OK; + for (const helper_check & c : checks) { + if (c.expected >= float(1 << 24)) { // the counts have to stay exact in f32 + ok = fail("expected result %.0f does not fit into the f32 mantissa, lower n_rounds", c.expected); + break; + } + } + + if (ok == TEST_OK) { + ok = run_and_check(sched.get(), g, backends, checks, tensor_len); + } + + 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 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 the copy started. +// Note: currently only reproducible on async H2D copy (=pinned memory). +static test_result test_inputless_splits_scheduling(ggml_backend_t gpu, ggml_backend_t 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 = { gpu, cpu }; + + const size_t graph_size = 64; + + ggml_backend_sched_ptr sched = create_test_scheduler(backends, 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_ptr ctx_cpu(ggml_init(params_static)); + + ggml_tensor * zero_cpu = ggml_new_tensor_1d(ctx_cpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(zero_cpu, "zero_cpu"); + + ggml_tensor * val55_cpu = ggml_new_tensor_1d(ctx_cpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(val55_cpu, "val55_cpu"); + + ggml_tensor * val66_cpu = ggml_new_tensor_1d(ctx_cpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(val66_cpu, "val66_cpu"); + + ggml_backend_buffer_ptr buf_cpu(ggml_backend_alloc_ctx_tensors(ctx_cpu.get(), cpu)); + GGML_ASSERT(buf_cpu != nullptr); + + ggml_context_ptr ctx_gpu(ggml_init(params_static)); + + ggml_tensor * zero_gpu = ggml_new_tensor_1d(ctx_gpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(zero_gpu, "zero_gpu"); + + ggml_tensor * one_gpu = ggml_new_tensor_1d(ctx_gpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(one_gpu, "one_gpu"); + + ggml_tensor * val44_gpu = ggml_new_tensor_1d(ctx_gpu.get(), GGML_TYPE_F32, tensor_len); + ggml_set_name(val44_gpu, "val44_gpu"); + + ggml_backend_buffer_ptr buf_gpu(ggml_backend_alloc_ctx_tensors(ctx_gpu.get(), gpu)); + GGML_ASSERT(buf_gpu != nullptr); + + 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)); + 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_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), graph_size); + + // GPU: add then sleep - occupies the GPU while later splits run + ggml_tensor * delayed = g.add(ggml_add(ctx_compute.get(), zero_gpu, one_gpu), GPU); + ggml_set_name(delayed, "to_delay"); + delayed = g.add(ggml_sleep(ctx_compute.get(), delayed, sleep_us), GPU); + ggml_set_name(delayed, "delayed"); + + // CPU: 0 + 55 -> 55 + ggml_tensor * v55 = g.add(ggml_add(ctx_compute.get(), zero_cpu, val55_cpu), CPU); + ggml_set_name(v55, "v55"); + + // GPU: 55 + 44 -> 99 + ggml_tensor * v99 = g.add(ggml_add(ctx_compute.get(), v55, val44_gpu), GPU); + ggml_set_name(v99, "v99"); + + // CPU: 0 + 66 -> 66 + ggml_tensor * v66 = g.add(ggml_add(ctx_compute.get(), zero_cpu, val66_cpu), CPU); + ggml_set_name(v66, "v66"); + + // GPU: increment both of the previous split outputs + ggml_tensor * out99 = g.add(ggml_add(ctx_compute.get(), v99, one_gpu), GPU); + ggml_set_name(out99, "out99"); + + ggml_tensor * out66 = g.add(ggml_add(ctx_compute.get(), v66, one_gpu), GPU); + ggml_set_name(out66, "out66"); + + ggml_set_output(delayed); + ggml_set_output(out99); + ggml_set_output(out66); + + const test_result ok = run_and_check(sched.get(), g, backends, { + { delayed, 1.0f }, + { out99, 100.0f }, + { out66, 67.0f }, + }, tensor_len); + + // 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); + + return ok; +} + +// Test that all async backends transmit their activations to the following async backend correctly. No user inputs tested. +static test_result test_chain_all_backends(const std::vector & backends, int64_t tensor_len, int32_t sleep_us, + bool use_device_host_buft) { + + 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); + } + } + } + + const size_t n_nodes = seq.size(); + const size_t n_sleeps = sleep_us > 0 ? n_nodes : 0; // one sleep per split + const size_t graph_size = n_nodes + n_sleeps + 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_ptr 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_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), 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.get(), k == 0 ? consts.zero[b] : out, consts.one[b]), b); + + // the sleep keeps the backend busy, so that the next one has to wait for this split + if (sleep_us > 0) { + out = g.add(ggml_sleep(ctx_compute.get(), out, sleep_us), b); + } + } + ggml_set_name(out, "out"); + ggml_set_output(out); + + const test_result ok = run_and_check(sched.get(), g, backends, {{ out, float(n_nodes) }}, tensor_len); + + return ok; +} + +// Tests data transfer between all combinations of backend pairs +// The receiving split takes one activation from the sender plus n_inputs user inputs +static test_result 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); + + ggml_backend_sched_ptr 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_ptr 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.get(), GGML_TYPE_F32, tensor_len); + ggml_format_name(in, "in%d", k); + ggml_set_input(in); + inputs.push_back(in); + } + + ggml_backend_buffer_ptr buf_inputs(ggml_backend_alloc_ctx_tensors(ctx_inputs.get(), backends[b_inputs])); + GGML_ASSERT(buf_inputs != nullptr); + + 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)); + } + + 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_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), graph_size); + + ggml_tensor * out = g.add(ggml_add(ctx_compute.get(), consts.zero[b_send], consts.one[b_send]), b_send); + + for (ggml_tensor * in : inputs) { + out = g.add(ggml_add(ctx_compute.get(), out, in), b_recv); + } + ggml_set_name(out, "out"); + ggml_set_output(out); + + test_result ok = run_and_check(sched.get(), g, backends, {{ out, float(1 + n_inputs) }}, tensor_len); + + for (ggml_tensor * in : inputs) { + if (ok != TEST_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; + } + } + } + + return ok; +} + +// 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 backend_a that receives one activation from each of the two backend_b splits. +static test_result test_y_shaped_graph(const std::vector & backends, int backend_a, int backend_b, int64_t tensor_len, + int32_t sleep_us, bool use_device_host_buft) { + + const size_t graph_size = 64; + + backend_consts consts(backends, tensor_len); + + ggml_backend_sched_ptr 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_ptr ctx_compute(ggml_init(params_compute)); + + helper_graph g(ctx_compute.get(), graph_size); + + ggml_tensor * seed_a = g.add(ggml_add(ctx_compute.get(), consts.zero[backend_a], consts.one[backend_a]), backend_a); + ggml_set_name(seed_a, "seed_a"); + + // only lane a sleeps, so that lane b is scheduled while it is still running + if (sleep_us > 0) { + seed_a = g.add(ggml_sleep(ctx_compute.get(), seed_a, sleep_us), backend_a); + ggml_set_name(seed_a, "seed_a_delayed"); + } + + ggml_tensor * lane_a = g.add(ggml_add(ctx_compute.get(), seed_a, consts.one[backend_b]), backend_b); + ggml_set_name(lane_a, "lane_a"); + + if (sleep_us > 0) { + lane_a = g.add(ggml_sleep(ctx_compute.get(), lane_a, sleep_us), backend_b); + ggml_set_name(lane_a, "lane_a_delayed"); + } + + ggml_tensor * seed_b = g.add(ggml_add(ctx_compute.get(), 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.get(), seed_b, consts.one[backend_b]), backend_b); + ggml_set_name(lane_b, "lane_b"); + + // neither lane ends on backend_a, so both activations are inputs of the last split + ggml_tensor * out = g.add(ggml_add(ctx_compute.get(), lane_a, lane_b), backend_a); + ggml_set_name(out, "out"); + ggml_set_output(out); + + const test_result ok = run_and_check(sched.get(), g, backends, {{ out, 4.0f }}, tensor_len); + + return ok; +} + +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) { + + // 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); + 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); + // Exclude ACCEL devices as they run in tandem with the CPU backend and only implement matmul, the tests here require ADD + if (type == GGML_BACKEND_DEVICE_TYPE_CPU || type == GGML_BACKEND_DEVICE_TYPE_ACCEL) { + 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.emplace_back(backend); + } + + if (backends.empty()) { + printf("no non-CPU backend found, skipping\n"); + return false; + } + + // ggml_backend_sched_new requires the CPU backend to be the last one + backends.emplace_back(ggml_backend_init_by_type(GGML_BACKEND_DEVICE_TYPE_CPU, nullptr)); + GGML_ASSERT(backends.back() != nullptr); + + for (const ggml_backend_ptr & backend : backends) { + ggml_backend_dev_t dev = ggml_backend_get_device(backend.get()); + + log_maybe("backend: %-10s host_buft = %d (%s)\n", ggml_backend_name(backend.get()), + ggml_backend_dev_host_buffer_type(dev) != nullptr, ggml_backend_dev_description(dev)); + } + log_maybe("\n"); + return true; +} + + +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; + + // nothing to test with synchronous CPU backend only + if (!initialize_backends(backends)) { + return 0; + } + + // the tests only borrow the backends, the ownership stays with this vector + std::vector backend_ptrs; + for (const ggml_backend_ptr & backend : backends) { + backend_ptrs.push_back(backend.get()); + } + + ggml_backend_t cpu = backend_ptrs.back(); + const size_t n_non_cpu_backends = backend_ptrs.size() - 1; + + // Test cases for single backends against CPU backend + for (size_t b_id = 0; b_id < n_non_cpu_backends; b_id++) { + ggml_backend_t backend = backend_ptrs[b_id]; + const char * name_backend = ggml_backend_name(backend); + + for (bool use_device_host_buft : { false, true }) { + // without a host buffer type this case is a duplicate of the pageable one + if (use_device_host_buft && ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend)) == nullptr) { + continue; + } + + log_maybe("=== Backend: %s, CPU sched buft: %s ===\n\n", name_backend, use_device_host_buft ? "device_host" : "pageable"); + + 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, cpu, tensor_len, sleep_us, use_device_host_buft)); + } + } + + 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 && ggml_backend_dev_host_buffer_type(ggml_backend_get_device(backend_ptrs[0])) == nullptr) { + 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("stress_test_dag n_lanes = %2d, n_rounds = %2d, tensor_len = %4d", n_lanes, n_rounds, tensor_len); + case_end(stress_test_dag(backend_ptrs, n_lanes, n_rounds, tensor_len, use_device_host_buft)); + } + } + } + + log_maybe("\n"); + + for (int tensor_len : { 2, 4096 }) { + for (int32_t sleep_us : { 0, 80 }) { + case_begin("test_chain_all_backends sleep_us = %4d, tensor_len = %4d", sleep_us, tensor_len); + case_end(test_chain_all_backends(backend_ptrs, tensor_len, sleep_us, 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(backend_ptrs, 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 backend_a = 0; backend_a < backend_ptrs.size(); backend_a++) { + for (size_t backend_b = 0; backend_b < backend_ptrs.size(); backend_b++) { + if (backend_a == backend_b) { + continue; + } + + const char * name_a = ggml_backend_name(backend_ptrs[backend_a]); + const char * name_b = ggml_backend_name(backend_ptrs[backend_b]); + + for (bool inputs_on_sender : { false, true }) { + for (bool parallel : { false, true }) { + for (int tensor_len : { 1, 4096 }) { + case_begin("test_pair_user_inputs %-8s -> %-8s inputs on %-8s parallel = %d, tensor_len = %4d", + name_a, name_b, inputs_on_sender ? "sender" : "receiver", parallel, tensor_len); + case_end(test_pair_user_inputs(backend_ptrs, (int) backend_a, (int) backend_b, tensor_len, + /*n_inputs =*/ 4, inputs_on_sender, parallel, use_device_host_buft)); + } + } + } + + for (int tensor_len : { 1, 4096 }) { + for (int32_t sleep_us : { 0, 80 }) { + case_begin("test_y_shaped %-8s -> %-8s sleep_us = %4d, tensor_len = %4d", name_a, name_b, sleep_us, tensor_len); + case_end(test_y_shaped_graph(backend_ptrs, (int) backend_a, (int) backend_b, tensor_len, sleep_us, use_device_host_buft)); + } + } + } + } + + log_maybe("\n"); + } + + if (print_log || n_ok + n_skip != n_test) { + printf("%d/%d tests passed", n_ok, n_test); + if (n_skip > 0) { + printf(", %d skipped", n_skip); + } + printf("\n"); + } + + return n_ok + n_skip == n_test ? 0 : 1; +} diff --git a/tests/test-sleep.cpp b/tests/test-sleep.cpp new file mode 100644 index 000000000000..748e02923bfc --- /dev/null +++ b/tests/test-sleep.cpp @@ -0,0 +1,161 @@ +// Tests for GGML_OP_SLEEP, on every backend that supports it. Three properties are checked: +// - the op only delays: input = output +// - 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. +// Known limitation: Vulkan does not specify the tick rate of the shader clock, see ggml_vk_sleep. A driver +// whose clock runs slower than assumed could undershoot the request and fail this bound. + +#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_alias = out->data == x->data; + 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 ") alias: %-4s data: %-4s time: %s\n", + desc, t_us, t_min_us, ok_alias ? "OK" : "FAIL", 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_alias && 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); + + 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 be a valid no-op, 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; +}