diff --git a/.gitignore b/.gitignore index d3f2e448..517deaf9 100644 --- a/.gitignore +++ b/.gitignore @@ -90,6 +90,9 @@ shim/perfagent-gpu-fpless # `make -C shim nvidia-workload`, needs the CUDA toolkit, and has no extension # for the *.so / *.o rules above to catch. shim/nvidia/testdata/cuda_workload +# The CONCURRENT CUDA workload the PC-sampling overhead benchmark is measured +# against: built by `make -C shim nvidia-concurrent`, same story. +shim/nvidia/testdata/cuda_concurrent bench/cmd/scenario/scenario bench/cmd/report/report diff --git a/.superpowers/sdd/task-12-overhead-report.md b/.superpowers/sdd/task-12-overhead-report.md new file mode 100644 index 00000000..d85575ed --- /dev/null +++ b/.superpowers/sdd/task-12-overhead-report.md @@ -0,0 +1,384 @@ +# Task 12 — Overhead: what to measure, against what, and what sinks Tier A + +Branch `feat/pc-sampling-overhead`, one commit on `origin/main` (`6f7cdf08`, the Task 11 merge). + +**This task produced the instrument, not the verdict.** The implementer has `CapEff: 0` and no +GPU. Every number the plan asks for is outstanding and **the tier decision is not yet made**. What +is here is a `bench/` scenario built so that running it on the RTX 3090 yields a decision rather +than a discussion, plus the offline half of its own correctness: the parsers, the medians, the +ratio and all four pre-committed threshold clauses are unit-tested, and the skip path is +exercised. + +No Go file under `gpu/` or `gpuprobe/` changed, and nothing under `shim/` changed except the new +`nvidia-concurrent` target in `shim/Makefile` and the new `.cu` it builds. Tier A, Tier B, tier +selection, the adapter and the ABI are untouched; no `.o` churn. + +--- + +## The workload, and why it is a second one + +`shim/nvidia/testdata/cuda_concurrent.cu`, new, built by `make -C shim nvidia-concurrent` with the +same `-O2 -g -lineinfo -arch=sm_86 -rdynamic` as `cuda_workload`. + +**I added a second workload rather than extending `cuda_workload.cu`, and the reason is the +measurement rather than tidiness.** `cuda_workload.cu` is a two-kernel *serial* loop on the +default stream with a `usleep` between iterations; its kernels are 64k elements of one FMA each. +That shape is exactly right for what it exists for — it is the fixture the adapter, the join, the +kernel-name table and the `-lineinfo` source resolution are all proven against, and the hardware +half of the phase gate asserts against **its** source lines. Changing its shape would change what +those gates measure. + +It is also the wrong instrument here, for the reason the plan states: a stream of trivial kernels +exaggerates per-launch costs and **understates serialization costs**, because serialization hurts +in proportion to the concurrency it destroys and a serial loop has none. A Tier A number measured +on it would be small for the same reason a stopped clock is right twice a day. + +The new workload: + +- **S streams** (default 4), created `cudaStreamNonBlocking` — spelled out rather than defaulted, + because the default (blocking) flag makes every stream serialize against the legacy default + stream and the workload would then have no concurrency at all, which is precisely the failure + the file exists to avoid; +- **non-trivial kernels.** One kernel is `rounds` dependent FMAs up then `rounds` dependent FMAs + down. The duration comes from the *dependency chain*, not from a large grid — deliberately, so + the kernel occupies a fraction of the device (16 blocks × 256 threads by default on an 82-SM + GA102) and several kernels genuinely co-reside. A big-grid kernel would saturate the device and + leave nothing for a second stream to overlap with, quietly turning this back into the serial + workload; +- **fixed work**, not fixed time: `iters × streams` kernels, whatever it takes; +- a **warm-up phase outside the timed region**, so module load, JIT, first-touch allocation and + the adapter's own steady state are not charged to the measurement; +- a **device sync every `sync_every` iterations**, which is not housekeeping: it bounds queue + depth and forces the pipeline to **refill** after every drain. Refill cost is exactly what makes + Tier A's damage outlast its burst, which is what the cost-over-duty ratio exists to detect; +- an **exact** self-check. Every kernel is the identity on its buffer (`+1` × rounds then `−1` × + rounds, with `rounds ≤ 2^23` so every partial sum is exactly representable in float), so + `max_abs_err` must be **exactly 0.0**, not "small". An arm that perturbed the computation fails + the run instead of contributing a number. + +**The loops survive the compiler, and that is checked rather than assumed.** `up` and `down` are +runtime kernel arguments and `#pragma unroll 1` blocks unroll-and-reassociate. Verified offline +with `cuobjdump -sass`: the SASS for `_Z20perfagent_conc_chainPfiiff` contains exactly **two +`FFMA`s, each inside its own backward branch** — two rolled loops of one dependent FMA. Nothing +was folded into a closed form, and `cuobjdump -elf` shows the `-lineinfo` `debug_line` sections +are present. + +**The realism is measured, not claimed.** The harness computes, from the profile the adapter +produces: + +``` +concurrency = Σ(exec EndNs − StartNs) / (max EndNs − min StartNs) +meanKernelUs = Σ(exec EndNs − StartNs) / count +``` + +and **fails the whole run** if the *baseline* arm's concurrency is below `--gpu-min-concurrency` +(1.5) or its mean kernel duration below `--gpu-min-kernel-us` (50). A workload that turned out to +be serial, or whose kernels turned out to be microseconds long, would otherwise produce a small, +green, meaningless Tier A cost — the benchmark-shaped instance of this project's standing defect. +Both spans are taken over the executions the profile *retained*, so ring eviction cannot skew the +ratio: numerator and denominator describe the same interval. + +--- + +## The baseline, and what it is not + +**The baseline arm is the shipping Phase 4 configuration with PC sampling off** — shim injected +via `CUDA_INJECTION64_PATH`, RUNTIME + RESOURCE callbacks, `CONCURRENT_KERNEL` activity, 100 ms +drain, launch sampling, and a `gpuprobe` consumer attached so the probe semaphores are armed and +the producer actually does its work. §9.1 already measured injection at 0.0% and the activity path +at −10.0%; those costs are paid whether or not PC sampling is on, so charging them to PC sampling +would be measuring the wrong thing twice. + +An **uninjected** run is taken as well, once, *before* the arms. It is labelled `calibration` in +the output and in the JSON, it is never used as a baseline, and it exists for two reasons: it +warms the device clocks before the first arm, and it proves the fixed work is sized for the device +in front of it (the run fails, with the retuning advice, if it falls outside 10–120 s). + +The consumer runs **without a symbolizer**, identically on every arm. Resolving a sampled launch's +stack is agent-side work that costs the workload nothing, and it cannot succeed anyway once the +workload has exited. The producer still samples launches and still captures stacks in BPF on every +arm, exactly as the shipping configuration does; only the agent-side resolution — a variance +source that cancels — is left out. + +--- + +## The arms, and what each proves it ran + +Five arms, **five interleaved runs each** (every arm once per round, not five of one then five of +the next, so thermal drift and boost state fall on all arms alike), medians, fixed work. + +| arm | env handed to the producer | duty | +| --- | --- | ---: | +| baseline | `PERFAGENT_GPU_PC_SAMPLING=off` | — | +| tier B continuous | `=continuous` | — | +| tier A 50 ms / 450 ms | `=serialized`, `BURST_MS=50`, `MAX_DUTY_PERMILLE=100`, `MAX_GAP_MS=450` | 10% | +| tier A 50 ms / 950 ms | `=serialized`, `BURST_MS=50`, `MAX_DUTY_PERMILLE=50`, `MAX_GAP_MS=950` | 5% | +| tier A 50 ms / 1950 ms | `=serialized`, `BURST_MS=50`, `MAX_DUTY_PERMILLE=25`, `MAX_GAP_MS=1950` | 2.5% | + +**The gaps are pinned from both sides, and that is deliberate.** The burst controller clamps its +gap into `[min_gap, max_gap]` where `min_gap = burst × (1/max_duty − 1)` — so 450 / 950 / 1950 ms +after a 50 ms burst *are* 10% / 5% / 2.5%. Setting the ceiling gives the minimum; setting +`MAX_GAP_MS` to the same value gives the maximum; the interval collapses to a point and whatever +the closed loop computes, the gap it returns is the arm's gap. With only the ceiling set, a +workload producing a high pair rate would have the loop lengthen the gap, the arm would run at a +duty nobody asked for, and the ratio would be divided by a denominator that never happened. The +tier is written **explicitly on every arm including the off one**, so an exported +`PERFAGENT_GPU_PC_SAMPLING` in the operator's shell cannot turn the baseline into a serializing +arm and make every other arm look free. + +**Each arm proves it ran in the mode it claims, from two ends that can disagree.** The producer's +own report line on stderr (`PERFAGENT_GPU_LOG=stderr`) and the consumer's counters are parsed +separately and both are recorded in the JSON, because an arm proved only by the consumer would +look identical whether the producer ran in the right mode and lost records or ran in the wrong +mode and had nothing to send. + +Every clause below **fails the run** when it does not hold: + +| arm | asserted | +| --- | --- | +| *all* | `ExecutionsSeen > 0` and the adapter printed a report line at all — an arm where the pipeline never ran satisfies every negative clause below perfectly and would be the cheapest arm in the table | +| baseline | producer `tier=off`; `PCSamplesDecoded == 0`; `SamplingWindowsDecoded == SamplingWindowsReceived == 0`; `bursts == 0`; `ExecutionsSerialized == 0` | +| tier B | producer `tier=continuous`; `pc_records > 0`; `PCSamplesDecoded > 0`; **zero** windows and **zero** bursts (a `CONTINUOUS` producer announcing a window would be claiming a perturbation it did not cause) | +| tier A | producer `tier=serialized`; `bursts ≥ --gpu-min-bursts` (4); `windows == 2N` or `2N−1`; `start_failed == stop_failed == 0`; `graph_execs == graph_refused == 0`; `PCSamplesDecoded > 0`; `SamplingWindowsReceived > 0`; `Snapshot.PCSampling == serialized`; and **`ExecutionsSerialized > 0`** | +| tier A | achieved duty inside `[configured/2, (burst+2·tick)/(burst+2·tick+gap)]` — the upper bound is *derived* from the burst timer's `burst/5` tick plus two ticks of scheduling slack, not a magic tolerance | +| cross-arm | a lower duty must open **strictly fewer** bursts than a higher one over the same fixed work | + +The two load-bearing ones are the last two rows and the `ExecutionsSerialized > 0` clause. Bursts +that overlapped no kernel serialized nothing, so an arm with zero serialized executions measured +the cost of starting and stopping CUPTI and not the cost of serialization — and would report a +beautifully small number for it. And if the three Tier A arms all opened the same number of +bursts, the duty environment did not take, they are one arm under three names, and their three +different ratios are fiction. + +On failure the run still writes its JSON — the arms are the diagnostic — but the `decision` object +stays at its zero value so nothing in the file can be read as a verdict, stderr says so in as many +words, and the exit code is 3. + +--- + +## The threshold logic + +`cost% = (arm median − baseline median) / baseline median × 100`, from the workload's own +fixed-work `elapsed_ms` (which excludes CUDA init, allocation and warm-up; the whole-process time +is recorded beside it, because a divergence between the two would mean the cost moved into +startup where the fixed-work number would hide it). + +`cost ÷ duty` is computed against the **configured** duty, and that is a choice, not arithmetic. +It is the number the thresholds name ("Tier A at 10% duty"), and it is the conservative one: the +achieved duty can only overshoot the configured one (the burst timer's granularity rounds bursts +up, never down), so dividing by the configured value can only make the ratio **larger** — +pessimistic for Tier A, never flattering. The achieved duty is reported beside it and separately +asserted to be inside the derived bound. + +The four clauses are evaluated **independently** and every one that fires is recorded, because +more than one can fire and the combination is itself information. The verdict is then resolved by +severity, stated explicitly rather than left to the order the code happens to be written in: + +``` +TIER_A_UNSHIPPABLE > TIER_A_DEEP_DIVE_ONLY > TIER_A_SHIPS_AT_A_SMALLER_DUTY > TIER_A_SHIPS_OPT_IN +``` + +| clause id | fires when | verdict it argues for | +| --- | --- | --- | +| `tier-b-cost-over-5pct` | tier B cost > 5% | `TIER_B_NOT_ALWAYS_ON` | +| `tier-a-10pct-duty-within-5pct-and-ratio-within-2` | headline duty within both bars | `TIER_A_SHIPS_OPT_IN` | +| `tier-a-cost-over-duty-above-2-at-every-duty` | ratio > 2 at every duty tested | `TIER_A_DEEP_DIVE_ONLY` | +| `tier-a-2.5pct-duty-cost-over-5pct` | lowest duty still over 5% | `TIER_A_UNSHIPPABLE` | + +**Two verdicts the plan does not name, and why they exist rather than a silent pick:** + +- `TIER_A_SHIPS_AT_A_SMALLER_DUTY` — the headline duty is over the wall-clock bar but a smaller + one is within **both** bars. Duty-cycling works here; it just needs turning down. The harness + names the largest qualifying duty. Promoting this to `SHIPS_OPT_IN` would claim a duty that was + measured over budget; demoting it to `DEEP_DIVE_ONLY` would discard a tier that is demonstrably + tunable. +- `TIER_A_INDETERMINATE` — nothing qualifies at any duty and neither harsh clause fired, which + means cost does not fall with duty the way serialization says it must, i.e. the measurement is + unsound. It prints *"do not read this as a pass"*. It is a named outcome for the same reason + `gpu_serialized` has an `"unknown"` that must never degrade to `"false"`. + +### A finding about the thresholds themselves + +`cost ÷ duty > 2` is the same statement as `cost% > 200 × duty`. At **2.5% duty that is exactly +`cost% > 5%`** — the wall-clock bar. So on the plan's own three duties, clause 3 **strictly +implies** clause 4, and `TIER_A_DEEP_DIVE_ONLY` is unreachable: the unshippable verdict always +outranks it. The two clauses separate only below 2.5% duty. + +This is a property of the thresholds, not of the hardware, and it is not something I resolved by +adjusting a threshold. The harness applies the clauses as written, and **prints the coincidence +whenever both fire**, so a controller reading a result does not conclude the deep-dive branch was +considered and rejected: + +> `NOTE both harsh clauses fired, and at the lowest duty tested (2.50%) they are the SAME +> condition: cost/duty > 2.0 means cost > 5.00%, and the wall bar is 5.0%. They separate only +> below 2.50% duty, which this table does not test — so TIER_A_DEEP_DIVE_ONLY could not have been +> reached here whatever the numbers. Add a lower-duty arm if that distinction is wanted.` + +The reassuring half is pinned too: on the plan's duties the decision is **total**. At 2.5% duty +"within the wall bar" and "within the ratio bar" are the same condition, so the lowest-duty arm +either qualifies for both — giving opt-in or a smaller duty — or fails both, giving unshippable. +`TestOnThePlansDutiesTheDecisionIsTotal` sweeps eight cost tables and asserts none of them reaches +`INDETERMINATE`. + +Whether to add a 1%-duty arm to make clause 3 distinguishable is a decision for whoever runs this. +I did not add one: the plan's arm table is what was pre-committed, and quietly extending it would +be the same category of move as quietly moving a bar. + +--- + +## The exact command for the RTX 3090 + +```bash +cd /home/diego/github/perf-agent # on the branch feat/pc-sampling-overhead +export CGO_CFLAGS="-I /usr/include/bpf -I /usr/include/pcap -I /home/diego/github/blazesym/capi/include" +export CGO_LDFLAGS="-L/home/diego/github/blazesym/target/release -lblazesym_c" +export LD_LIBRARY_PATH=/home/diego/github/blazesym/target/release + +make -C shim nvidia nvidia-concurrent +make bench-build +sudo setcap cap_bpf,cap_perfmon,cap_checkpoint_restore+ep ./bench/cmd/scenario/scenario + +make bench-gpu-pc-overhead +``` + +Equivalently, by hand: + +```bash +./bench/cmd/scenario/scenario --scenario gpu-pc-overhead --runs 5 \ + --out bench-gpu-pc-overhead.json +``` + +`cap_sys_admin` is **not** in that set and must not be added: the capability constraint is +`cap_bpf,cap_perfmon,cap_checkpoint_restore`, and `getcap` on the binary showing no +`cap_sys_admin` is the standing Phase 1 assertion. Do not put the binary in `/tmp` — that mount is +`nosuid` and file caps do not survive exec. + +**Expect roughly 26 fixed-work runs** (one calibration + 5 arms × 5 rounds). At the default sizing +each is ~20–30 s plus attach and drain, so budget 15–25 minutes. + +**If the calibration pass says the work is mis-sized**, retune and re-run — the failure message +names the flags. `--gpu-rounds` sets kernel duration (the dependency chain length); +`--gpu-iters` sets how many. If the baseline arm fails the concurrency floor, raise +`--gpu-streams` or lower `--gpu-blocks` so each kernel occupies less of the device and more of +them co-reside. + +**Reading the result:** the table, the ratios, the `THRESHOLD FIRED` lines and the two `DECISION` +lines are printed to stdout and the whole thing is in the JSON under `gpu_pc_overhead`. Exit 0 +means the measurement completed — *whatever the verdict*; an honest `TIER_A_UNSHIPPABLE` is a +successful run of this benchmark. Exit 3 means an arm could not prove what it measured, and then +the numbers are **not** a tier decision. + +When the numbers exist, record them in +`docs/superpowers/plans/2026-08-25-gpu-pc-sampling.md` under Task 12 and apply the Tier A verdict +to Task 11's defaults, per the phase gate. + +--- + +## Verification actually run + +``` +go build ./... && go vet ./... clean +go test ./... -count=1 all pass +~/go/bin/golangci-lint run --timeout=5m 0 issues +make -C shim nvidia-concurrent exit 0 (nvcc 13.3, sm_86) +./bench/cmd/scenario/scenario --scenario gpu-pc-overhead + BENCH_SKIPPED: missing required capabilities (CAP_BPF, CAP_PERFMON, + CAP_CHECKPOINT_RESTORE); run: sudo setcap ... + exit 0, no output file written +``` + +The skip is clean: exit 0, one `BENCH_SKIPPED` line, and **no JSON file** — a partial file with no +numbers in it is the thing most likely to be mistaken for a result. The other scenarios' skip +behaviour is unchanged (`--scenario pid-large` still reports the larger capability set). + +The capability message names `CAP_BPF, CAP_PERFMON, CAP_CHECKPOINT_RESTORE` and deliberately not +`CAP_SYS_ADMIN`: this scenario needs gpuprobe's own set, and checking for the larger one that +`bench/cmd/scenario`'s other scenarios need would skip on a correctly-capped machine — a skip for +the wrong reason being indistinguishable from a skip for the right one. + +Static checks on the workload, offline: + +``` +cuobjdump -sass → 2 FFMA, each inside its own backward branch (two rolled loops + of one dependent FMA; nothing folded, nothing unrolled) +cuobjdump -elf → debug_line sections present (-lineinfo took) +argument validation: unknown flag, --rounds past 2^23, and a zero dimension all + exit 2 with the reason, before any CUDA call +``` + +New tests, all offline, all in `bench/cmd/scenario/gpupc_test.go` (45 cases): + +- **the four clauses**, each fired and each *not* fired, including exactly-at-the-bar (the plan + says `> 5%`, so 5.0 does not fire — a strict-versus-non-strict slip is the classic way a + pre-committed threshold quietly becomes a different one); +- the every-duty ratio clause is **universal, not headline** — one arm inside the bar stops it; +- the harsher verdict wins when clauses co-fire, and the coincidence is reported; +- the decision is total on the plan's duties; +- the recorded bars are the committed ones; +- **the medians and the ratio**: cost against the baseline, ratio against the configured duty, the + baseline not a cost against itself, and `medianFloat` not reordering its input; +- **the parsers**: a Tier A report, a Tier B report, an off report, and — the one that matters — + a producer that printed nothing leaves the tier **empty**, never `"off"`, because "the adapter + was loaded with sampling off" and "the adapter was never loaded" must not look the same; +- **the arm assertions**, each shown red: an arm that measured nothing, a baseline that was + sampling, a Tier A arm that serialized nothing, one with too few bursts, one whose windows do + not reconcile with its bursts, one at the wrong duty (and one legitimately overshooting, which + must pass), one in a CUDA-graph process, a Tier B arm that emitted a window, a Tier B arm that + sampled nothing; +- **the workload guards**: a serial baseline and a trivial-kernel baseline both refused, a + realistic one accepted; three Tier A arms with equal burst counts refused; +- **the concurrency measurement**: fully overlapping kernels → 4, serial kernels → 1, an empty + snapshot → 0 rather than NaN (NaN would sail straight past the floor comparison), degenerate + intervals ignored; +- **the arm table itself**: the duties are the plan's, and each gap equals the adapter's own + `burst × (1/max_duty − 1)` — if either side ever changes, this fails; +- **the environment**: every arm names its tier explicitly, the Tier A arms pin the gap from both + sides, and the two non-Tier-A arms carry no burst environment at all; +- **all four skip branches**, reachable via an injected-predicate variant, since only the first + can ever fire on this machine. + +--- + +## Cannot verify — every number is outstanding + +`CapEff: 0`, no GPU on this machine. **No arm of this benchmark has ever been executed.** No CUDA +kernel in `cuda_concurrent.cu` has ever run; it has been compiled and disassembled, nothing more. + +**The tier decision is not yet made, and no part of it may be made from this file.** + +Outstanding, all of it: + +1. **Every number in the arm table.** Baseline wall-clock, Tier B's cost, Tier A's cost at 10%, 5% + and 2.5% duty, every achieved kernel throughput, every cost ÷ duty ratio. None exists. +2. **Which threshold fires, and therefore whether Tier A ships at all**, and whether Tier B + remains a candidate for always-on. +3. **Whether the workload is actually concurrent on a 3090.** The design argues 16 blocks × 256 + threads leaves room for four kernels to co-reside on 82 SMs, and the SASS confirms the + dependency chain is real, but the achieved concurrency is a hardware measurement. The harness + asserts a floor rather than assuming it, so a negative answer is a loud failure with retuning + advice and not a quiet underestimate — but the answer is unknown. +4. **Whether the default sizing lands in the 10–120 s window** on a 3090. `--gpu-rounds 64000` + was estimated from a dependency-chain latency guess, not measured. The calibration pass exists + because this is unknown. +5. **Whether pinning `MAX_DUTY_PERMILLE` and `MAX_GAP_MS` to the same gap really pins the burst + controller.** It follows from `burst_next_gap_ns`'s clamp read as source, and + `core/burst_test.cc` proves the clamp against a fake clock, but this exact combination has + never been given to the adapter on hardware. The achieved-duty assertion is what would catch + it. +6. **Whether the adapter's report line parses as expected on a real run.** The parser was written + against the `logf` format strings in `cupti_adapter.cc` and tested against hand-written + fixtures reproducing them. A format drift would show up as an empty producer tier, which is a + failing assertion rather than a silent zero — but it has not been read off a real process. +7. **Whether 25 sequential `gpuprobe.Attach`/`Close` cycles are clean** — the enrollment + rendezvous binding and unbinding twenty-five times in one process has not been exercised. +8. **Whether the baseline arm is stable enough for a 5% bar to be meaningful.** The spread across + the five runs is recorded per arm for exactly this reason: a median whose spread is larger than + the effect is not a measurement. Whether that holds on the lab machine is unknown. +9. Everything Tasks 6, 10 and 11 could not verify remains unverified and is unchanged by this + task. In particular Task 10's item 2 — whether the sampling windows and the converted activity + timestamps really land in the same clock domain closely enough for the intersection to mean + what it says — is load-bearing *here* as well: the `ExecutionsSerialized > 0` assertion and the + serialized/not-serialized split in the arm evidence both rest on it. + +Not verifiable at all, on hardware or otherwise: MPS and cross-process contention for the sampling +hardware, which would perturb any of these arms invisibly. diff --git a/Makefile b/Makefile index 5831902f..532e9141 100644 --- a/Makefile +++ b/Makefile @@ -121,3 +121,27 @@ bench-self: bench-build test-workloads --cpu-budget 1.5 --resolution-budget 0.5 \ --out bench-self.json @echo "self-profile bench written to bench-self.json" + +# GPU PC-sampling overhead (plan Task 12): the marginal cost of Tier B and of +# Tier A at three duty fractions, against the shipping Phase 4 configuration +# with PC sampling off. Needs an NVIDIA GPU, the CUPTI adapter and the +# concurrent CUDA workload; reports BENCH_SKIPPED and exits 0 without any of +# them. +# +# The capability set is gpuprobe's own and is SMALLER than the one +# bench-scenarios needs. cap_sys_admin is deliberately not in it. +# +# Exit codes: 0 when the measurement completed (whatever the verdict — an +# honest TIER_A_UNSHIPPABLE is a successful run of this benchmark), 3 when an +# arm could not prove it ran in the mode it claims, in which case the numbers +# are not a tier decision and must not be recorded as one. +.PHONY: bench-gpu-pc-overhead +bench-gpu-pc-overhead: bench-build + @$(MAKE) -C shim nvidia nvidia-concurrent + @if ! getcap ./bench/cmd/scenario/scenario | grep -q cap_bpf; then \ + echo "*** scenario binary missing caps; run: sudo setcap cap_bpf,cap_perfmon,cap_checkpoint_restore+ep ./bench/cmd/scenario/scenario"; \ + exit 1; \ + fi + ./bench/cmd/scenario/scenario --scenario gpu-pc-overhead --runs 5 \ + --out bench-gpu-pc-overhead.json + @echo "gpu pc-sampling overhead written to bench-gpu-pc-overhead.json" diff --git a/bench/README.md b/bench/README.md index 9dae03e6..f8c12f06 100644 --- a/bench/README.md +++ b/bench/README.md @@ -19,6 +19,51 @@ Two-layer benchmark for `--unwind dwarf` startup cost. Companion to - `system-wide-mixed` — N processes across Go/Python/Rust/Node from `test/workloads/`, attached via `-a`. Measures `/proc/*` walk + per-PID maps parse + per-distinct-binary compile. +- `gpu-pc-overhead` — the marginal cost of GPU PC sampling, and the + pre-committed thresholds that turn it into a decision. See below. + +## `gpu-pc-overhead` + +Plan Task 12. Measures the **marginal** cost of PC sampling: the baseline arm +is the shipping Phase 4 configuration (shim injected, RUNTIME + RESOURCE +callbacks, `CONCURRENT_KERNEL` activity, 100 ms drain, consumer attached) with +PC sampling **off**, not an uninjected run. Spec §9.1 already measured +injection and the activity path, and those costs are paid either way. + +Five arms, five interleaved runs each, medians, fixed work rather than fixed +time: + +| arm | duty | +| --- | --- | +| baseline (PC sampling off) | — | +| Tier B continuous | — | +| Tier A 50 ms / 450 ms | 10% | +| Tier A 50 ms / 950 ms | 5% | +| Tier A 50 ms / 1950 ms | 2.5% | + +The workload is `shim/nvidia/testdata/cuda_concurrent.cu`: several streams, +non-trivial kernel durations, genuine overlap. It is a **second** workload, not +a change to `cuda_workload.cu` — the serial fixture the adapter and the phase +gate are proven against gives serialization nothing to destroy, so measuring +Tier A on it would understate its cost. The harness measures the achieved +concurrency out of the profile and **fails** if the baseline arm is near +serial or its kernels are microseconds long. + +Every arm proves it ran in the mode it claims, from both ends independently — +the adapter's own report line on stderr and the consumer's counters — and a +mismatch fails the run rather than contributing a number. Cross-arm, a lower +duty must open strictly fewer bursts, or the three Tier A arms are one arm +under three names. + +```bash +make -C shim nvidia nvidia-concurrent +make bench-build +sudo setcap cap_bpf,cap_perfmon,cap_checkpoint_restore+ep ./bench/cmd/scenario/scenario +make bench-gpu-pc-overhead +``` + +Exit codes: `0` when the measurement completed (whatever the verdict), `3` +when an arm could not prove what it measured. ## First-time setup @@ -55,13 +100,25 @@ The aggregator (`bench/cmd/report/`) reads JSON and produces markdown. ## Flags `bench/cmd/scenario`: -- `--scenario pid-large | system-wide-mixed` (required) +- `--scenario pid-large | system-wide-mixed | self | gpu-pc-overhead` (required) - `--processes N` (default 30) — fleet size for system-wide - `--runs N` (default 5) — iterations - `--drop-cache` (default off) — drop page cache between runs (warm-cache by default) - `--out PATH` — JSON output path - `--workloads-dir PATH` — auto-detected if not set +`gpu-pc-overhead` only: +- `--gpu-shim PATH` / `--gpu-workload PATH` — the adapter and the concurrent workload +- `--gpu-iters N` / `--gpu-rounds N` — the fixed work; the calibration pass says + when they need retuning for the device in front of you +- `--gpu-streams N` (4) / `--gpu-blocks N` (16) / `--gpu-threads N` (256) — + the concurrency; fewer blocks means more kernels co-reside +- `--gpu-sync-every N` (4) — device sync cadence; forces concurrency to refill +- `--gpu-min-concurrency` (1.5) / `--gpu-min-kernel-us` (50) — the guards that + refuse to report numbers from a microbenchmark +- `--gpu-min-bursts` (4) — the floor on bursts per Tier A arm +- `--gpu-min-calibration-sec` (10) / `--gpu-max-calibration-sec` (120) + `bench/cmd/report`: - `--in PATH` (repeatable) — summary mode - `--diff A.json --diff B.json` — diff mode diff --git a/bench/cmd/scenario/gpupc.go b/bench/cmd/scenario/gpupc.go new file mode 100644 index 00000000..3dfe2d81 --- /dev/null +++ b/bench/cmd/scenario/gpupc.go @@ -0,0 +1,1245 @@ +package main + +// The "gpu-pc-overhead" scenario: the marginal cost of GPU PC sampling, and +// the pre-committed thresholds that turn it into a decision. +// +// This file is the harness for plan Task 12. It cannot be completed on a +// machine with no GPU: what it produces here is the instrument, and what it +// produces on an RTX 3090 is the number that decides whether Tier A ships. +// +// What is measured, and against what +// ---------------------------------- +// The baseline arm is NOT "no injection". Spec §9.1 already measured +// injection at 0.0% and the activity path at −10.0%, and those costs are paid +// whether or not PC sampling is on. The question here is strictly the +// MARGINAL cost of PC sampling, so the baseline is the shipping Phase 4 +// configuration -- shim injected, RUNTIME + RESOURCE callbacks, +// CONCURRENT_KERNEL activity, drain at 100 ms, consumer attached -- with PC +// sampling OFF. An uninjected run is taken too, before the arms, but it is +// labelled "calibration" and is never used as the baseline. +// +// Why a concurrent workload and not the 393k launches/s ceiling +// ------------------------------------------------------------- +// §9.1's ceiling is the wrong instrument for this question. It exaggerates +// per-launch costs and UNDERSTATES serialization costs, because serialization +// hurts in proportion to the concurrency it destroys and a stream of trivial +// kernels has almost none. shim/nvidia/testdata/cuda_concurrent.cu exists for +// this measurement: several streams, non-trivial kernel durations, genuine +// overlap. See its header for why it is a second workload rather than a +// change to cuda_workload.cu. +// +// The concurrency is not assumed. It is measured out of the profile +// (sum of exec durations over the span they cover) and the run FAILS if the +// baseline arm's value is near 1 -- a workload that turned out to be serial +// would produce a small, green, meaningless Tier A cost, which is the same +// defect as a counter reading green when things are worst. +// +// Why every arm has to prove which mode it ran in +// ----------------------------------------------- +// The benchmark-shaped instance of this project's standing defect is an arm +// that did not actually enable the tier it claims. Such an arm reports a +// wonderfully small overhead. So every arm asserts, from BOTH ends +// independently -- the producer's own report line on stderr and the +// consumer's counters -- that it ran in the mode it says, and a mismatch +// fails the whole run rather than contributing a number. + +import ( + "bufio" + "bytes" + "context" + "errors" + "fmt" + "io" + "math" + "os" + "os/exec" + "path/filepath" + "sort" + "strconv" + "strings" + "time" + + "kernel.org/pub/linux/libs/security/libcap/cap" + + "github.com/dpsoft/perf-agent/bench/internal/schema" + "github.com/dpsoft/perf-agent/gpu" + "github.com/dpsoft/perf-agent/gpuprobe" +) + +// --------------------------------------------------------------------------- +// The pre-committed thresholds and the verdicts they produce. +// +// These numbers are from the plan and were fixed before any data existed. A +// threshold decided after seeing the data is not a threshold, which is why +// they are constants here and are echoed into the output JSON: a reader can +// see they were not adjusted to fit. +// --------------------------------------------------------------------------- + +const ( + // gpuPCMaxWallPercent: "> 5% wall-clock on a realistic workload". + gpuPCMaxWallPercent = 5.0 + // gpuPCMaxCostOverDuty: "cost ÷ duty ≤ 2". + gpuPCMaxCostOverDuty = 2.0 +) + +// The verdicts, as stable identifiers rather than prose, so the outcome is a +// decision a script can read and not a paragraph a reader can reinterpret. +const ( + verdictTierBAlwaysOnCandidate = "TIER_B_ALWAYS_ON_CANDIDATE" + verdictTierBExplicitOnly = "TIER_B_NOT_ALWAYS_ON" + + verdictTierAOptIn = "TIER_A_SHIPS_OPT_IN" + verdictTierASmallerDuty = "TIER_A_SHIPS_AT_A_SMALLER_DUTY" + verdictTierADeepDiveOnly = "TIER_A_DEEP_DIVE_ONLY" + verdictTierAUnshippable = "TIER_A_UNSHIPPABLE" + // verdictTierAIndeterminate is reachable only from a shape the plan's + // four clauses do not cover -- for instance a cost that does not + // increase with duty, which would mean the measurement itself is + // unsound. It is a named outcome rather than a silent fallthrough to + // the friendliest verdict, for the same reason gpu_serialized has an + // "unknown" that must never degrade to "false". + verdictTierAIndeterminate = "TIER_A_INDETERMINATE" +) + +// The threshold clause identifiers, in the plan's own order. +const ( + clauseTierBOverBudget = "tier-b-cost-over-5pct" + clauseTierAHeadlineWithin = "tier-a-10pct-duty-within-5pct-and-ratio-within-2" + clauseTierARatioOverAtAll = "tier-a-cost-over-duty-above-2-at-every-duty" + clauseTierASmallestOverBud = "tier-a-2.5pct-duty-cost-over-5pct" +) + +// gpuPCArmSpec is one arm's configuration. The three Tier A arms differ only +// in their gap, which is what makes the duty the single moving variable. +type gpuPCArmSpec struct { + Name string + Tier gpu.PCSamplingTier + BurstMs int + GapMs int +} + +// gpuPCArms is the arm table from the plan, in the order it runs them: +// baseline first (everything else is measured against it), then Tier B, then +// Tier A at decreasing duty. +// +// The Tier A gaps are not free parameters. The adapter derives its minimum gap +// from the duty ceiling -- min_gap = burst * (1/max_duty - 1) -- so 450 / 950 +// / 1950 ms after a 50 ms burst ARE 10% / 5% / 2.5%, and setting the maximum +// gap to the same value pins the burst controller's closed loop to exactly +// that gap instead of letting it tune. See gpuPCArmEnv. +var gpuPCArms = []gpuPCArmSpec{ + {Name: "baseline (pc sampling off)", Tier: gpu.PCSamplingOff}, + {Name: "tier B continuous", Tier: gpu.PCSamplingContinuous}, + {Name: "tier A 50ms/450ms (10% duty)", Tier: gpu.PCSamplingSerialized, BurstMs: 50, GapMs: 450}, + {Name: "tier A 50ms/950ms (5% duty)", Tier: gpu.PCSamplingSerialized, BurstMs: 50, GapMs: 950}, + {Name: "tier A 50ms/1950ms (2.5% duty)", Tier: gpu.PCSamplingSerialized, BurstMs: 50, GapMs: 1950}, +} + +// dutyConfigured is burst/(burst+gap): the fraction of wall-clock the arm asks +// to spend with kernels serialized. Zero for the two non-Tier-A arms. +func (a gpuPCArmSpec) dutyConfigured() float64 { + if a.Tier != gpu.PCSamplingSerialized || a.BurstMs+a.GapMs == 0 { + return 0 + } + return float64(a.BurstMs) / float64(a.BurstMs+a.GapMs) +} + +// gpuPCConfig is everything the scenario takes from the command line. +type gpuPCConfig struct { + ShimPath string + WorkloadPath string + Runs int + + Iters int + Warmup int + Streams int + Rounds int + Blocks int + Threads int + SyncEvery int + + // MinConcurrency and MinKernelUs are the two guards that keep this + // from silently measuring a microbenchmark. Both are asserted against + // the BASELINE arm, which is the only arm whose concurrency is + // supposed to be undisturbed. + MinConcurrency float64 + MinKernelUs float64 + // MinBursts is how many bursts the lowest-duty Tier A arm must have + // opened for its duty to mean anything. At 2.5% the cycle is 2 s, so + // this is also the real floor on how long the fixed work must take. + MinBursts uint64 + // MinCalibrationSec / MaxCalibrationSec bound the uninjected run, so + // a workload sized far too small or far too large announces itself + // with the retuning advice instead of producing noise. + MinCalibrationSec float64 + MaxCalibrationSec float64 +} + +// --------------------------------------------------------------------------- +// The skip path: the only thing this scenario can prove without hardware. +// --------------------------------------------------------------------------- + +// gpuPCSkipReason returns the reason this scenario cannot run here, or "" if +// it can. Every reason is a SKIP and not a failure: a machine with no GPU is +// not a broken machine, and a benchmark that failed there would be turned off +// in CI and would then never run anywhere. +// +// The capability set checked is gpuprobe's own -- CAP_BPF, CAP_PERFMON, +// CAP_CHECKPOINT_RESTORE -- and deliberately NOT the larger set +// bench/cmd/scenario's other scenarios need. Checking for CAP_SYS_ADMIN here +// would skip on a correctly-capped machine, and a skip for the wrong reason is +// indistinguishable from a skip for the right one. +func gpuPCSkipReason(cfg gpuPCConfig) string { + return gpuPCSkipReasonWith(hasGPUCaps(), hasNVIDIADevice(), cfg) +} + +// gpuPCSkipReasonWith is gpuPCSkipReason with the two environment probes +// injected, so all four skip branches are reachable from a unit test on a +// machine that can only ever produce the first. A skip path that cannot be +// exercised is a skip path nobody has read. +func gpuPCSkipReasonWith(caps, gpuPresent bool, cfg gpuPCConfig) string { + if !caps { + return "missing required capabilities (CAP_BPF, CAP_PERFMON, CAP_CHECKPOINT_RESTORE); " + + "run: sudo setcap cap_bpf,cap_perfmon,cap_checkpoint_restore+ep ./bench/cmd/scenario/scenario " + + "(not from /tmp: it is nosuid and file caps do not survive exec)" + } + if !gpuPresent { + return "no NVIDIA GPU on this machine (/dev/nvidiactl absent); " + + "this scenario measures CUPTI PC sampling and has nothing to measure without one" + } + if _, err := os.Stat(cfg.ShimPath); err != nil { + return fmt.Sprintf("CUPTI adapter not built at %s (%v); build it with: make -C shim nvidia", + cfg.ShimPath, err) + } + if _, err := os.Stat(cfg.WorkloadPath); err != nil { + return fmt.Sprintf("concurrent CUDA workload not built at %s (%v); "+ + "build it with: make -C shim nvidia-concurrent", cfg.WorkloadPath, err) + } + return "" +} + +// hasGPUCaps mirrors gpuprobe/gate_test.go's hasCaps: Permitted as well as +// Effective, because a setcap'd binary has not promoted Permitted yet, and +// never a bare Geteuid check. +func hasGPUCaps() bool { + if os.Geteuid() == 0 { + return true + } + set := cap.GetProc() + if set == nil { + return false + } + for _, w := range []cap.Value{cap.BPF, cap.PERFMON, cap.CHECKPOINT_RESTORE} { + ok := false + for _, flag := range []cap.Flag{cap.Permitted, cap.Effective} { + if have, err := set.GetFlag(flag, w); err == nil && have { + ok = true + break + } + } + if !ok { + return false + } + } + return true +} + +// hasNVIDIADevice is the cheapest honest test for "a CUDA context can be +// created here". It does not shell out to nvidia-smi: this runs on a machine +// where the answer is no, and a missing binary and a missing GPU must not be +// reported as the same thing. +func hasNVIDIADevice() bool { + if _, err := os.Stat("/dev/nvidiactl"); err == nil { + return true + } + if ents, err := os.ReadDir("/proc/driver/nvidia/gpus"); err == nil && len(ents) > 0 { + return true + } + return false +} + +// --------------------------------------------------------------------------- +// Running the arms. +// --------------------------------------------------------------------------- + +// runGPUPCOverhead runs the calibration pass and then the five arms, +// interleaved, five runs each, and fills doc.GPUPC. It returns false when any +// assertion failed; main exits non-zero on that, because a benchmark that +// could not prove what it measured must not look like one that did. +func runGPUPCOverhead(doc *schema.Document, cfg gpuPCConfig, out io.Writer) bool { + kernels := cfg.Iters * cfg.Streams + res := &schema.GPUPCOverhead{ + Workload: schema.GPUPCWorkload{ + Path: cfg.WorkloadPath, Iters: cfg.Iters, Warmup: cfg.Warmup, + Streams: cfg.Streams, Rounds: cfg.Rounds, Blocks: cfg.Blocks, + Threads: cfg.Threads, SyncEvery: cfg.SyncEvery, KernelsRun: kernels, + }, + } + doc.GPUPC = res + doc.Config.Runs = cfg.Runs + + // The calibration pass: the same fixed work with NO adapter injected. + // It warms the device clocks before the first arm and it proves the + // workload is sized sanely. It is not an arm and is never the + // baseline -- see this file's header. + calRun, err := runConcurrentWorkload(cfg, nil) + if err != nil { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: calibration run failed: %v\n", err) + return false + } + calRun.RunN = 0 + res.Calibration = schema.GPUPCArm{ + Name: "calibration (uninjected; NOT the baseline)", Tier: "none", + Runs: []schema.GPUPCRun{calRun}, MedianWallMs: calRun.WallMs, + MedianKernelsPerS: calRun.KernelsPerS, + } + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: calibration (uninjected, not an arm): "+ + "%.1f ms for %d kernels, %.0f kernels/s\n", calRun.WallMs, kernels, calRun.KernelsPerS) + if sec := calRun.WallMs / 1000; sec < cfg.MinCalibrationSec || sec > cfg.MaxCalibrationSec { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: FAILED: the fixed work takes %.1f s uninjected, "+ + "outside the sane window [%.0f s, %.0f s]. Retune with --gpu-rounds (kernel "+ + "duration) and --gpu-iters (how many), then re-run. Too short and the 2.5%%-duty "+ + "arm cannot open enough bursts for its duty to mean anything; too long and the "+ + "five interleaved runs take longer than the operator will wait.\n", + sec, cfg.MinCalibrationSec, cfg.MaxCalibrationSec) + return false + } + + arms := make([]schema.GPUPCArm, len(gpuPCArms)) + for i, spec := range gpuPCArms { + arms[i] = schema.GPUPCArm{ + Name: spec.Name, Tier: spec.Tier.String(), + BurstMs: spec.BurstMs, GapMs: spec.GapMs, + DutyConfigured: spec.dutyConfigured(), + } + } + + // INTERLEAVED, per §9.1's method: every arm once per round, rather + // than five of one arm then five of the next. Thermal drift, clock + // boost state and any other slow drift then fall on all five arms + // alike instead of on whichever ran last. + ok := true + for run := 1; run <= cfg.Runs; run++ { + for i, spec := range gpuPCArms { + r, err := measureGPUPCArm(cfg, spec) + if err != nil { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: FAILED: arm %q run %d: %v\n", spec.Name, run, err) + return false + } + r.RunN = run + if err := assertArmRanInItsMode(cfg, spec, r); err != nil { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: FAILED: arm %q run %d did not prove it "+ + "ran in the mode it claims: %v\n", spec.Name, run, err) + _, _ = fmt.Fprintf(out, " evidence: %+v\n", r.Evidence) + ok = false + } + arms[i].Runs = append(arms[i].Runs, r) + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: run %d/%d %-32s %8.1f ms %7.0f kern/s "+ + "conc %.2f kern %.0f us\n", run, cfg.Runs, spec.Name, r.WallMs, + r.KernelsPerS, r.Concurrency, r.MeanKernelUs) + } + // Fail after a COMPLETE round rather than at the first bad arm: + // one round gives every arm's evidence side by side, which is + // the useful diagnostic, and four more rounds of a + // known-unsound configuration is ten minutes of GPU time that + // tells nobody anything. + if !ok { + break + } + } + + // Recorded whether or not the assertions held. The arms are the + // evidence of what went wrong; what is withheld on failure is the + // DECISION, which stays at its zero value so nothing in the file can + // be read as a verdict. + summarizeGPUPCArms(arms) + res.Arms = arms + if !ok { + return false + } + + if err := assertBaselineIsRealistic(cfg, arms[0]); err != nil { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: FAILED: %v\n", err) + return false + } + if err := assertDutyKnobDidSomething(arms); err != nil { + _, _ = fmt.Fprintf(out, "gpu-pc-overhead: FAILED: %v\n", err) + return false + } + res.Decision = decideGPUPC(arms, gpuPCMaxWallPercent, gpuPCMaxCostOverDuty) + + renderGPUPCTable(out, res) + for _, line := range res.Decision.Lines { + _, _ = fmt.Fprintln(out, line) + } + return true +} + +// measureGPUPCArm attaches a consumer configured for the arm's tier, runs the +// fixed work under the adapter with the arm's environment, and collects both +// the timing and the evidence. +// +// A fresh attach per arm-run is deliberate. The tier is a property of the +// Timeline as well as of the producer, so one long-lived consumer could not +// carry five tiers; and per-run counters that start at zero are what make the +// evidence assertions exact rather than differential. +func measureGPUPCArm(cfg gpuPCConfig, spec gpuPCArmSpec) (schema.GPUPCRun, error) { + timeline := gpu.NewTimeline(gpu.TimelineConfig{PCSampling: spec.Tier}) + + // No symbolizer, on purpose, and identically on every arm. Resolving a + // sampled launch's stack is agent-side work that costs the WORKLOAD + // nothing, and it cannot succeed anyway once the workload has exited + // (its /proc//maps is gone). Leaving it out removes a variance + // source without removing any cost the arms differ in: the producer + // still samples launches and still captures stacks in BPF, on every + // arm, exactly as the shipping configuration does. + c, err := gpuprobe.Attach(gpuprobe.Config{ + ShimPath: cfg.ShimPath, + PID: 0, // the process that will map the adapter does not exist yet + Backend: gpu.BackendCUPTI, + Sink: timeline, + }) + if err != nil { + return schema.GPUPCRun{}, fmt.Errorf("attach: %w", err) + } + defer func() { _ = c.Close() }() + + ctx, cancel := context.WithCancel(context.Background()) + done := make(chan struct{}) + var runErr error + go func() { + defer close(done) + if err := c.Run(ctx); err != nil && !errors.Is(err, context.Canceled) { + runErr = err + } + }() + + run, err := runConcurrentWorkload(cfg, gpuPCArmEnv(cfg, spec)) + if err != nil { + cancel() + <-done + return schema.GPUPCRun{}, err + } + + // The adapter's atexit handler has already run cuptiActivityFlushAll + // and flushed both batches, so everything is in the ringbuf; this is + // for the consumer goroutine to drain the tail. + time.Sleep(500 * time.Millisecond) + cancel() + <-done + if runErr != nil { + return schema.GPUPCRun{}, fmt.Errorf("consumer: %w", runErr) + } + c.Flush() + + snap := timeline.Snapshot() + st := c.Stats() + run.Concurrency, run.MeanKernelUs = concurrencyOf(snap) + run.Evidence.PCSamplesDecoded = st.PCSamplesDecoded + run.Evidence.SamplingWindowsDecoded = st.SamplingWindowsDecoded + run.Evidence.SamplingWindowsReceived = snap.SamplingWindowsReceived + run.Evidence.ExecutionsSeen = len(snap.Executions) + run.Evidence.ExecutionsSerialized = snap.ExecutionsSerialized + run.Evidence.ExecutionsNotSerialized = snap.ExecutionsNotSerialized + run.Evidence.ExecutionsUnknown = snap.ExecutionsSerializationUnknown + run.Evidence.SnapshotTier = snap.PCSampling.String() + return run, nil +} + +// gpuPCArmEnv builds the producer environment for an arm. +// +// The Tier A gap is pinned rather than tuned, and that is the whole reason +// both duty knobs are set. The burst controller clamps its gap into +// [min_gap, max_gap] where min_gap = burst * (1/max_duty - 1). Setting +// max_duty so that min_gap is the arm's gap AND setting max_gap to the same +// value collapses the interval to a point: whatever the closed loop computes, +// the gap it returns is the arm's gap. Without the second knob a workload +// producing a high pair rate would have the loop lengthen the gap, the arm +// would run at a duty nobody asked for, and the cost-over-duty ratio would be +// computed against a denominator that never happened. +func gpuPCArmEnv(cfg gpuPCConfig, spec gpuPCArmSpec) []string { + env := []string{ + "CUDA_INJECTION64_PATH=" + cfg.ShimPath, + "PERFAGENT_GPU_LOG=stderr", + // Set EXPLICITLY on every arm including the off one, never left + // to be inherited: an exported PERFAGENT_GPU_PC_SAMPLING in the + // operator's shell must not turn the baseline arm into a + // serializing one, which would make every other arm look free. + gpu.PCSamplingEnvVar + "=" + spec.Tier.EnvValue(), + } + if spec.Tier != gpu.PCSamplingSerialized { + return env + } + // burst * (1/duty - 1) == gap => duty == burst / (burst + gap). + permille := int(math.Round(1000 * float64(spec.BurstMs) / float64(spec.BurstMs+spec.GapMs))) + return append(env, + fmt.Sprintf("PERFAGENT_GPU_PC_BURST_MS=%d", spec.BurstMs), + fmt.Sprintf("PERFAGENT_GPU_PC_MAX_DUTY_PERMILLE=%d", permille), + fmt.Sprintf("PERFAGENT_GPU_PC_MAX_GAP_MS=%d", spec.GapMs), + ) +} + +// runConcurrentWorkload runs one fixed-work pass. extraEnv nil means an +// UNINJECTED run (the calibration pass): no CUDA_INJECTION64_PATH, so the +// adapter is not loaded at all. +func runConcurrentWorkload(cfg gpuPCConfig, extraEnv []string) (schema.GPUPCRun, error) { + cmd := exec.Command(cfg.WorkloadPath, + fmt.Sprintf("--iters=%d", cfg.Iters), + fmt.Sprintf("--warmup=%d", cfg.Warmup), + fmt.Sprintf("--streams=%d", cfg.Streams), + fmt.Sprintf("--rounds=%d", cfg.Rounds), + fmt.Sprintf("--blocks=%d", cfg.Blocks), + fmt.Sprintf("--threads=%d", cfg.Threads), + fmt.Sprintf("--sync-every=%d", cfg.SyncEvery), + "--linger-ms=0", + ) + // os.Environ() carries whatever the operator exported; every variable + // this scenario cares about is appended AFTER it, and os/exec keeps + // the last occurrence of a duplicate key. + cmd.Env = append(os.Environ(), extraEnv...) + if extraEnv == nil { + // The calibration pass must be genuinely uninjected even if the + // operator has CUDA_INJECTION64_PATH exported. + cmd.Env = append(cmd.Env, "CUDA_INJECTION64_PATH=") + } + var stdout, stderr bytes.Buffer + cmd.Stdout, cmd.Stderr = &stdout, &stderr + + t0 := time.Now() + err := cmd.Run() + processMs := float64(time.Since(t0).Microseconds()) / 1000.0 + if err != nil { + return schema.GPUPCRun{}, fmt.Errorf("workload: %w\nstdout: %s\nstderr: %s", + err, stdout.String(), stderr.String()) + } + + run, err := parseConcurrentLine(stdout.String()) + if err != nil { + return schema.GPUPCRun{}, fmt.Errorf("%w\nstdout: %s", err, stdout.String()) + } + run.ProcessMs = processMs + // Exactly zero, not "small": every kernel is the identity on its + // buffer. A non-zero value means this arm perturbed the computation + // and its timing must not be used. + if run.MaxAbsErr != 0 { + return schema.GPUPCRun{}, fmt.Errorf( + "workload reported max_abs_err=%g, want exactly 0: the computation was "+ + "corrupted, so this run's timing means nothing", run.MaxAbsErr) + } + if extraEnv != nil { + run.Evidence = parseAdapterReport(stderr.String()) + } + return run, nil +} + +// parseConcurrentLine reads cuda_concurrent's single result line. +func parseConcurrentLine(stdout string) (schema.GPUPCRun, error) { + var line string + sc := bufio.NewScanner(strings.NewReader(stdout)) + for sc.Scan() { + if strings.HasPrefix(sc.Text(), "concurrent: iters=") { + line = sc.Text() + } + } + if line == "" { + return schema.GPUPCRun{}, errors.New("workload printed no \"concurrent: iters=...\" result line") + } + kv := newKeyValues(line) + elapsed, ok := kv.float("elapsed_ms") + if !ok { + return schema.GPUPCRun{}, fmt.Errorf("no elapsed_ms in %q", line) + } + kps, _ := kv.float("kernels_per_s") + errv, ok := kv.float("max_abs_err") + if !ok { + return schema.GPUPCRun{}, fmt.Errorf("no max_abs_err in %q", line) + } + return schema.GPUPCRun{WallMs: elapsed, KernelsPerS: kps, MaxAbsErr: errv}, nil +} + +// parseAdapterReport reads the producer's own account of the run off stderr. +// +// This is the producer half of the evidence and it is deliberately NOT +// derived from the consumer's counters: the two ends can disagree, and the gap +// between them is the loss. An arm proved only by the consumer would look +// identical whether the producer ran in the right mode and lost records, or +// ran in the wrong mode and had nothing to send. +func parseAdapterReport(stderr string) schema.GPUPCEvidence { + var ev schema.GPUPCEvidence + sc := bufio.NewScanner(strings.NewReader(stderr)) + sc.Buffer(make([]byte, 0, 64*1024), 1024*1024) + for sc.Scan() { + line := sc.Text() + switch { + // "perfagent-cupti: pc exit tier=serialized ... pc_records=N ..." + case strings.Contains(line, "perfagent-cupti: pc ") && strings.Contains(line, " tier=") && + strings.Contains(line, " pc_records="): + kv := newKeyValues(line) + if s, ok := kv.str("tier"); ok { + ev.ProducerTier = s + } + ev.ProducerPCRecords, _ = kv.uint("pc_records") + ev.ProducerGraphExecs, _ = kv.uint("graph_execs") + // "perfagent-cupti: pc_sampling=off tier_refused=N ..." + case strings.Contains(line, "perfagent-cupti: pc_sampling=off"): + ev.ProducerTier = gpu.PCSamplingNameOff + // "perfagent-cupti: tier A bursts=N burst_ns=N duty=F ..." + case strings.Contains(line, "perfagent-cupti: tier A bursts="): + kv := newKeyValues(line) + ev.ProducerBursts, _ = kv.uint("bursts") + ev.ProducerWindows, _ = kv.uint("windows") + ev.ProducerDuty, _ = kv.float("duty") + ev.ProducerStartFailed, _ = kv.uint("start_failed") + ev.ProducerStopFailed, _ = kv.uint("stop_failed") + ev.ProducerGraphRefuse, _ = kv.uint("graph_refused") + } + } + return ev +} + +// keyValues splits a log line into its key=value tokens. Later occurrences of +// a key win, which matches the log's own "the last word is the final state" +// shape. +type keyValues map[string]string + +func (kv keyValues) str(k string) (string, bool) { v, ok := kv[k]; return v, ok } + +func (kv keyValues) float(k string) (float64, bool) { + s, ok := kv[k] + if !ok { + return 0, false + } + f, err := strconv.ParseFloat(s, 64) + return f, err == nil +} + +func (kv keyValues) uint(k string) (uint64, bool) { + s, ok := kv[k] + if !ok { + return 0, false + } + u, err := strconv.ParseUint(s, 10, 64) + return u, err == nil +} + +func newKeyValues(line string) keyValues { + kv := keyValues{} + for _, tok := range strings.Fields(line) { + if i := strings.IndexByte(tok, '='); i > 0 { + kv[tok[:i]] = tok[i+1:] + } + } + return kv +} + +// concurrencyOf turns a snapshot into the two properties that decide whether +// this workload is the realistic one the plan requires: +// +// concurrency = sum(exec duration) / (max end - min start) +// meanKernelUs = sum(exec duration) / count +// +// Both are computed over the executions the profile RETAINED, which may be a +// suffix of the run if the ring evicted. That is fine and is why the span is +// taken from the retained set rather than from the workload's own clock: both +// numerator and denominator then describe the same interval. +func concurrencyOf(snap gpu.Snapshot) (concurrency, meanKernelUs float64) { + var busy uint64 + var lo, hi uint64 + n := 0 + for i := range snap.Executions { + e := snap.Executions[i].Exec + if e.EndNs <= e.StartNs { + continue // a zero or inverted interval carries no duration + } + busy += e.EndNs - e.StartNs + if n == 0 || e.StartNs < lo { + lo = e.StartNs + } + if e.EndNs > hi { + hi = e.EndNs + } + n++ + } + if n == 0 { + return 0, 0 + } + meanKernelUs = float64(busy) / float64(n) / 1000.0 + if hi > lo { + concurrency = float64(busy) / float64(hi-lo) + } + return concurrency, meanKernelUs +} + +// --------------------------------------------------------------------------- +// The assertions. Each arm proves it ran in the mode it says it ran in. +// --------------------------------------------------------------------------- + +// assertArmRanInItsMode is the answer to "a benchmark that silently measures +// the wrong thing". Every clause below can go red, and the red ones are the +// ones that would otherwise report a flatteringly small overhead. +func assertArmRanInItsMode(cfg gpuPCConfig, spec gpuPCArmSpec, r schema.GPUPCRun) error { + var problems []string + fail := func(format string, a ...any) { problems = append(problems, fmt.Sprintf(format, a...)) } + + // Non-vacuity FIRST, and on every arm. An arm where the pipeline + // never ran at all satisfies every negative assertion below + // perfectly, and would be the cheapest arm in the table. + if r.Evidence.ExecutionsSeen == 0 { + fail("no GPU executions reached the timeline: the adapter was not injected, " + + "the probes did not attach, or the workload ran no kernels — this arm " + + "measured nothing") + } + if r.Evidence.ProducerTier == "" { + fail("the adapter printed no report line on stderr: PERFAGENT_GPU_LOG did not " + + "take, or the adapter was never loaded into the workload") + } + + switch spec.Tier { + case gpu.PCSamplingOff: + if r.Evidence.ProducerTier != gpu.PCSamplingNameOff { + fail("producer reports tier %q, want %q", r.Evidence.ProducerTier, gpu.PCSamplingNameOff) + } + // Off means off. If any of these moved, the baseline is itself + // paying a PC-sampling cost and every other arm's margin is + // understated by exactly that amount. + if r.Evidence.PCSamplesDecoded != 0 { + fail("%d PC samples decoded with the tier off", r.Evidence.PCSamplesDecoded) + } + if r.Evidence.SamplingWindowsDecoded != 0 || r.Evidence.SamplingWindowsReceived != 0 { + fail("%d sampling windows decoded / %d received with the tier off", + r.Evidence.SamplingWindowsDecoded, r.Evidence.SamplingWindowsReceived) + } + if r.Evidence.ProducerBursts != 0 { + fail("producer opened %d bursts with the tier off", r.Evidence.ProducerBursts) + } + if r.Evidence.ExecutionsSerialized != 0 { + fail("%d executions marked serialized with the tier off", r.Evidence.ExecutionsSerialized) + } + + case gpu.PCSamplingContinuous: + if r.Evidence.ProducerTier != gpu.PCSamplingNameContinuous { + fail("producer reports tier %q, want %q", r.Evidence.ProducerTier, gpu.PCSamplingNameContinuous) + } + if r.Evidence.ProducerPCRecords == 0 { + fail("producer drained 0 PC records: Tier B was selected but never collected") + } + if r.Evidence.PCSamplesDecoded == 0 { + fail("0 PC samples reached the consumer: this arm is baseline with a different name") + } + // A CONTINUOUS producer announcing a window would be claiming a + // perturbation it did not cause. + if r.Evidence.SamplingWindowsDecoded != 0 || r.Evidence.ProducerBursts != 0 { + fail("Tier B emitted %d windows / opened %d bursts; it must do neither", + r.Evidence.SamplingWindowsDecoded, r.Evidence.ProducerBursts) + } + + case gpu.PCSamplingSerialized: + if r.Evidence.ProducerTier != gpu.PCSamplingNameSerialized { + fail("producer reports tier %q, want %q", r.Evidence.ProducerTier, gpu.PCSamplingNameSerialized) + } + if r.Evidence.ProducerBursts < cfg.MinBursts { + fail("producer opened %d bursts, want at least %d: too few for the duty "+ + "fraction to mean anything over this run", + r.Evidence.ProducerBursts, cfg.MinBursts) + } + // One open record and one closed record per burst, minus the + // close of a burst still open when the process exited. + if w, b := r.Evidence.ProducerWindows, r.Evidence.ProducerBursts; w != 2*b && w != 2*b-1 { + fail("producer emitted %d windows for %d bursts, want 2N or 2N-1", w, b) + } + if r.Evidence.ProducerStartFailed != 0 || r.Evidence.ProducerStopFailed != 0 { + fail("cuptiPCSamplingStart failed %d times, Stop %d times: some bursts did "+ + "not happen and the achieved duty is not the configured one", + r.Evidence.ProducerStartFailed, r.Evidence.ProducerStopFailed) + } + if r.Evidence.ProducerGraphRefuse != 0 || r.Evidence.ProducerGraphExecs != 0 { + fail("CUDA graph executions observed (%d) / Tier A refusals (%d): Tier A "+ + "stops bursting in such a process, so this arm did not run Tier A "+ + "for its whole length", + r.Evidence.ProducerGraphExecs, r.Evidence.ProducerGraphRefuse) + } + if r.Evidence.PCSamplesDecoded == 0 { + fail("0 PC samples reached the consumer despite %d bursts", r.Evidence.ProducerBursts) + } + if r.Evidence.SamplingWindowsReceived == 0 { + fail("no sampling window reached the agent: nothing in the resulting profile " + + "could say which executions ran perturbed") + } + // The load-bearing one. Bursts that overlapped no execution + // serialized nothing, so an arm with zero serialized executions + // measured the cost of starting and stopping CUPTI and not the + // cost of serialization. + if r.Evidence.ExecutionsSerialized == 0 { + fail("not one execution was marked gpu_serialized=\"true\": the bursts did " + + "not overlap any kernel, so this arm did not measure serialization") + } + if r.Evidence.SnapshotTier != gpu.PCSamplingNameSerialized { + fail("the agent's own snapshot reports tier %q", r.Evidence.SnapshotTier) + } + // The achieved duty against what the arm asked for. The upper + // bound is DERIVED, not a magic tolerance: the burst timer ticks + // at burst/5 by default, so a burst runs 50..60 ms, and two + // ticks of slack covers scheduling overshoot on a loaded + // machine. Anything above that and the arm ran at a duty nobody + // configured, which would make its ratio meaningless. + want := spec.dutyConfigured() + tick := float64(spec.BurstMs) / 5 + hi := (float64(spec.BurstMs) + 2*tick) / (float64(spec.BurstMs) + 2*tick + float64(spec.GapMs)) + if d := r.Evidence.ProducerDuty; d > hi || d < want/2 { + fail("producer achieved duty %.4f, outside [%.4f, %.4f] for a configured %.4f", + d, want/2, hi, want) + } + } + + if len(problems) == 0 { + return nil + } + return errors.New(strings.Join(problems, "; ")) +} + +// assertBaselineIsRealistic is the guard against measuring a microbenchmark. +// +// The plan rules the saturating launch-rate ceiling out as an instrument +// precisely because it understates serialization costs. A workload that turned +// out to be effectively serial, or whose kernels turned out to be +// microseconds long, would reproduce that error while looking like a proper +// measurement — so the two properties the plan requires are asserted rather +// than asserted-in-a-comment. +func assertBaselineIsRealistic(cfg gpuPCConfig, base schema.GPUPCArm) error { + if base.MedianConcurrency < cfg.MinConcurrency { + return fmt.Errorf( + "the baseline arm's kernel concurrency is %.2f, below the %.2f floor: this "+ + "workload has almost no concurrency for serialization to destroy, so "+ + "every Tier A number it produces would be an underestimate. Raise "+ + "--gpu-streams, or lower --gpu-blocks so each kernel occupies less of "+ + "the device and more of them co-reside", + base.MedianConcurrency, cfg.MinConcurrency) + } + if base.MedianKernelUs < cfg.MinKernelUs { + return fmt.Errorf( + "the baseline arm's mean kernel duration is %.1f us, below the %.1f us floor: "+ + "this is a launch-rate microbenchmark, which the plan rules out as the "+ + "instrument for this question. Raise --gpu-rounds", + base.MedianKernelUs, cfg.MinKernelUs) + } + return nil +} + +// assertDutyKnobDidSomething is the cross-arm proof that the three Tier A arms +// are three arms and not the same arm three times. +// +// Burst count over a fixed run length is inversely proportional to burst+gap, +// so 10% duty must open strictly more bursts than 5%, which must open strictly +// more than 2.5%. If the three came out equal, the duty environment did not +// take and all three "duties" are one duty — from which a cost-over-duty ratio +// would be pure fiction. +func assertDutyKnobDidSomething(arms []schema.GPUPCArm) error { + var a []schema.GPUPCArm + for _, arm := range arms { + if arm.Tier == gpu.PCSamplingNameSerialized { + a = append(a, arm) + } + } + if len(a) < 2 { + return nil + } + sort.Slice(a, func(i, j int) bool { return a[i].DutyConfigured > a[j].DutyConfigured }) + for i := 1; i < len(a); i++ { + hi := medianUint(burstCounts(a[i-1])) + lo := medianUint(burstCounts(a[i])) + if hi <= lo { + return fmt.Errorf( + "arm %q opened a median of %d bursts and the lower-duty arm %q opened %d: "+ + "a lower duty must open strictly fewer bursts over the same fixed "+ + "work. The duty environment did not take, so the three Tier A arms "+ + "are one arm under three names and their cost-over-duty ratios are "+ + "meaningless", + a[i-1].Name, hi, a[i].Name, lo) + } + } + return nil +} + +func burstCounts(arm schema.GPUPCArm) []uint64 { + out := make([]uint64, 0, len(arm.Runs)) + for _, r := range arm.Runs { + out = append(out, r.Evidence.ProducerBursts) + } + return out +} + +// --------------------------------------------------------------------------- +// Summarizing and deciding. Everything below is pure and unit-tested. +// --------------------------------------------------------------------------- + +// summarizeGPUPCArms fills in each arm's medians and, from the baseline's, +// each arm's cost and cost-over-duty ratio. +// +// Medians, not means, and fixed work rather than fixed time: §9.1's method, +// unchanged. A mean over five runs is moved by one slow run; a median is not, +// and one slow run on a shared machine is the common case. +func summarizeGPUPCArms(arms []schema.GPUPCArm) { + for i := range arms { + a := &arms[i] + a.MedianWallMs = medianFloat(field(a.Runs, func(r schema.GPUPCRun) float64 { return r.WallMs })) + a.MedianKernelsPerS = medianFloat(field(a.Runs, func(r schema.GPUPCRun) float64 { return r.KernelsPerS })) + a.MedianConcurrency = medianFloat(field(a.Runs, func(r schema.GPUPCRun) float64 { return r.Concurrency })) + a.MedianKernelUs = medianFloat(field(a.Runs, func(r schema.GPUPCRun) float64 { return r.MeanKernelUs })) + a.DutyAchieved = medianFloat(field(a.Runs, func(r schema.GPUPCRun) float64 { return r.Evidence.ProducerDuty })) + } + if len(arms) == 0 || arms[0].MedianWallMs <= 0 { + return + } + base := arms[0].MedianWallMs + for i := 1; i < len(arms); i++ { + a := &arms[i] + a.CostPercent = (a.MedianWallMs - base) / base * 100 + if a.DutyConfigured > 0 { + // Against the CONFIGURED duty, deliberately. It is the + // number the thresholds name ("Tier A at 10% duty"), and + // it is the conservative choice: the achieved duty is at + // or above the configured one (the burst timer's + // granularity can only overshoot), so dividing by the + // configured value can only make the ratio LARGER — + // pessimistic for Tier A, never flattering. The achieved + // duty is reported beside it and is separately asserted + // to be within a derived bound. + a.CostOverDuty = (a.CostPercent / 100) / a.DutyConfigured + } + } +} + +func field(runs []schema.GPUPCRun, f func(schema.GPUPCRun) float64) []float64 { + out := make([]float64, 0, len(runs)) + for _, r := range runs { + out = append(out, f(r)) + } + return out +} + +// medianFloat returns the median, or 0 for an empty slice. Even lengths take +// the mean of the two middle values. +func medianFloat(v []float64) float64 { + if len(v) == 0 { + return 0 + } + s := append([]float64(nil), v...) + sort.Float64s(s) + n := len(s) + if n%2 == 1 { + return s[n/2] + } + return (s[n/2-1] + s[n/2]) / 2 +} + +func medianUint(v []uint64) uint64 { + if len(v) == 0 { + return 0 + } + s := append([]uint64(nil), v...) + sort.Slice(s, func(i, j int) bool { return s[i] < s[j] }) + return s[len(s)/2] +} + +// decideGPUPC evaluates the four pre-committed threshold clauses against the +// arms and returns the verdict. +// +// Every clause is evaluated INDEPENDENTLY and every one that fires is +// recorded, because more than one can fire and the combination is itself +// information. The verdict is then resolved by severity, which is stated here +// rather than left to the order the clauses happen to be written in: +// +// unshippable > deep-dive-only > smaller-duty > opt-in +// +// The residual case — the headline duty is over budget but a smaller one is +// not, and the ratio holds — is a named verdict of its own rather than a +// fallthrough. The plan's four clauses do not cover it and a harness that +// quietly picked the friendliest neighbouring answer would be making the +// decision the thresholds exist to take out of anyone's hands. +func decideGPUPC(arms []schema.GPUPCArm, maxWallPercent, maxCostOverDuty float64) schema.GPUPCDecision { + d := schema.GPUPCDecision{ + MaxWallPercent: maxWallPercent, + MaxCostOverDut: maxCostOverDuty, + } + + var tierB *schema.GPUPCArm + var tierA []schema.GPUPCArm + for i := range arms { + switch arms[i].Tier { + case gpu.PCSamplingNameContinuous: + tierB = &arms[i] + case gpu.PCSamplingNameSerialized: + tierA = append(tierA, arms[i]) + } + } + sort.Slice(tierA, func(i, j int) bool { return tierA[i].DutyConfigured > tierA[j].DutyConfigured }) + + // ---- Tier B. + switch { + case tierB == nil: + d.TierB = verdictTierBExplicitOnly + d.Lines = append(d.Lines, "DECISION tier B: no continuous arm ran; "+ + "absent a measurement, Tier B does not ship as always-on") + case tierB.CostPercent > maxWallPercent: + d.Fired = append(d.Fired, clauseTierBOverBudget) + d.TierB = verdictTierBExplicitOnly + d.Lines = append(d.Lines, fmt.Sprintf( + "THRESHOLD FIRED %s: tier B costs %+.2f%% wall-clock, above the %.1f%% bar", + clauseTierBOverBudget, tierB.CostPercent, maxWallPercent)) + d.Lines = append(d.Lines, "DECISION tier B: "+verdictTierBExplicitOnly+ + " — Tier B does not ship as always-on; it becomes an explicitly-enabled "+ + "mode like Tier A") + default: + d.TierB = verdictTierBAlwaysOnCandidate + d.Lines = append(d.Lines, fmt.Sprintf( + "threshold %s not fired: tier B costs %+.2f%% wall-clock, within the %.1f%% bar", + clauseTierBOverBudget, tierB.CostPercent, maxWallPercent)) + d.Lines = append(d.Lines, "DECISION tier B: "+verdictTierBAlwaysOnCandidate+ + " — Tier B remains a candidate for always-on on this evidence") + } + + // ---- Tier A. + if len(tierA) == 0 { + d.TierA = verdictTierAIndeterminate + d.Lines = append(d.Lines, "DECISION tier A: "+verdictTierAIndeterminate+ + " — no serialized arm ran, so no threshold can be evaluated") + return d + } + headline := tierA[0] // the largest duty tested: the plan's "10% duty" + smallest := tierA[len(tierA)-1] + + headlineWithin := headline.CostPercent <= maxWallPercent && headline.CostOverDuty <= maxCostOverDuty + ratioOverEverywhere := true + for _, a := range tierA { + if a.CostOverDuty <= maxCostOverDuty { + ratioOverEverywhere = false + break + } + } + smallestOverBudget := smallest.CostPercent > maxWallPercent + + if headlineWithin { + d.Fired = append(d.Fired, clauseTierAHeadlineWithin) + d.Lines = append(d.Lines, fmt.Sprintf( + "THRESHOLD FIRED %s: %s costs %+.2f%% (bar %.1f%%) at cost/duty %.2f (bar %.1f)", + clauseTierAHeadlineWithin, headline.Name, headline.CostPercent, + maxWallPercent, headline.CostOverDuty, maxCostOverDuty)) + } + if ratioOverEverywhere { + d.Fired = append(d.Fired, clauseTierARatioOverAtAll) + d.Lines = append(d.Lines, fmt.Sprintf( + "THRESHOLD FIRED %s: cost/duty is above %.1f at every duty tested (%s)", + clauseTierARatioOverAtAll, maxCostOverDuty, ratioList(tierA))) + } + if smallestOverBudget { + d.Fired = append(d.Fired, clauseTierASmallestOverBud) + d.Lines = append(d.Lines, fmt.Sprintf( + "THRESHOLD FIRED %s: the lowest duty tested (%s) still costs %+.2f%%, above the %.1f%% bar", + clauseTierASmallestOverBud, smallest.Name, smallest.CostPercent, maxWallPercent)) + } + // The two harshest clauses can coincide arithmetically, and when they do + // the reader is told rather than left to notice. + // + // cost/duty > R is the same statement as cost% > 100*R*duty. At + // duty = R*... — concretely, with the plan's bars (R = 2, W = 5%) the + // two coincide EXACTLY at duty = W/(100*R) = 2.5%, which is the lowest + // duty the plan's arm table tests. So on that table "ratio above 2 at + // every duty" strictly IMPLIES "the lowest duty is over the wall bar", + // and the deep-dive-only verdict is unreachable: the unshippable one + // always outranks it. That is a property of the thresholds, not of the + // hardware, and it is stated here so a controller reading a result does + // not conclude that the deep-dive branch was considered and rejected. + if ratioOverEverywhere && smallestOverBudget { + d.Lines = append(d.Lines, fmt.Sprintf( + "NOTE both harsh clauses fired, and at the lowest duty tested (%.2f%%) they are "+ + "the SAME condition: cost/duty > %.1f means cost > %.2f%%, and the "+ + "wall bar is %.1f%%. They separate only below %.2f%% duty, which this "+ + "table does not test — so %s could not have been reached here whatever "+ + "the numbers. Add a lower-duty arm if that distinction is wanted.", + smallest.DutyConfigured*100, maxCostOverDuty, + maxCostOverDuty*smallest.DutyConfigured*100, maxWallPercent, + maxWallPercent/(100*maxCostOverDuty)*100, verdictTierADeepDiveOnly)) + } + + switch { + case smallestOverBudget: + d.TierA = verdictTierAUnshippable + d.Lines = append(d.Lines, "DECISION tier A: "+verdictTierAUnshippable+ + " — serialization costs more than the sampling window can explain and "+ + "duty-cycling has no remaining lever. Tier A is unshippable in this phase") + case ratioOverEverywhere: + d.TierA = verdictTierADeepDiveOnly + d.Lines = append(d.Lines, "DECISION tier A: "+verdictTierADeepDiveOnly+ + " — duty-cycling is not buying what it appears to. Tier A ships only as a "+ + "deliberate deep-dive mode, with Task 11's operator warning and no "+ + "suggestion that it suits continuous use") + case headlineWithin: + d.TierA = verdictTierAOptIn + d.Lines = append(d.Lines, "DECISION tier A: "+verdictTierAOptIn+ + " — ships as an opt-in tier, as planned") + default: + if best, ok := largestQualifyingDuty(tierA, maxWallPercent, maxCostOverDuty); ok { + d.TierA = verdictTierASmallerDuty + d.Lines = append(d.Lines, fmt.Sprintf( + "DECISION tier A: %s — the headline %.1f%% duty is over budget, but "+ + "%s is within both bars (%+.2f%%, cost/duty %.2f). Ship opt-in "+ + "with that as the default duty, not %.1f%%", + verdictTierASmallerDuty, headline.DutyConfigured*100, best.Name, + best.CostPercent, best.CostOverDuty, headline.DutyConfigured*100)) + } else { + d.TierA = verdictTierAIndeterminate + d.Lines = append(d.Lines, "DECISION tier A: "+verdictTierAIndeterminate+ + " — no duty tested is within both bars, yet neither the "+ + "every-duty ratio clause nor the lowest-duty clause fired. "+ + "That combination means cost does not rise with duty, which "+ + "means the measurement is unsound. Re-run before deciding "+ + "anything; do not read this as a pass") + } + } + return d +} + +// largestQualifyingDuty returns the highest-duty arm that is within BOTH bars. +func largestQualifyingDuty(tierA []schema.GPUPCArm, maxWallPercent, maxCostOverDuty float64) (schema.GPUPCArm, bool) { + for _, a := range tierA { // already sorted by descending duty + if a.CostPercent <= maxWallPercent && a.CostOverDuty <= maxCostOverDuty { + return a, true + } + } + return schema.GPUPCArm{}, false +} + +func ratioList(tierA []schema.GPUPCArm) string { + parts := make([]string, 0, len(tierA)) + for _, a := range tierA { + parts = append(parts, fmt.Sprintf("%.1f%% duty: %.2f", a.DutyConfigured*100, a.CostOverDuty)) + } + return strings.Join(parts, ", ") +} + +// renderGPUPCTable prints the arm table, the ratios and the evidence each arm +// produced. The evidence column is not decoration: it is what lets a reader +// see that the arm claiming to be Tier A at 2.5% duty really did open bursts +// and really did serialize executions. +func renderGPUPCTable(w io.Writer, res *schema.GPUPCOverhead) { + _, _ = fmt.Fprintln(w) + _, _ = fmt.Fprintf(w, "GPU PC-sampling overhead — %d kernels of fixed work per run, "+ + "%d streams, %d-round kernels\n", res.Workload.KernelsRun, res.Workload.Streams, + res.Workload.Rounds) + _, _ = fmt.Fprintln(w, "baseline = the shipping Phase 4 configuration with PC sampling OFF, "+ + "not an uninjected run") + _, _ = fmt.Fprintln(w) + _, _ = fmt.Fprintf(w, "| %-30s | %5s | %10s | %8s | %9s | %8s | %6s | %8s |\n", + "arm", "duty", "wall ms", "cost", "cost/duty", "kern/s", "conc", "kernel us") + _, _ = fmt.Fprintf(w, "|%s|%s|%s|%s|%s|%s|%s|%s|\n", strings.Repeat("-", 32), + strings.Repeat("-", 7), strings.Repeat("-", 12), strings.Repeat("-", 10), + strings.Repeat("-", 11), strings.Repeat("-", 10), strings.Repeat("-", 8), + strings.Repeat("-", 10)) + for i, a := range res.Arms { + duty, cost, ratio := "—", "—", "—" + if a.DutyConfigured > 0 { + duty = fmt.Sprintf("%.1f%%", a.DutyConfigured*100) + ratio = fmt.Sprintf("%.2f", a.CostOverDuty) + } + if i > 0 { + cost = fmt.Sprintf("%+.2f%%", a.CostPercent) + } + _, _ = fmt.Fprintf(w, "| %-30s | %5s | %10.1f | %8s | %9s | %8.0f | %6.2f | %8.0f |\n", + a.Name, duty, a.MedianWallMs, cost, ratio, a.MedianKernelsPerS, + a.MedianConcurrency, a.MedianKernelUs) + } + _, _ = fmt.Fprintln(w) + _, _ = fmt.Fprintln(w, "what each arm proved it ran (medians across runs):") + for _, a := range res.Arms { + _, _ = fmt.Fprintf(w, " %-30s producer tier=%s pc_records=%d bursts=%d windows=%d duty=%.4f | "+ + "consumer pc_samples=%d windows=%d execs=%d serialized=%d not=%d unknown=%d\n", + a.Name, + firstEvidence(a).ProducerTier, + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ProducerPCRecords })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ProducerBursts })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ProducerWindows })), + a.DutyAchieved, + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.PCSamplesDecoded })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.SamplingWindowsReceived })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return uint64(r.Evidence.ExecutionsSeen) })), //nolint:gosec // a count + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ExecutionsSerialized })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ExecutionsNotSerialized })), + medianUint(field64(a.Runs, func(r schema.GPUPCRun) uint64 { return r.Evidence.ExecutionsUnknown })), + ) + } + _, _ = fmt.Fprintln(w) +} + +func firstEvidence(a schema.GPUPCArm) schema.GPUPCEvidence { + if len(a.Runs) == 0 { + return schema.GPUPCEvidence{} + } + return a.Runs[0].Evidence +} + +func field64(runs []schema.GPUPCRun, f func(schema.GPUPCRun) uint64) []uint64 { + out := make([]uint64, 0, len(runs)) + for _, r := range runs { + out = append(out, f(r)) + } + return out +} + +// defaultShimPath and defaultConcurrentWorkload are resolved relative to the +// repository root when the flags are left unset, the same way the other +// scenarios auto-detect test/workloads. +func defaultShimPath() string { return filepath.Join("shim", "libperfagent-gpu-nvidia.so") } +func defaultConcurrentWorkload() string { + return filepath.Join("shim", "nvidia", "testdata", "cuda_concurrent") +} + +// runGPUPCScenario is the whole entry point for --scenario gpu-pc-overhead: +// the skip path, the arms, the JSON, and the exit code. +// +// The exit codes are the same shape the "self" scenario already uses: 0 when +// the measurement completed (whatever the verdict — an honest +// TIER_A_UNSHIPPABLE is a successful run of this benchmark), 3 when an +// assertion failed and the numbers therefore mean nothing. A skip is 0 with a +// single BENCH_SKIPPED line and no output file, because a partial file with no +// numbers in it is the thing most likely to be mistaken for a result. +func runGPUPCScenario(cfg gpuPCConfig, outPath string) { + if reason := gpuPCSkipReason(cfg); reason != "" { + _, _ = fmt.Fprintln(os.Stdout, "BENCH_SKIPPED: "+reason) + os.Exit(0) + } + + doc := &schema.Document{ + Scenario: "gpu-pc-overhead", + StartedAt: time.Now().UTC(), + Config: schema.Config{Runs: cfg.Runs}, + System: gatherSystemInfo(), + } + ok := runGPUPCOverhead(doc, cfg, os.Stdout) + + out := outPath + if out == "" { + out = fmt.Sprintf("bench-gpu-pc-overhead-%d.json", time.Now().Unix()) + } + f, err := os.Create(out) + if err != nil { + _, _ = fmt.Fprintf(os.Stderr, "create %s: %v\n", out, err) + os.Exit(2) + } + if err := schema.Write(f, doc); err != nil { + _ = f.Close() + _, _ = fmt.Fprintf(os.Stderr, "write %s: %v\n", out, err) + os.Exit(2) + } + if err := f.Close(); err != nil { + _, _ = fmt.Fprintf(os.Stderr, "close %s: %v\n", out, err) + os.Exit(2) + } + _, _ = fmt.Fprintf(os.Stdout, "wrote %s\n", out) + + if !ok { + _, _ = fmt.Fprintln(os.Stderr, + "gpu-pc-overhead: the run did not prove what it measured; the numbers in "+ + "the output file are NOT a tier decision and must not be recorded as one") + os.Exit(3) + } +} diff --git a/bench/cmd/scenario/gpupc_test.go b/bench/cmd/scenario/gpupc_test.go new file mode 100644 index 00000000..9f29e68b --- /dev/null +++ b/bench/cmd/scenario/gpupc_test.go @@ -0,0 +1,720 @@ +package main + +// The offline half of Task 12. +// +// The measurement needs an RTX 3090 and none of it can be run here. What CAN +// be run here is everything between the measurement and the decision: the +// parsers that turn a producer's log line and a workload's result line into +// evidence, the medians, the ratio, and the four pre-committed threshold +// clauses. Those are the parts that decide the outcome, and a harness whose +// decision logic is only exercised on hardware is a harness whose decision +// logic has never been tested at all. +// +// Every threshold test states the arm table it feeds in, so a reader can see +// the thresholds are being applied to the numbers the plan named rather than +// to numbers chosen to produce a preferred answer. + +import ( + "fmt" + "os" + "strconv" + "strings" + "testing" + + "github.com/stretchr/testify/assert" + "github.com/stretchr/testify/require" + + "github.com/dpsoft/perf-agent/bench/internal/schema" + "github.com/dpsoft/perf-agent/gpu" +) + +// armsWith builds an arm table: a baseline, a Tier B arm at tierBCost percent, +// and one Tier A arm per (duty, cost) pair. The ratio is derived, never +// supplied, so a test cannot accidentally assert against a ratio that does not +// follow from the cost it set. +func armsWith(tierBCost float64, duties, costs []float64) []schema.GPUPCArm { + if len(duties) != len(costs) { + panic("duties and costs must be the same length") + } + base := 1000.0 + arms := []schema.GPUPCArm{ + {Name: "baseline", Tier: gpu.PCSamplingNameOff, MedianWallMs: base}, + {Name: "tier B", Tier: gpu.PCSamplingNameContinuous, + MedianWallMs: base * (1 + tierBCost/100), CostPercent: tierBCost}, + } + for i, duty := range duties { + c := costs[i] + arms = append(arms, schema.GPUPCArm{ + Name: fmt.Sprintf("tier A %g%% duty", duty*100), + Tier: gpu.PCSamplingNameSerialized, + DutyConfigured: duty, + MedianWallMs: base * (1 + c/100), + CostPercent: c, + CostOverDuty: (c / 100) / duty, + }) + } + return arms +} + +// planDuties is the plan's own arm table: 10%, 5%, 2.5%. +var planDuties = []float64{0.10, 0.05, 0.025} + +// armsFor is armsWith over the plan's three duties. +func armsFor(tierBCost float64, tierACosts [3]float64) []schema.GPUPCArm { + return armsWith(tierBCost, planDuties, tierACosts[:]) +} + +func decide(arms []schema.GPUPCArm) schema.GPUPCDecision { + return decideGPUPC(arms, gpuPCMaxWallPercent, gpuPCMaxCostOverDuty) +} + +// --------------------------------------------------------------------------- +// The four pre-committed clauses. +// --------------------------------------------------------------------------- + +// Tier B over 5% wall-clock: Tier B does not ship as always-on. +func TestTierBOverFivePercentStopsBeingAlwaysOn(t *testing.T) { + d := decide(armsFor(7.5, [3]float64{1.0, 0.5, 0.25})) + assert.Equal(t, verdictTierBExplicitOnly, d.TierB) + assert.Contains(t, d.Fired, clauseTierBOverBudget) + assert.Contains(t, strings.Join(d.Lines, "\n"), "does not ship as always-on") +} + +// The mirror, and it matters as much: a Tier B that is within budget must not +// be quietly demoted, or the clause would be unfalsifiable. +func TestTierBWithinFivePercentStaysAnAlwaysOnCandidate(t *testing.T) { + d := decide(armsFor(1.2, [3]float64{1.0, 0.5, 0.25})) + assert.Equal(t, verdictTierBAlwaysOnCandidate, d.TierB) + assert.NotContains(t, d.Fired, clauseTierBOverBudget) +} + +// Exactly at the bar is within it: the plan says "> 5%" fires, so 5.0 does +// not. Pinned because a strict-versus-non-strict slip is the classic way a +// pre-committed threshold quietly becomes a different threshold. +func TestTierBExactlyAtTheBarDoesNotFire(t *testing.T) { + d := decide(armsFor(gpuPCMaxWallPercent, [3]float64{1.0, 0.5, 0.25})) + assert.Equal(t, verdictTierBAlwaysOnCandidate, d.TierB) +} + +// Tier A at 10% duty costs <= 5% AND cost/duty <= 2: ships as an opt-in tier. +// 4% cost at 10% duty is a ratio of 0.4. +func TestTierAWithinBothBarsShipsOptIn(t *testing.T) { + d := decide(armsFor(1.0, [3]float64{4.0, 2.0, 1.0})) + assert.Equal(t, verdictTierAOptIn, d.TierA) + assert.Contains(t, d.Fired, clauseTierAHeadlineWithin) + assert.Contains(t, strings.Join(d.Lines, "\n"), "ships as an opt-in tier") +} + +// cost/duty > 2 at every duty tested: deep-dive mode only. +// +// Reaching this verdict at all takes a duty BELOW 2.5%, and the reason is +// arithmetic rather than a quirk of this table: cost/duty > 2 is the same +// statement as cost > 200 x duty percent, which at 2.5% duty is exactly +// "cost > 5%" — the wall-clock bar. See +// TestTheTwoHarshClausesCoincideAtThePlansLowestDuty. So the duties here are +// 10%/5%/1%, and the costs (25%, 12%, 3%) put every ratio above 2 while +// leaving the lowest duty inside the 5% wall bar. +func TestTierARatioAboveTwoAtEveryDutyIsDeepDiveOnly(t *testing.T) { + d := decide(armsWith(1.0, []float64{0.10, 0.05, 0.01}, []float64{25.0, 12.0, 3.0})) + assert.Equal(t, verdictTierADeepDiveOnly, d.TierA) + assert.Contains(t, d.Fired, clauseTierARatioOverAtAll) + assert.NotContains(t, d.Fired, clauseTierASmallestOverBud) + joined := strings.Join(d.Lines, "\n") + assert.Contains(t, joined, "deliberate deep-dive mode") + assert.Contains(t, joined, "no suggestion that it suits continuous use") +} + +// The finding, pinned so it cannot be un-noticed: on the plan's own arm table +// the deep-dive-only clause is UNREACHABLE, because at 2.5% duty it is the +// same condition as the unshippable clause and the unshippable one outranks +// it. A harness that reported deep-dive-only from these duties would be +// reporting a verdict its own thresholds cannot produce. +func TestTheTwoHarshClausesCoincideAtThePlansLowestDuty(t *testing.T) { + // 2.5% duty, cost 6%: over the 5% wall bar, and ratio 2.4 > 2. The + // higher duties are pushed over the ratio bar too (25% at 10% duty is + // 2.5; 12% at 5% is 2.4), so the every-duty clause genuinely fires. + d := decide(armsFor(1.0, [3]float64{25.0, 12.0, 6.0})) + assert.Contains(t, d.Fired, clauseTierARatioOverAtAll) + assert.Contains(t, d.Fired, clauseTierASmallestOverBud) + assert.Equal(t, verdictTierAUnshippable, d.TierA, + "the strictly harsher verdict must win") + assert.Contains(t, strings.Join(d.Lines, "\n"), + "they are the SAME condition", + "the coincidence must be reported, not left for the reader to notice") +} + +// One duty with a ratio at or below 2 is enough to stop the every-duty clause +// firing. Asserted because "at every duty tested" is a universal, and a +// harness that treated it as "at the headline duty" would reach a harsher +// verdict than the plan committed to. +func TestTierARatioClauseIsUniversalNotHeadline(t *testing.T) { + // 25% at 10% duty is a ratio of 2.5, over the bar; 12% at 5% is 2.4, + // over it; 4% at 2.5% is 1.6, WITHIN it. One arm within the bar is + // enough to stop a clause that says "at every duty tested". + d := decide(armsFor(1.0, [3]float64{25.0, 12.0, 4.0})) + assert.NotContains(t, d.Fired, clauseTierARatioOverAtAll) + assert.NotEqual(t, verdictTierADeepDiveOnly, d.TierA) +} + +// Tier A at 2.5% duty still over 5% wall-clock: unshippable in this phase. +// This clause outranks every other verdict, so the table below is deliberately +// one in which the ratio clause ALSO fires — the harsher answer must win. +func TestTierASmallestDutyOverBudgetIsUnshippableAndOutranksEverything(t *testing.T) { + d := decide(armsFor(1.0, [3]float64{30.0, 18.0, 9.0})) + assert.Equal(t, verdictTierAUnshippable, d.TierA) + assert.Contains(t, d.Fired, clauseTierASmallestOverBud) + assert.Contains(t, d.Fired, clauseTierARatioOverAtAll) + assert.Contains(t, strings.Join(d.Lines, "\n"), "unshippable in this phase") +} + +// The residual the plan's four clauses do not name: the headline duty is over +// the wall-clock bar, but a smaller duty is within BOTH bars. Duty-cycling +// works here — it just needs to be turned down — so the honest outcome is a +// named verdict that says which duty to ship, not a silent promotion to +// "opt-in" and not a demotion to "deep dive". +// +// 8% at 10% duty (ratio 0.8, over the wall bar); 4% at 5% (ratio 0.8, within +// both); 2% at 2.5% (ratio 0.8, within both). +func TestTierAOverBudgetAtTheHeadlineButTunableNamesTheDuty(t *testing.T) { + d := decide(armsFor(1.0, [3]float64{8.0, 4.0, 2.0})) + assert.Equal(t, verdictTierASmallerDuty, d.TierA) + assert.NotContains(t, d.Fired, clauseTierAHeadlineWithin) + assert.NotContains(t, d.Fired, clauseTierASmallestOverBud) + joined := strings.Join(d.Lines, "\n") + // The LARGEST qualifying duty, not just any: 5%, not 2.5%. + assert.Contains(t, joined, "tier A 5% duty is within both bars") +} + +// A table where nothing qualifies at any duty and neither harsh clause fires +// is not a pass. It means cost does not fall with duty the way serialization +// says it must, so the measurement itself is unsound — and it gets a verdict +// that says so rather than the friendliest neighbouring answer. +// +// Duties 10%/5%/1%; costs 8%, 6%, 3%. The two larger duties are over the wall +// bar (so nothing qualifies there); at 1% duty 3% is INSIDE the wall bar but +// its ratio is 3.0, outside the ratio bar — so nothing qualifies anywhere, the +// every-duty ratio clause does not fire (10% duty is at 0.8), and the +// lowest-duty wall clause does not fire either. +func TestTierAWithNothingQualifyingAnywhereIsIndeterminateNotAPass(t *testing.T) { + d := decide(armsWith(1.0, []float64{0.10, 0.05, 0.01}, []float64{8.0, 6.0, 3.0})) + assert.Empty(t, d.Fired, "no Tier A clause fires on this table") + assert.Equal(t, verdictTierAIndeterminate, d.TierA) + assert.Contains(t, strings.Join(d.Lines, "\n"), "do not read this as a pass") +} + +// The counterpart, and the reassuring half: on the PLAN's own duties the +// decision is total. At 2.5% duty "within the wall bar" and "within the ratio +// bar" are the same condition, so the lowest-duty arm either qualifies for +// both — giving opt-in or a smaller duty — or fails both, giving unshippable. +// Indeterminate is unreachable there, which is why it exists as a guard for +// other duty tables rather than as an expected outcome. +func TestOnThePlansDutiesTheDecisionIsTotal(t *testing.T) { + for _, costs := range [][3]float64{ + {1, 0.5, 0.25}, {4, 2, 1}, {8, 4, 2}, {12, 7, 4}, + {25, 12, 6}, {30, 18, 9}, {0.1, 0.1, 0.1}, {60, 30, 15}, + } { + d := decide(armsFor(1.0, costs)) + assert.NotEqual(t, verdictTierAIndeterminate, d.TierA, + "costs %v must reach a real verdict", costs) + assert.NotEmpty(t, d.TierA) + } +} + +// A table with no serialized arm at all must not report a Tier A verdict. +func TestNoTierAArmIsIndeterminate(t *testing.T) { + arms := armsFor(1.0, [3]float64{1, 1, 1})[:2] + d := decide(arms) + assert.Equal(t, verdictTierAIndeterminate, d.TierA) + assert.Equal(t, verdictTierBAlwaysOnCandidate, d.TierB) +} + +// The thresholds are echoed into the decision so the output file itself shows +// which numbers were applied. A run whose recorded bars differ from the plan's +// is a run whose verdict was taken against different thresholds. +func TestDecisionEchoesTheCommittedThresholds(t *testing.T) { + d := decide(armsFor(1.0, [3]float64{1, 1, 1})) + assert.InDelta(t, 5.0, d.MaxWallPercent, 1e-9) + assert.InDelta(t, 2.0, d.MaxCostOverDut, 1e-9) +} + +// --------------------------------------------------------------------------- +// The summary arithmetic. +// --------------------------------------------------------------------------- + +// The cost is against the BASELINE arm — the shipping Phase 4 configuration +// with PC sampling off — and the ratio is against the CONFIGURED duty. Both +// are pinned here because both are choices, not arithmetic. +func TestSummarizeComputesCostAgainstBaselineAndRatioAgainstConfiguredDuty(t *testing.T) { + arms := []schema.GPUPCArm{ + {Name: "baseline", Tier: gpu.PCSamplingNameOff, + Runs: runsWithWall(100, 110, 90, 105, 95)}, + {Name: "tier A 10%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.10, + Runs: runsWithWall(120, 130, 110, 125, 115)}, + } + summarizeGPUPCArms(arms) + assert.InDelta(t, 100, arms[0].MedianWallMs, 1e-9, "median of 90,95,100,105,110") + assert.InDelta(t, 120, arms[1].MedianWallMs, 1e-9, "median of 110,115,120,125,130") + assert.InDelta(t, 20, arms[1].CostPercent, 1e-9) + assert.InDelta(t, 2.0, arms[1].CostOverDuty, 1e-9, "0.20 cost over 0.10 duty") + assert.Zero(t, arms[0].CostPercent, "the baseline is not a cost against itself") +} + +func runsWithWall(v ...float64) []schema.GPUPCRun { + out := make([]schema.GPUPCRun, 0, len(v)) + for i, w := range v { + out = append(out, schema.GPUPCRun{RunN: i + 1, WallMs: w}) + } + return out +} + +func TestMedianFloat(t *testing.T) { + assert.Zero(t, medianFloat(nil)) + assert.InDelta(t, 3.0, medianFloat([]float64{5, 1, 3}), 1e-9) + assert.InDelta(t, 2.5, medianFloat([]float64{1, 2, 3, 4}), 1e-9) + // The input must not be reordered under the caller. + in := []float64{3, 1, 2} + _ = medianFloat(in) + assert.Equal(t, []float64{3, 1, 2}, in) +} + +// --------------------------------------------------------------------------- +// The evidence parsers. These are what stand between "the arm ran in the mode +// it claims" and "the arm reported a flatteringly small overhead". +// --------------------------------------------------------------------------- + +func TestParseAdapterReportReadsATierARun(t *testing.T) { + stderr := strings.Join([]string{ + "perfagent-cupti: pc_sampling=on tier=A/kernel-serialized", + "perfagent-cupti: graph_execs=0 multi_device=0 devices=1", + "perfagent-cupti: tier A bursts=41 burst_ns=2091000000 duty=0.1043 gap_ns=450000000 " + + "windows=82 range_end_drains=41 start_failed=0 stop_failed=0 graph_refused=0 sampling_now=0", + "perfagent-cupti: pc exit tier=serialized period=8(=256 cycles) stall_reasons=38 " + + "ctx_seen=1 ctx_enabled=1 ctx_enable_failed=0 pc_records=1828 pcs=352 " + + "graph_execs=0 multi_device=0 finalize_seen=0", + }, "\n") + ev := parseAdapterReport(stderr) + assert.Equal(t, "serialized", ev.ProducerTier) + assert.Equal(t, uint64(1828), ev.ProducerPCRecords) + assert.Equal(t, uint64(41), ev.ProducerBursts) + assert.Equal(t, uint64(82), ev.ProducerWindows) + assert.InDelta(t, 0.1043, ev.ProducerDuty, 1e-9) + assert.Zero(t, ev.ProducerStartFailed) + assert.Zero(t, ev.ProducerStopFailed) + assert.Zero(t, ev.ProducerGraphRefuse) + assert.Zero(t, ev.ProducerGraphExecs) +} + +// The startup line spells the tier "A/kernel-serialized" and the exit report +// spells it "serialized". The parser must take it from the exit report, which +// is the line that also carries what the run actually did — reading the +// startup spelling would report a tier the run merely intended. +func TestParseAdapterReportPrefersTheExitReportSpelling(t *testing.T) { + ev := parseAdapterReport(strings.Join([]string{ + "perfagent-cupti: pc_sampling=on tier=B/continuous", + "perfagent-cupti: pc exit tier=continuous period=8(=256 cycles) pc_records=903 " + + "graph_execs=0 multi_device=0", + }, "\n")) + assert.Equal(t, "continuous", ev.ProducerTier) + assert.Equal(t, uint64(903), ev.ProducerPCRecords) + assert.Zero(t, ev.ProducerBursts, "Tier B does not burst") +} + +func TestParseAdapterReportReadsTheOffRun(t *testing.T) { + ev := parseAdapterReport(strings.Join([]string{ + "perfagent-cupti: pc_sampling=off tier=none", + "perfagent-cupti: graph_execs=0 multi_device=0 devices=1", + "perfagent-cupti: pc_sampling=off tier_refused=0 " + + "(set PERFAGENT_GPU_PC_SAMPLING=continuous or =serialized)", + }, "\n")) + assert.Equal(t, "off", ev.ProducerTier) + assert.Zero(t, ev.ProducerPCRecords) + assert.Zero(t, ev.ProducerBursts) +} + +// A producer that printed nothing must leave the tier EMPTY, never "off". +// This is the difference between "the adapter was loaded and PC sampling was +// off" and "the adapter was never loaded at all", and the second one is an arm +// that measured nothing while looking like the cheapest arm in the table. +func TestParseAdapterReportOnSilenceLeavesTheTierUnset(t *testing.T) { + ev := parseAdapterReport("some unrelated stderr from the workload\n") + assert.Empty(t, ev.ProducerTier) +} + +func TestParseConcurrentLine(t *testing.T) { + r, err := parseConcurrentLine( + "concurrent: iters=20000 warmup=64 streams=4 rounds=64000 blocks=16 threads=256 " + + "sync_every=4 kernels=80000 elapsed_ms=24310.750 kernels_per_s=3290.7 " + + "max_abs_err=0.000000000\n") + require.NoError(t, err) + assert.InDelta(t, 24310.750, r.WallMs, 1e-6) + assert.InDelta(t, 3290.7, r.KernelsPerS, 1e-6) + assert.Zero(t, r.MaxAbsErr) +} + +func TestParseConcurrentLineRefusesSilence(t *testing.T) { + _, err := parseConcurrentLine("concurrent: some other line\n") + require.Error(t, err) + assert.Contains(t, err.Error(), "no \"concurrent: iters=...\" result line") +} + +// --------------------------------------------------------------------------- +// The arm-proof assertions. Every one of these can go red, which is the point: +// the failure they exist to catch reports a small, green overhead number. +// --------------------------------------------------------------------------- + +func goodOffRun() schema.GPUPCRun { + return schema.GPUPCRun{Evidence: schema.GPUPCEvidence{ + ProducerTier: gpu.PCSamplingNameOff, ExecutionsSeen: 4000, + ExecutionsNotSerialized: 4000, SnapshotTier: gpu.PCSamplingNameOff, + }} +} + +func goodTierARun() schema.GPUPCRun { + return schema.GPUPCRun{Evidence: schema.GPUPCEvidence{ + ProducerTier: gpu.PCSamplingNameSerialized, ProducerPCRecords: 1828, + ProducerBursts: 41, ProducerWindows: 82, ProducerDuty: 0.104, + PCSamplesDecoded: 1828, SamplingWindowsDecoded: 82, + SamplingWindowsReceived: 82, ExecutionsSeen: 4000, + ExecutionsSerialized: 380, ExecutionsNotSerialized: 3620, + SnapshotTier: gpu.PCSamplingNameSerialized, + }} +} + +func testCfg() gpuPCConfig { + return gpuPCConfig{MinConcurrency: 1.5, MinKernelUs: 50, MinBursts: 4} +} + +func TestTheGoodArmsPass(t *testing.T) { + cfg := testCfg() + require.NoError(t, assertArmRanInItsMode(cfg, gpuPCArms[0], goodOffRun())) + require.NoError(t, assertArmRanInItsMode(cfg, gpuPCArms[2], goodTierARun())) +} + +// The single most important negative: an arm that ran nothing at all. It +// satisfies every "must be zero" clause perfectly and would be the fastest arm +// in the table. +func TestAnArmThatMeasuredNothingFails(t *testing.T) { + r := goodOffRun() + r.Evidence.ExecutionsSeen = 0 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[0], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "measured nothing") +} + +// The baseline must be genuinely off. A baseline that was itself sampling +// understates every other arm's margin by exactly its own cost. +func TestABaselineThatWasSamplingFails(t *testing.T) { + r := goodOffRun() + r.Evidence.PCSamplesDecoded = 17 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[0], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "17 PC samples decoded with the tier off") +} + +// A Tier A arm whose bursts overlapped no kernel measured the cost of starting +// and stopping CUPTI, not the cost of serialization — and would report a +// beautifully small number for it. +func TestATierAArmThatSerializedNothingFails(t *testing.T) { + r := goodTierARun() + r.Evidence.ExecutionsSerialized = 0 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[2], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "did not measure serialization") +} + +// A Tier A arm that never bursted is the baseline under another name. +func TestATierAArmWithTooFewBurstsFails(t *testing.T) { + r := goodTierARun() + r.Evidence.ProducerBursts, r.Evidence.ProducerWindows = 2, 4 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[2], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "want at least 4") +} + +// The window count must reconcile with the burst count: one open record and +// one closed record per burst, minus the close of a burst still open at exit. +func TestWindowsMustReconcileWithBursts(t *testing.T) { + r := goodTierARun() + r.Evidence.ProducerWindows = 60 // neither 82 nor 81 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[2], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "want 2N or 2N-1") + + r.Evidence.ProducerWindows = 81 // the hard-exit shape: 2N-1 + require.NoError(t, assertArmRanInItsMode(testCfg(), gpuPCArms[2], r)) +} + +// An arm that ran at a duty nobody configured has a meaningless ratio, because +// the ratio's denominator is the configured duty. +func TestATierAArmRunningAtTheWrongDutyFails(t *testing.T) { + r := goodTierARun() + r.Evidence.ProducerDuty = 0.31 // asked for ~0.10 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[2], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "achieved duty 0.3100") + + // The burst timer ticks at burst/5, so a 50 ms burst really runs + // 50..60 ms and the achieved duty legitimately overshoots the + // configured one. That must NOT fail. + r.Evidence.ProducerDuty = 0.118 + require.NoError(t, assertArmRanInItsMode(testCfg(), gpuPCArms[2], r)) +} + +// CUDA graphs make Tier A refuse to burst. An arm in such a process ran Tier A +// for part of its length and baseline for the rest, which is not an arm. +func TestATierAArmInAGraphProcessFails(t *testing.T) { + r := goodTierARun() + r.Evidence.ProducerGraphExecs = 3 + err := assertArmRanInItsMode(testCfg(), gpuPCArms[2], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "CUDA graph executions observed") +} + +// Tier B must emit no window at all: a CONTINUOUS producer announcing one +// would be claiming a perturbation it did not cause. +func TestATierBArmThatEmittedAWindowFails(t *testing.T) { + r := schema.GPUPCRun{Evidence: schema.GPUPCEvidence{ + ProducerTier: gpu.PCSamplingNameContinuous, ProducerPCRecords: 900, + PCSamplesDecoded: 900, SamplingWindowsDecoded: 2, ExecutionsSeen: 4000, + SnapshotTier: gpu.PCSamplingNameContinuous, + }} + err := assertArmRanInItsMode(testCfg(), gpuPCArms[1], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "it must do neither") +} + +// A Tier B arm that decoded no PC sample is the baseline with a different +// name, and would report ~0% overhead for a tier that did nothing. +func TestATierBArmThatSampledNothingFails(t *testing.T) { + r := schema.GPUPCRun{Evidence: schema.GPUPCEvidence{ + ProducerTier: gpu.PCSamplingNameContinuous, ExecutionsSeen: 4000, + SnapshotTier: gpu.PCSamplingNameContinuous, + }} + err := assertArmRanInItsMode(testCfg(), gpuPCArms[1], r) + require.Error(t, err) + assert.Contains(t, err.Error(), "baseline with a different name") +} + +// --------------------------------------------------------------------------- +// The workload guards and the cross-arm proof. +// --------------------------------------------------------------------------- + +// A serial workload produces a small, green and meaningless Tier A cost — the +// exact defect the plan rules the 393k launches/s ceiling out for. +func TestASerialBaselineFailsAsAMicrobenchmark(t *testing.T) { + err := assertBaselineIsRealistic(testCfg(), schema.GPUPCArm{ + MedianConcurrency: 1.02, MedianKernelUs: 300}) + require.Error(t, err) + assert.Contains(t, err.Error(), "almost no concurrency for serialization to destroy") +} + +func TestTrivialKernelsFailAsAMicrobenchmark(t *testing.T) { + err := assertBaselineIsRealistic(testCfg(), schema.GPUPCArm{ + MedianConcurrency: 3.4, MedianKernelUs: 3}) + require.Error(t, err) + assert.Contains(t, err.Error(), "launch-rate microbenchmark") +} + +func TestARealisticBaselinePasses(t *testing.T) { + require.NoError(t, assertBaselineIsRealistic(testCfg(), schema.GPUPCArm{ + MedianConcurrency: 3.4, MedianKernelUs: 300})) +} + +// If the duty environment did not take, the three Tier A arms are one arm +// under three names — and their three different ratios are pure fiction. +// Burst count over fixed work is inversely proportional to burst+gap, so a +// lower duty must open strictly fewer bursts. +func TestThreeTierAArmsWithTheSameBurstCountFail(t *testing.T) { + arms := []schema.GPUPCArm{ + {Name: "baseline", Tier: gpu.PCSamplingNameOff}, + {Name: "10%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.100, + Runs: runsWithBursts(41, 41, 41)}, + {Name: "5%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.050, + Runs: runsWithBursts(41, 41, 41)}, + } + err := assertDutyKnobDidSomething(arms) + require.Error(t, err) + assert.Contains(t, err.Error(), "one arm under three names") +} + +func TestDecreasingBurstCountsAcrossDutiesPass(t *testing.T) { + arms := []schema.GPUPCArm{ + {Name: "baseline", Tier: gpu.PCSamplingNameOff}, + {Name: "10%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.100, + Runs: runsWithBursts(48, 49, 48)}, + {Name: "5%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.050, + Runs: runsWithBursts(24, 25, 24)}, + {Name: "2.5%", Tier: gpu.PCSamplingNameSerialized, DutyConfigured: 0.025, + Runs: runsWithBursts(12, 12, 13)}, + } + require.NoError(t, assertDutyKnobDidSomething(arms)) +} + +func runsWithBursts(v ...uint64) []schema.GPUPCRun { + out := make([]schema.GPUPCRun, 0, len(v)) + for i, b := range v { + out = append(out, schema.GPUPCRun{RunN: i + 1, + Evidence: schema.GPUPCEvidence{ProducerBursts: b}}) + } + return out +} + +// --------------------------------------------------------------------------- +// The arm table itself, and the environment it produces. +// --------------------------------------------------------------------------- + +// The three Tier A gaps are not free numbers: they are exactly what the +// adapter's own duty ceiling produces for a 50 ms burst, which is why setting +// the ceiling pins the gap. If either side ever changes, this fails. +func TestTheArmTableIsThePlansAndItsDutiesAreExact(t *testing.T) { + require.Len(t, gpuPCArms, 5) + assert.Equal(t, gpu.PCSamplingOff, gpuPCArms[0].Tier) + assert.Equal(t, gpu.PCSamplingContinuous, gpuPCArms[1].Tier) + for i, want := range []float64{0.10, 0.05, 0.025} { + a := gpuPCArms[2+i] + assert.Equal(t, gpu.PCSamplingSerialized, a.Tier) + assert.Equal(t, 50, a.BurstMs, "the plan's burst length") + assert.InDelta(t, want, a.dutyConfigured(), 1e-9, "arm %q", a.Name) + // min_gap = burst * (1/max_duty - 1), the adapter's own formula. + assert.InDelta(t, float64(a.GapMs), float64(a.BurstMs)*(1/want-1), 1e-9) + } + assert.Zero(t, gpuPCArms[0].dutyConfigured(), "the off arm serializes nothing") + assert.Zero(t, gpuPCArms[1].dutyConfigured(), "tier B serializes nothing") +} + +// The tier is written to the producer's environment EXPLICITLY on every arm +// including the off one. An exported PERFAGENT_GPU_PC_SAMPLING in the +// operator's shell must not turn the baseline into a serializing arm, which +// would make every other arm look free. +func TestEveryArmWritesItsTierExplicitly(t *testing.T) { + cfg := gpuPCConfig{ShimPath: "/tmp/adapter.so"} + for _, spec := range gpuPCArms { + env := gpuPCArmEnv(cfg, spec) + assert.Contains(t, env, gpu.PCSamplingEnvVar+"="+spec.Tier.EnvValue(), + "arm %q must name its tier", spec.Name) + assert.Contains(t, env, "CUDA_INJECTION64_PATH=/tmp/adapter.so") + } +} + +// Both duty knobs, and that is the point: the ceiling sets the burst +// controller's MINIMUM gap and the max-gap sets its maximum, so the interval +// the loop clamps into collapses to a point and the arm runs at exactly the +// duty it claims. With only the ceiling set, a high pair rate would have the +// loop lengthen the gap and the arm would run at a duty nobody asked for. +func TestTierAArmsPinTheGapFromBothSides(t *testing.T) { + cfg := gpuPCConfig{ShimPath: "x"} + for i, permille := range []int{100, 50, 25} { + env := gpuPCArmEnv(cfg, gpuPCArms[2+i]) + joined := strings.Join(env, " ") + assert.Contains(t, joined, "PERFAGENT_GPU_PC_BURST_MS=50") + assert.Contains(t, joined, + "PERFAGENT_GPU_PC_MAX_DUTY_PERMILLE="+strconv.Itoa(permille)) + assert.Contains(t, joined, + "PERFAGENT_GPU_PC_MAX_GAP_MS="+strconv.Itoa(gpuPCArms[2+i].GapMs)) + } +} + +// The two non-Tier-A arms must set no burst knob at all: a baseline carrying +// Tier A's environment would be a baseline that bursts. +func TestNonTierAArmsCarryNoBurstEnvironment(t *testing.T) { + for _, spec := range gpuPCArms[:2] { + joined := strings.Join(gpuPCArmEnv(gpuPCConfig{ShimPath: "x"}, spec), " ") + assert.NotContains(t, joined, "PERFAGENT_GPU_PC_BURST_MS") + assert.NotContains(t, joined, "PERFAGENT_GPU_PC_MAX_DUTY_PERMILLE") + assert.NotContains(t, joined, "PERFAGENT_GPU_PC_MAX_GAP_MS") + } +} + +// --------------------------------------------------------------------------- +// The concurrency measurement, which is what proves the workload is the +// realistic one rather than a microbenchmark wearing its name. +// --------------------------------------------------------------------------- + +// Four kernels of 100 us each, all overlapping over a 100 us span: the +// concurrency is 4 and the mean duration is 100 us. This is the shape +// cuda_concurrent.cu is built to produce and the shape Tier A destroys. +func TestConcurrencyOfFullyOverlappingKernels(t *testing.T) { + snap := gpu.Snapshot{Executions: []gpu.ExecutionView{ + execAt(1000, 101000), execAt(1000, 101000), + execAt(1000, 101000), execAt(1000, 101000), + }} + c, us := concurrencyOf(snap) + assert.InDelta(t, 4.0, c, 1e-9) + assert.InDelta(t, 100.0, us, 1e-9) +} + +// The same four kernels run back to back instead: concurrency 1. This is the +// value the baseline guard refuses, and it is refused precisely because a +// serial workload gives Tier A nothing to destroy. +func TestConcurrencyOfSerialKernelsIsOne(t *testing.T) { + snap := gpu.Snapshot{Executions: []gpu.ExecutionView{ + execAt(0, 100000), execAt(100000, 200000), + execAt(200000, 300000), execAt(300000, 400000), + }} + c, us := concurrencyOf(snap) + assert.InDelta(t, 1.0, c, 1e-9) + assert.InDelta(t, 100.0, us, 1e-9) +} + +// A snapshot with no executions must report zero rather than dividing by +// zero: a run that produced nothing has no concurrency, and NaN would sail +// straight past the floor comparison. +func TestConcurrencyOfNothingIsZeroNotNaN(t *testing.T) { + c, us := concurrencyOf(gpu.Snapshot{}) + assert.Zero(t, c) + assert.Zero(t, us) +} + +// An inverted or zero-length interval carries no duration and must not be +// counted as one, nor drag the span backwards. +func TestConcurrencyIgnoresDegenerateIntervals(t *testing.T) { + snap := gpu.Snapshot{Executions: []gpu.ExecutionView{ + execAt(1000, 101000), + execAt(5000, 5000), // zero length + execAt(9000, 8000), // inverted + }} + c, us := concurrencyOf(snap) + assert.InDelta(t, 1.0, c, 1e-9) + assert.InDelta(t, 100.0, us, 1e-9) +} + +func execAt(start, end uint64) gpu.ExecutionView { + return gpu.ExecutionView{Exec: gpu.GPUKernelExec{StartNs: start, EndNs: end}} +} + +// --------------------------------------------------------------------------- +// The skip path — the only thing this scenario can prove without hardware, and +// therefore the one thing here that must be proven exhaustively. +// --------------------------------------------------------------------------- + +// All four branches, in order, each reachable. On the implementer's machine +// only the first can ever fire; a skip path whose other branches have never +// been executed is a skip path nobody has read. +func TestSkipReasonsAreReachableAndOrdered(t *testing.T) { + // Two paths that exist, so the last two branches are about the + // arguments and not about the filesystem. + present := t.TempDir() + "/exists" + require.NoError(t, os.WriteFile(present, []byte("x"), 0o600)) + missing := t.TempDir() + "/absent" + + full := gpuPCConfig{ShimPath: present, WorkloadPath: present} + + assert.Contains(t, gpuPCSkipReasonWith(false, false, full), + "missing required capabilities") + assert.Contains(t, gpuPCSkipReasonWith(false, false, full), "CAP_BPF", + "the message must name gpuprobe's own set, not the larger one the other scenarios need") + assert.NotContains(t, gpuPCSkipReasonWith(false, false, full), "CAP_SYS_ADMIN", + "checking for CAP_SYS_ADMIN here would skip on a correctly-capped machine") + + assert.Contains(t, gpuPCSkipReasonWith(true, false, full), "no NVIDIA GPU") + + assert.Contains(t, gpuPCSkipReasonWith(true, true, + gpuPCConfig{ShimPath: missing, WorkloadPath: present}), "make -C shim nvidia") + + assert.Contains(t, gpuPCSkipReasonWith(true, true, + gpuPCConfig{ShimPath: present, WorkloadPath: missing}), "make -C shim nvidia-concurrent") + + assert.Empty(t, gpuPCSkipReasonWith(true, true, full), + "with caps, a GPU and both binaries present, nothing may skip — otherwise the "+ + "scenario would skip on the one machine it exists to run on") +} diff --git a/bench/cmd/scenario/main.go b/bench/cmd/scenario/main.go index 0cf75f89..8644f4a6 100644 --- a/bench/cmd/scenario/main.go +++ b/bench/cmd/scenario/main.go @@ -37,7 +37,7 @@ func modeFromFlag(s string) dwarfagent.Mode { func main() { var ( - scenario = flag.String("scenario", "", "pid-large | system-wide-mixed | self (required)") + scenario = flag.String("scenario", "", "pid-large | system-wide-mixed | self | gpu-pc-overhead (required)") processes = flag.Int("processes", 30, "fleet size for system-wide-mixed") runs = flag.Int("runs", 5, "iterations per scenario") dropCache = flag.Bool("drop-cache", false, "drop page cache between runs (root-only)") @@ -49,6 +49,26 @@ func main() { selfDuration = flag.Duration("self-duration", 10*time.Second, "capture window for each perf-agent in the self scenario") cpuBudget = flag.Float64("cpu-budget", 0, "self scenario: max allowed CPU overhead ratio (agent samples / workload samples); 0 disables the gate") resolutionBudget = flag.Float64("resolution-budget", 0, "self scenario: min allowed kernel-symbol resolution rate; 0 disables the gate") + + // gpu-pc-overhead specific flags. Defaults are the plan's arms + // and the workload sizing the report calibrates; the controller + // tunes only --gpu-rounds and --gpu-iters, and only if the + // calibration pass says the fixed work is mis-sized for the + // device in front of it. + gpuShim = flag.String("gpu-shim", "", "gpu-pc-overhead: the CUPTI adapter .so (default ./shim/libperfagent-gpu-nvidia.so)") + gpuWorkload = flag.String("gpu-workload", "", "gpu-pc-overhead: the concurrent CUDA workload (default ./shim/nvidia/testdata/cuda_concurrent)") + gpuIters = flag.Int("gpu-iters", 20000, "gpu-pc-overhead: timed iterations; each launches --gpu-streams kernels") + gpuWarmup = flag.Int("gpu-warmup", 64, "gpu-pc-overhead: untimed warm-up iterations before the clock starts") + gpuStreams = flag.Int("gpu-streams", 4, "gpu-pc-overhead: concurrent CUDA streams — the concurrency Tier A destroys") + gpuRounds = flag.Int("gpu-rounds", 64000, "gpu-pc-overhead: FMA rounds per direction; this is what sets kernel duration") + gpuBlocks = flag.Int("gpu-blocks", 16, "gpu-pc-overhead: blocks per kernel; deliberately a fraction of the device so kernels co-reside") + gpuThreads = flag.Int("gpu-threads", 256, "gpu-pc-overhead: threads per block") + gpuSyncEvery = flag.Int("gpu-sync-every", 4, "gpu-pc-overhead: device sync every N iterations; bounds queue depth and forces concurrency to refill") + gpuMinConc = flag.Float64("gpu-min-concurrency", 1.5, "gpu-pc-overhead: minimum kernel concurrency the BASELINE arm must show, or the run fails as a microbenchmark") + gpuMinKernelUs = flag.Float64("gpu-min-kernel-us", 50, "gpu-pc-overhead: minimum mean kernel duration the BASELINE arm must show, in microseconds") + gpuMinBursts = flag.Uint64("gpu-min-bursts", 4, "gpu-pc-overhead: minimum bursts every Tier A arm must open for its duty to mean anything") + gpuMinCalSec = flag.Float64("gpu-min-calibration-sec", 10, "gpu-pc-overhead: shortest acceptable uninjected fixed-work time") + gpuMaxCalSec = flag.Float64("gpu-max-calibration-sec", 120, "gpu-pc-overhead: longest acceptable uninjected fixed-work time") ) flag.Parse() @@ -57,6 +77,32 @@ func main() { os.Exit(2) } + // gpu-pc-overhead is handled first and entirely on its own. It needs + // gpuprobe's capability set and NOT the larger one below (checking for + // CAP_SYS_ADMIN would skip on a correctly-capped machine, and a skip + // for the wrong reason is indistinguishable from a skip for the right + // one), it needs a GPU, and it uses none of test/workloads/ — so the + // auto-detect below must not be allowed to exit(2) before its own skip + // path has had a chance to speak. + if *scenario == "gpu-pc-overhead" { + gcfg := gpuPCConfig{ + ShimPath: *gpuShim, WorkloadPath: *gpuWorkload, Runs: *runs, + Iters: *gpuIters, Warmup: *gpuWarmup, Streams: *gpuStreams, + Rounds: *gpuRounds, Blocks: *gpuBlocks, Threads: *gpuThreads, + SyncEvery: *gpuSyncEvery, MinConcurrency: *gpuMinConc, + MinKernelUs: *gpuMinKernelUs, MinBursts: *gpuMinBursts, + MinCalibrationSec: *gpuMinCalSec, MaxCalibrationSec: *gpuMaxCalSec, + } + if gcfg.ShimPath == "" { + gcfg.ShimPath = defaultShimPath() + } + if gcfg.WorkloadPath == "" { + gcfg.WorkloadPath = defaultConcurrentWorkload() + } + runGPUPCScenario(gcfg, *outPath) + return + } + // The "self" scenario is pure orchestration — it spawns // perf-agent subprocesses which carry their own file caps. // Other scenarios use dwarfagent.NewProfilerWithMode in-process diff --git a/bench/internal/schema/schema.go b/bench/internal/schema/schema.go index 3153a7ab..75ff4678 100644 --- a/bench/internal/schema/schema.go +++ b/bench/internal/schema/schema.go @@ -20,6 +20,12 @@ type Document struct { System System `json:"system"` StartedAt time.Time `json:"started_at"` Runs []Run `json:"runs"` + + // GPUPC holds the gpu-pc-overhead scenario's whole result. It hangs + // off Document rather than off Run because that scenario's unit of + // result is the ARM (five of them, five interleaved runs each), not + // the run. Nil, and absent from the JSON, for every other scenario. + GPUPC *GPUPCOverhead `json:"gpu_pc_overhead,omitempty"` } type Config struct { @@ -63,16 +69,16 @@ type Run struct { // instead of "0x". A drop = blazesym kernel symbolization // broke (the original v1.2.0 lockdown class of bug). type SelfMetrics struct { - WorkloadPID int `json:"workload_pid"` - AgentPID int `json:"agent_pid"` - WorkloadCPUSamples int `json:"workload_cpu_samples"` - AgentCPUSamples int `json:"agent_cpu_samples"` - CPUOverheadRatio float64 `json:"cpu_overhead_ratio"` - KernelLocationsTotal int `json:"kernel_locations_total"` - KernelLocationsNamed int `json:"kernel_locations_named"` - KernelResolutionRate float64 `json:"kernel_resolution_rate"` - CPUOverheadBudgetMet bool `json:"cpu_overhead_budget_met"` - ResolutionRateBudgetMet bool `json:"resolution_rate_budget_met"` + WorkloadPID int `json:"workload_pid"` + AgentPID int `json:"agent_pid"` + WorkloadCPUSamples int `json:"workload_cpu_samples"` + AgentCPUSamples int `json:"agent_cpu_samples"` + CPUOverheadRatio float64 `json:"cpu_overhead_ratio"` + KernelLocationsTotal int `json:"kernel_locations_total"` + KernelLocationsNamed int `json:"kernel_locations_named"` + KernelResolutionRate float64 `json:"kernel_resolution_rate"` + CPUOverheadBudgetMet bool `json:"cpu_overhead_budget_met"` + ResolutionRateBudgetMet bool `json:"resolution_rate_budget_met"` } type Binary struct { @@ -118,3 +124,178 @@ func Read(r io.Reader) (*Document, error) { } return &d, nil } + +// --------------------------------------------------------------------------- +// The GPU PC-sampling overhead scenario (plan Task 12). +// +// Everything below is additive and every field is omitzero/omitempty, so a +// document from any other scenario is byte-identical to what it was and +// SchemaVersion does not move. +// +// The shape is deliberately "one row per arm, plus the decision", because the +// point of this scenario is not a measurement, it is a DECISION taken against +// thresholds committed before the data existed. A reader who has to compute +// the ratio themselves is a reader who can talk themselves into a different +// answer. +// --------------------------------------------------------------------------- + +// GPUPCOverhead is the whole result of the gpu-pc-overhead scenario: the +// workload it ran, the arms, and the verdict the pre-committed thresholds +// produce from them. +type GPUPCOverhead struct { + // Workload is the fixed work every arm ran, verbatim as configured. + Workload GPUPCWorkload `json:"workload"` + + // Calibration is one UNINJECTED run of the same fixed work, taken + // before the arms. It is NOT the baseline and is never used as one: + // the baseline is the shipping Phase 4 configuration with PC sampling + // off, because §9.1 already measured injection and the activity path + // and those costs are paid whether or not PC sampling is on. This row + // exists to prove the workload is sized sanely and to warm the device + // clocks before the first arm. + Calibration GPUPCArm `json:"calibration,omitzero"` + + // Arms are the five arms in the order they were run, baseline first. + Arms []GPUPCArm `json:"arms"` + + // Decision is what the thresholds say. See GPUPCDecision. + Decision GPUPCDecision `json:"decision"` +} + +// GPUPCWorkload records the fixed work, so a number can never be compared +// against one produced by a differently-sized run. +type GPUPCWorkload struct { + Path string `json:"path"` + Iters int `json:"iters"` + Warmup int `json:"warmup"` + Streams int `json:"streams"` + Rounds int `json:"rounds"` + Blocks int `json:"blocks"` + Threads int `json:"threads"` + SyncEvery int `json:"sync_every"` + KernelsRun int `json:"kernels_run"` +} + +// GPUPCArm is one arm: its configuration, its per-run measurements, the +// medians, and the evidence that it actually ran in the mode it claims. +type GPUPCArm struct { + // Name is the arm as the report table names it. + Name string `json:"name"` + // Tier is the PC-sampling tier as the agent spells it: off, + // continuous or serialized. + Tier string `json:"tier"` + // BurstMs and GapMs are zero except on the Tier A arms. + BurstMs int `json:"burst_ms,omitzero"` + GapMs int `json:"gap_ms,omitzero"` + // DutyConfigured is burst/(burst+gap) as configured: 0.10, 0.05, + // 0.025. It is the denominator the pre-committed thresholds name + // ("Tier A at 10% duty"), so it is what the ratio is computed + // against. DutyAchieved is what the producer reported it actually + // did, and the two disagreeing by more than the tolerance FAILS the + // arm rather than being reported as a smaller ratio. + DutyConfigured float64 `json:"duty_configured,omitzero"` + DutyAchieved float64 `json:"duty_achieved,omitzero"` + + // Runs holds every measurement, not just the median, so the spread + // is visible. A median across five runs whose spread is larger than + // the effect is not a measurement. + Runs []GPUPCRun `json:"runs"` + + // The medians across Runs. WallMs is the workload's own fixed-work + // wall clock (its printed elapsed_ms), which excludes CUDA + // initialization, allocation and the warm-up. + MedianWallMs float64 `json:"median_wall_ms"` + MedianKernelsPerS float64 `json:"median_kernels_per_s"` + // MedianConcurrency is sum(exec duration) / (max end - min start) + // over the executions the profile retained: how many kernels were + // resident at once, on average. It is the property Tier A destroys, + // and the baseline arm's value is asserted against a floor -- a + // workload that turned out to be serial would produce a small, green + // and meaningless Tier A cost. + MedianConcurrency float64 `json:"median_concurrency,omitzero"` + // MedianKernelUs is the mean GPU kernel duration in microseconds. + // Asserted against a floor on the baseline arm: the plan requires + // non-trivial kernel durations, and a microbenchmark's would not be. + MedianKernelUs float64 `json:"median_kernel_us,omitzero"` + + // CostPercent is (this arm's median - baseline's median) / baseline's + // median, as a percentage. Zero on the baseline arm itself. + CostPercent float64 `json:"cost_percent,omitzero"` + // CostOverDuty is CostPercent/100 divided by DutyConfigured. THE + // number that decides Tier A: serialization's damage does not stop + // when the burst does, because concurrency has to refill afterwards. + // Near 1 means duty-cycling works and the tier is tunable to any + // budget; far above 1 means a smaller duty will not rescue it. + CostOverDuty float64 `json:"cost_over_duty,omitzero"` +} + +// GPUPCRun is one execution of one arm. +type GPUPCRun struct { + RunN int `json:"run_n"` + // WallMs is the workload's own fixed-work elapsed time; ProcessMs is + // the whole child process including CUDA init and teardown. Both are + // recorded because a divergence between them means the cost moved + // into startup, which the fixed-work number would hide. + WallMs float64 `json:"wall_ms"` + ProcessMs float64 `json:"process_ms"` + KernelsPerS float64 `json:"kernels_per_s"` + MaxAbsErr float64 `json:"max_abs_err"` + Concurrency float64 `json:"concurrency,omitzero"` + MeanKernelUs float64 `json:"mean_kernel_us,omitzero"` + + // Evidence is what this run PROVED about the mode it ran in, from + // both ends independently: the producer's own report line and the + // consumer's counters. Recorded rather than merely asserted, so a + // reader of the JSON can re-check the assertions the harness made. + Evidence GPUPCEvidence `json:"evidence"` +} + +// GPUPCEvidence is the proof that an arm ran in the mode it says it ran in. +// +// This exists because the standing failure mode on this project is a check +// reading green exactly when things are worst, and its benchmark-shaped +// instance is an arm that did not actually enable the tier it claims. Such an +// arm reports a wonderfully small overhead. Every field here is asserted +// against the arm's expectations and a mismatch fails the whole run. +type GPUPCEvidence struct { + // From the producer's report line on stderr (PERFAGENT_GPU_LOG=stderr). + ProducerTier string `json:"producer_tier"` + ProducerPCRecords uint64 `json:"producer_pc_records"` + ProducerBursts uint64 `json:"producer_bursts,omitzero"` + ProducerWindows uint64 `json:"producer_windows,omitzero"` + ProducerDuty float64 `json:"producer_duty,omitzero"` + ProducerStartFailed uint64 `json:"producer_start_failed,omitzero"` + ProducerStopFailed uint64 `json:"producer_stop_failed,omitzero"` + ProducerGraphExecs uint64 `json:"producer_graph_execs,omitzero"` + ProducerGraphRefuse uint64 `json:"producer_graph_refused,omitzero"` + + // From the consumer: gpuprobe.Stats and gpu.Snapshot. + PCSamplesDecoded uint64 `json:"pc_samples_decoded"` + SamplingWindowsDecoded uint64 `json:"sampling_windows_decoded"` + SamplingWindowsReceived uint64 `json:"sampling_windows_received"` + ExecutionsSeen int `json:"executions_seen"` + ExecutionsSerialized uint64 `json:"executions_serialized"` + ExecutionsNotSerialized uint64 `json:"executions_not_serialized"` + ExecutionsUnknown uint64 `json:"executions_serialization_unknown"` + SnapshotTier string `json:"snapshot_tier"` +} + +// GPUPCDecision is the verdict the pre-committed thresholds produce. A +// threshold decided after seeing the data is not a threshold, so these are +// evaluated mechanically and reported as an outcome, never as a discussion. +type GPUPCDecision struct { + // Thresholds echoes the numbers that were committed in the plan, so + // a reader of the JSON can see they were not adjusted to fit. + MaxWallPercent float64 `json:"max_wall_percent"` + MaxCostOverDut float64 `json:"max_cost_over_duty"` + + // TierB and TierA are the two verdicts, as stable identifiers. See + // the constants in the scenario. + TierB string `json:"tier_b"` + TierA string `json:"tier_a"` + + // Fired lists every threshold clause that fired, in the plan's order. + Fired []string `json:"fired,omitempty"` + // Lines is the human-readable rendering the harness printed. + Lines []string `json:"lines,omitempty"` +} diff --git a/docs/superpowers/plans/2026-08-25-gpu-pc-sampling.md b/docs/superpowers/plans/2026-08-25-gpu-pc-sampling.md index 6e50b94d..a63ae4a8 100644 --- a/docs/superpowers/plans/2026-08-25-gpu-pc-sampling.md +++ b/docs/superpowers/plans/2026-08-25-gpu-pc-sampling.md @@ -547,6 +547,18 @@ Because Tier A perturbs the workload, `serialized` must also be refused unless e Record the numbers in this file when they exist. A threshold decided after seeing the data is not a threshold. +**The harness exists; the numbers do not.** `bench/cmd/scenario --scenario gpu-pc-overhead` implements everything above — the five arms, the interleaving, the medians, the ratio and the four clauses — and prints the verdict as a stable identifier rather than as prose. The workload is `shim/nvidia/testdata/cuda_concurrent.cu` (see `.superpowers/sdd/task-12-overhead-report.md` for why it is a second workload and not a change to `cuda_workload.cu`). On the RTX 3090: + +```bash +make -C shim nvidia nvidia-concurrent && make bench-build +sudo setcap cap_bpf,cap_perfmon,cap_checkpoint_restore+ep ./bench/cmd/scenario/scenario +make bench-gpu-pc-overhead +``` + +**Nothing here has been measured.** The arm table above is still expectations, not results, and the tier decision is not yet made. + +One finding from building it, which is a property of the thresholds rather than of the hardware: `cost ÷ duty > 2` is the same statement as `cost > 200 × duty` percent, so at **2.5% duty it is exactly `cost > 5%`** — the wall-clock bar. On the three duties above, the third clause therefore strictly implies the fourth and *"ships only as a deliberate deep-dive mode"* is unreachable; the unshippable verdict always outranks it. The two clauses separate only below 2.5% duty. The harness reports this whenever both fire rather than letting a reader conclude the deep-dive branch was considered and rejected. Adding a 1%-duty arm would make the distinction real; that is a decision for whoever runs it. + **Verification without a GPU or capabilities — partial.** The scenario harness builds and reports `BENCH_SKIPPED` without caps or GPU, in the shape `bench/cmd/scenario/main.go` already uses. That is all that can be proven offline. **Must be measured on the RTX 3090 afterwards:** every number above. **This task cannot be completed without hardware, and no part of the tier decision may be made without it.** diff --git a/shim/Makefile b/shim/Makefile index 440dc5d0..7de25b69 100644 --- a/shim/Makefile +++ b/shim/Makefile @@ -4,6 +4,8 @@ # test-tsan the concurrency case under ThreadSanitizer (clang++ only) # nvidia libperfagent-gpu-nvidia.so, the CUPTI adapter (needs the CUDA toolkit) # nvidia-workload the CUDA test workload the adapter is proven against +# nvidia-concurrent the CONCURRENT CUDA workload the PC-sampling overhead +# benchmark is measured against (bench/ scenario gpu-pc-overhead) # clean CXX ?= c++ CC ?= cc @@ -24,7 +26,7 @@ CUDA_ARCH ?= sm_86 CORE_SRC := core/batch.cc core/clock.cc core/cubin.cc core/cubinqueue.cc core/drain.cc core/enroll.cc core/sampler.cc CORE_OBJ := $(CORE_SRC:.cc=.o) -.PHONY: all test test-tsan check-cubin-defer nvidia nvidia-workload clean +.PHONY: all test test-tsan check-cubin-defer nvidia nvidia-workload nvidia-concurrent clean all: libperfagent-gpu-core.a libperfagent-gpu-core.a: $(CORE_OBJ) @@ -177,6 +179,21 @@ nvidia/testdata/cuda_workload: nvidia/testdata/cuda_workload.cu $(NVCC) -O2 -g -lineinfo -arch=$(CUDA_ARCH) \ -Xcompiler -fno-omit-frame-pointer -Xcompiler -rdynamic -o $@ $< +# The concurrent workload the PC-sampling overhead benchmark measures against. +# Same flags as cuda_workload for the same reasons (-lineinfo so a PC sample +# taken during a burst resolves to a line of it; -g and -rdynamic so a sampled +# launch's CPU stack names the submitting function; not stripped). +# +# It is a SECOND workload rather than a change to cuda_workload.cu because the +# two answer different questions: cuda_workload is the serial fixture the +# adapter, the join and the source resolution are proven against and the +# hardware gate asserts against ITS source lines, while a serialization +# measurement needs concurrency to destroy. See the header of the .cu. +nvidia-concurrent: nvidia/testdata/cuda_concurrent +nvidia/testdata/cuda_concurrent: nvidia/testdata/cuda_concurrent.cu + $(NVCC) -O2 -g -lineinfo -arch=$(CUDA_ARCH) \ + -Xcompiler -fno-omit-frame-pointer -Xcompiler -rdynamic -o $@ $< + # Task 5's structural assertion, run as a BUILD step rather than asserted in a # comment: the copy of a cubin must happen inside the vendor callback, so no # deferral may sit between callback entry and the memcpy. core/cubinqueue.h @@ -278,4 +295,5 @@ clean: perfagent-gpu-dlopen-host libperfagent-gpu-fpless.so \ stub/so_fpless_bridge.o stub/so_fpless_stub.o \ perfagent-gpu-fpless stub/fpless_bridge.o stub/fpless_main.o stub/fpless_stub.o \ - libperfagent-gpu-nvidia.so nvidia/testdata/cuda_workload + libperfagent-gpu-nvidia.so nvidia/testdata/cuda_workload \ + nvidia/testdata/cuda_concurrent diff --git a/shim/nvidia/testdata/cuda_concurrent.cu b/shim/nvidia/testdata/cuda_concurrent.cu new file mode 100644 index 00000000..a5635e5a --- /dev/null +++ b/shim/nvidia/testdata/cuda_concurrent.cu @@ -0,0 +1,287 @@ +// The CONCURRENT CUDA workload the PC-sampling overhead benchmark is measured +// against (plan Task 12). It is a second workload, not a modification of +// cuda_workload.cu, and the reason is the whole point of this file. +// +// Why not extend cuda_workload.cu +// ------------------------------- +// cuda_workload.cu is a two-kernel SERIAL loop on the default stream with a +// usleep between iterations. Its kernels are ~64k elements of one FMA each -- +// microseconds -- and nothing ever overlaps anything. That shape is exactly +// right for what it exists for: it is the fixture the adapter, the join, the +// kernel-name table and the -lineinfo source resolution are all proven +// against, and the hardware half of the phase gate asserts against ITS source +// lines. Changing its shape would change what those gates measure. +// +// It is also exactly the WRONG instrument for measuring Tier A. The plan says +// so directly: a saturating stream of trivial kernels exaggerates per-launch +// costs and UNDERSTATES serialization costs, because serialization hurts in +// proportion to the concurrency it destroys and a serial loop has none to +// destroy. A benchmark run on cuda_workload.cu would report a small Tier A +// number for the same reason a stopped clock reports the right time twice a +// day. +// +// What this workload does instead +// ------------------------------- +// S independent CUDA streams, each carrying its own buffer, each iteration +// launching one non-trivial kernel per stream with NO synchronization between +// the streams. The kernels are sized to occupy a FRACTION of the device +// (blocks x threads well under the SM count x occupancy), so several of them +// genuinely co-reside -- which is the property Tier A's KERNEL_SERIALIZED +// mode destroys, and therefore the property the measurement needs to have in +// the first place. +// +// The concurrency is not asserted here; it is MEASURED, by the benchmark, out +// of the profile the adapter produces: +// +// concurrency = sum(exec duration) / (max end - min start) +// +// and the benchmark refuses to report any overhead number if that comes out +// near 1 on the baseline arm. A workload that turned out to be serial anyway +// would otherwise produce a small, green, meaningless Tier A cost. +// +// A periodic device sync every --sync-every iterations is deliberate, not +// housekeeping: it bounds queue depth, and it makes the pipeline REFILL after +// every drain. Refill cost is precisely what makes Tier A's damage outlast its +// burst, which is what the cost-over-duty ratio exists to detect. +// +// The exact self-check +// -------------------- +// Each kernel adds +1 to every element `rounds` times and then adds -1 to +// every element `rounds` times. With rounds <= 2^23 every partial sum is +// exactly representable in float, so the composition is the IDENTITY: every +// element must come back bit-exact. `up` and `down` are runtime kernel +// arguments, so the compiler cannot fold the loops away, and `#pragma unroll +// 1` keeps them from being unrolled into a reassociated sum. +// +// max_abs_err must therefore be exactly 0.0, not "small". An arm that +// perturbed the computation fails the run rather than contributing a number. +// +// Built with -lineinfo -g like cuda_workload, so a PC sample taken during a +// burst resolves to a line of this file rather than to nothing. +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#define CHECK(call) \ + do { \ + const cudaError_t _e = (call); \ + if (_e != cudaSuccess) { \ + fprintf(stderr, "concurrent: %s failed: %s\n", #call, \ + cudaGetErrorString(_e)); \ + exit(2); \ + } \ + } while (0) + +// One kernel: `rounds` dependent FMAs up, then `rounds` dependent FMAs down. +// +// The dependency chain is what gives the kernel a duration measured in +// hundreds of microseconds out of a tiny grid -- the alternative (a huge grid +// of trivial work) would saturate the device and leave nothing for a second +// stream to overlap with, which would quietly turn this back into the serial +// workload it exists not to be. +__global__ void perfagent_conc_chain(float *buf, int n, int rounds, float up, + float down) { + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= n) return; + float v = buf[i]; +#pragma unroll 1 + for (int r = 0; r < rounds; r++) v = fmaf(v, 1.0f, up); +#pragma unroll 1 + for (int r = 0; r < rounds; r++) v = fmaf(v, 1.0f, down); + buf[i] = v; +} + +// noinline so a sampled launch's captured CPU stack has a frame naming the +// submitting function rather than one inlined blob of the loop body -- the +// same reason cuda_workload.cu marks its two launchers noinline. +static __attribute__((noinline)) void launch_chain(float *buf, int n, int rounds, + int blocks, int threads, + cudaStream_t s) { + perfagent_conc_chain<<>>(buf, n, rounds, 1.0f, -1.0f); +} + +static uint64_t mono_ns() { + struct timespec t; + clock_gettime(CLOCK_MONOTONIC, &t); + return (uint64_t)t.tv_sec * 1000000000ULL + (uint64_t)t.tv_nsec; +} + +// Same release contract as cuda_workload.cu and the stub: a sampled launch's +// CPU stack is symbolized against /proc//maps of THIS process, which the +// kernel destroys the instant it exits. The consumer releases us by closing +// our stdin; linger_ms is the backstop for a run by hand from a terminal. +// The overhead benchmark passes 0 -- it needs no stacks -- but a run by hand +// under cmd/gpu-cuda-profile does. +static void linger(unsigned linger_ms) { + if (!linger_ms) return; + struct pollfd p{}; + p.fd = STDIN_FILENO; + p.events = POLLIN; + const uint64_t deadline = mono_ns() + (uint64_t)linger_ms * 1000000ULL; + for (;;) { + const uint64_t now = mono_ns(); + if (now >= deadline) return; + const int left = (int)((deadline - now) / 1000000ULL) + 1; + const int rc = poll(&p, 1, left); + if (rc < 0) { + if (errno == EINTR) continue; + return; + } + if (rc == 0) return; + char buf[256]; + const ssize_t got = read(STDIN_FILENO, buf, sizeof(buf)); + if (got <= 0) return; + } +} + +struct Opts { + unsigned iters = 20000; // timed iterations; each launches `streams` kernels + unsigned warmup = 64; // untimed iterations before the clock starts + unsigned streams = 4; // concurrent streams == concurrent kernels + unsigned rounds = 64000; // FMA rounds per direction; sets kernel duration + unsigned blocks = 16; // grid; deliberately a FRACTION of the device + unsigned threads = 256; + unsigned sync_every = 4; // device sync every N iterations; bounds queue depth + unsigned linger_ms = 0; +}; + +static bool parse_uint(const char *arg, const char *name, unsigned *out) { + const size_t n = strlen(name); + if (strncmp(arg, name, n) != 0 || arg[n] != '=') return false; + char *end = nullptr; + const unsigned long v = strtoul(arg + n + 1, &end, 10); + if (end == arg + n + 1 || (end && *end)) { + fprintf(stderr, "concurrent: %s: not a number\n", arg); + exit(2); + } + *out = (unsigned)v; + return true; +} + +int main(int argc, char **argv) { + Opts o; + for (int i = 1; i < argc; i++) { + const char *a = argv[i]; + if (parse_uint(a, "--iters", &o.iters)) continue; + if (parse_uint(a, "--warmup", &o.warmup)) continue; + if (parse_uint(a, "--streams", &o.streams)) continue; + if (parse_uint(a, "--rounds", &o.rounds)) continue; + if (parse_uint(a, "--blocks", &o.blocks)) continue; + if (parse_uint(a, "--threads", &o.threads)) continue; + if (parse_uint(a, "--sync-every", &o.sync_every)) continue; + if (parse_uint(a, "--linger-ms", &o.linger_ms)) continue; + fprintf(stderr, + "concurrent: unknown argument %s\n" + "usage: cuda_concurrent [--iters=N] [--warmup=N] [--streams=N] " + "[--rounds=N] [--blocks=N] [--threads=N] [--sync-every=N] " + "[--linger-ms=N]\n", + a); + return 2; + } + if (!o.iters || !o.streams || !o.rounds || !o.blocks || !o.threads) { + fprintf(stderr, "concurrent: iters, streams, rounds, blocks and threads " + "must all be positive\n"); + return 2; + } + if (!o.sync_every) o.sync_every = 1; + // 2^23: past this the +1/-1 partial sums stop being exactly representable + // in float and the identity check below would start reporting a rounding + // artefact as a corrupted computation. + if (o.rounds > (1u << 23)) { + fprintf(stderr, "concurrent: --rounds=%u exceeds 2^23; the exact " + "identity check would no longer hold\n", o.rounds); + return 2; + } + + const int n = (int)(o.blocks * o.threads); + const size_t bytes = (size_t)n * sizeof(float); + + float **d = (float **)calloc(o.streams, sizeof(float *)); + cudaStream_t *s = (cudaStream_t *)calloc(o.streams, sizeof(cudaStream_t)); + if (!d || !s) { fprintf(stderr, "concurrent: out of memory\n"); return 2; } + + float *h = (float *)malloc(bytes); + if (!h) { fprintf(stderr, "concurrent: out of memory\n"); return 2; } + for (int i = 0; i < n; i++) h[i] = 1.0f; + + for (unsigned k = 0; k < o.streams; k++) { + // Non-blocking so a stream never implicitly synchronizes against the + // legacy default stream. With the default (blocking) flag every one of + // these streams would serialize against stream 0 and the workload + // would have no concurrency at all -- which is the exact failure this + // file exists to avoid, so it is spelled out rather than defaulted. + CHECK(cudaStreamCreateWithFlags(&s[k], cudaStreamNonBlocking)); + CHECK(cudaMalloc(&d[k], bytes)); + CHECK(cudaMemcpy(d[k], h, bytes, cudaMemcpyHostToDevice)); + } + + // Warm-up, OUTSIDE the timed region: module load, JIT, first-touch + // allocation and the driver's own lazy initialization all land here rather + // than in the fixed-work measurement. It also gives the adapter's drain + // timer and (in Tier A) its burst controller a few cycles to reach steady + // state before the clock starts. + for (unsigned i = 0; i < o.warmup; i++) + for (unsigned k = 0; k < o.streams; k++) + launch_chain(d[k], n, (int)o.rounds, (int)o.blocks, (int)o.threads, s[k]); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaGetLastError()); + + // FIXED WORK, not fixed time: iters x streams kernels, whatever it takes. + // The wall clock over this region is the number the benchmark compares + // across arms. + const uint64_t t0 = mono_ns(); + for (unsigned i = 0; i < o.iters; i++) { + for (unsigned k = 0; k < o.streams; k++) + launch_chain(d[k], n, (int)o.rounds, (int)o.blocks, (int)o.threads, s[k]); + if ((i % o.sync_every) == o.sync_every - 1) CHECK(cudaDeviceSynchronize()); + } + CHECK(cudaDeviceSynchronize()); + const uint64_t t1 = mono_ns(); + CHECK(cudaGetLastError()); + + const double elapsed_ms = (double)(t1 - t0) / 1e6; + const unsigned long long kernels = (unsigned long long)o.iters * o.streams; + + // Exactly zero, not "small". Every kernel is the identity on its buffer, + // so any non-zero here means the computation was corrupted and this run's + // timing must not be used. + double max_abs_err = 0.0; + for (unsigned k = 0; k < o.streams; k++) { + CHECK(cudaMemcpy(h, d[k], bytes, cudaMemcpyDeviceToHost)); + for (int i = 0; i < n; i++) { + const double e = (double)h[i] - 1.0; + const double a = e < 0 ? -e : e; + if (a > max_abs_err) max_abs_err = a; + } + } + + printf("concurrent: iters=%u warmup=%u streams=%u rounds=%u blocks=%u " + "threads=%u sync_every=%u kernels=%llu elapsed_ms=%.3f " + "kernels_per_s=%.1f max_abs_err=%.9f\n", + o.iters, o.warmup, o.streams, o.rounds, o.blocks, o.threads, + o.sync_every, kernels, elapsed_ms, + elapsed_ms > 0 ? (double)kernels * 1000.0 / elapsed_ms : 0.0, + max_abs_err); + fflush(stdout); + + linger(o.linger_ms); + + for (unsigned k = 0; k < o.streams; k++) { + CHECK(cudaFree(d[k])); + CHECK(cudaStreamDestroy(s[k])); + } + free(h); + free(d); + free(s); + // Not cudaDeviceReset(), for the reason cuda_workload.cu gives: it tears + // the context down before the adapter's atexit flush runs and the last + // activity records go with it. + return max_abs_err == 0.0 ? 0 : 1; +}