From 97e2bf27db3c02919b639c285de91eaa106ad565 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 21 Sep 2026 14:35:28 -0300 Subject: [PATCH 1/5] Add KI tests for record_event and wait_event Co-Authored-By: Claude Fable 5.1 --- lib/KernelInterface/Project.toml | 2 +- lib/KernelInterface/test/events.jl | 93 +++++++++++++++++++++++++++ lib/KernelInterface/test/testsuite.jl | 5 ++ 3 files changed, 99 insertions(+), 1 deletion(-) create mode 100644 lib/KernelInterface/test/events.jl diff --git a/lib/KernelInterface/Project.toml b/lib/KernelInterface/Project.toml index b88be8d6b..936a21459 100644 --- a/lib/KernelInterface/Project.toml +++ b/lib/KernelInterface/Project.toml @@ -1,7 +1,7 @@ name = "KernelInterface" uuid = "4ee993da-d684-4d17-a7dd-4e58e78d92bf" authors = ["Valentin Churavy and contributors"] -version = "0.2.1" +version = "0.2.2" [compat] julia = "1.10" diff --git a/lib/KernelInterface/test/events.jl b/lib/KernelInterface/test/events.jl new file mode 100644 index 000000000..b50b6111a --- /dev/null +++ b/lib/KernelInterface/test/events.jl @@ -0,0 +1,93 @@ +# Tests for `record_event`/`wait_event`, the backend side of the protocol behind +# `KernelAbstractions.@spawn`. The protocol is specified in KernelInterface terms, so +# it is replayed here by hand with `Threads.@spawn`: record in the spawning task, then +# `device!`, `wait_event` and `synchronize` in the child task. + +# Burns a fixed amount of work per work-item before writing `v`. The accumulator is a +# linear congruential step, which the compiler cannot fold away, and it feeds into the +# store so the loop cannot be dropped. +function slow_fill_kernel(A, v, ::Val{iters}) where {iters} + i = KI.get_global_id().x + acc = UInt32(i) + for k in UInt32(1):UInt32(iters) + acc = acc * 0x19660d + k + end + if i <= length(A) + @inbounds A[i] = ifelse(acc == 0x12345678, -v, v) + end + return +end + +function events_testsuite(backend) + b = backend() + dev = KI.device(b) + + # Few work-items and many dependent iterations: tens of milliseconds on a GPU, + # and still well under a second per launch on a CPU-backed queue. + N = 64 + iters = Val(2^22) + slow_fill(A, v) = KI.@kernel b numworkgroups = 1 workgroupsize = N slow_fill_kernel(A, v, iters) + + # Time one launch, after a warm-up that absorbs compilation, and queue enough of + # them back to back to keep the spawner's queue busy for about 100ms. + A = KI.zeros(b, Float32, N) + slow_fill(A, 1.0f0) + KI.synchronize(b) + slow_time = @elapsed begin + slow_fill(A, 1.0f0) + KI.synchronize(b) + end + launches = clamp(ceil(Int, 0.1 / slow_time), 4, 64) + + @testset "ordered across tasks" begin + # The child queues nothing but the wait, so its `synchronize` can only return + # once the spawner's queued work has drained. A backend that forgets + # `wait_event` for its event type fails with a MethodError here, and a + # `wait_event` that does nothing returns in about a millisecond. The clock + # starts before `record_event`, so a backend whose `record_event` is the + # default full `synchronize` passes just the same. The data check alone would + # not do: drivers that track hazards between command buffers (Metal, for its + # default buffers) order the readback after the fills without any wait. + for v in 1:launches + slow_fill(A, Float32(v)) + end + start = time_ns() + ev = KI.record_event(b) + task = Threads.@spawn begin + # `wait_event` acts on the active device's queue, so select it first. + KI.device!(b, dev) + KI.wait_event(b, ev) + KI.synchronize(b) + elapsed = (time_ns() - start) / 1.0e9 + elapsed, Array(A) + end + elapsed, result = fetch(task) + # Half the expected drain time leaves room for GPU clock ramp-up between the + # calibration launch and these. + @test elapsed >= launches * slow_time / 2 + @test all(==(Float32(launches)), result) + KI.synchronize(b) + end + + if KI.ndevices(b) > 1 + @testset "cross-device" begin + # `@spawn backend device=id` records on the spawner's device and waits on + # another one, so a multi-device backend must accept a foreign event. The + # ordering itself is not observable without peer access; check that the + # wait is accepted and that work on the other device still runs. + other = mod1(dev + 1, KI.ndevices(b)) + slow_fill(A, 1.0f0) + ev = KI.record_event(b) + task = Threads.@spawn begin + KI.device!(b, other) + KI.wait_event(b, ev) + B = KI.ones(b, Float32, N) + KI.synchronize(b) + Array(B) + end + @test all(==(1.0f0), fetch(task)) + KI.synchronize(b) + end + end + return nothing +end diff --git a/lib/KernelInterface/test/testsuite.jl b/lib/KernelInterface/test/testsuite.jl index f69841a6b..4457ef3df 100644 --- a/lib/KernelInterface/test/testsuite.jl +++ b/lib/KernelInterface/test/testsuite.jl @@ -27,12 +27,17 @@ end include("interface.jl") +include("events.jl") function testsuite(backend, backend_str, backend_mod, AT, DAT; skip_tests = Set{String}()) @conditional_testset "Interface" skip_tests begin interface_testsuite(backend, AT) end + @conditional_testset "Events" skip_tests begin + events_testsuite(backend) + end + return end From 5da28451629871ba899fe197db32b1a16a795471 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 21 Sep 2026 14:44:28 -0300 Subject: [PATCH 2/5] Formatting --- lib/KernelInterface/test/events.jl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lib/KernelInterface/test/events.jl b/lib/KernelInterface/test/events.jl index b50b6111a..52cef21f3 100644 --- a/lib/KernelInterface/test/events.jl +++ b/lib/KernelInterface/test/events.jl @@ -10,7 +10,7 @@ function slow_fill_kernel(A, v, ::Val{iters}) where {iters} i = KI.get_global_id().x acc = UInt32(i) for k in UInt32(1):UInt32(iters) - acc = acc * 0x19660d + k + acc = acc * 0x0019660d + k end if i <= length(A) @inbounds A[i] = ifelse(acc == 0x12345678, -v, v) From 5acb42cbbe618ce7f35a76799f4ddc9d95bba6c1 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 21 Sep 2026 15:11:41 -0300 Subject: [PATCH 3/5] More robust --- lib/KernelInterface/test/events.jl | 28 +++++++++++++++++----------- 1 file changed, 17 insertions(+), 11 deletions(-) diff --git a/lib/KernelInterface/test/events.jl b/lib/KernelInterface/test/events.jl index 52cef21f3..2fb13d124 100644 --- a/lib/KernelInterface/test/events.jl +++ b/lib/KernelInterface/test/events.jl @@ -28,22 +28,26 @@ function events_testsuite(backend) iters = Val(2^22) slow_fill(A, v) = KI.@kernel b numworkgroups = 1 workgroupsize = N slow_fill_kernel(A, v, iters) - # Time one launch, after a warm-up that absorbs compilation, and queue enough of - # them back to back to keep the spawner's queue busy for about 100ms. + # Time a launch, after a warm-up that absorbs compilation, and queue enough of + # them back to back to keep the spawner's queue busy for a couple hundred + # milliseconds. The minimum of a few runs discards one-off stalls such as a GC + # pause, which would otherwise inflate the launch count. A = KI.zeros(b, Float32, N) slow_fill(A, 1.0f0) KI.synchronize(b) - slow_time = @elapsed begin - slow_fill(A, 1.0f0) - KI.synchronize(b) + slow_time = minimum(1:3) do _ + @elapsed begin + slow_fill(A, 1.0f0) + KI.synchronize(b) + end end - launches = clamp(ceil(Int, 0.1 / slow_time), 4, 64) + launches = clamp(ceil(Int, 0.2 / slow_time), 4, 64) @testset "ordered across tasks" begin # The child queues nothing but the wait, so its `synchronize` can only return # once the spawner's queued work has drained. A backend that forgets # `wait_event` for its event type fails with a MethodError here, and a - # `wait_event` that does nothing returns in about a millisecond. The clock + # `wait_event` that does nothing returns in a few milliseconds. The clock # starts before `record_event`, so a backend whose `record_event` is the # default full `synchronize` passes just the same. The data check alone would # not do: drivers that track hazards between command buffers (Metal, for its @@ -62,11 +66,13 @@ function events_testsuite(backend) elapsed, Array(A) end elapsed, result = fetch(task) - # Half the expected drain time leaves room for GPU clock ramp-up between the - # calibration launch and these. - @test elapsed >= launches * slow_time / 2 - @test all(==(Float32(launches)), result) KI.synchronize(b) + drained = (time_ns() - start) / 1.0e9 + # Compare against the drain time of this very run rather than the calibration, + # so a stall during calibration cannot fail a correct backend. Half of it + # leaves room for a GC pause in the spawner's final `synchronize`. + @test elapsed >= drained / 2 + @test all(==(Float32(launches)), result) end if KI.ndevices(b) > 1 From 2cb82d00dd4f89fe832c4575fe9ac668de541d2a Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 21 Sep 2026 15:22:37 -0300 Subject: [PATCH 4/5] Try --- lib/KernelInterface/test/events.jl | 63 ++++++++++++++++-------------- 1 file changed, 34 insertions(+), 29 deletions(-) diff --git a/lib/KernelInterface/test/events.jl b/lib/KernelInterface/test/events.jl index 2fb13d124..56ae6e09c 100644 --- a/lib/KernelInterface/test/events.jl +++ b/lib/KernelInterface/test/events.jl @@ -3,14 +3,15 @@ # it is replayed here by hand with `Threads.@spawn`: record in the spawning task, then # `device!`, `wait_event` and `synchronize` in the child task. -# Burns a fixed amount of work per work-item before writing `v`. The accumulator is a -# linear congruential step, which the compiler cannot fold away, and it feeds into the -# store so the loop cannot be dropped. -function slow_fill_kernel(A, v, ::Val{iters}) where {iters} +# Burns `iters` dependent steps per work-item before writing `v`. The accumulator is +# a linear congruential step, which the compiler cannot fold away, and it feeds into +# the store so the loop cannot be dropped. `iters` is a run-time argument so the +# duration can be tuned below without recompiling. +function slow_fill_kernel(A, v, iters::UInt32) i = KI.get_global_id().x acc = UInt32(i) - for k in UInt32(1):UInt32(iters) - acc = acc * 0x0019660d + k + for k in UInt32(1):iters + acc = acc * 0x19660d + k end if i <= length(A) @inbounds A[i] = ifelse(acc == 0x12345678, -v, v) @@ -22,26 +23,28 @@ function events_testsuite(backend) b = backend() dev = KI.device(b) - # Few work-items and many dependent iterations: tens of milliseconds on a GPU, - # and still well under a second per launch on a CPU-backed queue. N = 64 - iters = Val(2^22) - slow_fill(A, v) = KI.@kernel b numworkgroups = 1 workgroupsize = N slow_fill_kernel(A, v, iters) - - # Time a launch, after a warm-up that absorbs compilation, and queue enough of - # them back to back to keep the spawner's queue busy for a couple hundred - # milliseconds. The minimum of a few runs discards one-off stalls such as a GC - # pause, which would otherwise inflate the launch count. A = KI.zeros(b, Float32, N) - slow_fill(A, 1.0f0) - KI.synchronize(b) - slow_time = minimum(1:3) do _ - @elapsed begin - slow_fill(A, 1.0f0) - KI.synchronize(b) + slow_fill(v, iters) = KI.@kernel b numworkgroups = 1 workgroupsize = N slow_fill_kernel(A, v, UInt32(iters)) + + # Time a launch as the minimum of a few runs: a backend's `synchronize` may run a + # GC or otherwise stall once in a while, and the minimum discards that. + function time_launch(iters) + return minimum(1:3) do _ + @elapsed begin + slow_fill(1.0f0, iters) + KI.synchronize(b) + end end end - launches = clamp(ceil(Int, 0.2 / slow_time), 4, 64) + + # Tune the kernel to about 10ms per launch, after a warm-up that absorbs + # compilation, and queue enough launches for a couple hundred milliseconds. + base = 2^20 + time_launch(base) + iters = clamp(round(Int, base * 0.01 / time_launch(base)), base, 2^30) + launches = 20 + expected = launches * time_launch(iters) @testset "ordered across tasks" begin # The child queues nothing but the wait, so its `synchronize` can only return @@ -52,8 +55,12 @@ function events_testsuite(backend) # default full `synchronize` passes just the same. The data check alone would # not do: drivers that track hazards between command buffers (Metal, for its # default buffers) order the readback after the fills without any wait. + # + # Collect beforehand so that a GC pause is unlikely to land inside the + # measurement and mask a missing wait. + GC.gc() for v in 1:launches - slow_fill(A, Float32(v)) + slow_fill(Float32(v), iters) end start = time_ns() ev = KI.record_event(b) @@ -67,11 +74,9 @@ function events_testsuite(backend) end elapsed, result = fetch(task) KI.synchronize(b) - drained = (time_ns() - start) / 1.0e9 - # Compare against the drain time of this very run rather than the calibration, - # so a stall during calibration cannot fail a correct backend. Half of it - # leaves room for a GC pause in the spawner's final `synchronize`. - @test elapsed >= drained / 2 + # A third of the calibrated drain time leaves room for the device clocking up + # between calibration and this run; a missing wait is far below that. + @test elapsed >= expected / 3 @test all(==(Float32(launches)), result) end @@ -82,7 +87,7 @@ function events_testsuite(backend) # ordering itself is not observable without peer access; check that the # wait is accepted and that work on the other device still runs. other = mod1(dev + 1, KI.ndevices(b)) - slow_fill(A, 1.0f0) + slow_fill(1.0f0, iters) ev = KI.record_event(b) task = Threads.@spawn begin KI.device!(b, other) From 91f952af75a8d71c6efeb8934990e98a19d160f6 Mon Sep 17 00:00:00 2001 From: Christian Guinard <28689358+christiangnrd@users.noreply.github.com> Date: Mon, 21 Sep 2026 15:23:05 -0300 Subject: [PATCH 5/5] Format --- lib/KernelInterface/test/events.jl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lib/KernelInterface/test/events.jl b/lib/KernelInterface/test/events.jl index 56ae6e09c..6ff22e089 100644 --- a/lib/KernelInterface/test/events.jl +++ b/lib/KernelInterface/test/events.jl @@ -11,7 +11,7 @@ function slow_fill_kernel(A, v, iters::UInt32) i = KI.get_global_id().x acc = UInt32(i) for k in UInt32(1):iters - acc = acc * 0x19660d + k + acc = acc * 0x0019660d + k end if i <= length(A) @inbounds A[i] = ifelse(acc == 0x12345678, -v, v)