Repository navigation
metal: dispatch/commit/sync/wait counters (gemma4-qat#181) - #3
Conversation
Adds mlx::core::metal::counters()/reset(): process-wide relaxed std::atomic<uint64_t> counts of kernel dispatches (CommandEncoder dispatch_threadgroups/dispatch_threads), command-buffer commits (CommandEncoder::commit) and blocking syncs (CommandEncoder::synchronize, waitUntilCompleted). Always compiled; no_metal.cpp returns zeros. Host waits on an array's event (eval/item) are not counted as syncs. Read only after the measured eval completes. doctest in tests/metal_counters_tests.cpp.
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 5447675087
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
…emma4-qat#181, Codex P2) Adds a fourth field `waits` to metal::Counters: host blocking waits on GPU completion, counted once per wait call (not blocking time) at every host-side wait in the Metal backend: EventImpl::wait (eval()/item() on an array still in flight, CPU streams waiting on a GPU event), the waitUntilCompleted in CommandEncoder::synchronize(), and the CPU-stream spin on a fast fence. `commits` includes automatic buffer splits and is not a host wait; `syncs` is documented as explicit stream synchronizations. counters()/reset() are independent loads/stores and only meaningful while no eval is in flight. tests/metal_counters_tests.cpp checks the new field (eval waits at most once, re-eval does not wait, synchronize adds exactly one wait). Verified on Apple Silicon: full doctest suite 280/280.
jeethu
changed the title
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: aaefa50058
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
…waits (gemma4-qat#181, Codex P2)
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2e5d6ccdf0
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
…ecovery syncs (gemma4-qat#181, Codex P2)
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 1fee32b6e6
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
1 participant
Part of numen-tech/gemma4-qat#178 (Phase 0), child numen-tech/gemma4-qat#181.
Stacks on #2 (
fork/86-affine-sym-0.32.2); retarget after #2 merges.Adds
mlx::core::metal::counters()/reset()(struct Counters { uint64_t dispatches, commits, syncs, waits; }, declared inmlx/backend/metal/metal.h):CommandEncoder::dispatch_threadgroups/dispatch_threads(every compute dispatch goes through these two).CommandEncoder::commit(the eval-timeneeds_commitsplit,gpu::finalize, andsynchronize).CommandEncoder::synchronize, i.e.mx.synchronize()).eval/item()on an array still in flight (event.cpp), a CPU stream waiting on a GPU event (fence.cpp, counted inside the CPU-stream lambda so it counts the spin, not the enqueue), and a sync'swaitUntilCompleted(so every sync is also a wait). MLX skips the wait when the array's event has already signaled, sowaitsmeasures stalls actually taken, never blocking time. Internal hook:metal::count_host_wait()(device.h), called immediately before each blocking wait.Relaxed
std::atomic<uint64_t>increments, always compiled;no_metal.cppreturns zeros. No happens-before with the GPU: snapshot only after the measured eval has completed and with no other thread encoding.Test:
tests/metal_counters_tests.cpp(doctest; matmul+eval bumps dispatches/commits with 0 syncs and ≤1 wait, a re-eval of an already-evaluated array adds no wait,synchronize()bumps syncs, commits and waits,reset()zeroes). Full C++ suite: 280/280 pass (Release,-DMLX_METAL_JIT=ON). No Python binding (not needed by the consumer; the C++ test covers it).Consumer: numen-tech/mlx-swift#4 (
prism-181-counters: extern "C" shim +MLX.GPUCounters), then gemma4-qat'sNumenDispatchTestsdispatch-budget gate.Review log: Codex round on gemma4-qat#181 (P2: host waits were not counted, and
syncsconflated explicit synchronizations with implicit host waits) →aaefa5005addswaitsand narrowssyncs./simplifypass: no change (already minimal).