trace: event tracer for the RPC backend, rpc-server and llama-server by danielhanchen · Pull Request #191 · unslothai/llama.cpp

danielhanchen · 2026-09-05T23:24:57Z

What this is

A timing and profiling system for the RPC backend and llama-server, so that the bottlenecks of a
two node layer split are measured rather than inferred. It replaces the ad hoc pieces used so far
(RPC command counters in the client, per group timers in the server) with one tracer that both
nodes write to and one tool that merges them.

Based on feature/pipeline-groups.

Design

ggml/include/ggml-trace.h and ggml/src/ggml-trace.cpp live in ggml-base, so every binary that
takes part (libllama, llama-server, the RPC backend, ggml-rpc-server) can raise events without a
new dependency.

Off by default, and free when off. The tracer writes nothing unless GGML_RPC_TRACE=<path>
is set, or ggml-rpc-server --trace <path>. Call sites read the exported flag directly:
```
const int64_t t0 = ggml_trace_flag ? ggml_trace_time_us() : 0;
```
so with tracing off the cost is one load and one branch; no clock is read, nothing is formatted
and no allocation happens. There is no change to any output byte with the tracer off or on.
JSON lines. One header object per file, then one object per event, one file per process.
A line is built in the calling thread and handed to the file under a mutex, and the file is
flushed every 128 lines so that a process stopped with a signal does not lose its tail.
Clock alignment. A new RPC_CMD_TRACE_SYNC carries a four timestamp exchange (client sends
at t1, peer stamps t2 on receive and t3 on reply, client stamps t4), done once per connection at
connect time and only while tracing is on; the peer always answers it. The resulting offset is
written into the client trace, and the merge tool uses it to put both nodes on one time line.
Measured offsets on the pair were about 1.8 s between the two boot clocks, with round trips in
the tens of microseconds.

What is recorded

RPC client, one record per command: command type, bytes sent and received, the tensor name or
graph uid it belongs to, the calling thread, the llama-server pipeline group (threads tag
themselves through a thread local), and timestamps at enqueue into the dispatcher, first byte
sent, last byte sent, our turn in the reply order, first byte of the reply, reply complete.

Scheduler, per split: the backend, the input count and the node count; and the staging path of
ggml_backend_tensor_copy broken into host allocation, read back and send. That is the
Spark 1 to CPU to Spark 2 hop of a layer split, visible as its own span.

rpc-server, one record per command served: receive start and end, execute start and end, reply
start and end, bytes each way, graph node count and device.

GPU. A backend submit only queues the work, so host timestamps around it say nothing about
when the kernels ran. The CUDA backend exports two entry points through the registry
(ggml_backend_cuda_trace_mark, ggml_backend_cuda_trace_poll), so ggml-base and the RPC backend
do not have to link against CUDA. They record CUDA events on the compute stream around
ggml_backend_graph_compute and report the completions on the host monotonic microsecond scale,
anchored by one event whose completion time on the host was measured once. Nothing ever waits on
the GPU: the completed spans are collected from points where the caller is idle anyway (the serve
loop between commands, the scheduler synchronize). The hooks are looked up per backend registry,
so a scheduler that runs over a CPU or RPC backend as well still gets the CUDA rows.

llama-server, per group and per iteration: batch build, submit (llama_decode),
llama_synchronize, post decode, sampling and result send, with slot counts. libllama, per
llama_decode: the scheduler split count, with the per split backend coming from the scheduler
events above.

Tools

scripts/rpc_trace/merge.py reads the files of both nodes, aligns them with the measured clock
offset and writes

a Chrome trace (chrome://tracing, or the Perfetto UI) with one row per node, thread and
pipeline group, the phases of each RPC command as nested slices, and one row per GPU carrying
the CUDA event timings, and
a text summary per decode step: local compute, transfer, peer compute, logits return, sampling,
the idle fraction of each GPU, the longest stretch of the step in which neither GPU was busy,
and the total idle time attributed to each host phase.

scripts/rpc_trace/cpu_check.sh is the CPU validation: two rpc-servers on the CPU backend and one
llama-server splitting layers over them, run once with the trace off and once with it on, with the
generated text compared. scripts/rpc_trace/gpu_trace.sh and scripts/rpc_trace/nonrpc_bracket.sh
run the cells below.

Overhead

Measured, not assumed: every configuration was run with the tracer off and on.

Layer split over two nodes, Qwen3.8-27B UD-Q4_K_XL, 32 concurrent, npp 128 / ntg 256,
--cache-ram 0, -c 16384, whole cell tok/s:

configuration	trace off	trace on
N=1 `--device CUDA0,RPC0`	95.91	96.21
N=1 `--device RPC0,CUDA0`	101.15	99.91
N=2 `--device RPC0,CUDA0`	138.96	141.62

That is +0.3, -1.2 and +1.9 percent, in both directions, so within run to run noise. The client
writes about 865 events per second (19 MB over a 170 s cell) and the peer about 60 per second.

Non-RPC and non-CUDA workloads

The tracer lives in ggml-base and adds two entry points to the CUDA backend's registry, so it is
on the path of workloads that never touch RPC. It costs them nothing.

Single GPU, no --rpc, llama-batched-bench -c 32768 -npp 512 -ntg 128 -npl 1,8,32 on the same
model, run base / new / base / new-with-the-trace-on, prompt and generation tok/s:

pass	npl 1	npl 8	npl 32
base	823.52, 11.72	829.56, 59.98	830.42, 114.97
new	821.85, 11.68	826.79, 58.49	679.02, 111.64
base (repeat)	675.46, 11.47	679.68, 56.18	680.51, 110.75
new, trace on	669.13, 11.34	678.69, 55.87	678.43, 110.57

The node dropped to a lower clock state between the second and the third pass (prompt throughput
falls from about 828 to about 679 for both builds), so the comparisons that hold are the ones
inside one clock state: new against base at npl 1 and 8 is -0.2 and -0.3 percent on prompt, new
against the base repeat at npl 32 is -0.2 percent on prompt and +0.8 on generation, and new with
the trace on against the base repeat is -0.3 and -0.2 percent. Nothing outside noise.

Bit exactness, five greedy prompts, 48 tokens each:

single GPU, no RPC: md5 74926c4ef135f3cc89ad20cd5ec7e445 for the base, for this branch with
the trace off, and for this branch with the trace on;
CPU only build, -DGGML_CUDA=OFF -DGGML_RPC=OFF: md5 177dc61e0703eba3bdaf7bf1131f0458 for
the same three.

Build configurations checked: -DGGML_RPC=OFF -DGGML_CUDA=ON and -DGGML_CUDA=OFF -DGGML_RPC=OFF both build clean, and the CPU only build still produces a usable trace of the
llama and scheduler events with no GPU rows.

Other backends. Vulkan has no get_proc_address at all and the registry returns NULL for that
case; Metal and SYCL fall through to NULL for a name they do not know. ggml_trace_gpu_begin
therefore gets no hook, returns 0 and records nothing, and no other line of those backends is
touched. ggml-cuda.cu is also the HIP and MUSA source, and the vendor headers did not map
cudaEventCreate, cudaEventQuery or cudaEventElapsedTime, so those two builds would not have
compiled; the three defines are added.

The three timelines

Per decode step, in milliseconds, from the traces. localGPU and peerGPU come from CUDA events
on the compute stream of each node, so they are what the GPUs did, not what the host queued.
idle both is the part of the step in which neither GPU was busy.

configuration	step	batch build	submit	synchronize	post decode	sampling	send	localGPU	peerGPU	idle both
N=1 CUDA0,RPC0	325.6	10.1	303.0	0.0	12.2	6.9	5.1	150.0	149.7	27.1
N=1 RPC0,CUDA0	312.3	10.2	159.2	129.7	12.8	6.9	5.7	140.0	154.9	24.3
N=2 RPC0,CUDA0 group 0	219.5	6.7	107.2	94.4	8.4	6.3	2.0	195.9	199.4	5.9
N=2 RPC0,CUDA0 group 1	221.4	8.6	110.4	92.7	6.7	3.8	2.8	194.6	199.3	6.6

GPU busy over the whole cell, and RPC bytes per step:

configuration	local GPU busy	peer GPU busy	bytes out	bytes in
N=1 CUDA0,RPC0	46.1%	46.0%	1035 kB	39441 kB
N=1 RPC0,CUDA0	44.8%	49.6%	1038 kB	11509 kB
N=2 RPC0,CUDA0	87.8%	89.7%	1219 kB per group	5445 kB per group

What the traces say.

N=1 with the output layer on the peer. 39.4 MB comes back over the link every step: the F32
logits of 32 rows over a 248320 token vocabulary. llama_decode does not return until the
whole step is finished (submit 303 ms, synchronize 0), because the blocking GET_TENSOR of the
logits sits inside it. The two GPUs never overlap: 150.0 and 149.7 ms of GPU work inside a
325.6 ms step, each of them idle through the other's stage. Named bottleneck: the serial chain
itself, with the logits return as the transport cost that pays for nothing.
N=1 with the output layer on the local node. The logits stay local and only 11.5 MB comes
back, and the step now splits into submit 159 ms and synchronize 130 ms. The step barely moves
(312.3 against 325.6) because the stages still alternate: 140.0 and 154.9 ms of GPU work in a
312.3 ms step. Named bottleneck: still the serial chain. The staged copy of the hidden state
through a host buffer (ggml_backend_tensor_copy, since the RPC backend has no
cpy_tensor_async) is 148.4 ms per step here, which is the span an asynchronous device to
device path would attack; that number includes the blocking wait for the peer's graph, so it is
an upper bound on the copy itself.
N=2 with the output layer on the local node. 195.9 and 199.4 ms of GPU work in a 220 ms
step, both GPUs at 88 to 90 percent, and only 5.9 to 6.6 ms per step in which neither GPU is
busy against 24 to 27 ms with one context. Named bottleneck: GPU compute. The host path is off
the critical path because one group runs it under the other group's GPU work.
Where the host stalls both GPUs, per step: with one context, post decode 12.1 to 12.8 ms
(sampling and streaming) and batch build 8.3 to 9.5 ms; with two groups, batch build 2.6 to
3.3 ms and post decode 0.1 to 0.3 ms. The single longest idle gap in every cell is a prefill
step's batch build, 0.9 to 1.3 s, which is a time to first token cost and not a decode cost.

Validation

CPU harness (scripts/rpc_trace/cpu_check.sh, two local CPU rpc-servers, three greedy prompts):
the generated text is byte identical with the trace off and on, the trace files parse, and the
measured clock offsets between the local processes are 0 and 1 us with a 14 to 16 us round trip.

Scope, relative to the base branch

This PR targets feature/pipeline-groups, and against that base it changes 20 files, 1835 insertions and 15 deletions.
Stated explicitly because a reviewer diffs against the base, not against the point in a stack a
branch was cut from, and a scope claim measured from the wrong reference is exactly the sentence a
reviewer trusts instead of checking.

…ma-server Adds one coherent timing system in place of the ad hoc counters, so the cost of a two node layer split is measured rather than inferred. ggml/include/ggml-trace.h, ggml/src/ggml-trace.cpp JSON line writer, off unless GGML_RPC_TRACE names a file. Call sites read the exported flag directly, so with tracing off the cost is one load and one branch and no clock is read. Also holds the GPU span helper: the timing hooks are resolved through the backend registry, so ggml-base does not link against any GPU runtime. ggml-cuda two entry points exported through the registry: record a CUDA event on the compute stream, and collect the events that have completed. The completions are reported on the host monotonic scale through one anchor event whose completion time was measured once, and nothing ever waits on the GPU. ggml-rpc client: one record per command with the bytes each way, the tensor or graph it belongs to, the thread, the llama-server group, and timestamps at enqueue, first byte sent, last byte sent, our turn in the reply order, first byte of the reply and reply complete. server: one record per command served with receive, execute and reply timestamps, the graph node count and payload size, and CUDA event timestamps around ggml_backend_graph_compute. new RPC_CMD_TRACE_SYNC: a four timestamp exchange at connect time, written into the client trace so the two nodes can be put on one time line. It is only sent while tracing is on. ggml-backend per scheduler split: backend, input count, node count and a GPU span around the submit; the staging path of ggml_backend_tensor_copy broken into host allocation, read back and send, which is the Spark 1 to CPU to Spark 2 cost of the split. llama-server and libllama per group and per iteration: batch build, submit, synchronize, post decode, sampling and result send with slot counts; per llama_decode the scheduler split count. scripts/rpc_trace/merge.py aligns the files with the measured clock offset and emits a Chrome trace with one row per node, thread and group plus a row per GPU, and a per step summary with the idle fraction of each GPU and the biggest idle gap. scripts/rpc_trace/cpu_check.sh two rpc-servers on the CPU backend, run with the trace off and on, output compared.

…erge by time A scheduler runs over several backends and only some of them offer the timing hooks. The probe was cached once for the process, so a split whose first backend was the CPU or the RPC backend disabled the GPU rows for the whole run, and a backend could have been handed to another registry's mark function. The answer is now kept per registry. merge.py walked every event of the file for every phase of every step. The events are indexed by start time and looked up with a bisect instead.

ggml-cuda.cu is also the HIP and MUSA source, and the GPU timing hook uses cudaEventCreate, cudaEventQuery and cudaEventElapsedTime, which the vendor headers did not map yet. Adds the three defines and a single GPU bracket script for the non RPC proof.

…r properly A traced process is normally stopped with a signal at the end of a run, and the tail of the stdio buffer was lost with it. The file is now flushed every 128 lines. The bench script launched the peer server through setsid, so the pid it captured was setsid's and the server, its child, survived every kill. The next cell then found the port taken, logged 'Failed to create server socket' and silently talked to the previous server, which is how a whole set of peer traces came back with nothing but a header. The launch no longer goes through setsid, ssh is given -n, the bind is checked, the port is waited on before and after each cell, and any leftover server on our port is matched with pgrep -x on the binary name and the port in /proc rather than with a pattern that would also match the remote shell.

…, drop the run artifacts The summary named only the single longest stretch of a step in which neither GPU was busy, which is always a prefill batch build and says nothing about the decode steps. It now also reports the total idle time attributed to each host phase, per step. Also removes the bench output and the driver script that were committed by mistake.

…ature/rpc-trace # Conflicts: # tools/server/server-context.cpp

danielhanchen · 2026-09-07T04:22:02Z

Added scripts/rpc_trace/device_idle.py to this PR, because the tracer answered the question it was built for and then could not answer the next one.

merge.py says where a decode step goes, per pipeline group. It does not say whose fault a given idle interval on one device is, and that distinction decides what a fix would have to look like: idle while the other device computes is a scheduling problem, idle while neither device computes is a host problem, and they have different fixes. The new tool intersects one device's idle with the other device's busy time to split exactly that, reusing merge.py's union/clip/gaps/Index.covered primitives rather than duplicating them.

It also splits the cell by phase, and that turned out to be the whole answer. Time counts as prefill when any group is inside an iteration that submitted more tokens than it had slots. This has to be a SET of intervals and not a range: a serving cell interleaves prompt batches with decode for its whole length, so "everything before the last prompt iteration" puts most of the decode inside the prefill phase and reports nothing.

What it found on a 27B layer split across two DGX Sparks at 128 concurrent rows with two pipeline groups, 313 s window, 557k events, both nodes clock-pinned:

                        busy     idle   idle while the     idle with
                                        other computes   neither computing
  local                89.16%   10.84       9.71              1.13
  peer                 91.55%    8.45       7.31              1.13

  prompt-batch iterations, 18.3% of the window:  peer busy 68.81%
  decode-only time,        81.6% of the window:  peer busy 96.68%

Nine tenths of the remaining idle is one device waiting on the other, and two thirds of it lives in the 18 percent of the window that carries a prompt batch. Median idle stretch on the peer is 6.1 ms in decode and 55.2 ms in a prompt iteration.

Non-RPC and other-path proof. This commit adds one new file under scripts/ and changes nothing else: no source file, no header, no CMake target, no build option, no runtime code path. git show --stat is a single added Python file. It cannot affect NVIDIA, AMD or CPU-only builds on Linux, Windows, WSL or Mac, cannot affect Vulkan, Metal, SYCL, HIP or MUSA, cannot affect KV caching or prefix caching, and cannot affect any workload, RPC or otherwise, because nothing in the library or the tools ever reads it. It is an offline reader of the JSONL files this PR already produces, and it runs on a laptop against a trace copied off the machine.

Resolve the conflicts in ggml/src/ggml-rpc/ggml-rpc.cpp and tools/server/server-context.cpp by keeping this branch's code and the base branch's comment wording. Verified with comment_tools check that only comments changed relative to 23d52c3.

danielhanchen · 2026-09-09T01:35:31Z

@codex review

danielhanchen · 2026-09-09T01:35:38Z

@codex security review

chatgpt-codex-connector · 2026-09-09T01:35:45Z

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-09T14:30:46.494975Z	`f98e338`	Manual request

ℹ️ 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" or "@codex security review".

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

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: e3bdc70900

ℹ️ 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".

chatgpt-codex-connector · 2026-09-09T01:41:29Z

    }
+    ggml_trace_open(nullptr, "rpc-client");
+    if (ggml_trace_flag) {
+        rpc_trace_sync(sock, endpoint);


Negotiate trace sync before sending the new RPC command

When GGML_RPC_TRACE is enabled against an older protocol-compatible 5.1 server, the HELLO check succeeds and this unconditionally sends RPC_CMD_TRACE_SYNC, which that server treats as an unknown command and closes the connection. rpc_trace_sync() merely logs the failed exchange, after which get_socket() caches the dead socket and the first real RPC operation fails. Advertise support through the connection capabilities or otherwise avoid sending the command unless the peer supports it.

Useful? React with 👍 / 👎.

Fixed in c9f1a11. Confirmed: this branch added RPC_CMD_TRACE_SYNC but left RPC_PROTO_MINOR_VERSION at 1 and had no record of the peer version at all, so a real 5.1 server passed the HELLO check and then closed the connection on the unknown command.

The minor goes to 2, negotiate_hello records the peer minor, and the sync is sent only when the peer is at least 2.

chatgpt-codex-connector · 2026-09-09T01:41:29Z

+    if (cuda_ctx->device != st.device) {
+        return;


Track CUDA trace state separately for each device

In any traced process using multiple CUDA devices, the first mark permanently selects one device and every mark for the other devices is silently discarded here. The higher-level tracer still returns nonzero tags because all CUDA devices share the supported registry, so multi-GPU llama-server and rpc-server traces omit work from every device except the first and produce incorrect utilization and idle summaries. Anchors and pending-event queues need to be maintained per CUDA device.

Useful? React with 👍 / 👎.

Fixed in 0405b69. The anchor and both queues are now per CUDA device, so the device check that used to drop marks is gone rather than made conditional.

chatgpt-codex-connector · 2026-09-09T01:41:29Z

+        st.device = cuda_ctx->device;
+        // every later mark is reported as anchor_us + elapsed(anchor, mark)
+        cudaEventRecord(st.anchor, cuda_ctx->stream());
+        cudaEventSynchronize(st.anchor);


Avoid synchronizing the stream while creating the trace anchor

On the first CUDA mark, this waits for the entire compute stream even though tracing is intended not to perturb execution. In the scheduler path, asynchronous input copies can already have been queued before ggml_trace_gpu_begin(), so enabling tracing forces those copies to complete before the graph is submitted, altering first-request latency and potentially destroying cross-device overlap in the very trace being measured. Establish the anchor asynchronously and resolve it during polling instead.

Useful? React with 👍 / 👎.

Fixed in 0405b69. The anchor is recorded and never waited on; poll takes its wall clock the first time it sees the anchor complete, which works because the anchor is recorded before any mark on that stream.

Stating the tradeoff plainly: absolute timestamps now carry an offset bounded by the poll interval, where before they were exact but the act of measuring drained the stream. The spacing between marks is unaffected either way, and that is what the idle summaries use.

…trace

…eature/rpc-trace

…ure/rpc-trace

…eline-groups into feature/rpc-trace

RPC_CMD_TRACE_SYNC was added without bumping the protocol minor, so a genuine 5.1 server passed the HELLO version check and then closed the connection on the unknown command. rpc_trace_sync only logged the failed exchange, after which get_socket cached the dead socket and the first real operation on it failed. The minor goes to 2, the peer's minor is recorded from its HELLO response, and the sync is sent only when the peer is at least 2.

Two problems in the GPU timing marks. The anchor and the pending queues were global and bound to whichever device marked first, and every mark from another device was dropped by the device check. The tracer above still handed out tags for those devices, so a multi-GPU trace silently lost all of their work and reported utilization and idle summaries computed from one device. State is now kept per CUDA device. The first mark also called cudaEventSynchronize on the anchor, which drains whatever the scheduler has already queued on that stream, so turning tracing on changed the execution it was supposed to observe. The anchor is only recorded now; its wall clock is taken in poll the first time it is seen complete. That leaves absolute timestamps with an offset bounded by the poll interval and leaves the spacing between marks exact.

…trace Takes the draft-context thread pools, the device-identity fix in the pipeline-group fit reservation, and the router child environment fix. No conflicts this time: they sit in regions the tracing spans do not touch. The merged tree builds llama-server and ggml-rpc-server before this commit was made.

danielhanchen · 2026-09-09T05:24:08Z

@codex review

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 3ab45d9c9d

ℹ️ 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".

chatgpt-codex-connector · 2026-09-09T05:37:13Z

+            if (cudaEventQuery(d.anchor) != cudaSuccess) {
+                continue;
+            }
+            d.anchor_us = ggml_time_us();


Anchor GPU timestamps to the event completion time

When the first poll happens after the anchor has already completed, this records the poll time rather than the anchor's actual completion time, then adds every event's elapsed offset to it. In the RPC server this is guaranteed for the first graph because ggml_backend_graph_compute() synchronizes before the next serve-loop poll, so the whole GPU timeline is shifted forward by roughly that graph's duration, corrupting cross-device overlap and idle attribution. Fresh evidence after the prior synchronization fix is that the replacement now derives the permanent wall-clock anchor solely from this delayed poll.

Useful? React with 👍 / 👎.

Agreed, and fixed in eee6abe.

The point stands on its own: anchor_us was ggml_time_us() taken at the first poll that observed the anchor complete, and every mark is reported as anchor_us + elapsed(anchor, mark), so the entire device timeline shifts forward by however long the anchor had already been complete. In the RPC server that delay is a whole graph, because ggml_backend_graph_compute() synchronizes before the serve loop polls again.

The constraint the original code was respecting is real and the fix keeps it: the anchor still is not waited on at record time, because synchronizing there would drain whatever the scheduler has already queued and change the execution the tracer is supposed to be observing.

Instead the wall clock is established in poll, and only when cudaStreamQuery reports the stream already idle. In that state a freshly recorded probe completes immediately, so cudaEventSynchronize returns at its completion and drains nothing, and cudaEventElapsedTime(anchor, probe) gives an exact tie between GPU and wall time. An idle stream is the normal condition at poll time in the serve loop, which is the case this item is about. A busy stream or a failed probe falls back to the previous approximation.

One thing I got wrong in the first draft and corrected before building, since it is the sort of thing worth recording: I originally allowed the anchor to be refined on a later poll if the first attempt fell back. That would move marks reported after the refinement relative to marks reported before it, putting a step in the middle of a single device timeline, which is harder to reason about than a consistent offset. The anchor is now frozen once established, on both paths.

What I did not run, stated plainly. This compiles in the CUDA build, but I did not measure it against a GPU. Demonstrating the numeric shift needs a device and, to be meaningful, a multi-GPU or two-node RPC trace; the pair lock on this machine was held continuously by other work for the whole window, and the standing rule here is to take it only where a device is genuinely required. So for this item the evidence is compilation and source-level reasoning, not measurement. The same caveat applies to the device-identity item. The other four are measured.

chatgpt-codex-connector · 2026-09-09T05:37:13Z

+    if (ggml_trace_flag) {
+        tls_srv.n_nodes = (int) n_nodes;
+        tls_srv.device  = (int) device;
+        tls_srv.gpu_tag = ggml_trace_gpu_begin(backends[device], "GRAPH_COMPUTE");


Preserve the RPC GPU's device identity

When one rpc-server exposes multiple CUDA devices, every device's span is named GRAPH_COMPUTE, and the emitted GPU event contains only that name and gpu_tag. merge.py uses the event name as its per-device row key, so it collapses all server GPUs into one row and unions them into one peer-utilization value. Fresh evidence after the prior per-device queue fix is that marks are now retained for every device but their device identity is discarded at this call site.

Useful? React with 👍 / 👎.

Confirmed and fixed in eee6abe.

The chain is exactly as described, and I checked the consuming end rather than assuming it: merge.py line 163 uses key = e.get("n", "gpu") as its per-device row key and labels the row GPU %s % key. Since the GPU event carries only the span name and the tag, and every device passed the fixed string GRAPH_COMPUTE, one rpc-server exposing several GPUs collapsed into a single row and a single unioned peer-utilization value.

The device index is now part of the span name, at both graph_compute and graph_recompute. Putting it in the name rather than adding a field is deliberate: it restores per-device rows and per-device labels with no change to the analysis scripts at all.

Checked before using a stack buffer: ggml_trace_gpu_begin stores the name as st.names[tag] = name, a std::string, so the value is copied and the buffer does not need to outlive the call.

Not measured, and worth being explicit about: this needs an rpc-server with more than one CUDA device to observe, which I did not have a lock window for. It compiles in the CUDA build. The verification here is source-level, on both the producing and the consuming side, not a captured multi-GPU trace.

chatgpt-codex-connector · 2026-09-09T05:37:13Z

 int llama_server(int argc, char ** argv) {
    std::setlocale(LC_NUMERIC, "C");

+    ggml_trace_open(nullptr, "llama-server");


Give router children separate trace files

When GGML_RPC_TRACE is enabled in router mode, the router opens this path and every spawned model child inherits the same environment variable through server_models::base_env. Each child then executes this call with the identical path, and ggml_trace_open() uses fopen(..., "wb"), so loading a child truncates the router or previous child's trace and concurrently running children overwrite one another through independent file offsets. Router tracing therefore needs per-process paths or must remove/replace the inherited variable before spawning children.

Useful? React with 👍 / 👎.

Confirmed and fixed in eee6abe.

Reproduced as a flip pair, both arms run end to end with a real child spawn, router plus one model, GGML_RPC_TRACE pointed at trace.jsonl.

Before, at 40d69fb: both the router and the child logged ggml_trace_open: tracing to .../trace.jsonl, the same path. One file survived, 4216 bytes, carrying a single header whose pid was the child. The router opened first and the child truncated it away, exactly as described.

After: two files. trace.jsonl, 120 bytes, holding the router header for pid 1671191 intact, and trace.stories260K.41219.jsonl, 8312 bytes, holding the child events for pid 1671650.

The router now replaces the inherited variable per child rather than appending to it, since execve keeps duplicate entries and getenv returns the first, so appending could not override an inherited value. The child path is derived from the base name with the suffix inserted before the extension, and it is keyed by port: a port is unique among live children, whereas a model name can carry / and : from an HF repo spec and is not filesystem-safe on every platform we build for, so the name is sanitised into the label rather than trusted as the key.

chatgpt-codex-connector · 2026-09-09T05:37:13Z

+            logits = [(e["t0"], e.get("t_recv1", e["t1"])) for e in cmds
+                      if e.get("n") == "GET_TENSOR" and e.get("bytes_in", 0) >= LOGITS_MIN_BYTES]


Identify logits by tensor rather than response size

In a layer split where a large hidden-state tensor is copied back from RPC for later local layers, its GET_TENSOR response also exceeds 64 KiB and is counted here as a logits return. Hidden-state and embedding transfers routinely exceed this threshold, so the summary's logits column can include unrelated staged transfers and materially misstate where return bandwidth is spent; use the recorded tensor subject or another explicit output marker instead of byte size.

Useful? React with 👍 / 👎.

Confirmed and fixed in eee6abe.

The identifying information was already in the trace, so no protocol change was needed. ggml_backend_rpc_buffer_get_tensor calls ggml_trace_set_subject(tensor->name, 0) immediately before the send, and that is emitted as subj. Logits are always result_output; embeddings are result_embd, result_embd_pooled or result_norm. The classification now keys on that rather than on bytes_in >= 64 KiB.

Reproduced with a synthetic trace holding one 10 ms step with both shapes present: a staged hidden state GET_TENSOR subj="l_out-15" of 8388608 bytes, and the real output read GET_TENSOR subj="result_output" of 513024 bytes.

case logits before logits after

staged hidden state plus real output 3.5 ms 0.5 ms

control, real output only 0.5 ms 0.5 ms, row byte-identical

control, trace predating subj 3.5 ms 3.5 ms, with an explicit note that it fell back

The transfer column stays at 3.5 ms throughout, so the staged copy is still accounted for in return bandwidth, it is simply no longer attributed to logits. Traces written before subj existed keep the old size heuristic and say so rather than silently reporting a different number.

chatgpt-codex-connector · 2026-09-09T05:37:13Z

+        g = e.get("grp", 0)
+        n_slots = e.get("n1", 0)
+        toks = [n for a, b, n in subs.get(g, []) if a >= e["t0"] and b <= e["t1"]]
+        (pre if (toks and max(toks) > max(n_slots, 1)) else dec)[g].append((e["t0"], e["t1"]))


Distinguish speculative decode batches from prefill

With speculative decoding enabled, a normal decode iteration submits several drafted tokens per processing slot, so max(toks) > n_slots is true and this classifies every such decode iteration as prefill. Consequently t_pre_end can move to the end of the run and the reported decode-phase idle and group-offset statistics become empty or incorrect. The phase decision needs an explicit prompt/decode signal rather than the tokens-per-slot heuristic.

Useful? React with 👍 / 👎.

Confirmed and fixed in eee6abe, and you were right that the heuristic was the wrong instrument rather than a badly tuned one.

Checking what was available: nothing in the trace stated the phase. iteration carried only n0 = grp.id and n1 = n_slots_processing, and submit only the batch token count. So rather than replace one inference with another, the server now emits the phase explicitly. server_trace_scope gained an optional prompt field and the iteration span sets it from the slot states read before pre_decode(), which is the code that acts on them: a slot in SLOT_STATE_STARTED or SLOT_STATE_PROCESSING_PROMPT is what puts prompt tokens in this batch.

Verified on a real trace rather than a synthetic one: a 24-token request against a CPU model produced exactly 1 iteration with "prompt":1, lasting 14 ms, followed by 22 with "prompt":0, each about 1 ms.

device_idle.py reads that flag first. Reproduced the original defect on a synthetic 2-prefill, 10-speculative-decode trace submitting 5 tokens for 1 processing slot:

before: prefill ends 0.240 s in (100.0% of the window) prefill iterations 12, decode iterations 0 no decode phase section emitted at all group 0: 0 decode iterations, covering 0.0% of the decode phase after: prefill ends 0.040 s in (16.7% of the window) prefill iterations 2, decode iterations 10 decode phase 0.200 s, local busy 40.00% idle 60.00%, peer busy 27.50% idle 72.50% group 0: 10 decode iterations, median 20.0 ms, covering 100.0% of the decode phase

Controls, both byte-identical before and after: an ordinary non-speculative trace, and a multimodal one where an image chunk decodes inside batch_build, which is the case most likely to be misread by any drafting-based signal. Both stay at 2 prefill and 10 decode.

A tokens-per-slot fallback is kept only for traces written before the flag existed, and with no drafting present it reduces to the old rule exactly.

One related gap I did not fix, flagging it rather than leaving it buried: the speculative accept path wraps neither common_sampler_sample_and_accept_n nor the process_token loop in a server_trace_scope, so under speculation there are no sampling or result_send spans and that generation work is invisible in the trace. Happy to add those in this PR if you would prefer it here rather than as a follow-up.

chatgpt-codex-connector · 2026-09-09T05:37:13Z

+        for host, values in by_host.items():
+            if host == f.host or f.host.startswith(host) or host.startswith(f.host):
+                cand = values


Match peer traces without relying on hostname text

With two or more remote RPC peers addressed by IP or DNS aliases that differ from their gethostname() values, none of the synchronized endpoint names match the corresponding trace headers here and the single-peer fallback is disabled. Those server files retain offset_us = 0, leaving events on unrelated monotonic clocks and corrupting merged overlap and idle calculations; Windows is worse because its blank header hostname prefix-matches the first endpoint for every peer. Record an unambiguous endpoint identity or otherwise associate each trace with its own clock exchange.

Useful? React with 👍 / 👎.

Confirmed and fixed in eee6abe, including the Windows case, which turned out to be the worse of the two exactly as you said.

Reproduced both, with two peers, A running 500 ms ahead and B 300 ms behind, and a true peer-GPU coverage of 40% of the window:

case before after

hosts gpu-a / gpu-b, endpoints 10.0.0.11:50052 / 10.0.0.12:50052 two "no clock offset" warnings, offsets +0.000 / +0.000 ms, peer GPU busy 0.0% offsets +499.999 / -300.001 ms, peer GPU busy 40.0%

Windows, both header hosts blank no warning at all, both peers bound to the first endpoint at +499.999 ms, peer GPU busy 20.0% +499.999 / -300.001 ms, peer GPU busy 40.0%

control, single peer, IP endpoint, no TRACE_SYNC in the trace +499.999 ms, 20.0% unchanged, fallback still fires

The Windows row is the one worth noting: it produced no diagnostic. host.startswith("") is true, so a blank header matched the first endpoint for every peer and the tool reported a confidently wrong number rather than declining to answer.

The fix stops using hostname text as the primary key. The clock_offset record already carries the peer sample t2/t3 taken on the peer clock, and the peer traced that same command as a TRACE_SYNC span with t_recv0/t_send1 on that clock, so a record belongs to the peer whose span brackets [t2, t3]. That is exact and needs no hostname. Hostname matching is now the second choice and was tightened: it requires a non-empty host, a dot-boundary match so gpu-a still matches gpu-a.lan while gpu-ab no longer does, a unique endpoint hit, and no rival peer on the same host. Buckets are keyed by the whole endpoint, so two peers on one host differing only by port no longer collapse, and [::1]:port parses correctly.

Where this is still not airtight, and what would close it. The bracket match falls through to the weaker paths if the peer is on a build without RPC_CMD_TRACE_SYNC, if the peer trace was opened after the sync ran, or if two peers on genuinely identical clocks both bracket the same interval. The clean fix is on the C++ side: ggml_trace_open writes only role, host, pid, t_open_us and wall_us, and the rpc-server already knows its bind string. Emitting that endpoint plus a random per-process id, with the client echoing the id in the clock_offset record, would make this a direct key lookup and retire the heuristics entirely. I left that out of this commit because it changes the trace header format and felt like your call. Say the word and I will add it.

Separately: the blank Windows host is itself a small fix. gethostname is skipped under _WIN32 because it is a Winsock call needing WSAStartup and ws2_32, but GetComputerNameExW(ComputerNameDnsHostname, ...) from windows.h gives the DNS host name with neither dependency. That is worth doing regardless, though it would not fix this item on its own, since a gethostname value still is not an endpoint.

…ature/rpc-trace Takes the change that stops dividing n_ctx when the KV is unified, along with the narrowed --ctx-size divisibility check and the README correction. No conflicts, and the merged tree builds llama-server before this commit was made.

Each was reproduced before being changed, and each has a control showing the unaffected path is untouched. The three script fixes are driven from synthetic traces built to contain the specific shape that breaks them. ggml-cuda: anchor GPU timestamps to the anchor's completion, not to the poll that noticed it. Marks are reported as anchor_us + elapsed(anchor, mark), and anchor_us was ggml_time_us() taken at the first poll that saw the anchor complete, so the whole device timeline shifted forward by however long it had already been complete. In the RPC server that is a full graph, because ggml_backend_graph_compute() synchronizes before the serve loop polls again, which is exactly the data cross-device overlap and idle attribution are computed from. The anchor still is not waited on at record time, since that would drain queued work and change the execution being observed. Instead poll records a probe and measures back to the anchor, but only when cudaStreamQuery reports the stream already idle, in which case the probe completes immediately and the wait drains nothing. That is the normal state at poll time in the serve loop. Busy stream or failed probe falls back to the old approximation. The anchor is frozen once established either way: refining it later would move marks reported after the change relative to those reported before, putting a step in the middle of one device's timeline. ggml-rpc: keep the RPC GPU's device identity. The GPU event carries only the span name and tag, and every device used the fixed name GRAPH_COMPUTE, so one rpc-server exposing several GPUs collapsed into a single row and a single peer-utilization figure. The device index is now part of the name, at both the compute and recompute sites. server: give each router child its own trace file. ggml_trace_open() falls back to GGML_RPC_TRACE and opens with "wb", and children inherit that variable through base_env, so every process opened the identical path. Measured on this branch: with one router and one model, a single 4216-byte file survived carrying only the child's header, the router's having been truncated away. Children now get a path derived from the base name and keyed by port. Erase before push, because execve keeps duplicate entries and getenv returns the first. server: state the phase in the iteration span instead of leaving it to be inferred. Readers cannot tell prefill from decode by tokens per slot, because under speculative decoding an ordinary decode submits several drafted tokens per slot. The span now carries an explicit prompt flag. Verified on a real trace: a 24-token request produced 1 prompt iteration of 14 ms followed by 22 decode iterations of about 1 ms. merge.py: identify logits by tensor subject rather than by response size. A GET_TENSOR over 64 KiB was counted as a logits return, but a layer split pulls staged hidden states back the same way and an 8 MiB activation was charged to the logits column. The subject is already emitted. On the staged-copy trace the logits column drops from 3.5 ms to 0.5 ms while the transfer column stays at 3.5 ms, so the copy is still accounted for, just not as an output. merge.py: associate peer traces with their own clock exchange rather than by hostname text. With peers addressed by IP or DNS alias nothing matched, the single-peer fallback stayed disabled, and both peers kept offset_us = 0 with their events on foreign clocks: peer GPU busy read 0.0% where the truth was 40.0%. On Windows the header host is blank because gethostname is skipped there, and a blank host prefix-matched the first endpoint, silently binding every peer to it and reading 20.0%. Records are now matched by which peer's TRACE_SYNC span brackets the exchange, which needs no hostname at all. Both cases now recover +499.999 and -300.001 ms and report 40.0%. Hostname matching remains as a second choice, now requiring a non-empty host and a dot-boundary match. device_idle.py: stop classifying speculative decode as prefill. max(toks) > n_slots is true for every speculative decode iteration, so t_pre_end slid to the end of the run and the decode-phase and group-offset sections came out empty. On a 2-prefill, 10-speculative-decode trace the phase split went from 12 prefill and 0 decode, with no decode section at all, to 2 and 10 with the decode section populated. It now prefers the explicit prompt flag above, and keeps a heuristic only for traces written before that flag existed.

danielhanchen · 2026-09-09T12:36:25Z

@codex review

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: eee6abed91

ℹ️ 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".

chatgpt-codex-connector · 2026-09-09T12:47:49Z

+    if (++s.n_unflushed >= TRACE_FLUSH_EVERY) {
+        s.n_unflushed = 0;
+        fflush(s.f);


Flush records before terminating the RPC server

When ggml-rpc-server is stopped with SIGTERM, as the supplied trace scripts do, the process has no signal handler and therefore never runs ggml_trace_close() or the static destructor. Consequently, every tail shorter than 128 events remains in the stdio buffer; a short capture can contain only its explicitly flushed header, while longer captures silently lose their final steps and produce incomplete timing summaries. Flush each record or add a graceful shutdown path that closes the trace.

Useful? React with 👍 / 👎.

Confirmed and fixed in f98e338, though not by the first of the two options you offered, and the reasoning is worth setting out.

The diagnosis is right: neither ggml-rpc-server nor the capture scripts install a handler, so SIGTERM runs neither ggml_trace_close() nor the static destructors, and a record count alone is the wrong bound. A tail shorter than 128 is dropped however long the run was, and a capture that never reaches 128 keeps only its explicitly flushed header.

I did not add a graceful shutdown path, because a handler that closes the trace is not safe here. fflush and fclose are not async-signal-safe, and emit() takes a mutex: a signal arriving while a worker holds it would deadlock the handler against the thread it interrupted, turning a lost tail into a hang. That is a worse failure than the one being fixed, and a handler also does nothing for SIGKILL.

Nor did I flush every record. cpu_check.sh exists to assert that tracing does not perturb what it measures, and a write syscall inside every timed interval is exactly the kind of thing that would quietly stop being true.

So the flush is bounded in time as well as in records: also flush once 100 ms have elapsed. The loss window becomes a fixed 100 ms rather than "however many events were left over", for any signal including ones nothing can catch, and it adds about ten fflush calls a second, which cannot move the numbers this tracer reports. Both bounds are kept, so a high-rate run still flushes on the record count and a slow or short one no longer waits indefinitely.

If you would prefer the guaranteed-lossless version for SIGTERM specifically, the safe shape is a flag set by the handler plus an EINTR-aware accept loop in ggml_backend_rpc_start_server, so the close runs on the normal path. That is a larger change reaching into the server loop, so I did not make it uninvited. Say the word and I will.

chatgpt-codex-connector · 2026-09-09T12:47:49Z

+            }
+        }
+        server_trace_scope span_iter("iteration", grp.id, n_slots_processing);
+        span_iter.prompt = n_slots_prompt > 0 ? 1 : 0;


Preserve mixed prompt and decode iterations

When continuous batching combines a prompting slot with already-generating slots, this group-wide Boolean marks the entire iteration as prompt; it can even do so when a started prompt slot is not admitted because the batch is full. device_idle.py::phase_windows() then classifies all generation work in that iteration as prefill and advances t_pre_end whenever a staggered request arrives, corrupting the decode-phase idle and group-offset statistics. Record the phase of the submitted batch, including a mixed state or separate prompt/decode counts, rather than whether any group slot is prompting.

Useful? React with 👍 / 👎.

Agreed on both counts, and fixed in f98e338. You are right that the group-wide boolean was the wrong shape, and right about the second case too: a slot can be in SLOT_STATE_STARTED and contribute nothing to this batch because the batch filled before it was admitted, which reading slot states before pre_decode() cannot see.

The span now carries the phase of the batch that was actually submitted, taken after batch.render() from the is_prompt flag each server_batch::token already carries, as prompt and decode token counts. Both cases you raise fall out of that: a slot that was not admitted contributes no tokens and does not colour the iteration, and a genuinely mixed batch reports both counts instead of collapsing to one.

device_idle.py prefers those counts and treats a batch containing any generation as decode, so a staggered arrival can no longer drag t_pre_end forward. Only a batch that is entirely prompt tokens is prefill, and an iteration that submitted nothing is neither. The single-count and no-field paths remain for traces written before this.

What I could not demonstrate, said plainly. I tried three times to capture a genuinely mixed iteration on a CPU build, with staggered concurrent requests, then a tighter four-request burst, then a long prompt chunked at -b 256 against an already-generating slot. Every capture came out as clean prefill iterations plus clean decode iterations: 2 prefill / 318 decode, 4 / 1594, and 1 / 599, with zero mixed and zero empty batches in all three. So on this build and these configurations the prompting slot appears to get its own iteration, and I have verified the counts are emitted and correct for the pure cases (for example prefill iterations of 77 and 41 prompt tokens with 0 decode) but not that the mixed path produces the numbers I expect, because I could not make one occur.

I would rather flag that than imply the mixed case was tested. The change is still strictly more faithful than what it replaces, since it reports what was submitted rather than inferring from state, and that part is what fixes the not-admitted case regardless. If you know a configuration that reliably produces a mixed batch here I will capture it and post the numbers.

chatgpt-codex-connector · 2026-09-09T12:47:49Z

+        for a, b, e in ix.overlapping(g0, g1):
+            cov = min(b, g1) - max(a, g0)
+            if cov > best_cov:
+                best_cov = cov
+                best = e


Attribute every traced phase in an idle gap

When a no-GPU gap spans multiple consecutive host phases, this loop retains only the single event with the largest overlap. The caller charges that event's coverage to its phase and labels the entire remainder unattributed, even when the remainder is fully covered by other traced phases such as batch_build, submit, and synchronize; the resulting per-step idle breakdown therefore systematically overstates untraced host time. Accumulate coverage for every phase across the gap while resolving any overlapping scopes.

Useful? React with 👍 / 👎.

Confirmed and fixed in f98e338. Keeping only the largest-overlap event meant the caller charged that one event to its phase and called everything else in the gap untraced, so a gap walked by batch_build, submit and synchronize in sequence reported most of itself as untraced host time.

Reproduced on a synthetic step whose real untraced time is known by construction: one 10 ms step with the GPU busy only between 6000 and 8000 us, a first gap covered by batch_build, submit, graph_compute and synchronize with 500 us genuinely untraced, and a second covered by post_decode, sampling and result_send with 600 us untraced. Ground truth 1.1 ms.

reported untraced

before 5.4 ms (4.9x overstated)

after 1.1 ms, exact

The full breakdown afterwards is batch_build 2.0, synchronize 1.3, graph_compute 1.2, unattributed 1.1, submit 0.8, post_decode 0.6, split 0.2, summing to exactly the 8.0 ms of gap. On the existing synthetic cases the per-step figure goes from 4.3 ms to 0.8 ms.

On overlap, since that is the way to get this wrong. The scopes nest and straddle: a sched/split inside a submit, a graph_compute crossing the boundary between submit and synchronize. Each point of the gap is charged to exactly one scope, the innermost covering it, which is the shortest, ties broken by later start then by name. It is a sweep over the sorted span edges, so the attributed total cannot exceed the gap length by construction rather than by assertion. Checked with a randomized property test over 20000 gap and scope configurations: no violation of that invariant.

Controls: a gap covered by exactly one phase with no nesting is byte-identical before and after, and across all 11 synthetic cases the only lines that change are the idle-breakdown ones. Every table column, clock offset and whole-window percentage is unchanged, since those come from covered() rather than from this path. The printed breakdown was widened from four entries to six, because gaps now decompose into more names and unattributed, which is the number a reader most needs, could otherwise fall off the end.

Honest note on the real-hardware evidence: on an actual CPU trace the change shows up as a decomposition rather than a large drop, submit 1.3 becoming submit 0.7 + copy_stage 0.5 + split 0.1 with unattributed unchanged at 0.2, because that trace's gaps happen to be dominated by one phase. The large overstatement is demonstrated on synthetic traces, where the true answer is known.

chatgpt-codex-connector · 2026-09-09T12:47:49Z

+    curl -s http://127.0.0.1:$PORT/completion -H 'Content-Type: application/json' \
+      -d "{\"prompt\":\"$p\",\"n_predict\":32,\"temperature\":0,\"top_k\":1,\"seed\":1}" \
+      | python3 -c 'import json,sys; print(json.load(sys.stdin)["content"])' >> "$OUT/$tag.out.txt"


Fail the smoke test when completions fail

When a completion request returns an HTTP error, invalid JSON, or a response without content, this pipeline fails but cell() continues because the script enables neither errexit nor pipefail and never checks the command status. If this happens in both the traced and untraced cells, both output files can be empty and cmp reports PASS, so the advertised non-perturbation smoke test succeeds without comparing any generated text. Propagate request/parser failures and assert that every expected response was recorded.

Useful? React with 👍 / 👎.

Confirmed and fixed in f98e338. This one deserves the sharpest wording of the five: the non-perturbation smoke test passed while comparing zero bytes of generated text.

Reproduced three ways, each with both cells failing so cmp compared two empty files:

failure injected before after

nothing listening on the port, curl exits 7 PASS, exit 0, 0 B compared exit 1, request failed, curl exit 7 (is anything listening on 8791?) and only 0 of 3 completions were recorded

HTTP 500 with a JSON error body PASS, exit 0, 0 B exit 1, returned HTTP 500, body: {"error":...}

HTTP 200, valid JSON, no content key PASS, exit 0, 0 B exit 1, response has no "content" (keys: model, stop, tokens_predicted)

All three printed output identical with the trace off and on: PASS. The causes are exactly as you set out: no pipefail, curl -s exits 0 on any HTTP status, and cell() ended on an array assignment that always succeeds, so no status ever propagated.

set -o pipefail is now on. errexit is deliberately not, because teardown is full of kill and grep -q that are expected to fail, and turning it on would have traded a false pass for a flaky failure. Every fallible step is checked explicitly instead, and the header says so, so the next person does not add set -e and quietly break the teardown. Each request checks the curl exit, requires HTTP 200, and validates the body for parseable JSON, an object, and a present and non-empty content. cell() requires all expected completions and a non-empty output file and returns its status, both call sites check it, and before cmp runs both files must be non-empty and carry exactly the expected number of response markers.

The control is the part that matters, since making a test always fail would look identical in the table above. A real end-to-end run against stories260K.gguf with two real rpc-servers still passes, and now says what it actually did: compared 3 completions, 271 bytes of generated text, then output identical with the trace off and on: PASS, exit 0, with merge.py accepting the resulting traces (96 decode steps, 5571 chrome events). It is not just always-fail, and it now reports its evidence rather than asserting a conclusion.

chatgpt-codex-connector · 2026-09-09T12:47:49Z

+        return 0;
+    }
+    s.role   = role != nullptr ? role : "unknown";
+    s.t_open = ggml_time_us();


Initialize the timer before opening a trace on Windows

When GGML_RPC_TRACE is set on Windows, llama_server() calls ggml_trace_open() before llama_backend_init(), which is the later call that initializes ggml's timer. This line therefore reaches ggml_time_us() while the Windows timer_freq is still zero, causing division by zero during startup instead of launching the traced server. Initialize the timer inside ggml_trace_open() or move every startup trace-open call after ggml_time_init().

Useful? React with 👍 / 👎.

Confirmed by inspection and fixed in f98e338. ggml_trace_open() takes a timestamp, both llama_server() and rpc-server call it before llama_backend_init(), and on Windows ggml_time_us() divides by timer_freq, which is what that later call sets. So setting GGML_RPC_TRACE on Windows divides by zero during startup.

ggml_time_init() is now called at the top of ggml_trace_open(). It is guarded by InitOnceExecuteOnce, so it is idempotent and costs nothing when the backend has already initialized it, and it fixes both call sites at once rather than relying on every future caller to be ordered correctly.

Stated plainly: I have no Windows machine here, so this is reasoned from the source and compiled, not executed. The non-Windows build is unaffected, since ggml_time_init() is a no-op there and the timestamp is unchanged. Someone with a Windows box should confirm that a traced server now starts, because the failure this fixes is a startup crash and I can only argue it away rather than show it.

…d report phase from the batch Five more review items. ggml-trace: initialize the timer before using it. Both llama_server() and rpc-server call ggml_trace_open() before llama_backend_init(), and ggml_trace_open() takes a timestamp. On Windows ggml_time_us() divides by timer_freq, which that later call is what sets, so setting GGML_RPC_TRACE divided by zero during startup instead of launching a traced server. ggml_time_init() is idempotent, so calling it at the top of ggml_trace_open() costs nothing when the backend is already up. This is reasoned from the source and compiled, not executed: I have no Windows machine here. ggml-trace: bound the flush window in time, not only in records. Neither ggml-rpc-server nor the capture scripts install a signal handler, so a SIGTERM runs neither ggml_trace_close() nor the static destructors and whatever sits in the stdio buffer is lost. Flushing every 128 records drops any tail shorter than that no matter how long the run was, and a capture short enough never to reach 128 keeps only its explicitly flushed header. The buffer is now also flushed once 100 ms have passed, so the loss window is bounded for any signal including SIGKILL, which no shutdown path could catch. A handler that closed the trace was the obvious alternative and is deliberately not used: fflush and fclose are not async-signal-safe, and a signal arriving while emit() holds its mutex would deadlock the handler against the thread it interrupted. Ten flushes a second cannot perturb what this tracer measures, whereas a flush per record would put a write syscall inside the very intervals being timed, and non-perturbation is a property this tooling is supposed to have. server: report the phase of the batch that was submitted. The previous flag was derived from slot states read before pre_decode(), which is wrong twice over: a slot can be in a prompt state and contribute nothing because the batch filled before it was admitted, and continuous batching mixes one slot's prompt with other slots' generation in a single batch. The iteration span now carries prompt and decode token counts taken from the rendered batch, using the is_prompt flag each token already carries. device_idle.py prefers those counts, treats a batch that is entirely prompt as prefill, treats any batch containing generation as decode so a late arrival cannot drag the end of the prefill phase forward, and ignores an iteration that submitted nothing. merge.py: attribute every phase in an idle gap. name_gap() kept only the single event with the largest overlap, so the caller charged that one event's coverage to its phase and called the whole remainder untraced, even where the rest of the gap was covered by other traced phases. Coverage is now accumulated for every phase, with each point of the gap charged to exactly one scope, the innermost covering it, so the total can never exceed the gap. On a synthetic step whose real untraced time is 1.1 ms, the reported figure goes from 5.4 ms to 1.1 ms, and on the existing cases from 4.3 ms to 0.8 ms per step. A randomized property test over 20000 gap and scope configurations found no case where the attributed coverage exceeds the gap. Gaps covered by a single phase are byte-identical, and every other column of the report is unchanged. cpu_check.sh: fail when the completions fail. The script set neither pipefail nor errexit and never checked the status of the request pipeline, and cell() ended on an assignment that always succeeds. A curl that could not connect, an HTTP 500, or a response without content therefore left an empty output file, and with both cells empty cmp compared nothing and the script printed PASS and exited 0. Demonstrated three ways: nothing listening, an HTTP 500, and a 200 whose body has no content field. All three printed the non-perturbation PASS having generated and compared zero bytes. Every request is now checked for curl failure, non-200 status, unparseable JSON and missing or empty content, both cells must record every expected completion, both files must be non-empty and carry exactly the expected number of responses before cmp runs, and the script reports how much text it actually compared. A real run against the model still passes and now says so with its evidence: 3 completions, 271 bytes of generated text.

danielhanchen · 2026-09-09T14:18:05Z

@codex review

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: f98e338f24

ℹ️ 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".

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+                d.anchor_us    = ggml_time_us();
+                d.anchor_fixed = true;


Wait for an idle anchor stream before fixing its wall time

With --pipeline-groups > 1 sharing a CUDA device, one group can call ggml_backend_sched_synchronize() and globally poll the tracer while the stream that recorded this device's anchor is still executing another group's graph. Fresh evidence after the earlier anchor fix is that this fallback still permanently sets anchor_us to that poll's wall time when cudaStreamQuery() reports the stream busy, shifting every subsequent span on the device and corrupting cross-group overlap and idle results. Leave the marks pending until the anchor stream is idle rather than freezing the approximation.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+                    best, bestc = None, 0.0
+                    for name, ix in host.items():
+                        cv = ix.covered(d0, d1)
+                        if cv > bestc:
+                            bestc, best = cv, name


Accumulate every host phase in device-idle gaps

When a neither-GPU interval spans consecutive phases such as batch_build, submit, and synchronize, this loop retains only the phase with the largest coverage; lines 176–180 then report the rest as untraced even though other phases cover it. Thus device_idle.py still overstates host/untraced idle and misattributes the gap, despite merge.py::attribute_gap() now splitting equivalent gaps across all applicable phases.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z


            llama_token id;
            {
+                server_trace_scope span("sampling", grp.id, slot.id);


Instrument the speculative post-decode path

When speculative decoding has a nonempty draft, the slot returns at lines 4485–4486 before entering this sampling scope, then performs common_sampler_sample_and_accept_n() and all corresponding process_token() calls later without either sampling or result_send scopes. Consequently speculative runs—the same runs supported by the new phase classification—systematically underreport both columns and classify that time as untraced; add equivalent scopes around the speculative branch.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+    if (++s.n_unflushed >= TRACE_FLUSH_EVERY || now - s.t_flush >= TRACE_FLUSH_EVERY_US) {
+        s.n_unflushed = 0;
+        s.t_flush     = now;
+        fflush(s.f);


Flush an idle trace tail without waiting for another event

When ggml-rpc-server emits its final records less than 100 ms after the previous flush and then sits idle until the supplied scripts send SIGTERM, no later call to emit() occurs to reevaluate this elapsed-time condition, so the unflushed tail is still discarded by immediate process termination. Fresh evidence after the earlier flush fix is that the timer is checked only while writing a new record; it does not actually bound how long existing buffered records remain unwritten. Use a graceful shutdown or an independent periodic flush so the final decode steps are preserved.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+            for (const auto & tok : batch.tokens) {
+                n_prompt += tok.is_prompt ? 1 : 0;
+            }
+            span_iter.n_prompt = n_prompt;
+            span_iter.n_decode = (int) batch.tokens.size() - n_prompt;


Count decodes submitted while building multimodal batches

When an iteration processes only a multimodal chunk, process_mtmd_chunk() submits the image decode inside pre_decode(), but that work is not represented in batch.tokens; these counts are therefore recorded as prompt: 0, decode: 0. device_idle.py::phase_windows() gives the explicit counts precedence and skips such an iteration entirely, excluding its GPU work from the prefill window and distorting multimodal idle reports. Include the tokens decoded by the MTMD path in the recorded prompt count.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+for pass in base_1 new base_2 new_traced; do
+  case $pass in base_*) BIN=$B0;; new*) BIN=$B1;; esac
+  tenv=()
+  [ "$pass" = new_traced ] && tenv=(GGML_RPC_TRACE=$O/nonrpc_trace.jsonl)


Actually enable tracing for the batched-bench trace pass

The new_traced batched-bench pass only exports GGML_RPC_TRACE, but llama-batched-bench never calls ggml_trace_open() and this non-RPC workload cannot reach the fallback call in get_socket(); repository-wide call-site inspection shows only llama-server, rpc-server, and the RPC connection path open the tracer. Consequently ggml_trace_flag remains zero and this pass measures another untraced run, so the script cannot validate the tracer's non-RPC overhead as advertised.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-09-09T14:30:42Z

+        ggml_trace_eventf("rpc.client", "GRAPH_COMPUTE", t_enq, t_send1,
+                          "\"t_send0\":%lld,\"t_send1\":%lld,\"bytes_out\":%zu,\"bytes_in\":0,"
+                          "\"reply\":0,\"ok\":1,\"n_nodes\":%d,\"dev\":%u,\"serialize_us\":%lld",
+                          (long long) t_ser, (long long) t_send1,
+                          input.size() + RPC_CMD_HEADER_BYTES, cgraph->n_nodes, rpc_ctx->device,


Release the RPC send lock before emitting trace records

With concurrent pipeline groups sharing an RPC connection, this trace emission still occurs while the mtx_send lock acquired at line 930 is held. Formatting the record, contending on the trace-file mutex, and occasionally calling fflush() therefore block every other group from submitting its next RPC command; because t_send1 is captured before that work, the added delay appears only as queueing in the following command and can change the overlap the trace is intended to measure. Capture the fields under the connection lock, release it, and emit the record afterward; the no-response send_rpc_cmd() helper has the same lock lifetime.

Useful? React with 👍 / 👎.

danielhanchen added 8 commits September 5, 2026 14:46

trace: script for the traced layer split cells on the pair

d6693f8

Merge remote-tracking branch 'origin/feature/pipeline-groups' into fe…

abe16e4

…ature/rpc-trace # Conflicts: # tools/server/server-context.cpp

trace: decompose one device's idle into scheduling and host stalls

988e09e

danielhanchen added 2 commits September 8, 2026 00:56

trace: reduce comment volume in the event tracer

23d52c3

chatgpt-codex-connector Bot reviewed Sep 9, 2026

View reviewed changes

Merge the review fixes from feature/pipeline-groups into feature/rpc-…

9debb3e

…trace

unslothai deleted a comment from chatgpt-codex-connector Bot Sep 9, 2026

danielhanchen added 6 commits September 8, 2026 19:22

Merge the parent-task cleanup fix from feature/pipeline-groups into f…

e5f7c6d

…eature/rpc-trace

Merge the router child_env fix from feature/pipeline-groups into feat…

6bf85af

…ure/rpc-trace

Merge the fit reservation and per-group thread pools from feature/pip…

aee4120

…eline-groups into feature/rpc-trace

chatgpt-codex-connector Bot reviewed Sep 9, 2026

View reviewed changes

danielhanchen added 2 commits September 8, 2026 22:40

chatgpt-codex-connector Bot reviewed Sep 9, 2026

View reviewed changes

		logits = [(e["t0"], e.get("t_recv1", e["t1"])) for e in cmds
		if e.get("n") == "GET_TENSOR" and e.get("bytes_in", 0) >= LOGITS_MIN_BYTES]

case	logits before	logits after
staged hidden state plus real output	3.5 ms	0.5 ms
control, real output only	0.5 ms	0.5 ms, row byte-identical
control, trace predating `subj`	3.5 ms	3.5 ms, with an explicit note that it fell back

case	before	after
hosts `gpu-a` / `gpu-b`, endpoints `10.0.0.11:50052` / `10.0.0.12:50052`	two "no clock offset" warnings, offsets +0.000 / +0.000 ms, peer GPU busy 0.0%	offsets +499.999 / -300.001 ms, peer GPU busy 40.0%
Windows, both header hosts blank	no warning at all, both peers bound to the first endpoint at +499.999 ms, peer GPU busy 20.0%	+499.999 / -300.001 ms, peer GPU busy 40.0%
control, single peer, IP endpoint, no TRACE_SYNC in the trace	+499.999 ms, 20.0%	unchanged, fallback still fires

failure injected	before	after
nothing listening on the port, curl exits 7	PASS, exit 0, 0 B compared	exit 1, `request failed, curl exit 7 (is anything listening on 8791?)` and `only 0 of 3 completions were recorded`
HTTP 500 with a JSON error body	PASS, exit 0, 0 B	exit 1, `returned HTTP 500, body: {"error":...}`
HTTP 200, valid JSON, no `content` key	PASS, exit 0, 0 B	exit 1, `response has no "content" (keys: model, stop, tokens_predicted)`

Uh oh!

Conversation

danielhanchen commented Sep 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What this is

Design

What is recorded

Tools

Overhead

Non-RPC and non-CUDA workloads

The three timelines

Validation

Scope, relative to the base branch

Uh oh!

danielhanchen commented Sep 7, 2026

Uh oh!

danielhanchen commented Sep 9, 2026

Uh oh!

danielhanchen commented Sep 9, 2026

Uh oh!

chatgpt-codex-connector Bot commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codex Review Summary

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

danielhanchen commented Sep 9, 2026

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

danielhanchen commented Sep 9, 2026

danielhanchen commented Sep 5, 2026 •

edited

Loading

chatgpt-codex-connector Bot commented Sep 9, 2026 •

edited

Loading