From 8da7e74e142be7761c1c58210f6282fda2541dca Mon Sep 17 00:00:00 2001 From: Scott Cutler Date: Wed, 22 Apr 2026 21:57:59 -0700 Subject: [PATCH] ggml-cuda: move GGML_CUDA_AR_WATCHDOG from CMake option to local define MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The watchdog is development-only; a global CMake option is overkill. Move the toggle to a #define at the top of allreduce.cu (set to 0 by default) and remove the option from ggml/CMakeLists.txt and the CUDA CMakeLists.txt add_compile_definitions block. Co-Authored-By: Claude Sonnet 4.6 fences, sets completion flag, then all threads exit - Watchdog thread simply polls ring head counters every 1ms and prints any new complete records — no CUDA event queries, no mutex, no queue - Zero overhead on the dispatch path (no queue posting, no memset) - Watchdog shutdown returns within ~1ms (atomic bool, no drain) - On bailout the kernel skips Phase 3 entirely and exits cleanly Verified: 20/20 prefill soak test clean at ~1112 t/s, no hangs. Co-Authored-By: Claude Sonnet 4.6 P32, tensors <= 256 KB. Notes in NOTES-allreduce.md. Co-Authored-By: Claude Sonnet 4.6 --- ggml/CMakeLists.txt | 1 - ggml/src/ggml-cuda/CMakeLists.txt | 4 ---- ggml/src/ggml-cuda/allreduce.cu | 26 ++++++++++++++++---------- 3 files changed, 16 insertions(+), 15 deletions(-) diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index 820053a96c..c721d76f18 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -209,7 +209,6 @@ option(GGML_CUDA_FA_ALL_QUANTS "ggml: compile all quants for FlashA option(GGML_CUDA_GRAPHS "ggml: use CUDA graphs (llama.cpp only)" ${GGML_CUDA_GRAPHS_DEFAULT}) option(GGML_CUDA_NCCL "ggml: use NVIDIA Collective Comm. Library" ON) option(GGML_CUDA_NCCL_STATIC "ggml: link NCCL statically (ON) or dynamically (OFF)" OFF) -option(GGML_CUDA_AR_WATCHDOG "ggml: enable internal AllReduce hang watchdog (debug)" OFF) set (GGML_CUDA_COMPRESSION_MODE "size" CACHE STRING "ggml: cuda link binary compression mode; requires cuda 12.8+") set_property(CACHE GGML_CUDA_COMPRESSION_MODE PROPERTY STRINGS "none;speed;balance;size") diff --git a/ggml/src/ggml-cuda/CMakeLists.txt b/ggml/src/ggml-cuda/CMakeLists.txt index 8f6c54d6cc..1d94873c1a 100644 --- a/ggml/src/ggml-cuda/CMakeLists.txt +++ b/ggml/src/ggml-cuda/CMakeLists.txt @@ -194,10 +194,6 @@ if (CUDAToolkit_FOUND) endif() endif() - if (GGML_CUDA_AR_WATCHDOG) - add_compile_definitions(GGML_CUDA_AR_WATCHDOG) - endif() - set(CUDA_CXX_FLAGS "") set(CUDA_FLAGS -use_fast_math -extended-lambda) diff --git a/ggml/src/ggml-cuda/allreduce.cu b/ggml/src/ggml-cuda/allreduce.cu index c65449c490..4225b03790 100644 --- a/ggml/src/ggml-cuda/allreduce.cu +++ b/ggml/src/ggml-cuda/allreduce.cu @@ -4,7 +4,13 @@ #include #include -#ifdef GGML_CUDA_AR_WATCHDOG +// Set to 1 to enable the AllReduce spin-limit watchdog (development only). +// When enabled, the debug kernel bails out after GGML_CUDA_AR_MAX_SPIN +// iterations and writes a record to a per-GPU ring buffer that the +// background watchdog thread prints. +#define GGML_CUDA_AR_WATCHDOG 0 + +#if GGML_CUDA_AR_WATCHDOG #include #include #include @@ -123,7 +129,7 @@ static __global__ void ggml_cuda_ar_f32_kernel( // and then sets the completion flag so the host watchdog thread can safely // read the record. // --------------------------------------------------------------------------- -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG // One debug record written by the kernel on spin-limit bailout. struct ggml_cuda_ar_debug_record { @@ -254,7 +260,7 @@ static constexpr int GGML_CUDA_AR_POOL_SIZE = 128; // preventing false-sharing stalls on the polling GPU. static constexpr size_t GGML_CUDA_AR_ARRIVAL_STRIDE = 128; -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG // Watchdog poll interval in milliseconds. static constexpr int GGML_CUDA_AR_WDOG_POLL_MS = 1; #endif @@ -279,7 +285,7 @@ struct ggml_cuda_ar_pipeline { // Use ggml_cuda_ar_arrival_ptr() to index. char * arrival; -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG // Per-GPU debug ring buffers in pinned host memory. Written by the debug // kernel on spin-limit bailout, read by the background watchdog thread. ggml_cuda_ar_debug_ring * debug_ring[GGML_CUDA_MAX_DEVICES]; @@ -301,7 +307,7 @@ static int * ggml_cuda_ar_arrival_ptr(const ggml_cuda_ar_pipeline * p, int slot, // this thread polls the ring head counters every 1ms and prints any new // complete records. Zero overhead on the dispatch path (no queue, no events). // --------------------------------------------------------------------------- -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG static void ggml_cuda_ar_wdog_thread(ggml_cuda_ar_pipeline * p) { int last_seen[GGML_CUDA_MAX_DEVICES] = {}; @@ -352,7 +358,7 @@ ggml_cuda_ar_pipeline * ggml_cuda_ar_pipeline_init( p->streams[i] = nullptr; p->ev_pool[i] = nullptr; } -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) { p->debug_ring[i] = nullptr; } @@ -409,7 +415,7 @@ ggml_cuda_ar_pipeline * ggml_cuda_ar_pipeline_init( memset(p->host_buf[i], 0, max_bytes); } -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG // Per-GPU debug ring buffers: written by the kernel on spin-limit bailout, // polled by the background watchdog thread. Each ring is pinned host // memory accessed only by its owning GPU (single-GPU host atomics OK). @@ -507,7 +513,7 @@ void ggml_cuda_ar_pipeline_free(ggml_cuda_ar_pipeline * p) { return; } -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG // Stop the watchdog thread first — it only reads pinned host memory, // no GPU resources, so this is safe and returns within ~1ms. p->wdog_stop.store(true); @@ -544,7 +550,7 @@ void ggml_cuda_ar_pipeline_free(ggml_cuda_ar_pipeline * p) { if (p->arrival) { cudaFreeHost(p->arrival); } -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG for (int i = 0; i < p->n_devices; ++i) { if (p->debug_ring[i]) { cudaFreeHost(p->debug_ring[i]); @@ -624,7 +630,7 @@ bool ggml_cuda_ar_allreduce( CUDA_CHECK(cudaEventRecord(ev.app, cuda_ctx->stream())); CUDA_CHECK(cudaStreamWaitEvent(p->streams[i], ev.app)); -#ifdef GGML_CUDA_AR_WATCHDOG +#if GGML_CUDA_AR_WATCHDOG ggml_cuda_ar_f32_kernel_dbg<<streams[i]>>>( static_cast(tensors[i]->data), static_cast(tensors[i]->data),