You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
PR #3002's evidence records the board at 71.3% busy pooled (63.2% on the
64-token legs, 76.3% on the 128-token legs). That figure was read as the
strongest lead in the run: a compute-bound decode should not be idle a third of
the time, which points at host-side stalls or serialization.
It is a windowing artifact. The 71.3% is a mean over the whole leg, and a
leg contains process start, model load and teardown as well as the four timed
generations. job.log records generate_start_unix and generate_end_unix for
every run to the microsecond, and the clock sampler runs at --interval 0.25,
so the samples can be windowed to the generation intervals exactly.
Windowed to the generate intervals, with the cold run 1 dropped as the fold
drops it, over all 8 legs of /mnt/nas_share/rc/strix-arm-2933/out/:
leg
samples
busy mean
sclk mean
n128-r1..r4
285 / 281 / 279 / 284
100.0 each
2877 / 2872 / 2870 / 2866
n64-r1..r4
143 / 144 / 144 / 143
100.0 each
2889 / 2875 / 2870 / 2871
pooled
1703
100.0
~2872 MHz
gpu_busy_percent is integer-valued and 1703 of 1703 samples read exactly
100. Minimum, p10, median and maximum are all 100.
The same treatment applied to llama.cpp
The survey's llama.cpp legs (/mnt/nas_share/rc/strix-survey-2497/out/survey-20260904/clock-llamacpp-r*.jsonl)
have no per-run timestamps, but the generation phase is unambiguous in the
series -- a single contiguous run of >= 95 samples in every leg. Over that
span:
leg
samples
span
busy mean
sclk mean
r1
52
12.8 s
99.3
2675
r2
54
13.3 s
99.3
2637
r3
54
13.3 s
99.3
2675
r4
54
13.3 s
99.2
2656
We hold the board at a higher occupancy and a 210 MHz higher shader clock than
llama.cpp does, and we are still 2.26x slower. That is the opposite of the
signature a host-side stall leaves. It is the signature of a shader-bound
kernel: llama.cpp's lower sclk is what a DRAM-bound decode looks like on an
APU whose SMU drops the shader clock while the memory controller is the
critical path, and our pinned 2872 MHz is what a decode looks like when the
shaders are the critical path.
What this rules out, and what it does not
The counter is a firmware activity average sampled every 0.25 s. All samples
reading the integer 100 bounds serialized host-side idle at under about 0.5%
of decode wall time, so per-token host work -- the hipStreamSynchronize in greedy_argmax_host (src/vllm/v1/sample/sampler.cpp:275), host
detokenization (src/vllm/v1/engine/detokenizer.cpp:448-456), the depth-1
engine loop from max_concurrent_batches=1 -- cannot be more than a percent or
so of the gap at this model size. It does not resolve microsecond gaps
between the ~700 eager kernel launches per token, which no ROCm build can
collapse today (src/vllm/platforms/rocm.cpp:91-97 leaves support_static_graph_mode() false). This tree's own same-binary A/B on
gfx1200 measured graph capture at 0-2%, sys time unchanged
(.agents/specs/rocm-decode-graph.md D7), which is consistent.
So at least 98% of the 103.5 ms/token excess is GPU kernel time.
What is asked
Window the clock record to the generate interval in the arm harness's
fold, or publish both numbers with the whole-leg one labelled as including
load. The current figure has already misdirected one investigation.
Correctness state carried, unchanged, per #2921: TOKEN_GATE=FAIL, 3 of 6
prompts divergent against llama.cpp b10451, 5 of 6 against vLLM compiled,
every divergence a near-tie at about 0.125 nats, no deterministic denominator
established on this path. Nothing here is a parity claim.
Row:
BACKEND-GATE-ROCM-LLAMACPPPR #3002's evidence records the board at 71.3% busy pooled (63.2% on the
64-token legs, 76.3% on the 128-token legs). That figure was read as the
strongest lead in the run: a compute-bound decode should not be idle a third of
the time, which points at host-side stalls or serialization.
It is a windowing artifact. The 71.3% is a mean over the whole leg, and a
leg contains process start, model load and teardown as well as the four timed
generations.
job.logrecordsgenerate_start_unixandgenerate_end_unixforevery run to the microsecond, and the clock sampler runs at
--interval 0.25,so the samples can be windowed to the generation intervals exactly.
Windowed to the generate intervals, with the cold run 1 dropped as the fold
drops it, over all 8 legs of
/mnt/nas_share/rc/strix-arm-2933/out/:gpu_busy_percentis integer-valued and 1703 of 1703 samples read exactly100. Minimum, p10, median and maximum are all 100.
The same treatment applied to llama.cpp
The survey's llama.cpp legs (
/mnt/nas_share/rc/strix-survey-2497/out/survey-20260904/clock-llamacpp-r*.jsonl)have no per-run timestamps, but the generation phase is unambiguous in the
series -- a single contiguous run of
>= 95samples in every leg. Over thatspan:
We hold the board at a higher occupancy and a 210 MHz higher shader clock than
llama.cpp does, and we are still 2.26x slower. That is the opposite of the
signature a host-side stall leaves. It is the signature of a shader-bound
kernel: llama.cpp's lower
sclkis what a DRAM-bound decode looks like on anAPU whose SMU drops the shader clock while the memory controller is the
critical path, and our pinned 2872 MHz is what a decode looks like when the
shaders are the critical path.
What this rules out, and what it does not
The counter is a firmware activity average sampled every 0.25 s. All samples
reading the integer 100 bounds serialized host-side idle at under about 0.5%
of decode wall time, so per-token host work -- the
hipStreamSynchronizeingreedy_argmax_host(src/vllm/v1/sample/sampler.cpp:275), hostdetokenization (
src/vllm/v1/engine/detokenizer.cpp:448-456), the depth-1engine loop from
max_concurrent_batches=1-- cannot be more than a percent orso of the gap at this model size. It does not resolve microsecond gaps
between the ~700 eager kernel launches per token, which no ROCm build can
collapse today (
src/vllm/platforms/rocm.cpp:91-97leavessupport_static_graph_mode()false). This tree's own same-binary A/B ongfx1200 measured graph capture at 0-2%,
systime unchanged(
.agents/specs/rocm-decode-graph.mdD7), which is consistent.So at least 98% of the 103.5 ms/token excess is GPU kernel time.
What is asked
fold, or publish both numbers with the whole-leg one labelled as including
load. The current figure has already misdirected one investigation.
rocprofv3 --kernel-traceon gfx1151 for this artifact, bothengines, one lease, identical workload, as
.agents/porting.mdrequires.No per-kernel attribution exists on this board; KQuantGemmK strides 32 lanes over nsb=K/256 superblocks: half the warp idles on 75% of decode calls, and it is 54% of ROCm decode #1910's table is gfx1200 on
a 9B model. Without it, the split of the ~102 ms/token among the kernel
families named in the linked issues stays a byte-weighted estimate.
Correctness state carried, unchanged, per #2921:
TOKEN_GATE=FAIL, 3 of 6prompts divergent against llama.cpp
b10451, 5 of 6 against vLLM compiled,every divergence a near-tie at about 0.125 nats, no deterministic denominator
established on this path. Nothing here is a parity claim.
Related: #2497, #2921, #2940, #3002, #2109, #1910, #2964, #332.