ggml: add a scheduler sanitizer

[am17an]

Overview

Add a sanitizer which catches "happens-before" races in the scheduler. To run you can use GGML_SCHED_SANITIZE=1 which would crash if finds a race.

How it works - it takes the idea from Vector Clocks to identify happens-before relations. The idea is there is a race condition when there is an operation on a memory range M -

1. previous operation was a write OR current operation is a write
2. it is by different "actors" (logical ggml_backend_t)
3. the current actor's version of the previous actor's clock is less than the previous actor's actual clock (i.e. clock_view[current][prev] < clock[prev])

Additional information

1. ggml_cpy_async is inconsistent in ggml, we need to clarify the contract there.
2. It seems to correctly identify the race detected in ggml: fix backend split scheduler race condition #26040.

ggml-sched-sanitize: RACE (write-after-read) on Vulkan_Host[278941760, 279064640)
ggml-sched-sanitize:   read  Vulkan0#1    @1    split 1    model.input_embed (cpy_async src)
ggml-sched-sanitize:   write CPU          @2    split 2    conv_states-0 (compute)
ggml-sched-sanitize:   no happens-before edge: CPU knows Vulkan0#1@0, needs >=1

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES, paired with codex.

[ggerganov]

I am testing this branch with the following patch in order to cause a synchronization issue:

diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index e73a7b890..b9f449ac1 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -2468,10 +2468,10 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_
             CUDA_CHECK(cudaEventCreateWithFlags(&cuda_ctx_src->copy_event, cudaEventDisableTiming));
         }
 
-        CUDA_CHECK(cudaEventRecord(cuda_ctx_src->copy_event, cuda_ctx_src->stream()));
+        //CUDA_CHECK(cudaEventRecord(cuda_ctx_src->copy_event, cuda_ctx_src->stream()));
 
         // wait on dst stream for the copy to complete
-        CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx_dst->stream(), cuda_ctx_src->copy_event, 0));
+        //CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx_dst->stream(), cuda_ctx_src->copy_event, 0));
     } else {
         // src and dst are on the same backend
         CUDA_CHECK(cudaMemcpyAsync(dst->data, src->data, ggml_nbytes(dst), cudaMemcpyDeviceToDevice, cuda_ctx_src->stream()));

Command:

make -j && GGML_SCHED_SANITIZE=2 GGML_CUDA_DEVICES=4 ./bin/llama-perplexity -hf ggml-org/qwen3-0.6b-gguf:Q8_0 -f ./wikitext-2-raw/wiki.test.raw --chunks 16 -sm layer -dev CUDA0,CUDA1,CUDA2,CUDA3

The sanitizer does not report any issue. Is this expected?

[am17an]

Yes that's expected because the backend after this patch is incorrectly implementing the cpy_async. This sanitizer cannot catch those cases, it assumes the functions are implemented correctly. It will be able to catch all the event sync logic used in ggml_backend

[am17an]

I was able to introduce a race in #21067 and the sanitizer was able to catch it.

[ggerganov]

The main concern for me is that implementing it like this, it is very prone for errors in the future where we forget to add the respective sanitizer call. Instead, it should be seamless - we don't have to remember to call the sanitizer every time we use the backend interface.

[am17an]

I see, for actually seamless integration we would need to override backend->iface calls which do the sanitize and then call the actual function. Let me see how feasible that is

[ggerganov]

we would need to override backend->iface calls

Or wrap them

[am17an]

wrapping is not seamless, as the user can still call backend->iface directly?

[ggerganov]

If we want just the scheduler logic to be sanitized, the scheduler can wrap them. If we want to sanitize even user calls, then we probably need a "sanitize backend" that wraps a backend.

[JohannesGaessler]

While it should be fine from a llama.cpp standpoint, from a ggml standpoint I think it is undesirable to have a global state like this.

[JohannesGaessler]

I have not yet been able to do a full review but based on what I've read so far I think the code should be organized differently:

• Move the ggml backend scheduler to a dedicated .cpp file (preferably in a dedicated PR).
• Add e.g. functions like ggml_sched_backend_synchronize that wrap ggml_backend_synchronize. Wrap things like ggml_backend_t in structs like ggml_sched_backend_t to make accidental misuse of the wrong API impossible.
• Make the sanitizer a part of the backend scheduler so that there is no global state and the sanitizer is not exposed to the rest of the codebase.
• Remove the concept of an "actor" and directly use the scheduler's backend IDs instead.

[ggerganov]

Yes good suggestions by @JohannesGaessler. Also, design the code in such a way that in the future we can implement other sanitizers. To achieve that, have a minimal internal scheduler callback API (f.ex declared in ggml-backend-sched.h with default implementation in ggml-backend-sched.cpp) that the sanitizers implement.