From 7eff378a08c7975b7355a75c4bf480492d9616ed Mon Sep 17 00:00:00 2001 From: Phil-amd Date: Mon, 17 Aug 2026 11:52:17 +0800 Subject: [PATCH] [Feature] Add architecture-general ISA resource diff tool Compare per-kernel register, spill, scratch, and LDS usage between two FLYDSL_DUMP_IR dump directories or JSON snapshots, exposing resource regressions that functional tests do not surface. To work outside CDNA, read register counts from the per-kernel `.set .num_vgpr`/`.num_agpr` symbols that LLVM emits on every AMDGPU target, rather than the CDNA-only `.agpr_count` metadata field, and count LDS traffic under both the `ds_read` and the gfx11+ `ds_load` spelling. Take the processor and the feature set from the target ID, normalizing the triple's environment field, which is spelled either empty or `unknown` for one and the same target. Report each metric as a value, as not applicable, or as unparsed, and exit 0, 1, or 2 for no regression, a regression, or an untrustworthy result. Fail closed on anything that would otherwise answer from a partial comparison: an unparsed or impossible metric, a dump file that does not parse or decode, a kernel entry with no identity or a duplicated one, and two sides whose targets are not provably the same. Covered by a backend-agnostic test over all three parser axes and exposed to agents as the `isa-resource-diff` skill. --- .claude/skills/isa-resource-diff/SKILL.md | 221 +++ .claude/skills/kernel-trace-analysis/SKILL.md | 5 + docs/kernel_tuning_guide.md | 23 + docs/testing_benchmarking_guide.md | 1 + scripts/isa_resource_table.py | 1331 +++++++++++++++++ tests/unit/test_isa_resource_table.py | 378 +++++ 6 files changed, 1959 insertions(+) create mode 100644 .claude/skills/isa-resource-diff/SKILL.md create mode 100644 scripts/isa_resource_table.py create mode 100644 tests/unit/test_isa_resource_table.py diff --git a/.claude/skills/isa-resource-diff/SKILL.md b/.claude/skills/isa-resource-diff/SKILL.md new file mode 100644 index 000000000..0592c14c6 --- /dev/null +++ b/.claude/skills/isa-resource-diff/SKILL.md @@ -0,0 +1,221 @@ +--- +name: isa-resource-diff +description: > + Detect per-kernel GPU resource regressions (VGPR, SGPR, register spills, scratch, + static LDS) by diffing the final ISA before and after a change, using + scripts/isa_resource_table.py. Compile-only: needs no GPU and no profiler run, so + it works on any target the compiler supports and runs in seconds. Use when asked + whether a change increased register pressure, caused spilling, or hurt occupancy, + when reviewing a kernel change for resource impact, or as a fast pre-check before + spending a profiling run. + Usage: /isa-resource-diff [] [--arch ] +allowed-tools: Read Write Bash Grep Glob +--- + +# ISA Resource Diff + +Compare per-kernel register, spill, scratch, and LDS usage between two builds to +catch resource regressions that functional tests do not surface. + +## Pick the right skill first + +| Question | Skill | +|---|---| +| Did my change increase registers / cause spills / grow LDS? | **this skill** (compile-only, seconds, no GPU) | +| *Why* is this kernel slow — which instructions stall, and on what? | `/kernel-trace-analysis` (needs a GPU run + rocprofv3 ATT trace) | +| Which commit made it slow? | `/bisect-perf-regression` (needs a runnable benchmark) | +| How do I collect a trace at all? | `/capture-kernel-trace` | + +This skill measures **resources, not time**. A clean result here does not mean +performance is unchanged — it means register/LDS/spill pressure is unchanged. +A regression here is a strong, cheap signal that is usually worth acting on +before profiling, because spilling and occupancy cliffs dominate most kernel +slowdowns. See §7 of `docs/kernel_tuning_guide.md` for what to do about one. + +## Arguments + +| Argument | Required | Description | +|---|---|---| +| `` | No | What to run to produce dumps. Defaults to asking the user. Example: `pytest tests/kernels/test_softmax.py -q` | +| `--arch ` | No | Target for compile-only runs, e.g. `gfx950`. Omit to use local hardware | + +If the user already has two dump directories or two JSON snapshots, skip to Step 3. + +## Workflow + +### Step 1 — Capture the "before" side + +Check out or stash to the baseline state first, then: + +```bash +FLYDSL_DUMP_IR=1 FLYDSL_DUMP_DIR=/tmp/isa-before FLYDSL_RUNTIME_ENABLE_CACHE=0 \ + python3 -m pytest tests/kernels/test_softmax.py -q +``` + +`FLYDSL_RUNTIME_ENABLE_CACHE=0` is **required, not optional** — see Pitfalls. + +For a target without local hardware, add `ARCH= COMPILE_ONLY=1`: + +```bash +ARCH=gfx950 COMPILE_ONLY=1 \ +FLYDSL_DUMP_IR=1 FLYDSL_DUMP_DIR=/tmp/isa-before FLYDSL_RUNTIME_ENABLE_CACHE=0 \ + python3 -m pytest tests/kernels/test_softmax.py -q +``` + +### Step 2 — Capture the "after" side + +Apply the change, then rerun **the identical command** into a *fresh* directory +(`/tmp/isa-after`). Same test, same parameters, same arch, same cache setting. + +### Step 3 — Diff + +```bash +python3 scripts/isa_resource_table.py diff /tmp/isa-before /tmp/isa-after +``` + +The tool requires **Python 3.10+**. If `python3` is older it exits 2 with a clear +message; use `python3.10 scripts/isa_resource_table.py …` instead. + +Both sides may independently be a dump directory or a `.json` snapshot, so a +baseline can be captured once and reused: + +```bash +python3 scripts/isa_resource_table.py summarize /tmp/isa-before --json baseline.json +python3 scripts/isa_resource_table.py diff baseline.json /tmp/isa-after +``` + +## Reading the output + +``` +* = regression trigger; other columns are informational. + vgpr = total (arch+acc, LLVM's occupancy number); arch_vgpr/agpr are its split -- do not add them. +kernel *vgpr arch_vgpr agpr ... *lds_static_bytes lds_read ... +------------------------------------------------------------------------------------------------- +gemm::d128_fmha_fwd_kernel_0 942->960(+18) 942->960(+18) 0 ... 212992->229376(+16384) 12 + +compared 1 of 1 kernels; 0 unchanged; 1 changed; worsened: 2; improved: 0 +RESULT: REGRESSION +``` + +Only changed and problematic kernels are printed. The last two stdout lines are +always a count line and a `RESULT:` verdict that matches the exit code exactly. + +**Columns marked `*` are regression triggers**; the rest are context. + +| Column | Trigger | Read from | What it is | +|---|---|---|---| +| `vgpr` | yes | `.vgpr_count` metadata | Total VGPRs, arch + accumulator — LLVM's own occupancy number | +| `arch_vgpr` | no | `.set` symbol `num_vgpr` | The arch half of that total | +| `agpr` | no | `.set` symbol `num_agpr` | The accumulator half of that total | +| `sgpr` | yes | `.sgpr_count` metadata | SGPRs including the fixed extras (VCC, XNACK, FLAT_SCRATCH) | +| `numbered_sgpr` | no | `.set` symbol `numbered_sgpr` | SGPRs without those extras; tells a real increase from VCC becoming live | +| `vgpr_spill` | yes | `.vgpr_spill_count` metadata | VGPRs the allocator spilled | +| `sgpr_spill` | yes | `.sgpr_spill_count` metadata | SGPRs the allocator spilled | +| `scratch_bytes` | yes | `.private_segment_fixed_size` metadata | Private segment per work-item; exact on every target | +| `lds_static_bytes` | yes | `.group_segment_fixed_size` metadata | Statically allocated LDS per work-group | +| `lds_read` / `lds_write` | no | instruction count | `ds_read`/`ds_load` and `ds_write`/`ds_store` sites | +| `scratch_store` / `scratch_load` | no | instruction count | `scratch_*` sites; `n/a` where spilling goes through `buffer_*` | +| `matrix_ops` | no | instruction count | MFMA / WMMA sites | + +Three things are easy to misread: + +- **Do not add `arch_vgpr` and `agpr` to `vgpr`.** `vgpr` is already the total + (arch + accumulator) and is the only VGPR-family trigger. The other two are its + split, shown so you can tell *which half* moved. Moving accumulators into AGPRs + — which the tuning guide recommends — deliberately does not count as a regression. +- **`n/a` is not `0`.** It means the quantity does not exist on this target, e.g. + `scratch_store`/`scratch_load` on a target that spills through `buffer_*`. Use + `scratch_bytes`, which is exact everywhere. +- **`?` means unparsed** — the tool could not read something it reports. Any `?` + forces exit 2. Never treat it as unchanged. + +### Exit codes + +| Code | stdout | Meaning | What an agent should do | +|---|---|---|---| +| `0` | `RESULT: OK` | Everything comparable, no trigger increased | Proceed | +| `1` | `RESULT: REGRESSION` | Everything comparable, a trigger increased | Investigate — this is a real finding | +| `2` | `RESULT: NOT TRUSTWORTHY` | The tool cannot answer | **Fix the inputs and rerun. Do not report "no regression"** | + +Exit `1` is a claim about the code under test; exit `2` is a claim about the +tool's own confidence. Everything that would leave the answer partial reports `2` +and never `1`: a crash, an empty dump directory, a dump file that does not parse +or does not decode, a metadata entry with no kernel identity or a duplicated one, +a negative resource count, a kernel present on only one side, and a target that +differs — or that the tool cannot name — on either side. + +For scripting, `-q/--quiet` prints only the verdict line, and `--json PATH` +writes the full comparison as machine-readable JSON: + +```bash +python3 scripts/isa_resource_table.py diff /tmp/isa-before /tmp/isa-after -q --json report.json +case $? in + 0) echo "no resource regression" ;; + 1) echo "REGRESSION — see report.json" ;; + *) echo "inconclusive — inputs are bad, do not claim a clean result" ;; +esac +``` + +## Acting on a regression + +Map the column that moved to a cause, then follow `docs/kernel_tuning_guide.md`: + +| Column increased | Usual cause | +|---|---| +| `vgpr` | More live values; larger tiles or deeper pipelining/unrolling | +| `vgpr_spill` / `sgpr_spill` / `scratch_bytes` | Register pressure crossed the budget — normally the most damaging of these signals | +| `lds_static_bytes` | Bigger shared tiles or added double buffering; may cross an occupancy step | +| `sgpr` alone, by 2, with `numbered_sgpr` flat | Usually just VCC becoming live — rarely meaningful | + +If a resource regression is confirmed but the kernel is not actually slower, +say so rather than "fixing" it: these are proxies for occupancy, not timings. +Confirm with `/kernel-trace-analysis` before reworking a kernel. + +## Pitfalls + +These silently produce a *confident wrong answer* if ignored: + +- **Dump directories are keyed by kernel name only, with no specialization key.** + Every JIT specialization of one kernel writes to the same directory and + overwrites the previous one, so a parametrized test leaves only its last + variant. Diff one shape at a time when the answer must be exact. +- **A cache hit produces no dump at all.** Hence `FLYDSL_RUNTIME_ENABLE_CACHE=0` + on both sides. Without it a kernel can silently vanish from one side, which + the tool reports as `ONLY IN BEFORE` and exit 2. +- **`lds_static_bytes` is static LDS only.** A kernel using + `SharedAllocator(static=False)` reports `0` no matter how much LDS it takes at + dispatch; the tool prints a `dynamic LDS in use` note for it. LDS regressions + in those kernels are invisible here. +- **The stage number in `NN_final_isa.s` varies between runs.** Nothing cleans the + dump directory, so reusing one can leave two files side by side. The tool uses + the highest-numbered one and warns; prefer a fresh directory per run. +- **Diffs across targets are refused** (exit 2). Register files and LDS banking + differ between architectures, and `xnack`/`sramecc` change code generation + within one, so both sides must report the same processor *and* the same target + features. A target the tool cannot name is refused for the same reason. The + triple's environment field is normalized, so `amdgcn-amd-amdhsa--gfx942` and + `amdgcn-amd-amdhsa-unknown-gfx942` are the same target. +- **Warnings on stderr never change the exit code.** They describe the input tree + and are worth reading before trusting a clean result. + +## Verifying the tool itself + +`tests/unit/test_isa_resource_table.py` is backend-agnostic and needs no build: + +```bash +python3 -m pytest tests/unit/test_isa_resource_table.py -q +``` + +It generates its ISA input in `make_isa()` rather than checking dumps in, so the +test file states exactly which parts of LLVM's output the parser relies on. The +shapes cover both sides of each axis that has already broken it once: whether the +target emits `.agpr_count`, which LDS mnemonic spelling it uses, and whether the +target ID spells the triple's environment as empty or as `unknown`. The rest of +the file pins the fail-closed verdicts: a mismatched processor, a mismatched +feature set, an unnameable target, an unreadable or undecodable dump file, a +kernel entry with no identity or a duplicated one, and a negative count all have +to reach exit 2. + +If the tool reports `?` on a dump that looks healthy, LLVM's assembly format has +probably drifted — update `make_isa()` to match the new shape rather than +loosening the parser, since the failure is deliberately loud. diff --git a/.claude/skills/kernel-trace-analysis/SKILL.md b/.claude/skills/kernel-trace-analysis/SKILL.md index c67447ceb..1be0d7776 100644 --- a/.claude/skills/kernel-trace-analysis/SKILL.md +++ b/.claude/skills/kernel-trace-analysis/SKILL.md @@ -18,6 +18,11 @@ an optimization plan. All analysis is done programmatically via `hotspot_analyzer.py` + `code.json`. Do **not** use GUI tools. +> If the question is only "did my change increase register pressure, spills, or +> LDS?", use `/isa-resource-diff` first — it is compile-only, needs no GPU or +> profiler run, and answers in seconds. Come here when you need to know *why* a +> kernel is slow rather than *what resources it uses*. + ## Arguments | Argument | Description | diff --git a/docs/kernel_tuning_guide.md b/docs/kernel_tuning_guide.md index dd0362dd6..28622d623 100644 --- a/docs/kernel_tuning_guide.md +++ b/docs/kernel_tuning_guide.md @@ -445,6 +445,29 @@ WHERE ks.KernelName LIKE '%target_kernel%' LIMIT 5; the binding limiter, so you know whether to cut VGPR, shrink LDS per block, or reduce SGPR pressure. +For a lightweight before/after check without collecting a profiler trace, dump +the final ISA of each build and diff the per-kernel **total VGPR (with its +arch/accumulator split), SGPR, register spills, scratch bytes, and static LDS** +with `scripts/isa_resource_table.py`. It is compile-only, so it needs no GPU: + +```bash +FLYDSL_DUMP_IR=1 FLYDSL_DUMP_DIR=/tmp/isa-before FLYDSL_RUNTIME_ENABLE_CACHE=0 \ + python3 -m pytest tests/kernels/test_softmax.py -q # repeat for /tmp/isa-after +python3 scripts/isa_resource_table.py diff /tmp/isa-before /tmp/isa-after +``` + +It exits `0` for no regression, `1` for one, and `2` when it cannot produce a +trustworthy answer — a crash never reports as `1`. The table marks regression +triggers with `*` and prints a legend above itself; `/isa-resource-diff` carries +the full column reference. Set `FLYDSL_RUNTIME_ENABLE_CACHE=0` on both runs: a +cache hit emits no dump at all, which would silently drop a kernel from one side. + +Because arch VGPRs and accumulator VGPRs share one register budget on targets +with a unified register file (gfx90a and later MFMA-capable parts), that +tool treats the **total** as the regression signal and reports `arch_vgpr` and +`agpr` alongside it for information only. Moving accumulators into AGPRs, as +recommended below, therefore does **not** register as a resource regression. + **Do not** use `maxnreg` to force `accum_vgpr=0` — it spills MFMA results through arch_vgpr via `v_accvgpr_read` (measured ~4.5× regression). diff --git a/docs/testing_benchmarking_guide.md b/docs/testing_benchmarking_guide.md index 82b9cf93b..a0f32363c 100644 --- a/docs/testing_benchmarking_guide.md +++ b/docs/testing_benchmarking_guide.md @@ -395,6 +395,7 @@ bash scripts/dumpir.sh | `scripts/run_tests.sh` | Full test runner (pytest + examples + FileCheck) | | `scripts/run_benchmark.sh` | Benchmark harness with configurable shapes | | `scripts/dumpir.sh` | IR dump helper script | +| `scripts/isa_resource_table.py` | Diff per-kernel ISA resource usage (compile-only) | | `tests/conftest.py` | Pytest fixtures (MLIR context, module, insert point) | | `tests/test_common.py` | `perftest()`, `checkAllclose()`, `verify_output()` | | `tests/utils.py` | `pertoken_quant()`, `shuffle_weight()` | diff --git a/scripts/isa_resource_table.py b/scripts/isa_resource_table.py new file mode 100644 index 000000000..10a29f05e --- /dev/null +++ b/scripts/isa_resource_table.py @@ -0,0 +1,1331 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 FlyDSL Project Contributors +"""Summarize and compare per-kernel resources in FlyDSL ISA dumps. + +The tool is self-contained and only consumes ``*final_isa.s`` files produced by +``FLYDSL_DUMP_IR``: + + python3 scripts/isa_resource_table.py diff /tmp/before /tmp/after + python3 scripts/isa_resource_table.py summarize --json snapshot.json + python3 scripts/isa_resource_table.py diff before.json after.json + +``diff`` accepts either two dump directories or two JSON snapshots, and the two +sides may be mixed. Every metric is scoped to one kernel, even when an ISA file +contains multiple kernels. + +The tool is architecture-general. Register counts come from the per-kernel +``.set .num_vgpr`` / ``.num_agpr`` / ``.numbered_sgpr`` assembler symbols, +which LLVM emits on every AMDGPU target, rather than from the ``.agpr_count`` +metadata field, which LLVM emits only for MFMA-capable targets. LDS traffic is +counted under either mnemonic spelling (``ds_read``/``ds_write`` on gfx9 and gfx10, +``ds_load``/``ds_store`` on gfx11 and later). + +Each metric cell is one of three states: a parsed value, ``n/a`` when the quantity +does not exist on the target, or ``?`` when the tool could not read something it +claims to report. Only the third is a failure: the tool exits 1 for a resource +regression and 2 whenever it cannot produce a trustworthy answer. +""" + +import sys + +# The whole module is parsed before anything runs, so this guard only protects +# against *runtime* use of newer syntax -- which is why the file deliberately +# avoids `from __future__ import annotations`, the walrus operator and `match`. +# PEP 585/604 annotations below are never evaluated on an old interpreter. +if sys.version_info < (3, 10): # pragma: no cover - exercised via subprocess + sys.stderr.write( + "scripts/isa_resource_table.py requires Python 3.10+ (running %s).\n" + "This repository targets 3.10+ (CONTRIBUTING.md; ruff target-version=py310).\n" + "Try: python3.10 scripts/isa_resource_table.py ...\n" % sys.version.split()[0] + ) + raise SystemExit(2) + +import argparse # noqa: E402 - imports must follow the interpreter guard +import json # noqa: E402 +import re # noqa: E402 +import traceback # noqa: E402 +from bisect import bisect_left # noqa: E402 +from dataclasses import dataclass, field, replace # noqa: E402 +from pathlib import Path # noqa: E402 +from typing import TextIO # noqa: E402 + +SCHEMA_VERSION = 2 +TOOL_NAME = "isa_resource_table" + + +class SnapshotError(ValueError): + """An ISA snapshot cannot be parsed or compared safely.""" + + +# -------------------------------------------------------------------------------------- +# Metric cells: a value, "not applicable here", or "could not read it" +# -------------------------------------------------------------------------------------- + +VALUE = "value" +NA = "n/a" +UNPARSED = "unparsed" + + +@dataclass(frozen=True) +class Cell: + """One metric for one kernel, in one of three explicitly distinct states. + + Keeping ``NA`` separate from ``UNPARSED`` is the point of the design: a target + with no accumulator register file legitimately has no AGPR count, and treating + that as a parse failure made the tool unusable on every target without MFMA. + """ + + state: str + value: int | None = None + reason: str = "" + + @staticmethod + def of(value: int) -> "Cell": + return Cell(VALUE, value) + + @staticmethod + def na(reason: str) -> "Cell": + return Cell(NA, None, reason) + + @staticmethod + def unparsed(reason: str) -> "Cell": + return Cell(UNPARSED, None, reason) + + def to_json(self) -> dict: + if self.state == VALUE: + return {"state": VALUE, "value": self.value} + return {"state": self.state, "reason": self.reason} + + @staticmethod + def from_json(raw: object, where: str) -> "Cell": + if not isinstance(raw, dict): + raise SnapshotError(f"{where} must be an object") + state = raw.get("state") + if state == VALUE: + value = raw.get("value") + if not isinstance(value, int) or isinstance(value, bool): + raise SnapshotError(f"{where}.value must be an integer") + if value < 0: + # A snapshot is just the other way to build a `Cell`, and the invariant + # belongs to the cell rather than to the assembly parser: every metric + # here is a count or a byte size. Enforcing it on one path only lets a + # saved snapshot carry a value the tool would never have written, and a + # diff down to it reads as an improvement. + raise SnapshotError(f"{where}.value is negative: {value}") + return Cell.of(value) + if state in (NA, UNPARSED): + reason = raw.get("reason", "") + if not isinstance(reason, str): + raise SnapshotError(f"{where}.reason must be a string") + return Cell(state, None, reason) + raise SnapshotError(f"{where}.state must be one of {VALUE!r}, {NA!r}, {UNPARSED!r}") + + +# -------------------------------------------------------------------------------------- +# The metric table +# -------------------------------------------------------------------------------------- + +TRIGGER = "trigger" +INFO = "info" + +METADATA = "metadata" +SYMBOL = "symbol" +INSTRUCTION = "instruction" + + +@dataclass(frozen=True) +class Metric: + key: str + kind: str + source: str + field: str + + +# Exactly one metric per physical quantity is a TRIGGER. In particular `vgpr` is the +# total (arch + accumulator) that LLVM itself uses for occupancy; `arch_vgpr` and `agpr` +# are its informational split. Counting all three would report a single AGPR change as +# two regressions, and would flag moving accumulators into AGPRs -- which the kernel +# tuning guide recommends -- as a regression. +METRICS = ( + Metric("vgpr", TRIGGER, METADATA, "vgpr_count"), + Metric("arch_vgpr", INFO, SYMBOL, "num_vgpr"), + Metric("agpr", INFO, SYMBOL, "num_agpr"), + Metric("sgpr", TRIGGER, METADATA, "sgpr_count"), + Metric("numbered_sgpr", INFO, SYMBOL, "numbered_sgpr"), + Metric("vgpr_spill", TRIGGER, METADATA, "vgpr_spill_count"), + Metric("sgpr_spill", TRIGGER, METADATA, "sgpr_spill_count"), + Metric("scratch_bytes", TRIGGER, METADATA, "private_segment_fixed_size"), + Metric("lds_static_bytes", TRIGGER, METADATA, "group_segment_fixed_size"), + Metric("lds_read", INFO, INSTRUCTION, "lds_read"), + Metric("lds_write", INFO, INSTRUCTION, "lds_write"), + Metric("scratch_store", INFO, INSTRUCTION, "scratch_store"), + Metric("scratch_load", INFO, INSTRUCTION, "scratch_load"), + Metric("matrix_ops", INFO, INSTRUCTION, "matrix_ops"), +) +KEYS = tuple(metric.key for metric in METRICS) +BY_KEY = {metric.key: metric for metric in METRICS} +TRIGGERS = frozenset(metric.key for metric in METRICS if metric.kind == TRIGGER) + +KNOWN_SET_SUFFIXES = frozenset({"num_vgpr", "num_agpr", "numbered_sgpr", "private_seg_size"}) + +LEGEND = ( + "* = regression trigger; other columns are informational.\n" + " vgpr = total (arch+acc, LLVM's occupancy number); arch_vgpr/agpr are its split " + "-- do not add them." +) + + +# -------------------------------------------------------------------------------------- +# Architecture +# -------------------------------------------------------------------------------------- + +RE_PROC = re.compile(r"^gfx([0-9a-f]{3,})$") +# FeatureArchitectedFlatScratch on gfx9: these spill through scratch_*, gfx90a does not. +GFX9_FLAT_SCRATCH = frozenset({"gfx940", "gfx941", "gfx942", "gfx950"}) +HIGHEST_KNOWN_GEN = 13 + + +@dataclass(frozen=True) +class Arch: + target_id: str = "" + processor: str | None = None + gen: int | None = None + # `xnack+` / `sramecc-` and friends, sorted, so two spellings of one set compare equal. + features: tuple = () + + @property + def known(self) -> bool: + return self.processor is not None + + @property + def extrapolated(self) -> bool: + return self.gen is not None and self.gen > HIGHEST_KNOWN_GEN + + def spills_via_scratch(self) -> bool | None: + """True/False when known, None when the target could not be identified.""" + if self.processor is None: + return None + if self.gen is not None and self.gen >= 11: + return True + return self.processor in GFX9_FLAT_SCRATCH + + def to_json(self) -> dict: + return { + "target_id": self.target_id, + "processor": self.processor, + "gen": self.gen, + "features": list(self.features), + } + + @staticmethod + def from_json(raw: object) -> "Arch": + if raw is None: + return Arch() + if not isinstance(raw, dict): + raise SnapshotError("arch must be an object") + processor = raw.get("processor") + gen = raw.get("gen") + if processor is not None and not isinstance(processor, str): + raise SnapshotError("arch.processor must be a string or null") + if gen is not None and (not isinstance(gen, int) or isinstance(gen, bool)): + raise SnapshotError("arch.gen must be an integer or null") + features = raw.get("features") or () + if not isinstance(features, (list, tuple)) or any(not isinstance(f, str) for f in features): + raise SnapshotError("arch.features must be an array of strings") + return Arch(str(raw.get("target_id") or ""), processor, gen, tuple(features)) + + +def parse_target_id(target_id: str) -> Arch: + """Split ``amdgcn-amd-amdhsa-unknown-gfx950:sramecc+:xnack+`` into its parts. + + ``AMDGPUTargetID::toString()`` emits ``----`` + and then the feature suffixes, and the environment is spelled either as empty + (``amdgcn-amd-amdhsa--gfx950``) or as ``unknown`` + (``amdgcn-amd-amdhsa-unknown-gfx950``) depending on how the triple reached the + backend. Both are in circulation and this repo's own dumps use the second, so + take the processor from the last ``-`` separated field. Keying off a literal + ``--`` recognizes only the first spelling and silently reports every real dump + as an unidentified target. + """ + if not target_id: + return Arch() + head, _, tail = target_id.partition(":") + # Sorted, because the feature set is what matters and LLVM's order is not a contract. + features = tuple(sorted(f for f in tail.split(":") if f)) + processor = head.rsplit("-", 1)[-1] + m = RE_PROC.match(processor) + if not m: + return Arch(target_id, None, None, features) + return Arch(target_id, processor, int(m.group(1)[:-2]), features) + + +def _target_blockers(before: Arch, after: Arch) -> tuple: + """Why two targets are not comparable, or ``()`` when they are. + + Resource counts mean the same thing only on the same processor *and* under the + same target features: `xnack` and `sramecc` change code generation and register + allocation, so a diff taken across them measures the build flags rather than the + change under test. The triple's environment field is spelled either empty or + `unknown` for one and the same target, so it is normalized away here by comparing + the parsed processor and feature set rather than the raw target ID string. + """ + unnamed = [side for side, arch in (("before", before), ("after", after)) if not arch.known] + if unnamed: + return ( + f"target not identified on the {' and '.join(unnamed)} side " + f"({before.target_id or 'no target directive'} vs {after.target_id or 'no target directive'}); " + "without a processor neither the scratch/spill applicability nor the arch match is decidable", + ) + if before.processor != after.processor: + return ( + f"architecture differs ({before.processor} vs {after.processor}); " "resource counts are not comparable", + ) + if before.features != after.features: + return ( + f"target features differ ({':'.join(before.features) or 'none'} vs " + f"{':'.join(after.features) or 'none'}); they change code generation, so this " + "diff would measure the build flags rather than the change under test", + ) + return () + + +# -------------------------------------------------------------------------------------- +# Parsing +# -------------------------------------------------------------------------------------- + +# A label is at column 0 and may carry a trailing comment: `my_kernel: ; @my_kernel`. +RE_LABEL = re.compile(r'^(?P"[^"]*"|[A-Za-z_.$][\w.$]*):[ \t]*(?:;.*)?$') +# `.size , .Lfunc_endN-` names the kernel's terminator explicitly. +RE_SIZE = re.compile( + r'^[ \t]*\.size[ \t]+(?P"[^"]*"|[A-Za-z_.$][\w.$]*)' r"[ \t]*,[ \t]*(?P\.Lfunc_end\d+)-(?P=name)[ \t]*$" +) +RE_SET = re.compile(r"^[ \t]*\.set[ \t]+(?P[A-Za-z_.$][\w.$]*)[ \t]*,[ \t]*(?P-?\d+)[ \t]*$") +RE_SET_ANY = re.compile(r"^[ \t]*\.set[ \t]+(?P[A-Za-z_.$][\w.$]*)[ \t]*,") +# Kernel-level metadata keys sit at column 4, or on the ` - ` sequence-dash line. +# Argument keys sit at column 6 or 8, so both branches must be literal. +RE_MD_KEY = re.compile(r"^(?: - | )\.(?P[A-Za-z_]\w*):[ \t]*(?P.*?)[ \t]*$") +RE_MNEMONIC = re.compile(r"^[ \t]+(?P[a-z][a-z0-9_]*)\b") +RE_AMDGCN_TARGET = re.compile(r'^[ \t]*\.amdgcn_target[ \t]+"([^"]*)"[ \t]*$') +RE_MD_TARGET = re.compile(r"^amdhsa\.target:[ \t]*(?:\"([^\"]*)\"|'([^']*)'|(\S+))[ \t]*$") + +# `(?:_|2)` keeps ds_read2_b32 / ds_write2_b32 while rejecting gfx11+'s +# ds_storexchg_rtn_b32 -- the renamed gfx9 atomic ds_wrxchg_rtn_b32, which a bare +# `^ds_store` would miscount as an LDS write. +_CATEGORIES = ( + ("ds_read", re.compile(r"^ds_read(?:_|2)")), + ("ds_load", re.compile(r"^ds_load(?:_|2)")), + ("ds_write", re.compile(r"^ds_write(?:_|2)")), + ("ds_store", re.compile(r"^ds_store(?:_|2)")), + ("scratch_store", re.compile(r"^scratch_store")), + ("scratch_load", re.compile(r"^scratch_load")), + ("matrix_ops", re.compile(r"^v_(?:mfma|wmma)")), +) + +METADATA_BEGIN = ".amdgpu_metadata" +METADATA_END = ".end_amdgpu_metadata" +DESCRIPTOR_BEGIN = ".amdhsa_kernel" +DESCRIPTOR_END = ".end_amdhsa_kernel" + + +def _unquote(name: str) -> str: + if len(name) >= 2 and name[0] == name[-1] == '"': + return name[1:-1] + return name + + +def clean_md_value(value: str) -> str: + """Strip a YAML tag and surrounding quotes: ``!str n`` -> ``n``.""" + v = value.strip() + if v.startswith("!"): + v = v.split(" ", 1)[1].strip() if " " in v else "" + if len(v) >= 2 and v[0] == v[-1] and v[0] in "\"'": + v = v[1:-1] + return v + + +def _categorize(mnemonic: str) -> str | None: + for name, pattern in _CATEGORIES: + if pattern.match(mnemonic): + return name + return None + + +@dataclass +class _Scan: + labels: dict = field(default_factory=dict) + sizes: dict = field(default_factory=dict) + symbols: dict = field(default_factory=dict) + entries: list = field(default_factory=list) + instr_lines: list = field(default_factory=list) + instr_cats: list = field(default_factory=list) + target_id: str | None = None + md_target_id: str | None = None + saw_metadata: bool = False + + +def _record_set(line: str, symbols: dict) -> None: + m = RE_SET.match(line) + if m: + sym, value = m.group("sym"), int(m.group("val")) + else: + loose = RE_SET_ANY.match(line) + if not loose: + return + # e.g. `.set k.num_vgpr, max(32, .Lhelper.num_vgpr)` when the kernel calls a + # function: record it as present-but-unreadable rather than dropping it. + sym, value = loose.group("sym"), None + base, _, suffix = sym.rpartition(".") + if not base or suffix not in KNOWN_SET_SUFFIXES: + return + if base.startswith(".L"): + # 90 of 233 real dumps emit `.set .L.num_vgpr` even for .globl kernels, + # because the local-linkage test differs from the one behind metadata `.name`. + base = base[2:] + symbols.setdefault(base, {})[suffix] = value + + +def _scan_text(text: str) -> _Scan: + """One pass over the file; no regex ever runs over the whole text.""" + scan = _Scan() + in_metadata = False + in_descriptor = False + + for index, line in enumerate(text.splitlines()): + if not line: + continue + if line[0] not in " \t": + if in_metadata: + if line.startswith("amdhsa.target:"): + m = RE_MD_TARGET.match(line) + if m: + scan.md_target_id = next(g for g in m.groups() if g is not None) + continue + m = RE_LABEL.match(line) + if m: + scan.labels.setdefault(_unquote(m.group("name")), index) + continue + + stripped = line.lstrip() + # `.end_*` must be tested first: it is a prefix collision with `.amdgpu_metadata`. + if stripped.startswith(METADATA_END): + in_metadata = False + continue + if stripped.startswith(METADATA_BEGIN): + in_metadata = True + scan.saw_metadata = True + continue + if stripped.startswith(DESCRIPTOR_END): + in_descriptor = False + continue + if stripped.startswith(DESCRIPTOR_BEGIN): + in_descriptor = True + continue + if in_descriptor: + continue + if in_metadata: + m = RE_MD_KEY.match(line) + if m: + if line.startswith(" - "): + scan.entries.append({}) + if scan.entries: + scan.entries[-1][m.group("key")] = clean_md_value(m.group("val")) + continue + if stripped.startswith(".size"): + m = RE_SIZE.match(line) + if m: + scan.sizes[_unquote(m.group("name"))] = m.group("end") + continue + if stripped.startswith(".set"): + _record_set(line, scan.symbols) + continue + if stripped.startswith(".amdgcn_target"): + m = RE_AMDGCN_TARGET.match(line) + if m: + scan.target_id = m.group(1) + continue + if stripped[0] in ".;": + continue + m = RE_MNEMONIC.match(line) + if m: + category = _categorize(m.group("m")) + if category: + scan.instr_lines.append(index) + scan.instr_cats.append(category) + return scan + + +@dataclass(frozen=True) +class KernelRecord: + name: str + arch: Arch + metrics: dict + source: str = "" + notes: tuple = () + problems: tuple = () + + @property + def unparsed_keys(self) -> tuple: + return tuple(key for key in KEYS if self.metrics[key].state == UNPARSED) + + +def _kernel_name(entry: dict) -> tuple: + """(name, problems) -- metadata `.name`, cross-checked against `.symbol`.""" + name = entry.get("name") or "" + symbol = entry.get("symbol") or "" + from_symbol = symbol[:-3] if symbol.endswith(".kd") else symbol + if name and from_symbol and name != from_symbol: + return name, (f"metadata .name {name!r} disagrees with .symbol {symbol!r}",) + if name: + return name, () + if from_symbol: + return from_symbol, () + return "", ("metadata entry has neither .name nor .symbol",) + + +def _metadata_cell(entry: dict, metric: Metric) -> Cell: + raw = entry.get(metric.field) + if raw is None: + return Cell.unparsed(f"metadata field .{metric.field} is absent") + try: + value = int(raw) + except ValueError: + return Cell.unparsed(f"metadata field .{metric.field} is not an integer: {raw!r}") + if value < 0: + # Every metric here is a count or a byte size. A negative one is impossible, so + # the dump is malformed; treating it as a value makes the drop toward it read as + # an improvement, which is the one verdict this tool must never invent. + return Cell.unparsed(f"metadata field .{metric.field} is negative: {value}") + return Cell.of(value) + + +def _symbol_cell(symbols: dict, name: str, metric: Metric) -> Cell: + if name not in symbols or metric.field not in symbols[name]: + return Cell.unparsed(f"assembler symbol .set {name}.{metric.field} is absent") + value = symbols[name][metric.field] + if value is None: + return Cell.unparsed( + f".set {name}.{metric.field} is an unresolved expression " + "(the kernel calls a function); use the metadata counts" + ) + if value < 0: + return Cell.unparsed(f".set {name}.{metric.field} is negative: {value}") + return Cell.of(value) + + +def _instruction_cells(counts: dict, arch: Arch, body_error: str | None) -> tuple: + """Returns (cells, notes, problems) for the five instruction-derived metrics.""" + if body_error is not None: + cells = {metric.key: Cell.unparsed(body_error) for metric in METRICS if metric.source == INSTRUCTION} + return cells, (), () + + problems = [] + if counts.get("ds_read", 0) and counts.get("ds_load", 0): + problems.append("mixed ds_read/ds_load spellings in one body; the dump is inconsistent") + if counts.get("ds_write", 0) and counts.get("ds_store", 0): + problems.append("mixed ds_write/ds_store spellings in one body; the dump is inconsistent") + + cells = { + "lds_read": Cell.of(counts.get("ds_read", 0) + counts.get("ds_load", 0)), + "lds_write": Cell.of(counts.get("ds_write", 0) + counts.get("ds_store", 0)), + "matrix_ops": Cell.of(counts.get("matrix_ops", 0)), + } + + scratch = arch.spills_via_scratch() + if scratch is False: + reason = ( + f"{arch.processor} spills through buffer_* instructions; " + "a scratch instruction count is not a spill signal on this target" + ) + cells["scratch_store"] = Cell.na(reason) + cells["scratch_load"] = Cell.na(reason) + else: + cells["scratch_store"] = Cell.of(counts.get("scratch_store", 0)) + cells["scratch_load"] = Cell.of(counts.get("scratch_load", 0)) + return cells, (), tuple(problems) + + +def parse_isa(path: str | Path) -> tuple: + """Parse one ``*final_isa.s`` into ``({kernel name: KernelRecord}, file problems)``. + + The second element carries what went wrong with the file as a whole rather than + with one kernel: an entry that cannot be identified, a kernel declared twice, a + byte that is not valid UTF-8. None of those belong to a record -- and dropping + them, as an earlier version did, is what let a file lose a kernel and still be + reported as trustworthy. + """ + path = Path(path) + data = path.read_bytes() + file_problems = [] + try: + text = data.decode("utf-8") + except UnicodeDecodeError as exc: + # Assembly is ASCII in practice, so a byte that will not decode means the file + # is damaged. Decoding leniently keeps the run alive instead of aborting with a + # traceback, but a replacement character can erase a mnemonic and quietly change + # an instruction count, so the file is reported rather than counted. + text = data.decode("utf-8", errors="replace") + file_problems.append( + f"{path}: not valid UTF-8 at byte {exc.start}; a damaged byte can erase a " + "mnemonic, so the instruction counts from this file are not trustworthy" + ) + scan = _scan_text(text) + + if not scan.saw_metadata: + return {}, tuple(file_problems) + + arch = parse_target_id(scan.target_id or scan.md_target_id or "") + arch_problems = [] + if scan.target_id and scan.md_target_id and scan.target_id != scan.md_target_id: + arch_problems.append(f".amdgcn_target {scan.target_id!r} disagrees with amdhsa.target {scan.md_target_id!r}") + + records = {} + for entry in scan.entries: + name, name_problems = _kernel_name(entry) + if not name: + # `_kernel_name` already said why. Dropping the entry silently removes a + # kernel from this side of the comparison without leaving a trace, which is + # exactly the partial-comparison case the exit-2 verdict exists for. + file_problems.extend(name_problems) + continue + if name in records: + # Last-wins would let a stale duplicate hide a changed counter. Keep the + # first and say so; `collect()` already treats a duplicate key across + # directories the same way. + file_problems.append(f"metadata declares kernel {name!r} more than once in {path}") + continue + + metrics = {} + for metric in METRICS: + if metric.source == METADATA: + metrics[metric.key] = _metadata_cell(entry, metric) + elif metric.source == SYMBOL: + metrics[metric.key] = _symbol_cell(scan.symbols, name, metric) + + counts, body_error = _body_counts(scan, name) + instruction_cells, notes, instruction_problems = _instruction_cells(counts, arch, body_error) + metrics.update(instruction_cells) + + notes = list(notes) + lds_static = metrics["lds_static_bytes"] + lds_traffic = metrics["lds_read"].value or 0 + lds_traffic += metrics["lds_write"].value or 0 + if lds_static.state == VALUE and lds_static.value == 0 and lds_traffic > 0: + notes.append("dynamic LDS in use; a static size of 0 does not bound this kernel's LDS") + + cross = _cross_check_scratch(metrics, scan.symbols.get(name, {})) + records[name] = KernelRecord( + name=name, + arch=arch, + metrics=metrics, + source=str(path), + notes=tuple(notes), + problems=tuple(name_problems) + tuple(arch_problems) + tuple(instruction_problems) + cross, + ) + return records, tuple(file_problems) + + +def _cross_check_scratch(metrics: dict, symbols: dict) -> tuple: + """`.private_seg_size` and `.private_segment_fixed_size` measure the same quantity.""" + cell = metrics["scratch_bytes"] + other = symbols.get("private_seg_size") + if cell.state != VALUE or other is None or other == cell.value: + return () + return (f"scratch bytes disagree: metadata {cell.value} vs .set private_seg_size {other}",) + + +def _body_counts(scan: _Scan, name: str) -> tuple: + """Instruction counts for one kernel body, or (empty, reason) when unbounded.""" + if name not in scan.labels: + return {}, f"no code label {name!r} in the file" + terminator = scan.sizes.get(name) + if terminator is None: + return {}, f"no .size line for {name!r}, so the end of its body is unknown" + if terminator not in scan.labels: + return {}, f"terminator {terminator!r} for {name!r} is never defined" + + lo = scan.labels[name] + 1 + hi = scan.labels[terminator] + if hi <= lo: + return {}, f"body of {name!r} is empty or inverted ({lo}..{hi})" + + start = bisect_left(scan.instr_lines, lo) + stop = bisect_left(scan.instr_lines, hi) + counts: dict = {} + for category in scan.instr_cats[start:stop]: + counts[category] = counts.get(category, 0) + 1 + return counts, None + + +# -------------------------------------------------------------------------------------- +# Snapshots +# -------------------------------------------------------------------------------------- + +RE_STAGE_PREFIX = re.compile(r"^(\d+)_") + + +@dataclass(frozen=True) +class Snapshot: + kernels: dict = field(default_factory=dict) + warnings: tuple = () + problems: tuple = () + + def __len__(self) -> int: + return len(self.kernels) + + def __bool__(self) -> bool: + return bool(self.kernels) + + @property + def trustworthy(self) -> bool: + return not self.problems and not any(record.unparsed_keys for record in self.kernels.values()) + + def to_json(self) -> dict: + return { + "schema": SCHEMA_VERSION, + "tool": TOOL_NAME, + "trustworthy": self.trustworthy, + "warnings": list(self.warnings), + "kernels": { + key: { + "source": record.source, + "arch": record.arch.to_json(), + "notes": list(record.notes), + "problems": list(record.problems), + "metrics": {k: record.metrics[k].to_json() for k in KEYS}, + } + for key, record in sorted(self.kernels.items()) + }, + } + + +def _pick_stage_file(paths: list) -> tuple: + """Newest final-ISA file in one directory, plus a warning naming the ignored ones.""" + if len(paths) == 1: + return paths[0], None + + def sort_key(path): + m = RE_STAGE_PREFIX.match(path.name) + stage = int(m.group(1)) if m else -1 + try: + mtime = path.stat().st_mtime + except OSError: + mtime = 0.0 + return (stage, mtime, path.name) + + ordered = sorted(paths, key=sort_key) + chosen = ordered[-1] + ignored = ", ".join(p.name for p in ordered[:-1]) + warning = f"{chosen.parent}: {len(paths)} final-ISA files, using {chosen.name} (ignoring stale {ignored})" + return chosen, warning + + +def collect(root: str | Path) -> Snapshot: + """Summarize every ``*final_isa.s`` under ``root``.""" + root = Path(root) + if not root.exists(): + raise SnapshotError(f"dump directory not found: {root}") + if not root.is_dir(): + raise SnapshotError(f"not a directory: {root}") + + by_directory: dict = {} + for path in sorted(root.rglob("*final_isa.s")): + by_directory.setdefault(path.parent, []).append(path) + if not by_directory: + raise SnapshotError(f"no *final_isa.s under {root} (is FLYDSL_DUMP_IR set?)") + + kernels: dict = {} + warnings: list = [] + problems: list = [] + files_read = 0 + + for directory in sorted(by_directory): + chosen, warning = _pick_stage_file(by_directory[directory]) + if warning: + warnings.append(warning) + files_read += 1 + + records, file_problems = parse_isa(chosen) + problems.extend(file_problems) + if not records: + # Skipping this file quietly makes the comparison silently partial: the + # kernels it should have contributed simply never appear, and the count + # line then reports full coverage of what is left. A file the tool + # cannot read is exactly the case the exit-2 verdict exists for. + problems.append( + f"{chosen}: not an LLVM AMDHSA dump (no {METADATA_BEGIN}); " + "its kernels are missing from this side, so the comparison would be partial" + ) + continue + + relative_dir = chosen.parent.relative_to(root).as_posix() + relative_dir = "" if relative_dir == "." else relative_dir + if relative_dir and Path(relative_dir).name not in records and len(records) == 1: + warnings.append( + f"{chosen.parent}: directory name does not match kernel " + f"{next(iter(records))!r}; keys are path-derived" + ) + for name, record in records.items(): + # The key is a pure function of one file's location and one kernel's own + # name. It must never depend on what else happens to be in the tree, or + # adding one kernel to the "after" run renames every pre-existing key. + key = f"{relative_dir}::{name}" + if key in kernels: + problems.append(f"duplicate kernel key {key!r}") + continue + kernels[key] = replace(record, source=str(chosen.relative_to(root))) + problems.extend(f"{key}: {p}" for p in record.problems) + + if records: + arch = next(iter(records.values())).arch + if not arch.known: + warnings.append( + f"{chosen}: architecture unknown; scratch instruction counts are raw " "traffic, not proven spills" + ) + elif arch.extrapolated: + warnings.append( + f"{chosen}: arch {arch.processor} is newer than this tool knows " + f"(gen {arch.gen}); its bucket is extrapolated" + ) + + if not kernels: + raise SnapshotError(f"no kernels found in {files_read} file(s) under {root}") + + warnings.append( + "FLYDSL_DUMP_IR writes one directory per kernel *name* with no specialization key, " + "so different JIT specializations overwrite each other and cache hits emit no dump. " + "Compare only trees from a single run with FLYDSL_RUNTIME_ENABLE_CACHE=0." + ) + return Snapshot(kernels=kernels, warnings=tuple(warnings), problems=tuple(problems)) + + +_V1_ALIASES = { + "vgpr": "vgpr", + "sgpr": "sgpr", + "agpr": "agpr", + "vgpr_spill": "vgpr_spill", + "sgpr_spill": "sgpr_spill", + "scratch_bytes": "scratch_bytes", + "lds_bytes": "lds_static_bytes", + "scratch_store": "scratch_store", + "scratch_load": "scratch_load", + "ds_read": "lds_read", +} + + +def _load_v1(raw: dict, path: Path) -> Snapshot: + kernels = {} + for kernel, metrics in raw.items(): + if not isinstance(kernel, str) or not isinstance(metrics, dict): + raise SnapshotError(f"{path} contains an invalid kernel entry") + cells = {key: Cell.unparsed("absent from this v1 snapshot") for key in KEYS} + for old_key, new_key in _V1_ALIASES.items(): + if old_key not in metrics: + continue + value = metrics[old_key] + if value is None: + # v1 could not tell "not applicable" from "failed to parse". + cells[new_key] = Cell.unparsed("null in a v1 snapshot; regenerate it") + elif isinstance(value, int) and not isinstance(value, bool): + cells[new_key] = Cell.of(value) + else: + raise SnapshotError(f"{path}: {kernel}.{old_key} must be an integer or null") + kernels[kernel] = KernelRecord(name=kernel, arch=Arch(), metrics=cells, source=str(path)) + return Snapshot( + kernels=kernels, + warnings=(f"{path} is a v1 snapshot; regenerate it with this version of the tool",), + ) + + +def load_snapshot(path: str | Path) -> Snapshot: + path = Path(path) + raw = json.loads(path.read_text(encoding="utf-8")) + if not isinstance(raw, dict): + raise SnapshotError(f"{path} must contain a JSON object") + if "schema" not in raw: + return _load_v1(raw, path) + + if raw.get("schema") != SCHEMA_VERSION: + raise SnapshotError(f"{path}: unsupported schema {raw.get('schema')!r}") + entries = raw.get("kernels") + if not isinstance(entries, dict): + raise SnapshotError(f"{path}: 'kernels' must be an object") + + kernels = {} + for kernel, body in entries.items(): + if not isinstance(kernel, str) or not isinstance(body, dict): + raise SnapshotError(f"{path} contains an invalid kernel entry") + raw_metrics = body.get("metrics") + if not isinstance(raw_metrics, dict): + raise SnapshotError(f"{path}: {kernel}.metrics must be an object") + cells = {} + for key in KEYS: + if key not in raw_metrics: + cells[key] = Cell.unparsed(f"absent from {path.name}") + else: + cells[key] = Cell.from_json(raw_metrics[key], f"{path}: {kernel}.{key}") + kernels[kernel] = KernelRecord( + name=kernel, + arch=Arch.from_json(body.get("arch")), + metrics=cells, + source=str(body.get("source") or ""), + notes=tuple(body.get("notes") or ()), + problems=tuple(body.get("problems") or ()), + ) + + problems = [] + if raw.get("trustworthy") is False: + problems.append(f"{path} was written from data the tool could not fully parse") + return Snapshot( + kernels=kernels, + warnings=tuple(raw.get("warnings") or ()), + problems=tuple(problems), + ) + + +def load_input(path: str | Path) -> Snapshot: + """Load a JSON snapshot or summarize a FLYDSL_DUMP_IR directory.""" + path = Path(path) + if path.is_dir(): + return collect(path) + if not path.exists(): + raise SnapshotError(f"input not found: {path}") + if path.suffix != ".json": + raise SnapshotError(f"expected a dump directory or a .json snapshot: {path}") + return load_snapshot(path) + + +# -------------------------------------------------------------------------------------- +# Comparison +# -------------------------------------------------------------------------------------- + +ONLY_IN_BEFORE = "ONLY IN BEFORE" +ONLY_IN_AFTER = "ONLY IN AFTER" + + +@dataclass(frozen=True) +class DiffRow: + name: str + before: KernelRecord | None + after: KernelRecord | None + blocked: tuple = () + + @property + def side(self) -> str | None: + """The one label for a one-sided row -- defined once, used by every caller.""" + if self.before is None: + return ONLY_IN_AFTER + if self.after is None: + return ONLY_IN_BEFORE + return None + + @property + def comparable(self) -> bool: + return self.side is None and not self.blocked + + def delta(self, key: str) -> int | None: + if not self.comparable: + return None + before_cell = self.before.metrics[key] + after_cell = self.after.metrics[key] + if before_cell.state != VALUE or after_cell.state != VALUE: + return None + return after_cell.value - before_cell.value + + @property + def changed(self) -> bool: + """The single definition of 'this kernel changed', used by every caller.""" + return any(self.delta(key) for key in KEYS) + + +@dataclass(frozen=True) +class Comparison: + rows: tuple = () + problems: tuple = () + warnings: tuple = () + notes: tuple = () + total_kernels: int = 0 + + @property + def compared_kernels(self) -> int: + return sum(1 for row in self.rows if row.comparable) + + @property + def unchanged_kernels(self) -> int: + return sum(1 for row in self.rows if row.comparable and not row.changed) + + @property + def changed_kernels(self) -> int: + return sum(1 for row in self.rows if row.comparable and row.changed) + + @property + def worsened_metrics(self) -> int: + return self._count_triggers(lambda d: d > 0) + + @property + def improved_metrics(self) -> int: + return self._count_triggers(lambda d: d < 0) + + def _count_triggers(self, predicate) -> int: + total = 0 + for row in self.rows: + if not row.comparable: + continue + for key in TRIGGERS: + delta = row.delta(key) + if delta is not None and predicate(delta): + total += 1 + return total + + @property + def exit_code(self) -> int: + if self.problems: + return 2 + return 1 if self.worsened_metrics else 0 + + @property + def verdict(self) -> str: + return {0: "RESULT: OK", 1: "RESULT: REGRESSION", 2: "RESULT: NOT TRUSTWORTHY"}[self.exit_code] + + +def _match_keys(before: Snapshot, after: Snapshot) -> tuple: + """Exact keys first, then pair leftovers by bare kernel name across layout drift.""" + common = before.kernels.keys() & after.kernels.keys() + pairs = {key: key for key in common} + notes: list = [] + + left = sorted(before.kernels.keys() - common) + right = sorted(after.kernels.keys() - common) + if left and right: + + def by_kernel(keys): + out: dict = {} + for key in keys: + out.setdefault(key.rpartition("::")[2], []).append(key) + return out + + left_by, right_by = by_kernel(left), by_kernel(right) + for kernel, left_keys in left_by.items(): + right_keys = right_by.get(kernel, []) + if len(left_keys) == 1 and len(right_keys) == 1: + pairs[left_keys[0]] = right_keys[0] + notes.append( + "matched by kernel name across differing layouts:\n" + f" before {left_keys[0]}\n" + f" after {right_keys[0]}" + ) + return pairs, tuple(notes) + + +def compare_snapshots(before: Snapshot, after: Snapshot) -> Comparison: + problems = list(before.problems) + list(after.problems) + if not before or not after: + problems.append(f"one or both inputs contain no kernels (before={len(before)}, after={len(after)})") + + pairs, notes = _match_keys(before, after) + matched_after = set(pairs.values()) + + rows = [] + for key in sorted(before.kernels.keys() | after.kernels.keys()): + if key in pairs: + before_record = before.kernels[key] + after_record = after.kernels[pairs[key]] + elif key in after.kernels and key not in matched_after: + rows.append(DiffRow(key, None, after.kernels[key])) + problems.append(f"{key}: {ONLY_IN_AFTER}") + continue + elif key in after.kernels: + continue # already consumed as the right-hand side of a drift match + else: + rows.append(DiffRow(key, before.kernels[key], None)) + problems.append(f"{key}: {ONLY_IN_BEFORE}") + continue + + blocked = list(_target_blockers(before_record.arch, after_record.arch)) + + for metric_key in KEYS: + before_state = before_record.metrics[metric_key].state + after_state = after_record.metrics[metric_key].state + if UNPARSED in (before_state, after_state): + reason = ( + before_record.metrics[metric_key].reason + if before_state == UNPARSED + else after_record.metrics[metric_key].reason + ) + blocked.append(f"unparsed metric {metric_key}: {reason}") + elif (before_state == NA) != (after_state == NA): + blocked.append( + f"metric {metric_key} is {before_state} before and {after_state} after; " "applicability changed" + ) + + rows.append(DiffRow(key, before_record, after_record, tuple(blocked))) + problems.extend(f"{key}: {reason}" for reason in blocked) + + total = len(before.kernels.keys() | after.kernels.keys()) - ( + len(pairs) - len(before.kernels.keys() & after.kernels.keys()) + ) + return Comparison( + rows=tuple(rows), + problems=tuple(problems), + warnings=tuple(before.warnings) + tuple(after.warnings), + notes=tuple(notes), + total_kernels=total, + ) + + +# -------------------------------------------------------------------------------------- +# Rendering +# -------------------------------------------------------------------------------------- + +MIN_NAME_WIDTH = 20 +MAX_NAME_WIDTH = 60 +NAME_TAIL = 12 +MIN_CELL_WIDTH = 5 + + +def _abbreviate(name: str, width: int) -> str: + """Keep the tail: real kernel names differ in their trailing config suffix.""" + if len(name) <= width: + return name + if width <= NAME_TAIL + 1: + return name[-width:] + head = width - NAME_TAIL - 1 + return name[:head] + "~" + name[-NAME_TAIL:] + + +def _name_column(names: list) -> dict: + if not names: + return {} + longest = max(len(name) for name in names) + width = min(max(MIN_NAME_WIDTH, longest), MAX_NAME_WIDTH) + shown = {name: _abbreviate(name, width) for name in names} + if len(set(shown.values())) != len(shown): + # Rows must never be indistinguishable; drop the cap rather than the identity. + return {name: name for name in names} + return shown + + +def _format_table(headers: list, rows: list) -> list: + """Fixed-width table whose columns are sized from the data. Nothing is truncated.""" + widths = [] + for index, header in enumerate(headers): + longest = max([len(header)] + [len(row[index]) for row in rows]) if rows else len(header) + widths.append(max(longest, MIN_CELL_WIDTH if index else len(header))) + + def line(cells): + parts = [cells[0].ljust(widths[0])] + parts.extend(cell.rjust(widths[i]) for i, cell in enumerate(cells[1:], start=1)) + return " ".join(parts).rstrip() + + header_line = line(headers) + out = [header_line, "-" * max(len(header_line), *(len(line(row)) for row in rows), 1)] + out.extend(line(row) for row in rows) + return out + + +def _header_cells() -> list: + return ["kernel"] + [("*" + key if key in TRIGGERS else key) for key in KEYS] + + +def _cell_text(cell: Cell) -> str: + if cell.state == VALUE: + return str(cell.value) + return "n/a" if cell.state == NA else "?" + + +def render_snapshot(snapshot: Snapshot, output: TextIO, error: TextIO | None = None) -> None: + error = error if error is not None else sys.stderr + names = _name_column(sorted(snapshot.kernels)) + rows = [] + for key in sorted(snapshot.kernels): + record = snapshot.kernels[key] + rows.append([names[key]] + [_cell_text(record.metrics[k]) for k in KEYS]) + + print(LEGEND, file=output) + for line in _format_table(_header_cells(), rows): + print(line, file=output) + + for key in sorted(snapshot.kernels): + for note in snapshot.kernels[key].notes: + print(f" note: {names[key]}: {note}", file=output) + + print(f"\n{len(snapshot)} kernels", file=output) + _render_problems(snapshot.warnings, snapshot.problems, output, error) + + +def _render_problems(warnings, problems, output: TextIO, error: TextIO) -> None: + for warning in warnings: + print(f"warning: {warning}", file=error) + if problems: + print(f"{len(problems)} kernel(s)/metric(s) not trustworthy (see stderr)", file=output) + print(f"\nNOT TRUSTWORTHY ({len(problems)} problem(s)):", file=error) + for problem in problems[:20]: + print(f" {problem}", file=error) + if len(problems) > 20: + print(f" ... and {len(problems) - 20} more", file=error) + + +def render_comparison(comparison: Comparison, output: TextIO, error: TextIO) -> None: + printable = [row for row in comparison.rows if not row.comparable or row.changed] + names = _name_column([row.name for row in printable]) + + rows = [] + for row in printable: + if row.side is not None: + rows.append([names[row.name], row.side] + [""] * (len(KEYS) - 1)) + continue + if row.blocked: + rows.append([names[row.name], "BLOCKED: " + row.blocked[0]] + [""] * (len(KEYS) - 1)) + continue + cells = [names[row.name]] + for key in KEYS: + delta = row.delta(key) + if delta: + before_value = row.before.metrics[key].value + after_value = row.after.metrics[key].value + cells.append(f"{before_value}->{after_value}({delta:+d})") + else: + cells.append(_cell_text(row.after.metrics[key])) + rows.append(cells) + + print(LEGEND, file=output) + for line in _format_table(_header_cells(), rows): + print(line, file=output) + + for note in comparison.notes: + print(f"\nnote: {note}", file=output) + + print( + f"\ncompared {comparison.compared_kernels} of {comparison.total_kernels} kernels; " + f"{comparison.unchanged_kernels} unchanged; {comparison.changed_kernels} changed; " + f"worsened: {comparison.worsened_metrics}; improved: {comparison.improved_metrics}", + file=output, + ) + _render_problems(comparison.warnings, comparison.problems, output, error) + print(comparison.verdict, file=output) + + +# -------------------------------------------------------------------------------------- +# Commands +# -------------------------------------------------------------------------------------- + + +def do_diff( + before_path: str | Path, + after_path: str | Path, + *, + output: TextIO | None = None, + error: TextIO | None = None, + json_path: str | Path | None = None, + quiet: bool = False, +) -> int: + output = output if output is not None else sys.stdout + error = error if error is not None else sys.stderr + comparison = compare_snapshots(load_input(before_path), load_input(after_path)) + if json_path is not None: + Path(json_path).write_text( + json.dumps(_comparison_json(comparison), indent=2, sort_keys=True) + "\n", + encoding="utf-8", + ) + if quiet: + _render_problems(comparison.warnings, comparison.problems, output, error) + print(comparison.verdict, file=output) + else: + render_comparison(comparison, output, error) + return comparison.exit_code + + +def _comparison_json(comparison: Comparison) -> dict: + return { + "schema": SCHEMA_VERSION, + "tool": TOOL_NAME, + "verdict": comparison.verdict, + "exit_code": comparison.exit_code, + "compared_kernels": comparison.compared_kernels, + "total_kernels": comparison.total_kernels, + "worsened_metrics": comparison.worsened_metrics, + "improved_metrics": comparison.improved_metrics, + "problems": list(comparison.problems), + "warnings": list(comparison.warnings), + "rows": [ + { + "kernel": row.name, + "side": row.side, + "blocked": list(row.blocked), + "deltas": {key: row.delta(key) for key in KEYS if row.delta(key)}, + } + for row in comparison.rows + if not row.comparable or row.changed + ], + } + + +def do_summarize( + dump_dir: str | Path, + *, + json_path: str | Path | None = None, + output: TextIO | None = None, + error: TextIO | None = None, + quiet: bool = False, +) -> int: + output = output if output is not None else sys.stdout + error = error if error is not None else sys.stderr + snapshot = collect(dump_dir) + if json_path is not None: + Path(json_path).write_text(json.dumps(snapshot.to_json(), indent=2, sort_keys=True) + "\n", encoding="utf-8") + if quiet: + _render_problems(snapshot.warnings, snapshot.problems, output, error) + else: + render_snapshot(snapshot, output, error) + + # Fail closed on unreadable metrics, not merely on an empty directory: a snapshot + # whose every metric is unparsed must not be reported -- or persisted -- as success. + if not snapshot.trustworthy: + unparsed = sorted({key for record in snapshot.kernels.values() for key in record.unparsed_keys}) + if unparsed: + print( + f"\nNOT TRUSTWORTHY: unparsed metric(s): {', '.join(unparsed)}", + file=error, + ) + print("RESULT: NOT TRUSTWORTHY", file=output) + return 2 + print("RESULT: OK", file=output) + return 0 + + +def _create_parser() -> argparse.ArgumentParser: + parser = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + commands = parser.add_subparsers(dest="command", required=True) + + summarize = commands.add_parser("summarize", help="summarize a FLYDSL_DUMP_IR directory") + summarize.add_argument("dump_dir", type=Path) + summarize.add_argument("--json", type=Path, metavar="PATH", help="write a reusable JSON snapshot") + summarize.add_argument("-q", "--quiet", action="store_true", help="suppress the table") + + diff = commands.add_parser("diff", help="compare two dump directories or JSON snapshots") + diff.add_argument("before", type=Path, help="before dump directory or JSON snapshot") + diff.add_argument("after", type=Path, help="after dump directory or JSON snapshot") + diff.add_argument("--json", type=Path, metavar="PATH", help="write the comparison as JSON") + diff.add_argument("-q", "--quiet", action="store_true", help="suppress the table") + return parser + + +def main(argv: list[str] | None = None) -> int: + args = _create_parser().parse_args(argv) + try: + if args.command == "diff": + return do_diff(args.before, args.after, json_path=args.json, quiet=args.quiet) + return do_summarize(args.dump_dir, json_path=args.json, quiet=args.quiet) + except SnapshotError as exc: + print(f"error: {exc}", file=sys.stderr) + print("RESULT: NOT TRUSTWORTHY") + return 2 + except Exception: # noqa: BLE001 - a crash must never be reported as a regression + traceback.print_exc() + print("error: internal error; the result is not trustworthy", file=sys.stderr) + print("RESULT: NOT TRUSTWORTHY") + return 2 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tests/unit/test_isa_resource_table.py b/tests/unit/test_isa_resource_table.py new file mode 100644 index 000000000..716dd6aa9 --- /dev/null +++ b/tests/unit/test_isa_resource_table.py @@ -0,0 +1,378 @@ +#!/usr/bin/env python3 + +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 FlyDSL Project Contributors + +"""Tests for scripts/isa_resource_table.py. + +The ISA text is generated here rather than checked in, so the file states exactly +which parts of LLVM's output the parser depends on: the metadata indentation, the +`.set .*` symbols, the `.size` terminator, and the mnemonic spelling. + +Targets differ along three axes the parser has to survive, so every assertion runs +over one shape from each side of them: + +* whether the target has an accumulator register file — LLVM emits `.agpr_count` + only when `STM.hasMAIInsts()`, so an MFMA-capable target has the field and a + target without MFMA does not; +* how LDS access is spelled — `ds_read`/`ds_write` through gfx10, renamed to + `ds_load`/`ds_store` from gfx11 on; +* how the target ID spells the triple's environment field — + `amdgcn-amd-amdhsa--gfx942` when it is empty, `amdgcn-amd-amdhsa-unknown-gfx942` + when it is `unknown`. This repo's own dumps use the second, and an earlier + version of this file exercised only the first, so the tool shipped unable to + identify the architecture of any dump it would actually be pointed at. + +The first two are ISA properties, not product families, and they do not track the +CDNA / RDNA marketing split: gfx1250 is a CDNA-generation part that LLVM places in +the GFX12 family (wave32, no `FeatureMAIInsts`), so it takes the second value on +both. An earlier version of this tool was tested only against a gfx9-shaped sample +and so stayed green while rejecting every kernel on gfx12xx. The third is a +property of how the triple reached the backend rather than of the target itself, +so it is pinned per shape only to keep both spellings under test. + +The generated shape was checked against real dumps: its metadata keys are a subset +of the keys LLVM emits, in the same sorted order, and every structural marker the +parser keys off (`.amdgcn_target`, the kernel label, `.Lfunc_endN`, `.size`, the +`.set` block, the ` - `/4-space/6-space indent levels) appears as it does there. +The keys left out are ones the parser never reads. +""" + +import json + +import pytest + +from scripts import isa_resource_table as irt +from scripts.isa_resource_table import VALUE, parse_isa + +pytestmark = [pytest.mark.l0_backend_agnostic] + +LDS_READS = 3 + + +def make_isa(arch, kernel, *, mfma, ds_load_spelling, num_vgpr, num_agpr, vgpr_count, env): + """One kernel's final ISA, in the shape LLVM emits for the given target. + + `mfma` selects whether the target has an accumulator register file, which is + what decides both the `.agpr_count` metadata field and the matrix mnemonic. + `ds_load_spelling` selects the gfx11+ LDS naming. `env` is the triple's + environment field, which `AMDGPUTargetID::toString()` writes out verbatim + between the OS and the processor. Everything else is common. + """ + target_id = f"amdgcn-amd-amdhsa-{env}-{arch}" + read = "ds_load_b128" if ds_load_spelling else "ds_read_b64" + write = "ds_store_b32" if ds_load_spelling else "ds_write_b32" + # Operands are v-registers on both sides: gfx90a+ unified the register file, and + # a target without an accumulator file has no `a[...]` to write to at all. + matrix = ( + "v_mfma_f32_16x16x16_f16 v[0:3], v4, v5, v[0:3]" + if mfma + else "v_wmma_f32_16x16x32_bf16 v[0:7], v[8:15], v[16:23], 0" + ) + lines = [ + f'\t.amdgcn_target "{target_id}"', + "\t.text", + f"{kernel}: ; @{kernel}", + *[f"\t{read} v[0:1], v2"] * LDS_READS, + f"\t{write} v3, v4", + f"\t{matrix}", + "\ts_endpgm", + # The parser must skip the descriptor block and bound the body at the + # `.size` terminator, not at s_endpgm. + '\t.section\t.rodata,"a",@progbits', + f"\t.amdhsa_kernel {kernel}", + f"\t\t.amdhsa_next_free_vgpr {num_vgpr}", + "\t.end_amdhsa_kernel", + "\t.text", + ".Lfunc_end0:", + f"\t.size\t{kernel}, .Lfunc_end0-{kernel}", + # Register counts come from these symbols, which every target emits. + f"\t.set {kernel}.num_vgpr, {num_vgpr}", + f"\t.set {kernel}.num_agpr, {num_agpr}", + f"\t.set {kernel}.numbered_sgpr, 53", + f"\t.set {kernel}.private_seg_size, 0", + "\t.set amdgpu.max_num_vgpr, 0", # module-level decoy, must not bind to a kernel + "\t.amdgpu_metadata", + "amdhsa.kernels:", + ] + # Metadata indentation is a contract: " - " opens an entry, kernel keys sit + # at column 4, and argument keys are deeper so they must not be mistaken for + # kernel keys. Keys are emitted in sorted order, as LLVM does. + lines += [f" - .agpr_count: {num_agpr}", " .args:"] if mfma else [" - .args:"] + lines += [ + " - .offset: 0", + " .size: 8", + " .value_kind: global_buffer", + " .group_segment_fixed_size: 39936", + f" .name: {kernel}", + " .private_segment_fixed_size: 0", + " .sgpr_count: 59", + " .sgpr_spill_count: 0", + f" .symbol: {kernel}.kd", + f" .vgpr_count: {vgpr_count}", + " .vgpr_spill_count: 0", + f"amdhsa.target: {target_id}", + ".end_amdgpu_metadata", + ] + return "\n".join(lines) + "\n" + + +TARGETS = { + # Accumulators in use, gfx9 LDS spelling. `.vgpr_count` is the arch+acc total, + # which is why it, and not its split, is the VGPR regression trigger. + "gfx942": dict( + arch="gfx942", + kernel="gemm_0", + mfma=True, + ds_load_spelling=False, + num_vgpr=256, + num_agpr=29, + vgpr_count=285, + env="unknown", + ), + # No accumulator register file, so no `.agpr_count` anywhere; gfx11+ LDS + # spelling. The AGPR count must still come through, from the `.set` symbol. + "gfx1250": dict( + arch="gfx1250", + kernel="fmha_0", + mfma=False, + ds_load_spelling=True, + num_vgpr=942, + num_agpr=0, + vgpr_count=942, + env="", + ), +} + + +def dump_tree(root, spec, **overrides): + return write_tree(root, make_isa(**{**spec, **overrides})) + + +def write_tree(root, text): + """A one-directory dump tree holding exactly the given ISA text.""" + (root / "k").mkdir(parents=True, exist_ok=True) + (root / "k" / "21_final_isa.s").write_text(text) + return root + + +def with_second_entry(text, mutate): + """`text` with a second kernel entry spliced into the metadata list. + + A single-entry file cannot express "one entry was dropped": losing the only + kernel already raises, which hides the case where a healthy kernel keeps a + damaged sibling company and the snapshot still claims to be complete. + """ + head, target = text.split("\namdhsa.target:", 1) + entry = head.split("amdhsa.kernels:\n", 1)[1] + return head + "\n" + mutate(entry).rstrip("\n") + "\namdhsa.target:" + target + + +@pytest.mark.parametrize("target", sorted(TARGETS), ids=sorted(TARGETS)) +def test_resources_are_read_and_diffed_on_both_target_shapes(tmp_path, target): + spec = TARGETS[target] + before = dump_tree(tmp_path / "before", spec) + + records, file_problems = parse_isa(before / "k" / "21_final_isa.s") + assert file_problems == () + (record,) = records.values() + assert record.name == spec["kernel"], "identity comes from the kernel, not its first argument" + + # Nothing may be unreadable on a healthy dump from either shape. The AGPR count + # in particular must come from the `.set` symbol, which every target emits, + # rather than from `.agpr_count`, which only MFMA-capable targets carry. + assert record.unparsed_keys == () + assert record.metrics["agpr"] == irt.Cell.of(spec["num_agpr"]) + assert record.metrics["vgpr"] == irt.Cell.of(spec["vgpr_count"]) + assert record.metrics["arch_vgpr"] == irt.Cell.of(spec["num_vgpr"]) + + # LDS traffic is counted under whichever spelling this target uses; a parser + # that knew only one of them would silently report 0 on the other. + assert record.metrics["lds_read"] == irt.Cell.of(LDS_READS) + assert record.metrics["lds_write"].state == VALUE + + # An unchanged pair is clean; a higher VGPR total is a regression; a missing + # input is untrustworthy and must never be reported as either of the first two. + assert irt.main(["diff", str(before), str(dump_tree(tmp_path / "same", spec))]) == 0 + worse = dump_tree(tmp_path / "worse", spec, vgpr_count=spec["vgpr_count"] + 8) + assert irt.main(["diff", str(before), str(worse)]) == 1 + assert irt.main(["diff", str(before), str(tmp_path / "nonexistent")]) == 2 + + +def test_target_id_is_parsed_under_both_environment_spellings(): + """The processor is the last `-` separated field, not whatever follows a `--`. + + `AMDGPUTargetID::toString()` writes the triple's environment out verbatim, so an + empty environment gives `...amdhsa--gfx942` and an `unknown` one gives + `...amdhsa-unknown-gfx942`. Recognizing only the first leaves `processor` unset on + every dump this repo actually produces, which disables the architecture-mismatch + guard and the scratch-versus-buffer spill classification at the same time, with + nothing but a warning to show for it. + """ + for target_id in ( + "amdgcn-amd-amdhsa--gfx942", + "amdgcn-amd-amdhsa-unknown-gfx942", + "amdgcn-amd-amdhsa-unknown-gfx942:sramecc+:xnack-", + ): + arch = irt.parse_target_id(target_id) + assert arch.processor == "gfx942", target_id + assert arch.gen == 9, target_id + assert arch.spills_via_scratch() is True, target_id + + # A target ID with no concrete processor must stay unidentified rather than take + # its last field on faith: `gfx11-generic` really does end in `generic`. + assert irt.parse_target_id("amdgcn-amd-amdhsa-unknown-gfx11-generic").processor is None + assert irt.parse_target_id("amdgcn-amd-amdhsa").processor is None + + +def test_architecture_mismatch_is_never_reported_as_comparable(tmp_path): + """Two different targets are exit 2, whether or not the tool can name them. + + Register counts from different architectures are not comparable quantities, and + an unidentified processor is not evidence that the two sides match -- comparing + them anyway is how a gfx942-against-gfx950 run printed RESULT: OK. + """ + spec = TARGETS["gfx942"] + before = dump_tree(tmp_path / "before", spec) + assert irt.main(["diff", str(before), str(dump_tree(tmp_path / "same", spec))]) == 0 + assert irt.main(["diff", str(before), str(dump_tree(tmp_path / "gfx950", spec, arch="gfx950"))]) == 2 + + # A target the tool cannot name blocks on its own: without a processor it can + # decide neither that the two sides match nor whether the scratch metrics apply. + # `gfx11-generic` really does end in `generic`, so this is reachable, and the fix + # if it ever matters is to teach `RE_PROC` about it, not to loosen the verdict. + generic = dump_tree(tmp_path / "generic11", spec, arch="gfx11-generic") + assert irt.parse_isa(generic / "k" / "21_final_isa.s")[0].popitem()[1].arch.known is False + assert irt.main(["diff", str(generic), str(dump_tree(tmp_path / "generic11b", spec, arch="gfx11-generic"))]) == 2 + assert irt.main(["diff", str(generic), str(dump_tree(tmp_path / "generic12", spec, arch="gfx12-generic"))]) == 2 + + +def test_an_unreadable_final_isa_blocks_the_comparison(tmp_path): + """A discovered dump file the parser cannot read must not degrade to a warning. + + Skipping it quietly leaves a comparison that is partial but still prints a + full-coverage count line and RESULT: OK, which is the single answer the exit-2 + verdict exists to prevent. + """ + spec = TARGETS["gfx942"] + before, after = tmp_path / "before", tmp_path / "after" + for side in (before, after): + dump_tree(side, spec) + (side / "truncated").mkdir(parents=True) + (side / "truncated" / "21_final_isa.s").write_text("") + + # The healthy half compares cleanly on its own, so the verdict below is + # attributable to the unreadable file and to nothing else in the tree. + assert irt.main(["diff", str(dump_tree(tmp_path / "a", spec)), str(dump_tree(tmp_path / "b", spec))]) == 0 + assert irt.main(["diff", str(before), str(after)]) == 2 + + +def test_target_features_are_part_of_comparability(tmp_path): + """Same processor is not the same target, and the same target has two spellings. + + `xnack` and `sramecc` change code generation and register allocation, so a diff + taken across them reports the build flags rather than the change under test. The + triple's environment field, by contrast, is empty or `unknown` for one and the + same target, so it must not read as a difference. + """ + spec = TARGETS["gfx942"] + base = make_isa(**spec) + feature = lambda suffix: write_tree(tmp_path / f"f{suffix}", base.replace("gfx942", f"gfx942:{suffix}")) + + # Environment spelling: same target, so still comparable. + empty_env = write_tree(tmp_path / "empty_env", base.replace("amdhsa-unknown-gfx942", "amdhsa--gfx942")) + assert irt.main(["diff", str(dump_tree(tmp_path / "plain", spec)), str(empty_env)]) == 0 + + assert irt.main(["diff", str(feature("xnack+")), str(feature("xnack+"))]) == 0 + assert irt.main(["diff", str(feature("xnack+")), str(feature("xnack-"))]) == 2 + assert irt.main(["diff", str(feature("xnack+")), str(feature("sramecc+:xnack+"))]) == 2 + + # Neither side declares a target at all: there is no evidence the two runs are the + # same target, and matching absences are not evidence of a match. + strip = lambda t: "".join( + line for line in t.splitlines(keepends=True) if ".amdgcn_target" not in line and "amdhsa.target:" not in line + ) + assert ( + irt.main(["diff", str(write_tree(tmp_path / "n1", strip(base))), str(write_tree(tmp_path / "n2", strip(base)))]) + == 2 + ) + + +def test_a_negative_resource_count_is_never_reported_as_an_improvement(tmp_path): + """Counts and byte sizes cannot go below zero, so a negative one means a bad dump. + + Accepting it is worse than dropping the kernel: the fall toward it is rendered as + a resource win, which is the one verdict this tool must never invent. + """ + spec = TARGETS["gfx942"] + base = make_isa(**spec) + before = dump_tree(tmp_path / "before", spec) + negative = write_tree(tmp_path / "negative", base.replace(".vgpr_count: 285", ".vgpr_count: -1")) + assert irt.main(["diff", str(before), str(negative)]) == 2 + + # Same for the `.set` symbols, which feed the informational split. + symbol = write_tree(tmp_path / "symbol", base.replace("gemm_0.num_agpr, 29", "gemm_0.num_agpr, -29")) + assert irt.main(["diff", str(before), str(symbol)]) == 2 + + +def test_a_kernel_entry_without_or_with_a_duplicate_identity_is_a_problem(tmp_path): + """An entry the tool cannot key must not vanish from the snapshot. + + Both mutations leave a healthy kernel behind, so the run still produces a table; + what it must not do is call that table complete. + """ + spec = TARGETS["gfx942"] + base = make_isa(**spec) + + anonymous = with_second_entry( + base, + lambda e: e.replace(" .name: gemm_0\n", "").replace(" .symbol: gemm_0.kd\n", ""), + ) + snapshot = irt.collect(write_tree(tmp_path / "anonymous", anonymous)) + assert len(snapshot) == 1 and not snapshot.trustworthy + + # A second entry under a name already seen would otherwise overwrite the first, + # so a stale duplicate could hide a changed counter behind an unchanged one. + duplicated = with_second_entry(base, lambda e: e.replace(".vgpr_count: 285", ".vgpr_count: 999")) + snapshot = irt.collect(write_tree(tmp_path / "duplicated", duplicated)) + assert snapshot.kernels["k::gemm_0"].metrics["vgpr"] == irt.Cell.of(spec["vgpr_count"]) + assert not snapshot.trustworthy + + +def test_an_undecodable_byte_makes_the_instruction_counts_untrustworthy(tmp_path): + """A damaged byte can erase a mnemonic, and an erased mnemonic is not a zero. + + Lenient decoding is still the right call -- a traceback would report worse than + exit 2 does -- but the replacement has to be admitted rather than counted. + """ + spec = TARGETS["gfx942"] + path = dump_tree(tmp_path / "corrupt", spec) / "k" / "21_final_isa.s" + raw = path.read_bytes() + cut = raw.index(b"ds_read_b64") + path.write_bytes(raw[: cut + 3] + b"\xff" + raw[cut + 4 :]) + + records, file_problems = parse_isa(path) + assert file_problems, "an undecodable byte has to be reported, not absorbed" + assert records["gemm_0"].metrics["lds_read"] == irt.Cell.of(LDS_READS - 1), "the count really did change" + assert irt.main(["diff", str(dump_tree(tmp_path / "clean", spec)), str(tmp_path / "corrupt")]) == 2 + + +def test_snapshot_loading_holds_the_same_domain_as_the_assembly_parser(tmp_path): + """A JSON snapshot is the other way to build a `Cell`, so it owes the same checks. + + The non-negativity of a count belongs to the metric, not to the parser that read + it. Enforcing it only where the assembly is read lets a saved snapshot carry a + value the tool would never have written, and a diff down to it reads as a win. + """ + snapshot = irt.collect(dump_tree(tmp_path / "dump", TARGETS["gfx942"])).to_json() + key = next(iter(snapshot["kernels"])) + + def written(name, cell): + copy = json.loads(json.dumps(snapshot)) + copy["kernels"][key]["metrics"]["vgpr"] = cell + (tmp_path / name).write_text(json.dumps(copy)) + return str(tmp_path / name) + + ten = written("ten.json", {"state": VALUE, "value": 10}) + assert irt.main(["diff", ten, ten]) == 0, "a healthy snapshot pair still compares" + assert irt.main(["diff", ten, written("negative.json", {"state": VALUE, "value": -1})]) == 2