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.