idle gaps are not compute
One NCCL configuration looked ordinary until the workload ran.
config isolated benchmark workload
tree-ll #6 (150.5 us) #8 (11,969 us/token, IQR 578)The isolated benchmark put tree-ll sixth among eight configurations. The decode-like workload put it last. Its median time per token was 45 percent above the winner and 23 percent above tree-simple, the other forced tree configuration nearest to it. Its spread was also much larger than the rows around it.
Those numbers came from an earlier four-A100 campaign. They are the observation that started this investigation. They are not a fresh measurement from the machine on which this revision was written.
That distinction became important when I tried to verify the explanation attached to the table. The public repository contains the experiment code, manifests, analyzers, and several later campaign identities. It does not contain the raw attempt archive from the campaign that produced this particular table. The repository’s own artifact note says that the old numeric claims cannot be independently regenerated from the released files.
The table is therefore historical reported evidence. Its arithmetic can be checked. Its experiment design can be inspected. Its exact medians cannot be recomputed from raw samples in the current checkout.
Then a more serious problem appeared. The explanation said that the workload overlapped every collective with the next matrix multiplication. The preserved workload source did not.
def _run_layer(torch, dist, activation, weight, comm_buffer):
torch.matmul(activation, weight)
dist.all_reduce(comm_buffer, op=dist.ReduceOp.SUM)There is no async_op=True. In PyTorch 2.4.1, the version named by the experiment, the default is synchronous. The Python wrapper obtains a work object and calls work.wait(). The NCCL backend makes PyTorch’s current CUDA stream wait for the NCCL completion event. The next matrix multiplication can be submitted by the host, but it cannot pass that event on the current stream.
The replay that improved the ranking did use async_op=True. It placed matrix multiplications between collective issue and wait. That replay certainly created compute and communication overlap. The preserved ground-truth program did not express the same schedule.
The result might still be useful. The old causal story is not.
The strange row had led below microbenchmarks and traces into a harder question: what must be true before a replay ranking counts as evidence about the workload rather than evidence about the replay?
the table survived longer than its evidence
The original campaign ran on one node of the LSU Rostam cluster with four A100 PCIe 40 GB GPUs and no NVLink. It compared eight configurations:
NCCL 2.19.3 default
NCCL 2.20.5 default
NCCL 2.20.5, Ring, LL
NCCL 2.20.5, Ring, LL128
NCCL 2.20.5, Ring, Simple
NCCL 2.20.5, Tree, LL
NCCL 2.20.5, Tree, LL128
NCCL 2.20.5, Tree, SimpleRing and Tree choose the broad collective algorithm. LL, LL128, and Simple choose protocol families inside NCCL. A protocol changes how data is staged, how much synchronization it requires, and how it trades latency against bandwidth. The two default rows let NCCL select an algorithm and protocol rather than forcing one combination.
NCCL warns that its debugging and tuning environment variables should not be left forced in production without a reason. That warning is easy to read as operational caution. The table shows the deeper reason. A setting can win one context and lose another because an algorithm is not intrinsically fast. It is fast for a message size, topology, arrival pattern, concurrent workload, and software version.
The reported ranking was:
config isolated workload
ring-ll 1 (107.5 us) 5 (8,466 us/token)
ring-ll128 2 (110.6 us) 3 (8,258)
2.20.5-default 3 (132.1 us) 2 (8,246)
ring-simple 4 (132.1 us) 4 (8,453)
2.19.3-default 5 (133.1 us) 1 (8,238)
tree-ll 6 (150.5 us) 8 (11,969, IQR 578)
tree-ll128 7 (157.7 us) 6 (9,006)
tree-simple 8 (203.8 us) 7 (9,768)The benchmark winner, ring-ll, was 2.768 percent slower than the workload winner:
(8,466 us/token / 8,238 us/token - 1) * 100 = 2.768 percenttree-ll was 45.290 percent slower than the winner:
(11,969 / 8,238 - 1) * 100 = 45.290 percentIt was 22.533 percent slower than tree-simple:
(11,969 / 9,768 - 1) * 100 = 22.533 percentThese calculations need only the printed medians. They do not validate the medians. Arithmetic and measurement provenance are different layers of evidence.
An archive capable of validating them would need more than a CSV copied from an analysis directory. It would need the frozen campaign definition, every terminal attempt, the selection that chose exactly one attempt for each expected cell, the completeness verdict, the raw measurement artifacts, and hashes joining those objects. It would also need enough environment identity to show which code, wheel, NCCL library, driver, node, and command produced each attempt.
The current CommCanary harness implements much of that chain. It freezes a manifest before execution. Retries append terminal attempt records instead of replacing the old ones. A selection names one attempt for every cell. A completeness pass refuses missing, duplicate, failed, stale, or unexpected cells. The analyzer emits rankings only from selected evidence after completeness passes. A raw archive descriptor binds the archive’s byte size and SHA-256 digest to the exact run, campaign, repository, manifest, selection, and verdict.
The old table predates that complete local archive. A later harness cannot retroactively create bytes that were not retained.
This does not require deleting the table. It requires making its status visible. Reported evidence can motivate a new investigation. It cannot silently become a reproducible result because source code exists nearby.
twenty eight decisions hide inside eight ranks
Eight configurations create 28 unordered pairs:
8 * 7 / 2 = 28For each pair, a proxy makes a decision. Configuration A is faster, configuration B is faster, or the two are tied under a declared tolerance. A complete ordering compresses those pairwise decisions into ranks, but the pairwise view matters because one misplaced row can reverse several decisions at once.
The printed isolated and workload orders agree on 18 of 28 pairs:
18 / 28 * 100 = 64.285714 percentRounded to one decimal place, that is 64.3 percent. The ten disagreements are:
2.19.3-default | 2.20.5-default
2.19.3-default | ring-ll
2.19.3-default | ring-ll128
2.19.3-default | ring-simple
2.20.5-default | ring-ll
2.20.5-default | ring-ll128
ring-ll | ring-ll128
ring-ll | ring-simple
tree-ll | tree-ll128
tree-ll | tree-simpleThis direct calculation treats every printed order as strict. The campaign analyzer used a more cautious relation. It compared the median difference with the larger of the two interquartile ranges. If the difference was smaller, it called the pair tied.
For configurations a and b, let m_a and m_b be their medians and let q_a and q_b be their IQR values. The campaign relation was:
if abs(m_a - m_b) < max(q_a, q_b):
relation = tied
elif m_a < m_b:
relation = a_faster
else:
relation = b_fasterThat policy is part of the result. Change it and the agreement percentage can change. The public table prints the IQR only for tree-ll, so the exact tie-aware relation matrix cannot be regenerated from the article alone. The fact that the strict rank calculation also yields 64.3 percent is a check on one displayed number, not a reconstruction of the analysis.
Pairwise agreement also discards magnitude. Reversing two rows separated by 0.1 percent counts the same as reversing two rows separated by 45 percent. A deployment decision needs both:
direction: did the proxy choose the same side?
magnitude: how expensive was the disagreement?
uncertainty: could sampling noise explain it?Kendall-style concordance adds information about ordering and ties. It still does not know that one wrong pair violates a latency SLO while another changes nothing. Statistical ranking is a necessary layer. It is not the final decision model.
The old experiment used median token latency as the workload verdict and collective latency as the microbenchmark verdict. Those quantities do not even have the same unit of work. One is microseconds per collective invocation. The other is microseconds per token after 32 layers of compute and communication. Ranking makes them comparable only in the narrow sense that smaller is preferred under both. It does not make their magnitudes interchangeable.
This is why the initial problem is a ranking problem rather than a calibration problem. Nobody should expect 150.5 microseconds per collective to equal 11,969 microseconds per token. The surprise is that a configuration which appears safe relative to seven alternatives becomes the worst relative choice once embedded in the program.
the run order can manufacture a winner
Eight configurations cannot be measured at the same instant on the same four GPUs. They are scheduled one after another. That schedule can become a hidden input.
Suppose the machine begins cold and gradually warms. Running every repetition of configuration A, then every repetition of B, produces:
A A A A A B B B B B C C C C C ...Configuration identity is now correlated with time. If clocks, temperature, another tenant, filesystem state, or background daemons drift, the measurement assigns that drift to configurations.
Reversing the order does not solve the problem. It changes which configuration inherits the drift.
A round-robin schedule is better:
A B C D E F G H
B C D E F G H A
C D E F G H A B
...Each configuration appears in different positions. Randomizing each round is better when the random seed and realized order are recorded. Blocking by round preserves a local comparison: every configuration is measured once before the next round begins.
The old campaign notes describe round-robin interleaving on one physical node. That design reduces one confounder. It does not eliminate time.
The analysis should retain round identity:
round 0:
A=...
B=...
...
round 1:
A=...
B=...
...Then comparisons can be paired within rounds. If round 3 is uniformly slow, a paired difference between A and B can remain stable even when their raw times rise.
For round r, define:
d_r = latency_candidate_r - latency_baseline_rThe distribution of d_r answers a direct rollout question. Positive values mean the candidate was slower than the baseline in that local block. It is often more informative than comparing two independent collections of absolute measurements.
Randomization does not excuse missing environment controls. Each attempt should record:
physical node identity
GPU UUIDs and PCIe topology
driver and CUDA versions
loaded NCCL library path and version
PyTorch wheel identity
environment variables affecting NCCL
GPU clocks and power limits when readable
CPU affinity
NUMA placement
process count and rank mapping
message sizes
warmup count
measurement count
profiler state
other visible GPU processesThe historical setup verified the loaded NCCL version rather than trusting an environment variable. That is the right instinct. A command that asks the loader to use one library is not proof that the process mapped it. /proc/self/maps and ncclGetVersion observe two different parts of the identity: the actual shared object mapping and the version reported by the library.
Warmup also needs a definition. “One warmup” can mean:
one collective
one token
one complete 256-token run
one communicator initialization
one pass after clocks stabilizeThese have different costs and effects. NCCL communicator initialization, CUDA context creation, kernel module loading, allocator growth, and autotuning can all make early iterations unlike steady state.
Discarding the first observation by habit can hide a production cost. Model startup and the first request matter when instances scale from zero. A campaign should separate:
cold-start metric
warm steady-state metric
transition lengthThe unit of replication is another trap. One 256-token job yields 8,192 collective events. Those are not 8,192 independent experiments. They share the same node, process lifetime, communicator, temperature trajectory, rank mapping, and configuration.
Treating every collective as an independent sample would produce an enormous sample count and an unjustifiably narrow confidence interval. This is pseudoreplication.
The physical job is closer to the independent unit when each job reinitializes the relevant state. If five jobs run per configuration, the strongest unconstrained claim usually has five independent repetitions, not 40,960.
Events inside a job remain useful. They reveal message-size effects, within-run drift, and tails. Their dependence must be preserved in uncertainty calculations. A hierarchical bootstrap can resample jobs first and events within selected jobs second. Resampling all events from one flat pool destroys the correlation structure.
The same problem applies to tokens. A p99 computed from thousands of tokens in one process may be precise about that process lifetime and poor at predicting a fresh job on another node.
The central table uses medians across runs. Without the raw attempt archive, none of these design checks can be repeated. The source and manifest show intent. Only attempt records show whether execution followed it.
five repetitions do not make rank eight certain
The interquartile range is:
IQR = 75th percentile - 25th percentileIt describes the middle half of observed values. It is resistant to one extreme outlier. It is not a confidence interval for the median.
With five independent run medians, quantile estimates are coarse. Depending on the quantile convention, the 25th and 75th percentiles may land directly on individual ordered samples or interpolate between them:
x0 <= x1 <= x2 <= x3 <= x4There is not enough information to characterize a rare tail. Saying tree-ll had a 25 times larger IQR than neighbors describes the observed sample. It does not establish that the population IQR ratio is exactly 25.
Ranking uncertainty can be estimated without pretending medians are normally distributed. A bootstrap repeatedly resamples independent runs with replacement, recomputes medians, and rebuilds ranks.
The smallest implementation is:
import random
from statistics import median
def bootstrap_rank_probability(samples, repetitions=10000, seed=7):
rng = random.Random(seed)
names = sorted(samples)
wins = {name: 0 for name in names}
for _ in range(repetitions):
medians = {}
for name in names:
observed = samples[name]
resampled = [
observed[rng.randrange(len(observed))]
for _ in observed
]
medians[name] = median(resampled)
winner = min(names, key=lambda name: (medians[name], name))
wins[winner] += 1
return {
name: wins[name] / repetitions
for name in names
}This code is real and runnable. Its validity depends on its inputs being independent at the resampled level. Passing 8,192 correlated events from one job does not fix a five-job experiment.
A paired bootstrap for round-robin data resamples whole rounds. Every selected round contributes all eight configurations, preserving common drift.
The output should include more than a most-likely order:
probability each config is best
probability each config is worst
probability candidate beats baseline
distribution of candidate minus baseline
probability effect exceeds an operational thresholdThe operational threshold matters. A statistically detectable 0.1 percent difference may not pay for a rollout. A noisy 10 percent regression may demand action before a conventional significance threshold is reached.
Let delta be the largest candidate slowdown the deployment accepts:
delta = 2 percentThe useful probability is:
P(candidate slowdown > 2 percent | observed experiment)It is not merely:
P(candidate is slower by any amount)Twenty-eight pairwise comparisons create another selection problem. If each pair is tested independently at a 5 percent false-positive rate, the chance of at least one false positive across 28 truly equal pairs is:
1 - (1 - 0.05)^28 = 0.762That 76.2 percent value assumes independent tests, which pairwise comparisons sharing configurations are not. The calculation is still a warning. Repeating many comparisons and reporting only interesting inversions invites false discoveries.
A predeclared family-level procedure can control this. So can reporting confidence intervals and the full pair matrix rather than converting every small difference into a categorical win.
Tie policy is one such procedure. The campaign used the larger IQR as a tolerance. That is conservative in some cases and unconventional in others. An IQR is a spread statistic, not the sampling uncertainty of a median difference. The policy may be reasonable for a decision heuristic, but it should not be described as a formal significance test.
The correct wording is:
Under the declared IQR-based tie rule, these pairwise decisions agreed.Not:
The configurations are statistically identical.Cross-validation across hardware is the strongest check on ranking stability. A proxy discovered and judged on one node can overfit that node’s topology and noise. Freeze the proxy and tie policy, then test on:
another node of the same SKU
another A100 PCIe node
an NVLink-connected A100 system
a newer GPU architecture
a multi-node fabricThe expected result need not be the same absolute rank everywhere. The test is whether the proxy predicts each local workload decision better than its baselines without retuning itself after every failure.
one source line changes the schedule
The workload was intended to resemble tensor-parallel decode. Each token traversed 32 layers. Each layer executed a bf16 matrix multiplication and an all_reduce. Message sizes cycled through 64 KiB, 128 KiB, and 256 KiB. With 256 tokens, that produces:
32 collectives/token * 256 tokens = 8,192 collectivesThe loop is small enough to inspect without any machine-learning vocabulary:
def token_step(token_index):
for layer in range(args.layers):
index = token_index * args.layers + layer
comm_buffer = comm_buffers[index % len(comm_buffers)]
torch.matmul(activation, weight)
dist.all_reduce(comm_buffer, op=dist.ReduceOp.SUM)The matrix multiplication uses one input array and one weight array to produce GPU work. The reduction sums a communication buffer across the four processes and leaves the sum on every process. The exact numerical relationship between the matrix multiplication and buffer does not matter for this scheduling question. Their order does.
The earlier prose described this loop as if it did:
work = dist.all_reduce(comm_buffer, async_op=True)
torch.matmul(next_activation, next_weight)
work.wait()That is a different program.
The first program asks for the default async_op=False. The second asks for a work handle and decides later where to wait. Both may let the Python thread continue before every GPU instruction has physically retired. Host asynchrony is not the same as device overlap.
This distinction is easy to miss because the word “blocking” can name three different boundaries:
host blocking:
the CPU thread cannot return from the call
stream dependency:
later GPU work in one stream cannot pass an event
device synchronization:
the host waits until all earlier work on a stream or device finishesPyTorch’s synchronous NCCL collective does not need to spin the host until the GPU finishes every byte. It can insert a dependency into the current CUDA stream and return. The CPU is then free to enqueue the next matrix multiplication. The GPU still sees an ordering edge.
That edge is enough to defeat the old overlap claim.
waiting does not always stop the host
The exact PyTorch 2.4.1 Python implementation is short. After validating the tensor and constructing options, all_reduce reaches this branch:
work = group.allreduce([tensor], opts)
if async_op:
return work
else:
work.wait()The workload takes the second branch.
The NCCL work object’s wait calls synchronizeInternal. In that function, synchronizeStream obtains PyTorch’s current CUDA stream and makes it wait on the NCCL end event:
void ProcessGroupNCCL::WorkNCCL::synchronizeStream() {
auto currentStream = at::cuda::getCurrentCUDAStream(device_.index());
ncclEndEvent_->block(currentStream);
}
void ProcessGroupNCCL::WorkNCCL::synchronizeInternal(
std::chrono::milliseconds timeout) {
synchronizeStream();
// additional host waiting occurs only in selected modes
}NCCL normally runs communication on its own stream. The current compute stream and the NCCL stream are separate queues. Separate queues make concurrency possible, but events can add dependencies between them.
The sequence becomes:
current compute stream NCCL stream
GEMM 0
all_reduce 0
wait for NCCL end event
GEMM 1
all_reduce 1
wait for NCCL end event
GEMM 2The wait is recorded on the current compute stream. GEMM 1 sits after that wait. It cannot begin until all_reduce 0 has reached its end event.
The CUDA Programming Guide describes a stream as an ordered work queue. Operations in one stream execute in enqueue order. Work in different streams may execute concurrently if dependencies and resources permit it. “May” matters. Multiple streams express the possibility of concurrency. They do not guarantee that kernels overlap, and an event can deliberately remove the possibility.
The source therefore supports this statement:
the host may enqueue later work before the collective physically finishesIt does not support this statement:
the next matrix multiplication runs concurrently with that collectiveOnly a GPU timeline showing kernel intervals could establish actual overlap. A Python call trace is insufficient. An NCCL event trace is insufficient. Even async_op=True is insufficient because resource pressure or dependencies can serialize work that was eligible to overlap.
The local source probe records the boundary without importing PyTorch:
preserved workload
all_reduce line: 172
async_op keyword: absent
explicit overlap replayer
all_reduce line: 196
async_op keyword: true
explicit wait line: 201The probe also records SHA-256 digests for both files and the inspected CommCanary revision, a1744e2b7facb8b1890ebe77c1840e1d716c266c. The checkout had uncommitted editorial and experiment-harness changes, so the file digests matter in addition to the Git commit.
That last sentence is not housekeeping. A Git revision identifies committed bytes. It does not identify dirty files. Reproducibility requires the actual source identity that produced the claim.
a timestamp can only tell you when
A profiler trace might keep an event like this:
{
"op": "all_reduce",
"bytes": 65536,
"ranks": [0, 1, 2, 3],
"group": "tp",
"start_us": 4312.5
}The next event might begin 400 microseconds later. Subtracting the timestamps gives a gap:
next_start_us - current_start_us = 400 usThat number does not say what occupied the interval.
It could contain:
100 us of collective execution followed by 300 us of GEMM
300 us of GEMM followed by 100 us of collective execution
100 us of collective and 300 us of partially overlapping GEMM
host scheduling delay
another CUDA stream
a page fault in the process recording the trace
rank arrival skew
an explicit sleep
some mixture of all of themTimestamps constrain chronology. They do not identify causality.
The CommCanary Kineto adapter is appropriately narrow. It selects record_param_comms events, reconstructs operation names, element counts, dtypes, process-group ranks, and single-rank start timestamps. Its emitted note explicitly refuses to claim cross-rank arrival skew, compute overlap, or measured exposed latency. A single-rank collective trace does not contain those facts.
The trace schema can represent more when another capture path provides it:
start_us
gap_us
compute_before_us
compute_overlap_us
compute_pressure
rank_arrival_us
arrival_skew_us
observed_exposed_us
concurrent_groupsRepresentable is not observed. A field with a default value of zero is dangerous when zero can mean either “measured none” or “not captured.” The compiler refuses one explicit uncertainty state, arrival_skew_unknown, because compressing an uncalibrated cross-rank offset would turn missing knowledge into a number. Compute fields need the same discipline. Current code carries compute_fields_uncertain through compilation, which prevents an uncertain capture from quietly receiving the strongest behavioral claim.
There are at least four clocks in a distributed GPU trace:
CPU monotonic clock on rank 0
CPU monotonic clock on rank 1
CUDA event clock on GPU 0
CUDA event clock on GPU 1The values are not automatically comparable. CUDA event durations are meaningful within supported device timing relationships. CPU timestamps on different hosts or processes need calibration. Profiler timestamps may be rebased to a trace start, which preserves local differences while discarding an absolute epoch. Cross-rank arrival skew needs either a common timebase or an explicit alignment method with an error bound.
An event at 4312.5 us looks precise to a tenth of a microsecond. Precision in the printed decimal does not establish synchronization between clocks.
sixty four kibibytes is not a transport description
The smallest trace row says bytes: 65536. Even that field needs a definition.
For an in-place all_reduce, 65,536 bytes can mean the logical tensor size on each rank. It does not mean exactly 65,536 bytes crossed each physical link. A ring reduction splits the tensor into chunks, moves them through reduce-scatter steps, then circulates the reduced chunks through all-gather steps. A tree sends different traffic along parent and child edges. Protocol framing, alignment, channels, and retries add physical traffic not represented by the tensor byte count.
For world size p, the familiar idealized ring volume per rank is:
reduce-scatter:
((p - 1) / p) * tensor_bytes sent
all-gather:
((p - 1) / p) * tensor_bytes sent
total:
2 * ((p - 1) / p) * tensor_bytes sent per rankFor four ranks and a 65,536-byte tensor:
2 * (3 / 4) * 65,536 bytes
= 98,304 bytes per rankThat is an algorithmic volume model, not a PCIe counter prediction. It ignores protocol overhead and topology.
An adapter also needs element size. Kineto may report input and output element counts plus a dtype. The importer multiplies the larger element count by bytes per element. For bf16:
32,768 elements * 2 bytes/element = 65,536 bytesIf dtype metadata is absent or misparsed, the same element count could be interpreted as 131,072 bytes for float32. That can move an event across an NCCL algorithm threshold and change the very configuration choice being studied.
Collectives other than all_reduce complicate the mapping. An all_gather input shard and output tensor differ by world size. A reduce_scatter does the reverse. all_to_all can use uneven split sizes. Point-to-point operations need sender and receiver. Recording one undifferentiated byte count and replaying it as every operation type silently changes communication.
Process-group membership matters too:
all_reduce 64 KiB over ranks [0, 1, 2, 3]
all_reduce 64 KiB over ranks [0, 2]The rows share operation and size. Their topology, contention, and volume per participant differ. A compressor that groups them because their first three fields match corrupts the program.
The current compiler’s operation identity includes ordered ranks and group, along with operation, bytes, phase, and point-to-point details. PARAM export creates explicit process-group initialization entries because a collective referencing an unregistered group would either fail or execute against the wrong membership.
There is another hidden choice in “same message sizes.” The workload cycles:
64 KiB, 128 KiB, 256 KiB, 64 KiB, ...A microbenchmark can test those sizes with equal frequency and still lose their phase relationship. If every 256 KiB event follows a particular matrix multiplication or aligns with a queue burst, shuffling sizes preserves a histogram and changes execution.
A trace must therefore distinguish:
value distribution:
how often each message size occurs
sequence:
which size follows which
phase:
where sizes occur in prefill, decode, or another region
concurrency:
which groups and kernels are active at the same timeThe isolated benchmark preserved a small part of the first item. The timestamp trace preserved more of the second. Neither fact alone determined the workload rank.
four programs can wear the same trace
Take one 300 microsecond matrix multiplication and one 100 microsecond collective. Four replay schedules can be described with nearly the same operation list.
The preserved blocking workload:
time us 0 300 400 700 800
compute [ GEMM 0 ] [ GEMM 1 ]
comm [AR0] [AR1]
overlap 0 usA communication-only timestamp replay:
time us 0 100 400 500
comm [AR0] [AR1]
GPU idle [ 300 us ]
overlap 0 usA sequential compute-fill replay:
time us 0 100 400 500
comm [AR0] [AR1]
compute [ GEMM ]
overlap 0 usAn explicit overlap intervention:
time us 0 100 300 400 600
comm [AR0] [AR1]
compute [ GEMM 0 ][ GEMM 1 ]
overlap 200 us across two collectivesThe local interval probe calculates the same schedules:
preserved blocking workload:
makespan=800 us, compute/comm-overlap=0 us
timestamp-paced communication replay:
makespan=500 us, compute/comm-overlap=0 us
sequential compute fill:
makespan=500 us, compute/comm-overlap=0 us
explicit overlap intervention:
makespan=600 us, compute/comm-overlap=200 usThe numbers are illustrative, not GPU measurements. Their job is to expose execution order.
Calling all four “trace replay” hides the variable that changes. The trace might list GEMM and all_reduce in a familiar sequence while the replayer decides whether operations block, which stream receives them, where waits appear, and whether timestamps cause sleeping or merely annotate events.
A replay grammar therefore needs operational semantics, not just a schema.
For every operation, it must define:
when the host issues it
which device and stream receives it
which earlier events it waits for
whether the host blocks
whether the current stream blocks
where completion is observed
what the reported duration spans
what work can be in flight at the same timeWithout those rules, two replayers can consume byte-identical JSON and execute different programs.
the faithful replay made silence faithful
The old campaign reported 60.7 percent pairwise agreement for a timestamp-paced PARAM replay. It preserved collective order, message sizes, and recorded start gaps. It ranked tree-ll second rather than last.
The earlier explanation said that this failed because the gap represented overlapping compute as idle time. Source inspection changes the first half of that sentence.
The gap still was not idle. In the blocking workload, the interval between the start of one collective and the start of the next includes completion of the first collective and execution of the next matrix multiplication. A communication-only replayer can render the noncommunication portion as sleep. That changes GPU occupancy, power state, cache state, memory traffic, and the host’s rate of submission.
So “idle gaps are not compute” remains correct.
What no longer follows is that the missing compute was concurrent with the collective in the original source. It was preceding compute in a serial dependency chain.
This corrected mechanism is less dramatic and more general. A timestamp is not a substitute for the work that elapsed between timestamps. Even fully serialized applications can depend on the identity of that work because earlier kernels change:
clock and power state
temperature
cache contents
memory-controller queues
allocator state
CUDA context state
CPU submission timing
rank arrival timingReplacing 300 microseconds of GEMM with 300 microseconds of sleep preserves wall time while changing all of those state variables.
This is the distinction between temporal fidelity and behavioral fidelity.
Temporal fidelity asks whether event times match.
Behavioral fidelity asks whether the system reaches the same states that influence the decision being tested.
A replay can have low timestamp error and poor behavioral fidelity. It can also have visibly inaccurate timestamps yet preserve a configuration ranking if the omitted timing variation does not affect that decision.
Neither possibility can be decided from the trace alone. It takes controls.
the compute fill changed the chip before the collective
The next historical mode filled each inter-collective gap with calibrated bf16 matrix multiplications. The earlier campaign reported a calibration of 44.03 microseconds for one 1024 by 1024 bf16 matrix multiplication on that node. The replay rounded:
gemm_count = round(gap_us / 44.03 us)Then it ran those matrix multiplications and the collective sequentially.
The reported agreement fell to 50.0 percent. More strangely, collective times became roughly 4 to 14 percent faster across configurations after compute was added.
The old article attributed the speedup to GPU clocks. The inserted GEMMs, it argued, raised the A100 from an idle state before the collective began. That is plausible. It is not proven by the retained evidence.
To prove the clock explanation, the run would need a time-aligned record of at least:
SM clock
memory clock
power draw
temperature
GEMM kernel interval
NCCL kernel interval
collective durationThe experiment would also need a negative control that warms the GPU without applying the same memory traffic, and perhaps another that applies memory traffic without comparable tensor-core work. Otherwise “warm clock” remains entangled with cache state, memory-controller state, thermal behavior, and launch timing.
The observation, if the old aggregate is accepted, is:
sequential compute fill preceded lower measured collective durationsThe inference is:
the compute fill may have changed clock or power stateThe article cannot promote the inference into a measurement just because it sounds physically reasonable.
This is a useful failure even without the final cause. Sequential fill did not recreate a concurrent contention state. It created a new predecessor workload. A tool that adds “realistic compute” needs to specify which relationship it intends to preserve:
same elapsed time before communication
same kernel type before communication
same average GPU occupancy
same simultaneous resource pressure
same temperature and clocks
same rank arrival distributionThese are not equivalent.
The calibrated quantum also introduces quantization error. A 70 microsecond gap becomes two GEMMs, or 88.06 microseconds under the historical calibration. A 20 microsecond gap becomes zero. Rounding preserves neither total time nor short-gap structure unless an error accumulator carries residuals forward.
If the replayer uses:
count_i = round(gap_i / quantum)then the error for gap i is:
error_i = count_i * quantum - gap_iEach absolute error can approach half a quantum. Across 8,192 events, signed errors may cancel by luck while local burst timing changes substantially. Reporting only total replay duration would hide that.
the overlap proxy was an intervention
The explicit overlap replayer uses a trace grammar with issue and wait entries. Its central path is:
work = dist.all_reduce(
buf,
group=groups[entry["pg_id"]],
async_op=True,
)
pending[entry["req"]] = (work, start)
# later trace entries may run matrix multiplies here
work, start = pending.pop(entry["req"])
work.wait()The PARAM exporter places the wait for collective k after the compute entries derived from the next gap. If an overlap-aware replayer consumes that grammar, the collective can remain in flight while those matrix multiplications execute.
This is a coherent stress test. It asks how each NCCL configuration behaves when collective work competes with matrix multiplication for GPU resources.
The old campaign reported 67.9 percent agreement and Kendall tau 0.54 for this mode. It put tree-ll seventh. It also reported an 8.2 microsecond IQR for tree-ll against 0.5 microseconds or less for the other replayed configurations.
That result was the first replay mode in the campaign reported to beat the isolated benchmark on both ranking measures. It remains interesting.
But the correct name for the mode is not faithful workload overlap. It is an explicit overlap intervention.
An intervention deliberately changes one factor to see how the system responds. Here the factor is simultaneous compute pressure during communication. If tree-ll becomes slow and unstable under that pressure, the result supports:
tree-ll is sensitive to the pressure created by this replayIt does not yet support:
the original workload applies the same pressure in the same scheduleThose statements would meet only after a trace of the original workload showed overlapping GEMM and NCCL kernel intervals with comparable resource use.
The distinction changes how the result should guide engineering.
As a canary, an intervention can still be valuable. Crash tests do not reproduce ordinary driving. Fault injection does not claim that a network cable disconnects on every request. A stress proxy can reveal configuration fragility that a calm microbenchmark hides.
As a causal explanation of the original workload ranking, it needs a bridge.
That bridge could be:
an Nsight Systems timeline from the workload
CUDA kernel start and end intervals for every relevant stream
NCCL kernel names and stream identities
matrix multiplication kernel intervals
rank-aligned traces
evidence that the replay matches overlap duration and pressureWithout that bridge, the ranking improvement is evidence for a useful probe design, not proof of workload fidelity.
the profiler must show intervals, not utilization
GPU utilization cannot prove overlap.
A monitoring sample might report 100 percent utilization during a one-second interval. It means the device was considered busy according to that monitor’s sampling rule. It does not reveal which kernels ran, whether communication and compute coincided, or how much useful work retired.
Two schedules can both show 100 percent:
schedule A:
500 ms GEMM
500 ms NCCL
no overlap
schedule B:
1,000 ms GEMM
1,000 ms NCCL
complete temporal overlapThe first executes one second of interval work in one second. The second executes two seconds of interval work in one second, assuming enough resources and the simplified durations. Utilization alone does not distinguish them.
An Nsight Systems trace can expose kernel intervals and CUDA dependencies. The capture needs:
CUDA kernel launches
CUDA API calls
NVTX ranges for token, layer, and replay phase
NCCL activity
CPU thread scheduling
GPU metrics or clock samples where supported
process and rank identityEach layer can carry an NVTX label:
with torch.cuda.nvtx.range(f"token={token} layer={layer} gemm"):
torch.matmul(activation, weight)
with torch.cuda.nvtx.range(f"token={token} layer={layer} all_reduce"):
dist.all_reduce(buffer)An NVTX range marks host-side intent. It does not replace the device intervals. The analysis must join the range with the launched GEMM kernels, NCCL kernels, stream identifiers, and event dependencies.
For one GEMM interval [g_s, g_e) and one communication interval [c_s, c_e), temporal overlap is:
overlap_us =
max(0, min(g_e, c_e) - max(g_s, c_s))For many intervals, simply summing every pair can double-count time when several kernels overlap. A correct total constructs the intersection of the union of compute intervals with the union of communication intervals:
compute_union = union(all compute intervals)
comm_union = union(all communication intervals)
overlap = duration(intersection(compute_union, comm_union))Then report denominators:
communication overlap fraction =
overlap / total communication-active time
compute overlap fraction =
overlap / total compute-active time
wall overlap fraction =
overlap / observation wall timeOne unlabeled “overlap percent” is ambiguous.
Temporal intersection still does not measure contention. Two kernels can overlap in time while using mostly separate resources. A small NCCL kernel might wait on memory transactions while a GEMM uses tensor cores. Or both might compete for SM issue slots, registers, shared memory, L2, HBM bandwidth, copy engines, and power.
Nsight Compute can inspect individual kernels for:
achieved occupancy
SM instruction throughput
memory throughput
L2 hit rate
DRAM throughput
warp stall reasons
register use
shared-memory use
eligible warps
tensor-core activityProfiling every kernel can perturb execution and serialize launches. The safer design uses Nsight Systems for representative full timelines, then profiles selected kernels in controlled repetitions with Nsight Compute. The article must not combine the two as if they came from an unperturbed single run.
CUDA event timing needs equal care. In the overlap replayer, the start event is recorded before asynchronous issue and the end event after work.wait(). The elapsed interval can include:
queueing before the NCCL kernel starts
NCCL execution
delayed progress
time until the wait entry is reached
stream-dependency effects
measurement event placementIt is an exposed issue-to-completion interval under that replay schedule. It is not automatically NCCL kernel duration.
Per-operation cudaDeviceSynchronize would make timing easy and destroy the pipeline being measured. Collecting events and synchronizing once at the end preserves more host run-ahead. It can also allow a large queue of pending work to alter resource scheduling compared with the source.
The trace should therefore preserve three timings separately:
host issue time
device kernel interval
application wait exposureThey answer different questions.
The decisive plot for the source correction would place five rows on one shared time axis:
Python and C++ collective calls
current PyTorch compute stream
NCCL stream
GEMM kernels
NCCL kernelsFor the blocking workload, the expected pattern is an NCCL completion event followed by the next GEMM on the current stream. For the explicit asynchronous workload, the prediction is that the next GEMM becomes eligible before the wait. Actual overlap then depends on resources.
If the trace contradicts that prediction, the prediction loses. Source semantics define allowed ordering. The timeline shows the execution that occurred.
a better rank is not a causal proof
Suppose proxy A agrees with 18 of 28 workload pairs. Proxy B agrees with 20. Proxy B is better on this sample by two pairwise decisions.
Several explanations fit:
B preserves a missing causal mechanism
B accidentally favors the same configurations
B overfits this eight-configuration set
B changes two ties at the analysis boundary
B increases variance and happens to move medians
B shares an unrecognized confounder with the workloadOne result cannot distinguish them.
The overlap mode changed several properties at once:
collectives became asynchronous
wait placement changed
matrix multiplications became concurrent candidates
host run-ahead changed
timing boundaries changed from per-op synchronization to one pass-end sync
CUDA event intervals spanned issue to completion
pending work accumulated in a request mapEven if concurrency is the intended independent variable, the implementation has a bundle of differences. A stronger experiment would separate them.
One ablation could issue asynchronously and wait immediately:
work = dist.all_reduce(buffer, async_op=True)
work.wait()That changes the API path without placing compute between issue and wait.
Another could delay the wait with host work only:
work = dist.all_reduce(buffer, async_op=True)
cpu_only_work()
work.wait()Another could run a memory-bound GPU kernel, then a compute-bound kernel, then a tensor-core matrix multiplication. Their different resource footprints would test whether the signal comes from SM occupancy, memory bandwidth, tensor cores, scheduling, or simply delayed waiting.
The prediction should be written before running:
If SM contention causes tree-ll instability, a compute-heavy overlapping
kernel should increase its dispersion more than a host-only delay. A
memory-heavy kernel may produce a different ordering if bandwidth is the
shared bottleneck.Then a profiler can compare stall reasons, kernel overlap, achieved occupancy, memory throughput, and clocks.
The historical campaign did not provide those ablations. The next honest state is “mechanism unresolved,” not a more polished story.
the control killed timing precision, not every notion of fidelity
The strongest part of the earlier investigation was a control that deliberately destroyed exact gap timing.
A stratified baseline retained operation types, message sizes, phases, and order, but replaced each signature’s individual gaps with representative values. The old campaign reported:
faithful overlap replay: 67.9 percent agreement
stratified overlap replay: 71.4 percent agreement
faithful timestamp replay: 60.7 percent
stratified timestamp replay: 60.7 percent
faithful sequential fill: 50.0 percent
stratified sequential fill: 57.1 percentExact timing did not separate the faithful artifact from the stratified control in those reported decisions.
The justified conclusion is narrow:
Within this workload, configuration set, hardware, replay implementations,
sampling uncertainty, and ranking metric, per-event gap precision did not
improve the observed ranking decision.It does not prove that timing never matters. Timing could matter for:
another message-size distribution
arrival bursts
queue reset gaps
tail events
multi-node congestion
collective overlap
thermal transients
credit exhaustion
an SLO based on p99 rather than medianIt also does not prove that structure plus concurrency is the unique sufficient representation. The preserved source correction removes confidence in that stronger claim.
The control still falsifies a tempting universal premise: lower timestamp error is not automatically higher decision fidelity.
This matters because trace-compression work can optimize an internal metric and assume the downstream decision follows. If compression reduces maximum gap error from 20 microseconds to 2 microseconds while both artifacts choose the wrong NCCL configuration, the better compression score has no operational value for this decision.
The right test is downstream and adversarial:
Does the artifact preserve the pairwise configuration relations?
Does it preserve tail events?
Does it preserve queue waits?
Does it preserve exposed communication?
Does it preserve the regression verdict?
Does a simpler or deliberately damaged control do just as well?When the damaged control ties, the optimized property has failed to justify its cost for that case.
one shared trace removes one circularity
The earlier per-configuration replay design captured one trace while the workload ran under each candidate configuration, then replayed each trace under the same configuration.
That is useful for isolating replay semantics. It is operationally circular.
If eight full workload runs are required to create eight traces, the operator already has the workload measurements needed to rank eight configurations. Replay has not reduced the expensive part.
The shared-trace design captures once under one reference configuration and replays the same artifact across all candidates:
capture:
workload + reference config -> one trace
replay:
same trace + config 1
same trace + config 2
...
same trace + config 8The historical article reported 71.4 percent agreement and tau 0.55 for a shared overlap trace captured under 2.20.5-default. It reported tree-ll seventh with the largest dispersion.
That design removes trace-content variation across candidates. Every candidate sees the same operation order, message sizes, and gap values. Any ranking difference then arises from candidate execution under the replay, measurement noise, or uncontrolled environment differences.
It does not remove the source-schedule problem. If the shared replayer applies a pressure pattern absent from the source workload, it is still an intervention.
The current experiment repository records newer complete shared-replay and overlap campaigns. Their physical archive descriptors and selections belong to a newer repository identity than the core campaign needed for the original ground-truth join. The analyzer correctly refuses to combine them into one trusted ranking result.
This is frustrating in exactly the right way. A result should not become publishable merely because compatible-looking rows exist in two directories.
a single event can preserve two bits
A decision-only minimizer asks a dangerously small question:
Can this subset reproduce the same pairwise ranking relations?For two configurations there is one pair. Ignoring ties, its answer is one bit: left or right. Many unrelated traces can produce the same bit.
The current CommCanary source includes a deterministic synthetic example with 100 communication events. Five events have large rank-arrival skew and large compute-overlap fields. The rest are ordinary. Two simulated configurations invert:
isolated trace:
isolated-fast-no-overlap
workload-overlap-friendly
full synthetic workload:
workload-overlap-friendly
isolated-fast-no-overlapThis is a simulator result. It is not a GPU measurement.
The fresh local run produced:
source events: 100
two-sample canary:
behavioral status: failed
configuration ranking: failed
random baseline:
source verification: failed
configuration ranking: failed
frequency baseline:
source verification: failed
configuration ranking: failed
cluster baseline:
source verification: failed
configuration ranking: failed
lossless 32-sample canary:
behaviorally_verified
configuration ranking: pass
behavior search:
selected 16 timing samples
behaviorally_verified
configuration ranking: passThen decision-only ddmin reduced the same 100 events to one event in 11 oracle calls. It preserved the ranking relation used by its oracle.
original events: 100
reduced events: 1
oracle calls: 11
budget exhausted: falseThe reducer labels its result as a decision-preserving subset and explicitly says it is not source-verified. That label is the important part.
The behavioral verifier asks more. It compiles the full source trace and candidate, replays both across configurations, and compares:
event coverage
source fidelity
median latency
p95 latency
p99 latency
maximum latency
mean latency
communication hidden percentage
queue-wait distributions
phase breakdowns
operation breakdowns
tail-event recall
pairwise configuration rankings
capture uncertaintyThe strong status requires full source coverage, source verification, behavioral metric checks, ranking checks, and no unresolved capture uncertainty.
This does not make the simulator equivalent to hardware. It prevents one particular lie inside the simulator: calling a subset behaviorally faithful because it preserved a small decision by coincidence.
The physical 8,192 to 1 reduction and its 46.4 percent agreement remain historical reports because the old raw trace is absent locally. The 100 to 1 reduction is the fresh reproducible mechanism test.
compression can make a file larger honestly
Compression has at least three denominators:
event compression = source logical events / stored logical records
byte compression = source JSON bytes / canary JSON bytes
decision compression = source cost / artifact cost at equal decision qualityThey can disagree.
A repetitive 8,192-event trace might collapse to a short motif:
[64K all_reduce, 128K all_reduce, 256K all_reduce] repeated many timesBut preserving timing samples, source digests, error intervals, rank offsets, phase identities, and verification metadata adds bytes. A canary can store fewer logical records and still occupy more bytes than a bare event list.
The historical campaign reported a byte ratio of 0.36 for its faithful canary:
source bytes / canary bytes = 0.36That means:
canary bytes = source bytes / 0.36
= 2.78 * source bytesCalling 0.36 “compression” without showing the ratio direction would be misleading. The artifact was about 2.8 times larger.
Sampling baselines reported much larger ratios because they discarded information. A 175 times smaller file is not an improvement if it changes the configuration decision. A larger verified artifact can be the correct output when evidence costs bytes.
The more meaningful operational ratio might be:
full workload GPU seconds needed for decision
------------------------------------------------
capture GPU seconds + replay GPU secondsEven that requires equal decision quality. A proxy that saves 99 percent of runtime and selects a 45 percent slower configuration is expensive.
Compression is therefore not the product property by itself. The property is decision cost at a declared error and uncertainty level.
the queue begins before the collective
The original campaign measured a closed decode-style loop. Production inference adds another system before the model: requests arrive, wait, receive service, and depart.
An NCCL configuration can change service time. Queueing can amplify that change.
Let:
lambda = completed requests per second
W = average seconds spent in the system per request
L = average number of requests in the systemFor a stable observed interval with consistent boundaries, Little’s Law is:
L = lambda * WThe units verify the equation:
requests = (requests / second) * secondsThe local queue probe sends 1,000 requests at fixed 14 millisecond intervals. Nine of every ten service times are 10 milliseconds. Every tenth is 40 milliseconds. The average service time is:
(9 * 10 ms + 1 * 40 ms) / 10 = 13 msThe offered rate is:
1 / 0.014 s = 71.429 requests/sThe service capacity based on the mean is:
1 / 0.013 s = 76.923 requests/sThe system is stable in the long run, but the 40 millisecond events build temporary queues.
The deterministic run produced:
open loop, fixed arrivals
throughput: 71.296 requests/s
queue wait p50/p95/p99: 6/26/26 ms
response p50/p95/p99: 20/40/40 ms
average in system L: 1.618565521 requests
lambda * W: 1.618565521 requestsThe finite-run throughput is slightly below the offered rate because the observation horizon includes draining the final queued requests.
The exact accounting is:
lambda = 1000 requests / horizon_seconds
W = sum(response_seconds) / 1000 requests
lambda * W
= (1000 requests / horizon_seconds)
* (sum(response_seconds) / 1000 requests)
= sum(response_seconds) / horizon_seconds
= average requests in systemThe last equality follows because each request contributes one unit to the in-system count for its response interval. Summing response intervals gives the area under the number-in-system curve.
A communication trace that records only collective durations has no request arrival process. It cannot predict this queue. A token trace with no admission policy cannot predict it either.
Once queueing matters, a 2.8 percent service-time regression can produce more than a 2.8 percent p99 regression near capacity. The extra service time holds queue positions longer, which increases the number of waiting requests, which increases later response times.
Ranking configuration by median collective duration can therefore disagree with ranking it by SLO attainment.
the request generator can erase its own queue
The same local probe also runs one closed-loop client. It submits the next request only after the previous response and a 1 millisecond think time.
closed loop
throughput: 71.429 requests/s
queue wait p50/p95/p99: 0/0/0 ms
response p50/p95/p99: 10/40/40 msIt reports no queue wait because it never has more than one outstanding request. When a 40 millisecond response occurs, the generator stops offering work for those 40 milliseconds. The slow response reduces offered load at precisely the moment when a real independent arrival stream would build a queue.
This is coordinated omission. The measurement schedule coordinates itself with system completion and omits the waiting times that would have occurred under an external arrival schedule.
The HdrHistogram documentation describes correction using an expected sampling interval. If one recorded value exceeds that interval, corrected recording accounts for values that would have been sampled while the long operation was in progress.
Correction is not a substitute for a representative generator. It is a way to reason about one known omission pattern.
For replay, the workload model must state whether it is:
closed:
each client waits for completion before issuing its next request
open:
arrivals occur independently of completion
partly open:
many closed clients collectively approximate an external rateIt must also state:
arrival distribution
concurrency
admission control
batching
timeout policy
retry policy
warmup
measurement window
drain policyA replay that preserves GPU kernels and collective overlap while changing the arrival model can still rank production configurations incorrectly.
This is why idle gaps are not compute, and compute intervals are not requests, and request timestamps are not a queueing policy. Each abstraction discards state needed by the layer above it.
rank arrival is part of collective latency
An all_reduce needs participation from every rank in its group. If rank 3 arrives 200 microseconds after ranks 0, 1, and 2, the earlier ranks may show time that looks like communication but is partly waiting for a peer.
Define rank arrival offsets relative to the first arrival:
rank 0: 0 us
rank 1: 12 us
rank 2: 7 us
rank 3: 200 usArrival skew is:
max(offsets) - min(offsets) = 200 usNow separate:
arrival wait:
time an early rank waits before the collective can make equivalent progress
transport and reduction:
time spent moving and reducing data once participation permits progress
exposed latency:
portion not hidden behind useful workThese quantities depend on where timing starts and ends. A CPU timer around dist.all_reduce can include Python dispatch and host waits. A CUDA event on the compute stream can include dependencies. An event on NCCL’s stream can exclude time before NCCL launch. A profiler event named all_reduce may represent a CPU operator rather than the complete device interval.
A single-rank trace cannot reconstruct arrival offsets on other ranks. Joining rank traces by sequence number can fail if events are dropped, groups interleave, or different groups reuse local sequence counters. A sound join needs operation identity:
process group
ordered membership
collective kind
message shape or byte count
sequence number in that group
phase or channel when necessaryIt also needs to detect ambiguity. Guessing a join because two events are close in time turns timing noise into topology.
Arrival skew can change algorithm ranking. A tree may expose a late parent differently from a ring waiting on a neighbor. Chunked protocols can begin partial progress before the last rank reaches the same logical point. The exact behavior depends on implementation, topology, and launch order.
Any claim that a proxy preserves collective performance while omitting rank arrival needs an experiment showing that arrival is irrelevant for the tested decision. It cannot be assumed.
clocks disagree before ranks do
Imagine rank 0 records a collective start at 1,000 microseconds and rank 1 records 980 microseconds. That does not prove rank 1 arrived 20 microseconds earlier.
Their clocks could have a 30 microsecond offset:
rank 0 clock = real time + 0 us
rank 1 clock = real time - 30 usIf the real arrivals were rank 0 at 1,000 and rank 1 at 1,010, the recorded values would be:
rank 0: 1,000 us
rank 1: 980 usNaive subtraction reverses the order.
Clock calibration needs a model. A simple affine model is:
local_time = scale * reference_time + offsetoffset accounts for clocks starting at different values. scale accounts for drift. Calibration samples estimate both, with an error interval affected by message delay and jitter.
For a short single-node run, a constant offset might be adequate. For a long multi-node capture, drift can matter. For CPU and GPU clocks, conversion can involve profiler correlation records rather than network exchange.
Every derived arrival value should therefore carry uncertainty:
arrival_offset_us = 200
calibration_error_us = 8If two configurations differ by 5 microseconds while clock alignment error is 8, the trace cannot support a strict arrival-sensitive comparison.
The same principle applies to profiler overhead. Capturing shapes, stacks, and CUDA activities adds work and storage. If capture changes token latency, kernel launch spacing, or memory pressure, the trace describes the instrumented workload, not the uninstrumented one.
A capture-overhead experiment needs paired runs:
control: workload with capture disabled
treatment: identical workload with capture enabledInterleave them to avoid drift. Compare median, tails, throughput, kernel schedule, and memory use. Report the overhead distribution, not one percentage.
The trace should record that uncertainty. A replay compiler cannot recover fidelity that capture already destroyed.
the trace needs a contract for what it does not know
A useful trace contract is more than required JSON keys.
For each field, it should say:
measurement source
clock domain
unit
resolution
calibration method
known error
missing-value meaning
join key
privacy classification
versioned semanticsConsider compute_overlap_us.
It could mean:
intersection of NCCL kernel intervals with all non-NCCL GPU kernels
intersection with GEMM kernels only
time from asynchronous issue until wait, regardless of actual device overlap
simulator-estimated hidden communication
user-provided annotationThose definitions produce different numbers.
The current CommCanary compiler keeps separate fields for compute before a collective, overlap duration, compute pressure, observed exposed latency, rank offsets, and concurrent groups. It also computes fidelity error budgets. That is better than one overloaded “gap.”
The compiler derives and stores a SHA-256 digest of the canonical source trace. It groups events by an operation identity that includes phase, operation, bytes, ranks, group, and relevant point-to-point fields. It preserves source counts and digests through grouping. It can store exact timing patterns or bounded approximations.
None of those mechanisms proves that the incoming fields were measured correctly. Integrity begins after capture. A perfectly hashed wrong trace is still wrong.
The assurance states therefore need to remain separate:
schema-valid:
fields have permitted shapes and values
integrity-bound:
bytes and transformations match recorded digests
source-verified:
compressed representation stays within declared source error bounds
behaviorally verified:
replay metrics and rankings match within declared tolerances
hardware validated:
claims were tested on identified physical systems
production representative:
arrival, load, topology, and SLO match the intended deploymentNo lower state implies a higher one.
a hash cannot tell whether a campaign finished
SHA-256 detects changed bytes with overwhelming probability. It does not tell whether all expected bytes exist.
Suppose a campaign expects:
8 configurations * 5 repetitions = 40 cellsThirty-nine successful files can all have valid hashes. The missing fortieth cell can be the slowest configuration’s failed run. Aggregating the 39 produces a clean but biased table.
Completeness needs an expected matrix frozen before execution:
configuration 1, repetition 0
configuration 1, repetition 1
...
configuration 8, repetition 4For each cell, retries can create several terminal attempts:
attempt 1: failed
attempt 2: success
attempt 3: success after an environment changeThe analysis must not silently glob the newest file. A selection must explicitly choose one attempt and bind its hash. The selection policy should explain why attempt 2 or 3 is admissible. If the environment changed, they may not be members of the same experiment.
Then a completeness verdict checks:
every expected cell has exactly one selected attempt
every selected attempt exists
the selected record hash still matches
the selected status is successful
no selected measurement contradicts its manifest identityOnly after that should rankings appear.
The current claim builder returns withheld-incomplete when completeness fails. That is the desired failure mode. Diagnostic aggregates can exist, but they should be visibly unsuitable for publication.
The old public sentence said every hardware measurement regenerated from the open kit. That sentence failed this standard because the old raw archive is absent. Removing it is not an editorial downgrade. It makes the public claim match the available evidence.
two valid campaigns can still refuse to meet
The local evidence-join probe creates two complete evidence records. Each has:
campaign identity
repository commit
manifest SHA-256
selected attempt SHA-256
measurement valueWhen the repository identities match, the miniature join accepts them:
accepted join:
core + replay
repository_commit=aaaaaaaa...When they differ, it refuses:
refused join:
repository identity mismatch
aaaaaaaa... != dddddddd...Changing one selected value from 123 to 122 microseconds also changes its canonical SHA-256 digest:
original:
566c6f05726761e5563f1ba72ea90c091f319b6ac01cfd4245cd10c74603ffea
tampered:
94ecf6f52d650513ebd68eca59edec10bca35cf79ca27f7883105f21c8f06ea3The real CommCanary analyzer enforces a stronger join. Joined campaigns must agree on repository identity, expected site contract, analysis policy, configuration definitions, workload definitions, and shared input identities. It also refuses repeated manifests and repeated run or campaign identities.
This is why the newer complete shared-replay and overlap archives cannot simply be pasted beside the old core ranking. Their rows may each be valid within their campaigns. The cross-campaign claim requires a shared identity contract they do not currently satisfy.
There are two legitimate repairs:
rerun the missing core campaign under the newer frozen identity
or
construct and justify a versioned compatibility bridge before execution,
then make that bridge part of the frozen analysis policyEditing the analyzer after seeing results to permit the join would change the experiment in response to its output.
traces also contain customer data
A communication trace without tensor values can still reveal:
model layer count
hidden dimensions inferred from message sizes
tensor-parallel group size
pipeline boundaries
sequence-length changes
prefill and decode phases
batching patterns
request bursts
mixture-of-experts routing imbalance
topology and hostnames
software versions
customer traffic timingRemoving payload bytes is not equivalent to removing sensitive information.
Hashing a hostname is often inadequate. A stable hash preserves equality, so repeated traces can still be linked. A small known set of hostnames can be brute-forced. Exact timestamps can correlate with logs in another system.
A privacy-preserving export needs a threat model. It should decide which questions the artifact must answer and remove fields not required for those questions.
Possible transformations include:
replace physical ranks with local group indices
remove hostnames and job IDs
rebase timestamps to a local zero
quantize timestamps only if the decision survives
bucket message sizes only if the decision survives
separate secret environment details from public evidence
keep salted, access-controlled linkage keys outside the artifact
enforce retention limits
record every export in an audit logEvery transformation is another compression. It needs the same behavioral test. If message-size bucketing changes transport selection, the privacy transform has changed the experiment.
The safest public artifact may be a verified report plus content hashes held by the customer, rather than the raw trace. That reduces independent inspection, so the report must state which evidence remains private and who can audit it.
Privacy and reproducibility pull in opposite directions. Pretending otherwise usually means one of them has been ignored.
the production question is not which row wins
An operator does not deploy a rank. An operator deploys a configuration into a workload with constraints.
A useful decision record needs:
candidate identity
baseline identity
hardware and topology
workload version
traffic model
throughput
time to first token
time per output token
inter-token latency
p50, p95, and p99
SLO attainment
error rate
GPU memory use
power or cost where available
confidence interval
rollback triggerA candidate can have lower median token latency and worse p99. It can increase throughput while violating an inter-token latency SLO. It can rank well in steady state and stall during model loading. It can save GPU time while requiring an unsupported NCCL environment variable.
The rollout should therefore be reversible.
One safe sequence is:
offline evidence:
replay and full-workload confirmation on representative hardware
shadow:
observe decisions without serving candidate output
small canary:
a bounded traffic fraction with automatic rollback
expanded canary:
enough volume to estimate tails and rare failures
fleet:
gradual expansion while retaining the baselineThis is not a public reading path. It is the execution path by which a configuration claim becomes a production change.
Rollback thresholds must be declared before the canary:
p99 inter-token latency exceeds baseline by more than X
SLO violation rate exceeds Y
GPU error count becomes nonzero
throughput falls below Z
rank timeouts exceed a fixed countIf thresholds are chosen after seeing the candidate, the decision inherits the same post-selection bias as an experiment that edits its tie policy after examining medians.
A performance canary also needs a negative control. Replaying the baseline configuration through the new measurement path detects changes caused by the path itself. A randomized no-op configuration label can test whether analysis accidentally uses label order. Repeating the old known-bad configuration can show whether the canary still has enough sensitivity to detect a regression.
a canary needs a result that should not move
A canary that only sees candidate differences has no way to tell a candidate effect from a broken instrument.
The first negative control is baseline against itself:
baseline A, label left
baseline A, label rightBoth labels point to byte-identical configuration and software. The analysis should call them tied at the expected false-positive rate. If it consistently ranks one label faster, run order, label handling, caching, or analysis has leaked into the result.
The second is replay without the intended mechanism:
communication-only timestamp replay
same replay plus host sleep
same replay plus CPU work
same replay plus sequential GPU work
same replay plus explicitly overlapping GPU workThese modes reveal which change moves the ranking. If host sleep and overlapping GPU work produce the same result, simultaneous GPU pressure is not required by the evidence.
The third is a known perturbation with a predictable sign. Add a bounded delay before one rank enters every tenth collective:
if rank == delayed_rank and sequence % 10 == 0:
time.sleep(delay_seconds)The prediction is not that every algorithm slows equally. It is that arrival-skew metrics and exposed waits must increase in the affected windows. If the capture cannot see the injected skew, it cannot support claims about natural skew.
The fourth is trace corruption:
drop one rank shard
change one message size
duplicate one event
change a selected attempt byte
join a campaign from a different repository identityThe correct result is refusal, not a degraded score. These tests exercise the evidence boundary rather than GPU performance.
The fifth is a deliberately poor configuration whose workload regression has already been independently established under the same frozen environment. A canary should detect it before being trusted to clear a subtle candidate. If it misses the known regression, an apparent pass for a new candidate means little.
Controls also need blind labels. If analysis code knows which directory is “candidate,” a sign error or sorting fallback can favor it. Map opaque identifiers to configurations only after aggregate generation:
config-17
config-42The mapping remains in a frozen private manifest. Public reports can reveal names after verdicts are fixed.
Sensitivity has a cost. A threshold small enough to detect a 1 percent change may false-alarm on harmless node noise. A threshold large enough to avoid noise may miss an economically important regression at fleet scale.
The threshold belongs to the deployment objective:
minimum effect worth acting on
maximum false rollout rate
maximum missed-regression rate
cost of one full-workload confirmation
cost of one wrong fleet decisionThe canary can abstain:
pass
fail
insufficient evidenceForcing every noisy comparison into pass or fail converts uncertainty into confidence. “Insufficient evidence” can trigger more repetitions, a full workload run, or a smaller rollout.
Behavioral verification and physical validation must remain separate in that workflow. A simulator can prove that a compressed artifact preserves simulator behavior. A hardware canary can test a physical candidate. Neither alone shows production representativeness.
The most useful canary output is not a recommendation without context. It is an auditable claim:
Under manifest M, on hardware H, with workload projection W and replay
semantics R, candidate C beat baseline B by D, uncertainty interval U,
while controls N1 through N5 behaved as predicted. Production dimensions
P1 and P2 were not represented.That sentence is longer than “use tree.” It contains the boundaries needed to decide whether the recommendation travels.
a faster row becomes capacity only after division
Performance value must carry units.
Consider a hypothetical fleet. These are scenario inputs, not measured customer data:
demand: 1,000,000 tokens/s
baseline throughput: 8,000 tokens/(s * GPU)
candidate throughput: 9,200 tokens/(s * GPU)
GPU cost: $3/(GPU * hour)
operating time: 8,760 hours/yearThe throughput improvement is:
(9,200 / 8,000 - 1) * 100 = 15 percentBaseline capacity:
ceil(
1,000,000 tokens/s
/
8,000 tokens/(s * GPU)
) = 125 GPUsThe units reduce correctly:
(tokens/s) / (tokens/(s * GPU)) = GPUCandidate capacity:
ceil(1,000,000 / 9,200) = 109 GPUsAvoided capacity at exactly this demand:
125 - 109 = 16 GPUsGross annual capacity value:
16 GPU
* $3 / (GPU * hour)
* 8,760 hour / year
= $420,480 / yearThe units again cancel:
GPU * dollars/(GPU * hour) * hour/year = dollars/yearThis is not profit. It omits:
latency SLO failures
redundancy and headroom
utilization
power and cooling
network cost
engineering labor
rollout risk
support burden
reserved-capacity commitments
whether 16 GPUs can actually be removed from a fragmented fleetIf the fleet keeps the same GPU count, the value may appear as additional traffic rather than avoided cost. If demand is below committed capacity, the immediate cash value may be zero. If the candidate raises tail latency, the 15 percent throughput figure may be unusable.
This is why a replay tool cannot stop at “agreement improved from 64.3 to 71.4 percent.” It must connect the extra correct decisions to deployment outcomes and the cost of being wrong.
the tuner is still the missing row
The two default configurations ranked first and second in the reported workload. Every forced algorithm and protocol was worse.
None of the historical proxies reproduced that cleanly. The overlap modes improved the position of tree-ll, but the defaults remained a large part of the residual disagreement.
NCCL’s internal selection can depend on information absent from a simple trace proxy:
message size
topology
channel count
transport
protocol thresholds
collective algorithm cost models
runtime version
buffer registration
communicator initialization
observed launch and resource stateSome of these are fixed at communicator setup. Some vary by message. Some are implementation details that can change between NCCL versions.
Forcing NCCL_ALGO and NCCL_PROTO disables part of that selection. A replay that compares forced settings against defaults is testing both execution and selection policy.
The missing experiment is not another rank table from the same proxy. It is a path trace of the tuner’s choices:
for every collective:
chosen algorithm
chosen protocol
channel count
transport path
message size
topology context
measured durationThen compare workload and replay choices event by event. If the default chooses different protocols under replay, the proxy changed its input to the tuner. If choices match but timings do not, the omitted state lies later in execution.
Until that is measured, the default rows bound what the proxy family can claim.
the observation now has three answers
The opening row supports three different statements, each with a different evidence class.
The historical campaign reported:
tree-ll ranked sixth in an isolated benchmark
tree-ll ranked eighth in the decode-like workload
the strict displayed orders agree on 18 of 28 pairsThe current source inspection established:
the preserved workload calls all_reduce without async_op=True
PyTorch 2.4.1 inserts a current-stream dependency in that path
the next current-stream GEMM cannot pass the NCCL completion event
the explicit overlap replayer executes a different schedule
the old overlap explanation is unsupported by the preserved sourceThe fresh local mechanism tests established:
timestamp silence, sequential fill, and explicit overlap are different programs
one closed-loop generator can hide queue waits created by an open arrival stream
a decision-only reducer can compress 100 synthetic events to one while losing behavior
a behavioral verifier can refuse weak artifacts
content hashes detect changed bytes
repository-identity mismatches can and should refuse an evidence joinNone of the local tests remeasured tree-ll on an A100.
The next physical experiment has a sharper shape now. Run the preserved blocking workload, an explicitly asynchronous workload, the communication-only replay, sequential compute fill, and the overlap intervention under the same frozen campaign. Capture GPU timelines, clocks, rank arrival, NCCL choices, and request-level outcomes. Interleave configurations. Keep every raw attempt. Select before analysis. Refuse incomplete joins.
The prediction is no longer that the overlap replayer faithfully reproduces the preserved workload.
It is:
If simultaneous compute pressure exposes a real tree-ll sensitivity, the
explicitly asynchronous workload and the overlap intervention should share
the instability. The preserved blocking workload may not. If all three share
it, another state variable links them. If only the old reported table contains
it, the result itself needs replication before its mechanism is discussed.The original table can still be true. Its first explanation was not.
That is where the investigation stops. The idle interval in a communication trace still does not mean the GPU did nothing. It also does not tell us that compute overlapped communication. A gap is elapsed time with its causes removed. Replaying it faithfully begins by admitting what the trace no longer knows.