`;
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/static/dkao-shape-input.js b/metainfer/tasks/dcu_kernel_auto_opt/static/dkao-shape-input.js
index c40b4da7..a0240b49 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/static/dkao-shape-input.js
+++ b/metainfer/tasks/dcu_kernel_auto_opt/static/dkao-shape-input.js
@@ -24,9 +24,11 @@ function parseCsvInts(raw) {
var W8A8_M_VALUES = [2, 16, 3072];
// Large-prefill boundary added on 2026-08-06 for TP4; extended to the
// Hy3 / MiniMax M3 / GLM5.2 TP8 catalogs on 2026-08-27 via per-topology
-// mValues. DeepSeek TP1/TP8 keep the three original M values. The backend
-// validates shapes against the contract.
+// mValues, plus the M=8 short-decode boundary on 2026-09-20. DeepSeek TP1/TP8
+// keep the three original M values. The backend validates shapes against the
+// contract (MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES = (8, 4096)).
var W8A8_TP4_M_VALUES = [2, 16, 3072, 4096];
+var W8A8_MODEL_TP8_M_VALUES = [2, 8, 16, 3072, 4096];
// Model label → W8A8 GEMM workload. (M, K) @ (K, N) per weight name and TP
// size, mirroring the model cards in the operator docs. DeepSeek keeps the
@@ -94,7 +96,7 @@ var MODEL_WORKLOADS = {
},
{
tp_size: 8,
- mValues: [2, 16, 3072, 4096],
+ mValues: W8A8_MODEL_TP8_M_VALUES,
operators: [
{ operator: "qkv_proj", K: 4096, N: 1280 },
{ operator: "o_proj", K: 1024, N: 4096 },
@@ -129,7 +131,7 @@ var MODEL_WORKLOADS = {
},
{
tp_size: 8,
- mValues: [2, 16, 3072, 4096],
+ mValues: W8A8_MODEL_TP8_M_VALUES,
operators: [
{ operator: "qkv_proj", K: 6144, N: 1280 },
{ operator: "qkv_proj_and_indexer_qk", K: 6144, N: 1536 },
@@ -167,7 +169,7 @@ var MODEL_WORKLOADS = {
},
{
tp_size: 8,
- mValues: [2, 16, 3072, 4096],
+ mValues: W8A8_MODEL_TP8_M_VALUES,
operators: [
{ operator: "fused_qkv_a_proj", K: 6144, N: 2624 },
{ operator: "q_b_proj", K: 2048, N: 2048 },
@@ -186,7 +188,8 @@ function modelCatalog(modelLabel) {
var shapes = [];
(model.topologies || []).forEach(function (topology) {
// Per-topology M values override the defaults; TP4 and the Hy3/MiniMax/
- // GLM TP8 catalogs cover the M=4096 large-prefill boundary.
+ // GLM TP8 catalogs cover the M=8 short-decode and M=4096 large-prefill
+ // boundaries.
var mValues = topology.mValues || (topology.tp_size === 4
? W8A8_TP4_M_VALUES : W8A8_M_VALUES);
topology.operators.forEach(function (item) {
@@ -439,6 +442,24 @@ function ShapeCatalog(_a) {
`;
}
+// Quick select/deselect helper shown above subset (Selected shapes only)
+// catalogs. A subset task starts with NOTHING selected on purpose: the user
+// ticks the shapes that should enter this task (or hits "Select all").
+function SubsetActions(_a) {
+ var count = _a.count, total = _a.total, onAll = _a.onAll, onNone = _a.onNone;
+ return html`
+
+ ${count} / ${total} selected
+
+
+ Only ticked shapes enter this task.
+
+ ${count === 0 ? html`
No shape selected — the task starts empty. Tick at least one shape above (or Select all).
` : null}
+ `;
+}
+
function ApiDefaultsPreview(_a) {
var catalog = _a.catalog;
var tpSizes = catalog
@@ -450,8 +471,10 @@ function ApiDefaultsPreview(_a) {
${catalog.length} logical shapes.
Workloads are grouped by TP size and kept as separate task identities.
- Each operator is optimized at M=2, M=16 and the large-prefill boundary
- M=3072; TP4 additionally covers M=4096.
+ Every operator covers the decode and prefill boundaries of its own
+ catalog (per TP size above), e.g. DeepSeek M=2/16/3072 with TP4 adding
+ M=4096, and the Hy3 / MiniMax M3 / GLM5.2 TP8 catalogs adding the short
+ decode boundary M=8 and M=4096.
${tpSizes.map(function (tp) {
@@ -470,7 +493,16 @@ function ApiDefaultsPreview(_a) {
});
}
});
- var mValues = tp === 4 ? W8A8_TP4_M_VALUES : W8A8_M_VALUES;
+ // Derive the M list from the catalog itself so model-specific
+ // catalogs (Hy3 / MiniMax M3 / GLM5.2 TP8) show their real M
+ // values (including the M=8 and M=4096 boundaries) instead of the
+ // DeepSeek defaults.
+ var mValues = topoShapes
+ .map(function (shape) { return shape.M; })
+ .filter(function (value, index, all) {
+ return all.indexOf(value) === index;
+ })
+ .sort(function (a, b) { return a - b; });
return html`
@@ -502,7 +534,9 @@ function AiShapeSubset(_a) {
var catalog = modelCatalog(model);
var _b = useState(function () {
var parsed = parseShapeRecords(value || "");
- return selectedShapes(parsed, true, catalog, model);
+ // Start EMPTY (no fallback to the full catalog): only shapes the user
+ // ticks (or Select-all) enter this task.
+ return selectedShapes(parsed, false, catalog, model);
}), selected = _b[0], setSelected = _b[1];
var emit = useCallback(function (next) {
@@ -515,6 +549,10 @@ function AiShapeSubset(_a) {
useEffect(function () { emit(selected); }, []);
+ var count = catalog.filter(function (shape) {
+ return selected[shape.id];
+ }).length;
+
return html`
@@ -523,6 +561,15 @@ function AiShapeSubset(_a) {
shapes stay on the trusted fallback and are checked again before the
candidate is published.
+ <${SubsetActions}
+ count=${count}
+ total=${catalog.length}
+ onAll=${function () {
+ var next = {};
+ catalog.forEach(function (shape) { next[shape.id] = true; });
+ emit(next);
+ }}
+ onNone=${function () { emit({}); }} />
<${ShapeCatalog}
shapes=${catalog}
selected=${selected}
@@ -573,7 +620,8 @@ function ManualGpuAssignment(_a) {
var _b = useState(function () {
var parsedShapes = parseShapeRecords(value || "");
return subset
- ? selectedShapes(parsedShapes, true, catalog, model)
+ // Start EMPTY for a manual subset: only ticked shapes enter this task.
+ ? selectedShapes(parsedShapes, false, catalog, model)
: selectedMap(catalog.map(function (shape) {
return shape.id;
}));
@@ -620,6 +668,21 @@ function ManualGpuAssignment(_a) {
appears exactly once; empty GPU cards are allowed.
${subset && html`
+ <${SubsetActions}
+ count=${catalog.filter(function (s) { return selected[s.id]; }).length}
+ total=${catalog.length}
+ onAll=${function () {
+ var nextSelected = {};
+ var nextOwners = Object.assign({}, owners);
+ catalog.forEach(function (shape) {
+ nextSelected[shape.id] = true;
+ if (nextOwners[shape.id] == null) {
+ nextOwners[shape.id] = operatorGroupedOwners(catalog)[shape.id];
+ }
+ });
+ emitState(nextOwners, nextSelected);
+ }}
+ onNone=${function () { emitState(owners, {}); }} />
<${ShapeCatalog}
shapes=${catalog}
selected=${selected}
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/static/dkao.css b/metainfer/tasks/dcu_kernel_auto_opt/static/dkao.css
index 55f7e00f..94b628aa 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/static/dkao.css
+++ b/metainfer/tasks/dcu_kernel_auto_opt/static/dkao.css
@@ -907,3 +907,38 @@
color: var(--muted);
font-size: var(--fs-xs);
}
+
+/* Subset (Selected shapes only): quick select-all / clear bar */
+.dkao-subset-actions {
+ display: flex;
+ align-items: center;
+ gap: 10px;
+ flex-wrap: wrap;
+ margin: 10px 0 4px;
+ padding: 8px 10px;
+ border: 1px dashed var(--border);
+ border-radius: var(--radius-sm);
+ background: var(--bg-elev);
+}
+.dkao-subset-count {
+ font-weight: 700;
+ color: var(--fg);
+ font-size: var(--fs-xs);
+}
+.dkao-subset-actions .dkao-shape-tab {
+ padding: 4px 12px;
+ border: 1px solid var(--border);
+ border-radius: var(--radius-sm);
+ background: var(--panel-2);
+ color: var(--fg);
+ cursor: pointer;
+ font-size: var(--fs-xs);
+}
+.dkao-subset-actions .dkao-shape-tab:disabled {
+ opacity: 0.5;
+ cursor: not-allowed;
+}
+.dkao-subset-hint {
+ color: var(--muted);
+ font-size: var(--fs-xxs);
+}
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/conftest.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/conftest.py
index 31681ab9..2b693a7b 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/conftest.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/conftest.py
@@ -21,6 +21,20 @@ def client(app):
return TestClient(app)
+@pytest.fixture(autouse=True)
+def _no_measurement_gate(monkeypatch):
+ """Offline tests never wait on a real device.
+
+ Every timed benchmark and PMC profile now passes the admission gate
+ (VRAM <= 90% and HCU == 0), which on a busy machine means a 30-minute wait
+ per check. Unit tests exercise command construction and parsing, not
+ admission, so the gate is off unless a test turns it back on (see
+ ``test_gpu_preflight.py``).
+ """
+ monkeypatch.setenv("METAINFER_GPU_PREFLIGHT", "0")
+ yield
+
+
def register_dkao_task(
state_dir, workspace_dir, task_id: str = "dkao-1"
) -> TaskEntry:
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_ahe_child_embed.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_ahe_child_embed.py
new file mode 100644
index 00000000..fca34b5c
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_ahe_child_embed.py
@@ -0,0 +1,175 @@
+"""The AHE page embeds this view for one question of one iteration.
+
+Every harness_evolve question *is* a real DKAO run, but it is spawned by the HE
+orchestrator's CLI and is therefore absent from the WebUI task registry -- so
+``task_or_404`` cannot resolve it and the DKAO page could not render it. The AHE
+page passes the child's artifact directory instead, which these routes accept
+under two conditions, both checked here:
+
+ * the directory really is a question of a *registered* harness_evolve task
+ (otherwise any path on the box could be read through these routes), and
+ * it is read-only: a round's roster belongs to the orchestrator that spawned
+ it, not to whoever is watching it.
+"""
+
+from __future__ import annotations
+
+import json
+
+from metainfer.server import tasks as _tasks
+from metainfer.server.tasks import TaskEntry
+
+
+def _he_task_with_child(tmp_path, *, register_owner: bool = True,
+ task_id: str = "he-embed", question: str = "m16_wqkv_a"):
+ """One registered HE task with one evaluated question on disk."""
+ state_dir = tmp_path / "he-state"
+ workspace = tmp_path / "he-workspace"
+ child_state = (workspace / "children" / "iteration_003" / question
+ / "state")
+ child_ws = child_state.parent / "workspace"
+ child_state.mkdir(parents=True)
+ (child_ws / "workers" / "worker_1").mkdir(parents=True)
+ (child_state / "requirements.json").write_text(json.dumps({
+ "task_id": "ahe-it003-wqkv", "task_type": "dcu-kernel-auto-opt",
+ "answers": {"shapes": []},
+ }), encoding="utf-8")
+ (child_state / "run.json").write_text(json.dumps({
+ "task_id": "ahe-it003-wqkv", "task_type": "dcu-kernel-auto-opt",
+ "current_phase": "parallel_explore", "current_iteration": 2,
+ "last_update": 123.0,
+ }), encoding="utf-8")
+ (child_ws / "plan.json").write_text(json.dumps({
+ "shapes": [{"id": "m16_wqkv_a"}], "execution_mode": "gen-and-opt",
+ }), encoding="utf-8")
+ (child_ws / "workers" / "worker_1" / "status.json").write_text(
+ json.dumps({"worker_id": "worker_1", "gpu": 1, "state": "optimizing"}),
+ encoding="utf-8")
+
+ state_dir.mkdir(parents=True, exist_ok=True)
+ if register_owner:
+ _tasks.add_task(TaskEntry(
+ id=task_id, type="harness-evolve", label="AHE",
+ state_dir=str(state_dir), workspace_dir=str(workspace),
+ created_at=0.0))
+ return child_state, child_ws, workspace
+
+
+def test_a_registered_he_child_is_served_through_the_dkao_routes(
+ tmp_path, client):
+ """The embedded view gets the same payload the standalone page gets."""
+ child_state, _, _ = _he_task_with_child(tmp_path)
+
+ resp = client.get(
+ "/api/dcu-kernel-auto-opt/ignored/summary",
+ params={"state_dir": str(child_state)})
+ assert resp.status_code == 200, resp.text
+ body = resp.json()
+ assert body["run"]["current_phase"] == "parallel_explore"
+ assert body["plan"]["shapes"][0]["id"] == "m16_wqkv_a"
+ assert body["workers"][0]["worker_id"] == "worker_1"
+ assert body["gate"] is not None
+
+ lanes = client.get(
+ "/api/dcu-kernel-auto-opt/ignored/iterations",
+ params={"state_dir": str(child_state)})
+ assert lanes.status_code == 200, lanes.text
+ # one lane per device, and the one this child wrote shows its real state
+ by_id = {w["worker_id"]: w for w in lanes.json()["workers"]}
+ assert "worker_1" in by_id
+ assert by_id["worker_1"]["gpu"] == 1
+ assert by_id["worker_1"]["state"] == "optimizing"
+
+
+def test_the_task_id_in_the_path_is_irrelevant_for_an_embedded_child(
+ tmp_path, client):
+ """The AHE page keys the URL by question; only state_dir resolves it."""
+ child_state, _, _ = _he_task_with_child(tmp_path)
+ other = _he_task_with_child(tmp_path / "second", task_id="he-embed-2",
+ question="m16_other")
+ first = client.get("/api/dcu-kernel-auto-opt/aaa/summary",
+ params={"state_dir": str(child_state)})
+ second = client.get("/api/dcu-kernel-auto-opt/aaa/summary",
+ params={"state_dir": str(other[0])})
+ assert first.status_code == second.status_code == 200
+ assert first.json() == second.json() # same content, different paths
+
+
+def test_an_unregistered_directory_is_refused(tmp_path, client):
+ """Without a registered HE owner, any path could be read: refuse it."""
+ child_state, _, _ = _he_task_with_child(tmp_path, register_owner=False)
+ resp = client.get("/api/dcu-kernel-auto-opt/x/summary",
+ params={"state_dir": str(child_state)})
+ assert resp.status_code == 404
+ assert "harness-evolve" in resp.json()["detail"]
+
+
+def test_arbitrary_paths_are_refused(tmp_path, client):
+ """Not a child of an iteration -> 404, even if the task exists."""
+ _, _, workspace = _he_task_with_child(tmp_path)
+ probes = [
+ tmp_path / "he-state", # the HE task's own state
+ workspace / "children", # too shallow
+ workspace / "children" / "iteration_003", # the question's parent
+ workspace / "children" / "iteration_003" / "m16_wqkv_a", # not "state"
+ tmp_path / "elsewhere" / "state", # unrelated path
+ ]
+ for probe in probes:
+ probe.mkdir(parents=True, exist_ok=True)
+ resp = client.get("/api/dcu-kernel-auto-opt/x/summary",
+ params={"state_dir": str(probe)})
+ assert resp.status_code == 404, f"{probe} was accepted"
+
+
+def test_without_a_state_dir_the_routes_behave_exactly_as_before(
+ tmp_path, client):
+ """No query parameter -> the ordinary registry lookup."""
+ resp = client.get("/api/dcu-kernel-auto-opt/nope/summary")
+ assert resp.status_code == 404
+ assert resp.json()["detail"] == "no such task: nope"
+
+
+def test_the_embedded_view_cannot_write(tmp_path, client):
+ """Read-only for real: the UI hides these, the server refuses them."""
+ child_state, _, _ = _he_task_with_child(tmp_path)
+ writes = [
+ ("post", "/workers/worker_1/guidance", {"text": "steer me"}),
+ ("post", "/workers/worker_1/restart", {}),
+ ("post", "/skills/some-skill/publish", {}),
+ ("post", "/skills/sync", {}),
+ ("post", "/skills/fuse", {"skill_name": "s"}),
+ ("post", "/skills/some-skill/rollback", {}),
+ ("post", "/variants", {"shape_id": "m16"}),
+ ("post", "/rename-repo", {"new_name": "nope"}),
+ ]
+ for method, path, payload in writes:
+ resp = getattr(client, method)(
+ f"/api/dcu-kernel-auto-opt/ignored{path}",
+ params={"state_dir": str(child_state)}, json=payload)
+ assert resp.status_code == 403, f"{path} -> {resp.status_code}"
+ assert "read-only" in resp.json()["detail"]
+
+
+def test_reads_are_not_blocked_by_the_write_guard(tmp_path, client):
+ """The guard must not leak onto GETs, or the mirror would be blank."""
+ child_state, _, _ = _he_task_with_child(tmp_path)
+ for path in ("/summary", "/iterations", "/state-graph", "/skills",
+ "/variants"):
+ resp = client.get(f"/api/dcu-kernel-auto-opt/ignored{path}",
+ params={"state_dir": str(child_state)})
+ assert resp.status_code == 200, f"{path} -> {resp.status_code}"
+
+
+def test_a_registered_task_of_another_type_is_still_rejected(tmp_path, client):
+ """The embed path must not become a way around the type check."""
+ _he_task_with_child(tmp_path) # HE task exists
+ other_state = tmp_path / "dkao-state"
+ other_ws = tmp_path / "dkao-ws"
+ other_state.mkdir(parents=True)
+ other_ws.mkdir(parents=True)
+ _tasks.add_task(TaskEntry(
+ id="calc-1", type="calc-theoretical-value", label="calc",
+ state_dir=str(other_state), workspace_dir=str(other_ws), created_at=0.0))
+
+ resp = client.get("/api/dcu-kernel-auto-opt/calc-1/summary")
+ assert resp.status_code == 409 # wrong type, as before
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_config.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_config.py
index c7ccec16..46bf00d6 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_config.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_config.py
@@ -81,7 +81,26 @@ def test_dsh_framework_resolves_wrapper_model(monkeypatch):
assert cfg.claude_model == "deepseek/deepseek-v4-flash-0731"
-def test_dsh_default_model_when_agent_model_missing(monkeypatch):
+def test_dsh_flash_41_label_resolves_to_flash_model(monkeypatch):
+ """The new frontend label maps to deepseek/deepseek-flash (4.1)."""
+ monkeypatch.delenv("DSH_AGENT_MODEL", raising=False)
+ req = _req("""
+shapes:
+ - {id: m2, M: 2, N: 16, K: 32}
+assignments:
+ worker_0: {gpu: 0, shapes: [m2]}
+""")
+ req["answers"]["agent_framework"] = "dsh"
+ req["answers"]["agent_model"] = "deepseek-flash-4.1"
+
+ assert load_config(req).claude_model == "deepseek/deepseek-flash"
+ # Un-pinned, the older label still resolves to its own pinned build.
+ req["answers"]["agent_model"] = "deepseek-v4-flash"
+ assert load_config(req).claude_model == "deepseek/deepseek-v4-flash-0731"
+
+
+def test_dsh_model_pin_collapses_catalog(monkeypatch):
+ """DSH_AGENT_MODEL pins a mixed-version host to one build for every dsh label."""
monkeypatch.setenv("DSH_AGENT_MODEL", "deepseek/deepseek-v4-flash-0731")
req = _req("""
shapes:
@@ -90,8 +109,24 @@ def test_dsh_default_model_when_agent_model_missing(monkeypatch):
worker_0: {gpu: 0, shapes: [m2]}
""")
req["answers"]["agent_framework"] = "dsh"
+ req["answers"]["agent_model"] = "deepseek-flash-4.1"
- assert load_config(req).claude_model == DSH_DEFAULT_MODEL_ID
+ assert load_config(req).claude_model == "deepseek/deepseek-v4-flash-0731"
+
+
+def test_dsh_default_model_when_agent_model_missing(monkeypatch):
+ monkeypatch.delenv("DSH_AGENT_MODEL", raising=False)
+ req = _req("""
+shapes:
+ - {id: m2, M: 2, N: 16, K: 32}
+assignments:
+ worker_0: {gpu: 0, shapes: [m2]}
+""")
+ req["answers"]["agent_framework"] = "dsh"
+
+ cfg = load_config(req)
+ assert cfg.claude_model == DSH_DEFAULT_MODEL_ID
+ assert cfg.claude_model == "deepseek/deepseek-flash"
def test_dsh_rejects_ccb_only_model(monkeypatch):
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contention.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contention.py
new file mode 100644
index 00000000..a58b2fc3
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contention.py
@@ -0,0 +1,91 @@
+"""Discarding a worker round that was measured on a shared device."""
+
+from __future__ import annotations
+
+import json
+import os
+import time
+from pathlib import Path
+
+from ..orchestrator.contention import (
+ discard_contaminated_round, iteration_is_contaminated, read_contention,
+)
+
+
+def _history(path: Path, rounds: int) -> None:
+ path.parent.mkdir(parents=True, exist_ok=True)
+ with path.open("w", encoding="utf-8") as fh:
+ for i in range(1, rounds + 1):
+ fh.write(json.dumps({
+ "schema_version": 1, "worker_id": "worker_0", "iteration": i,
+ "shape_id": "op_a", "accepted": i % 2 == 1,
+ "metrics": {"median_us": 100.0 - i, "p90_us": 100.0 - i,
+ "passed": True, "mismatch_count": 0},
+ }) + "\n")
+
+
+def _window(path: Path, **fields) -> None:
+ path.parent.mkdir(parents=True, exist_ok=True)
+ path.write_text(json.dumps(fields), encoding="utf-8")
+
+
+def test_a_round_inside_the_window_is_discarded_and_quarantined(tmp_path):
+ state = tmp_path / "state"
+ history = tmp_path / "runs" / "op_a" / "experiments.jsonl"
+ iteration_dir = tmp_path / "runs" / "op_a" / "iteration3"
+ iteration_dir.mkdir(parents=True)
+ _history(history, 3)
+ now = time.time()
+ frozen_from, recovered_at = now - 60, now - 30
+ # the round's work happened while the device was shared
+ stamp = frozen_from + 10
+ os.utime(iteration_dir, (stamp, stamp))
+ _window(state / "contention.json", question="op_a", gpu=0, active=True,
+ frozen_from=frozen_from, recovered_at=recovered_at)
+
+ result = discard_contaminated_round(
+ experiments_path=history, iteration=3, state_dir=state,
+ iteration_dir=iteration_dir)
+
+ assert result["discarded"] is True
+ assert result["rounds_removed"] == 1
+ kept = [json.loads(l) for l in history.read_text().splitlines()]
+ assert [r["iteration"] for r in kept] == [1, 2] # round 3 removed
+ quarantined = json.loads(Path(result["quarantined"][0]).read_text())
+ assert quarantined["iteration"] == 3
+ assert quarantined["metrics"]["median_us"] == 97.0 # evidence preserved
+ assert quarantined["reason"] == "device_shared_during_measurement"
+
+
+def test_a_round_outside_the_window_is_kept(tmp_path):
+ state = tmp_path / "state"
+ history = tmp_path / "runs" / "op_a" / "experiments.jsonl"
+ iteration_dir = tmp_path / "runs" / "op_a" / "iteration1"
+ iteration_dir.mkdir(parents=True)
+ _history(history, 2)
+ now = time.time()
+ # the device was taken over AFTER this round finished
+ _window(state / "contention.json", frozen_from=now + 30, recovered_at=None)
+
+ result = discard_contaminated_round(
+ experiments_path=history, iteration=1, state_dir=state,
+ iteration_dir=iteration_dir)
+ assert result["discarded"] is False
+ assert len(history.read_text().splitlines()) == 2
+
+
+def test_no_contention_record_means_nothing_is_discarded(tmp_path):
+ history = tmp_path / "runs" / "op_a" / "experiments.jsonl"
+ _history(history, 2)
+ result = discard_contaminated_round(experiments_path=history, iteration=1,
+ state_dir=tmp_path / "state")
+ assert result["discarded"] is False
+ assert "no contention" in result["reason"]
+
+
+def test_an_open_window_contaminates_the_running_round(tmp_path):
+ """Contention still active: the round that was running is not trusted."""
+ window = {"frozen_from": time.time() - 10, "active": True}
+ contaminated, why = iteration_is_contaminated(1, window)
+ assert contaminated is True and "taken over" in why
+ assert read_contention(tmp_path / "missing") == {}
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contract_snapshot.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contract_snapshot.py
index ca271898..9a0324ef 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contract_snapshot.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_contract_snapshot.py
@@ -148,20 +148,90 @@ def test_live_contract_accepts_model_catalog_tp8_m4096_shapes():
def test_default_optimization_shapes_unchanged_by_model_tp8_boundary():
- # Adding the model-catalog TP8 M=4096 boundary must not change the
+ # Adding the model-catalog TP8 M=8/M=4096 boundaries must not change the
# DeepSeek-only default workload or its serial-validation fallback scope.
from ..api.int8w8a8gemm.int8_w8a8_gemm_api import (
DEFAULT_OPTIMIZATION_SHAPES,
MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES,
)
- assert MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES == (4096,)
+ assert MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES == (8, 4096)
assert len(DEFAULT_OPTIMIZATION_SHAPES) == 42
assert all(
str(shape["id"]).startswith("tp") for shape in DEFAULT_OPTIMIZATION_SHAPES
)
+def test_live_contract_accepts_model_catalog_tp8_m8_shapes():
+ # Model-catalog TP8 short-decode boundary (2026-09-20): Hy3 / MiniMax M3 /
+ # GLM5.2 TP8 operators are optimizable at M=8 ("Selected shapes only");
+ # DeepSeek TP8 defaults keep the original three M values.
+ contract = resolve_operator_api("Quantized GEMM", "INT8 W8A8")
+ shapes = {
+ "hy3_tp8_qkv_proj_m8": {
+ "tp_size": 8, "operator": "qkv_proj",
+ "M": 8, "N": 1280, "K": 4096,
+ },
+ "hy3_tp8_o_proj_m8": {
+ "tp_size": 8, "operator": "o_proj",
+ "M": 8, "N": 4096, "K": 1024,
+ },
+ "hy3_tp8_shared_gate_up_proj_m8": {
+ "tp_size": 8, "operator": "shared_gate_up_proj",
+ "M": 8, "N": 384, "K": 4096,
+ },
+ "hy3_tp8_shared_down_proj_m8": {
+ "tp_size": 8, "operator": "shared_down_proj",
+ "M": 8, "N": 4096, "K": 192,
+ },
+ "minimax_tp8_qkv_proj_m8": {
+ "tp_size": 8, "operator": "qkv_proj",
+ "M": 8, "N": 1280, "K": 6144,
+ },
+ "minimax_tp8_qkv_proj_and_indexer_qk_m8": {
+ "tp_size": 8, "operator": "qkv_proj_and_indexer_qk",
+ "M": 8, "N": 1536, "K": 6144,
+ },
+ "minimax_tp8_o_proj_m8": {
+ "tp_size": 8, "operator": "o_proj",
+ "M": 8, "N": 6144, "K": 1024,
+ },
+ "minimax_tp8_shared_gate_up_proj_m8": {
+ "tp_size": 8, "operator": "shared_gate_up_proj",
+ "M": 8, "N": 768, "K": 6144,
+ },
+ "minimax_tp8_shared_down_proj_m8": {
+ "tp_size": 8, "operator": "shared_down_proj",
+ "M": 8, "N": 6144, "K": 384,
+ },
+ "glm52_tp8_fused_qkv_a_proj_m8": {
+ "tp_size": 8, "operator": "fused_qkv_a_proj",
+ "M": 8, "N": 2624, "K": 6144,
+ },
+ "glm52_tp8_q_b_proj_m8": {
+ "tp_size": 8, "operator": "q_b_proj",
+ "M": 8, "N": 2048, "K": 2048,
+ },
+ "glm52_tp8_kv_b_proj_m8": {
+ "tp_size": 8, "operator": "kv_b_proj",
+ "M": 8, "N": 3584, "K": 512,
+ },
+ "glm52_tp8_o_proj_m8": {
+ "tp_size": 8, "operator": "o_proj",
+ "M": 8, "N": 6144, "K": 2048,
+ },
+ "glm52_tp8_shared_gate_up_proj_m8": {
+ "tp_size": 8, "operator": "shared_gate_up_proj",
+ "M": 8, "N": 512, "K": 6144,
+ },
+ "glm52_tp8_shared_down_proj_m8": {
+ "tp_size": 8, "operator": "shared_down_proj",
+ "M": 8, "N": 6144, "K": 256,
+ },
+ }
+ validate_contract_shapes(contract, shapes)
+
+
def test_live_contract_accepts_minimax_m3_tp4_shapes():
# Regression for minimaxm3-dsh-tp4-m4096-1-0c2f84a9: the MiniMax M3 TP4
# M=4096 workload stopped in prepare because TP=4 qkv_proj (K, N)
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_execution_profile.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_execution_profile.py
new file mode 100644
index 00000000..dd9996c1
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_execution_profile.py
@@ -0,0 +1,266 @@
+"""Production profile: priority in the device queue, but never a gate bypass.
+
+The operator's decision (2026-09-16, option 1): a production DKAO task outranks
+the evolving harness when devices are handed out, but it still only measures on
+a device that passes VRAM<=90% / HCU==0 — a shared card yields stale numbers no
+matter who asked for the run.
+"""
+
+from __future__ import annotations
+
+import json
+
+from ..orchestrator import config as cfg
+
+import importlib.util
+
+import pytest
+_GPU_BROKER_AVAILABLE = (
+ importlib.util.find_spec("metainfer.orchestrator.gpu_broker") is not None
+)
+requires_gpu_broker = pytest.mark.skipif(
+ not _GPU_BROKER_AVAILABLE,
+ reason="metainfer.orchestrator.gpu_broker is not available in this tree",
+)
+
+
+def _lease(monkeypatch, *, grants, profile, answers=None):
+ """Run lease_gpus_for_task against a fake broker and report the call."""
+ from metainfer.orchestrator import gpu_broker
+
+ calls = []
+
+ class _FakeBroker:
+ def acquire(self, holder, **kwargs):
+ calls.append(kwargs)
+ return [{"gpu": g, "holder": holder} for g in grants]
+
+ monkeypatch.setattr(gpu_broker, "GpuBroker", _FakeBroker)
+ assignments = [cfg.WorkerAssignment(worker_id="worker_0", gpu=0,
+ shape_ids=["s1"])]
+ merged = {"task_id": "t1"}
+ merged.update(answers or {})
+ rebuilt, info = cfg.lease_gpus_for_task(
+ assignments, merged, production=(profile == "production"))
+ return calls, rebuilt, info
+
+
+@requires_gpu_broker
+def test_production_leases_at_the_highest_priority(monkeypatch):
+ calls, rebuilt, info = _lease(monkeypatch, grants=[2], profile="production")
+ assert calls and calls[0]["priority"] == 100 # PRIORITY_PRODUCTION
+ assert calls[0]["meta"]["profile"] == "production"
+ assert info["profile"] == "production" and info["gpus"] == [2]
+ assert rebuilt[0].gpu == 2
+
+
+@requires_gpu_broker
+def test_default_profile_still_uses_production_priority_on_the_broker(monkeypatch):
+ """DKAO has always asked at priority 100; the profile makes it explicit."""
+ calls, _, info = _lease(monkeypatch, grants=[0], profile="default")
+ assert calls[0]["priority"] == 100
+ assert calls[0]["meta"]["profile"] == "default"
+ assert info["profile"] == "default"
+
+
+@requires_gpu_broker
+def test_production_waits_longer_for_a_device_that_passes_the_gate(monkeypatch):
+ """Nothing passing the gate: production keeps waiting instead of giving up."""
+ calls, _, _ = _lease(monkeypatch, grants=[0], profile="production",
+ answers={"gpu_lease_max_waits": "2",
+ "production_max_waits": "48",
+ "gpu_lease_wait_seconds": "0"})
+ assert calls[0]["priority"] == 100
+
+
+@requires_gpu_broker
+def test_no_device_after_the_budget_reports_the_gate_honestly(monkeypatch):
+ import pytest
+
+ class _EmptyBroker:
+ def acquire(self, holder, **kwargs):
+ return []
+
+ from metainfer.orchestrator import gpu_broker
+ monkeypatch.setattr(gpu_broker, "GpuBroker", _EmptyBroker)
+ assignments = [cfg.WorkerAssignment(worker_id="worker_0", gpu=0,
+ shape_ids=["s1"])]
+ with pytest.raises(RuntimeError) as excinfo:
+ cfg.lease_gpus_for_task(assignments, {"task_id": "t1",
+ "gpu_lease_max_waits": "2",
+ "gpu_lease_wait_seconds": "0"},
+ production=True)
+ message = str(excinfo.value)
+ assert "production" in message
+ # and it must say the gate was NOT skipped
+ assert "not skipped" in message and "shared card" in message
+
+
+@requires_gpu_broker
+def test_load_config_accepts_the_profile_and_forces_the_broker(monkeypatch):
+ """``production`` implies broker-mediated scheduling (that is the queue)."""
+ from metainfer.orchestrator import gpu_broker
+
+ seen = {}
+
+ class _FakeBroker:
+ def acquire(self, holder, **kwargs):
+ seen.update(kwargs)
+ return [{"gpu": 1, "holder": holder}]
+
+ monkeypatch.setattr(gpu_broker, "GpuBroker", _FakeBroker)
+ req = {
+ "task_id": "prod-1",
+ "answers": {
+ "execution_mode": "Generate & optimize (auto-create kernel repo)",
+ "target_repo_path": "11111",
+ "shape_config": ("model: m\nassignment_mode: ai\nshape_scope: subset\n"
+ "shapes:\n - {id: s1, M: 16, N: 64, K: 32, tp_size: 8,"
+ " operator: o_proj}\n"),
+ "max_iterations": "1",
+ "execution_profile": "production",
+ # deliberately ask for occupy: production must override it
+ "gpu_mode": "occupy",
+ },
+ }
+ config = cfg.load_config(req)
+ assert config.execution_profile == "production"
+ assert config.gpu_mode == "scheduled" # broker carries the priority
+ assert seen["priority"] == 100
+ assert config.gpu_leases.get("profile") == "production"
+
+
+def test_an_unknown_profile_is_rejected():
+ import pytest
+
+ req = {
+ "task_id": "bad-1",
+ "answers": {
+ "execution_mode": "Generate & optimize (auto-create kernel repo)",
+ "target_repo_path": "11111",
+ "shape_config": ("model: m\nassignment_mode: ai\nshape_scope: subset\n"
+ "shapes:\n - {id: s1, M: 16, N: 64, K: 32, tp_size: 8,"
+ " operator: o_proj}\n"),
+ "max_iterations": "1",
+ "execution_profile": "fastest",
+ },
+ }
+ with pytest.raises(ValueError):
+ cfg.load_config(req)
+
+
+@requires_gpu_broker
+def test_replace_assignments_keeps_the_profile(monkeypatch):
+ """Re-planning the workers must not silently downgrade a production task."""
+ from metainfer.orchestrator import gpu_broker
+
+ class _FakeBroker:
+ def acquire(self, holder, **kwargs):
+ return [{"gpu": 3, "holder": holder}]
+
+ monkeypatch.setattr(gpu_broker, "GpuBroker", _FakeBroker)
+ req = {
+ "task_id": "prod-2",
+ "answers": {
+ "execution_mode": "Generate & optimize (auto-create kernel repo)",
+ "target_repo_path": "11111",
+ "shape_config": ("model: m\nassignment_mode: ai\nshape_scope: subset\n"
+ "shapes:\n - {id: s1, M: 16, N: 64, K: 32, tp_size: 8,"
+ " operator: o_proj}\n"),
+ "max_iterations": "1",
+ "execution_profile": "production",
+ },
+ }
+ config = cfg.load_config(req)
+ replanned = cfg.replace_assignments(config, config.assignments)
+ assert replanned.execution_profile == "production"
+ assert replanned.gpu_leases.get("priority") == 100
+
+
+def test_gpu_leases_acquired_event_names_the_profile():
+ """The timeline event carries the rank, so the GPU view can label a holder."""
+ from ..orchestrator.gen_and_opt_pipeline import GenAndOptPipeline
+
+ events = []
+
+ class _Store:
+ task_dir = None
+
+ def append_timeline(self, name, payload):
+ events.append((name, payload))
+
+ pipeline = GenAndOptPipeline.__new__(GenAndOptPipeline)
+ pipeline.store = _Store()
+ pipeline._register_gpu_leases(type("C", (), {
+ "gpu_leases": {"holder": "dkao:prod-2", "gpus": [3], "mode": "scheduled",
+ "profile": "production", "priority": 100},
+ })())
+ acquired = dict(events)["gpu_leases_acquired"]
+ assert acquired["profile"] == "production"
+ assert acquired["priority"] == 100
+
+
+def test_summary_reports_the_profile_and_the_held_devices(tmp_path):
+ """The task view shows production mode + which devices the task leased."""
+ from ..server.routes import _execution_profile, _latest_gpu_lease
+
+ state_dir = tmp_path / "state"
+ state_dir.mkdir()
+ (state_dir / "requirements.json").write_text(json.dumps({
+ "task_id": "prod-3",
+ "answers": {"execution_profile": "production"},
+ }), encoding="utf-8")
+ (state_dir / "timeline.jsonl").write_text("\n".join([
+ json.dumps({"ts": 1.0, "type": "phase_start",
+ "payload": {"phase": "prepare"}}),
+ json.dumps({"ts": 2.0, "type": "gpu_leases_acquired", "payload": {
+ "holder": "dkao:prod-3", "gpus": [1, 2], "mode": "scheduled",
+ "profile": "production", "priority": 100}}),
+ json.dumps({"ts": 3.0, "type": "gpu_leases_released",
+ "payload": {"holder": "dkao:prod-3", "gpus": [1, 2]}}),
+ ]), encoding="utf-8")
+
+ assert _execution_profile(state_dir) == "production"
+ leases = _latest_gpu_lease(state_dir)
+ assert leases["gpus"] == [1, 2] and leases["priority"] == 100
+
+
+def test_summary_defaults_to_the_default_profile(tmp_path):
+ from ..server.routes import _execution_profile, _latest_gpu_lease
+
+ state_dir = tmp_path / "state"
+ state_dir.mkdir()
+ (state_dir / "requirements.json").write_text(json.dumps({
+ "task_id": "plain-1", "answers": {}}), encoding="utf-8")
+ assert _execution_profile(state_dir) == "default"
+ assert _latest_gpu_lease(state_dir) == {}
+
+
+def test_gate_status_summarises_the_blocked_checks(tmp_path):
+ """A task parked on the gate must say why, in its own state dir."""
+ from ..server.routes import _gate_status
+
+ state_dir = tmp_path / "state"
+ state_dir.mkdir()
+ rows = [
+ {"ts": 1.0, "event": "gate_blocked", "site": "generate_probe",
+ "gpu": 0, "attempt": 1, "max_waits": 48,
+ "reasons": ["device busy: HCU 98.0% > 0.0%"]},
+ {"ts": 1801.0, "event": "gate_ok", "site": "generate_probe", "gpu": 1,
+ "passing": [1], "reasons": []},
+ ]
+ (state_dir / "measurement_gate.jsonl").write_text(
+ "\n".join(json.dumps(r) for r in rows), encoding="utf-8")
+
+ gate = _gate_status(state_dir)
+ assert gate["checks"] == 2 and gate["blocked"] == 1
+ assert gate["last_event"] == "gate_ok" and gate["last_gpu"] == 1
+ assert gate["last_reasons"] == []
+
+
+def test_gate_status_is_empty_without_an_audit_trail(tmp_path):
+ from ..server.routes import _gate_status
+
+ state_dir = tmp_path / "state"
+ state_dir.mkdir()
+ assert _gate_status(state_dir)["checks"] == 0
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gate_policy.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gate_policy.py
new file mode 100644
index 00000000..8cb57b72
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gate_policy.py
@@ -0,0 +1,158 @@
+"""gates.yaml is wired: the runtime reads it, the manifest decides if it applies."""
+
+from __future__ import annotations
+
+import shutil
+
+import pytest
+import yaml
+
+from ..orchestrator import gate_policy as gp
+from ..orchestrator import planner
+from ..orchestrator.harness_io import default_harness_dir
+
+
+@pytest.fixture(autouse=True)
+def _clear_cache():
+ gp.reset_cache()
+ yield
+ gp.reset_cache()
+
+
+def _workspace(tmp_path, *, gates_wired=True, policy_wired=True):
+ root = tmp_path / "harness"
+ shutil.copytree(default_harness_dir(), root)
+ manifest = yaml.safe_load((root / "manifest.yaml").read_text(encoding="utf-8"))
+ manifest["components"]["gates"]["wired"] = gates_wired
+ manifest["components"]["planner_policy"]["wired"] = policy_wired
+ (root / "manifest.yaml").write_text(yaml.safe_dump(manifest, sort_keys=False),
+ encoding="utf-8")
+ return root
+
+
+def _edit_gates(root, **values):
+ data = yaml.safe_load((root / "gates.yaml").read_text(encoding="utf-8"))
+ for key, value in values.items():
+ node = data["gates"]
+ parts = key.split(".")
+ for part in parts[:-1]:
+ node = node.setdefault(part, {})
+ node[parts[-1]] = value
+ (root / "gates.yaml").write_text(yaml.safe_dump(data, sort_keys=False),
+ encoding="utf-8")
+
+
+def test_seed_matches_builtin_defaults():
+ """The seed file mirrors the historical constants (no surprise changes)."""
+ seed = yaml.safe_load(
+ (default_harness_dir() / "gates.yaml").read_text(encoding="utf-8")
+ )["gates"]
+ defaults = gp.builtin_gates()
+ assert seed["round_acceptance_improvement_percent"] == (
+ defaults["round_acceptance_improvement_percent"])
+ assert seed["shadow"]["min_improvement_percent"] == (
+ defaults["shadow"]["min_improvement_percent"])
+ assert seed["plateau"]["max_regression_percent"] == (
+ defaults["plateau"]["max_regression_percent"])
+ assert seed["isa_gate"]["required_valid_isa_guided_rounds"] == (
+ defaults["isa_gate"]["required_valid_isa_guided_rounds"])
+ assert seed["task_budget"]["default_max_iterations"] == (
+ defaults["task_budget"]["default_max_iterations"])
+
+
+def test_editing_gates_changes_runtime_values(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ _edit_gates(root, **{
+ "round_acceptance_improvement_percent": 2.5,
+ "p90_guard": "1.05",
+ "shadow.enabled": False,
+ "shadow.min_improvement_percent": 0.7,
+ "plateau.max_regression_percent": 3.5,
+ "isa_gate.required_valid_isa_guided_rounds": 4,
+ "task_budget.default_max_iterations": 6,
+ })
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+
+ assert gp.round_acceptance_improvement_percent() == 2.5
+ assert gp.p90_tolerance() == 1.05
+ assert gp.shadow_enabled() is False
+ assert gp.shadow_min_improvement_percent() == 0.7
+ assert gp.plateau_max_regression_percent() == 3.5
+ assert gp.isa_required_valid_rounds() == 4
+ assert gp.default_max_iterations() == 6
+ # the pipeline helpers see the same values
+ from ..orchestrator import w8a8_pipeline as W
+ assert W._gates.round_acceptance_improvement_percent() == 2.5
+
+
+def test_unwired_gates_component_is_ignored(tmp_path, monkeypatch):
+ """wired:false means the file is inert — the manifest is meaningful."""
+ root = _workspace(tmp_path, gates_wired=False)
+ _edit_gates(root, **{"round_acceptance_improvement_percent": 9.0})
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+ assert gp.round_acceptance_improvement_percent() == 1.0
+ assert gp.component_wired("gates") is False
+
+
+def test_unwired_planner_policy_falls_back_to_builtin(tmp_path, monkeypatch):
+ root = _workspace(tmp_path, policy_wired=False)
+ data = yaml.safe_load((root / "planner_policy.yaml").read_text(encoding="utf-8"))
+ data["fallback"]["m16"] = ["consolidate"] * 8
+ (root / "planner_policy.yaml").write_text(yaml.safe_dump(data, sort_keys=False),
+ encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ effective = planner.policy()
+ assert effective["fallback"]["m16"] != ["consolidate"] * 8
+ assert effective["fallback"] == planner._builtin_policy()["fallback"]
+
+
+def test_required_hip_rounds_rule_and_literal(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+ assert gp.required_hip_rounds(10) == 8 # documented expression
+ _edit_gates(root, **{"isa_gate.required_hip_rounds_rule": 3})
+ gp.reset_cache()
+ assert gp.required_hip_rounds(10) == 3 # literal override
+
+
+def test_p90_guard_mode_parsing(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+ assert gp.p90_guard_mode() == "no_p90_regression"
+ assert gp.p90_tolerance() == 1.0
+ _edit_gates(root, **{"p90_guard": "not-a-number"})
+ gp.reset_cache()
+ assert gp.p90_tolerance() == 1.0 # never loosens by accident
+
+
+def test_snapshot_records_what_a_run_used(tmp_path, monkeypatch):
+ """Children record their effective gates so the mechanism gate can verify."""
+ import json
+
+ root = _workspace(tmp_path)
+ _edit_gates(root, **{"round_acceptance_improvement_percent": 3.25})
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+ state = tmp_path / "state"
+ out = gp.snapshot(state)
+ assert out is not None and out.is_file()
+ payload = json.loads(out.read_text(encoding="utf-8"))
+ assert payload["gates"]["round_acceptance_improvement_percent"] == 3.25
+ assert payload["gates_wired"] is True
+ assert payload["harness_root"] == str(root)
+
+
+def test_snapshot_marks_unwired_gates(tmp_path, monkeypatch):
+ import json
+
+ root = _workspace(tmp_path, gates_wired=False)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ gp.reset_cache()
+ out = gp.snapshot(tmp_path / "state2")
+ payload = json.loads(out.read_text(encoding="utf-8"))
+ assert payload["gates_wired"] is False
+ assert payload["gates"]["round_acceptance_improvement_percent"] == 1.0
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_mode.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_mode.py
new file mode 100644
index 00000000..213135c2
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_mode.py
@@ -0,0 +1,108 @@
+"""DKAO GPU modes: occupy (default, pins devices) vs scheduled (broker leases)."""
+
+from __future__ import annotations
+
+import pytest
+
+import importlib.util
+
+_GPU_BROKER_AVAILABLE = (
+ importlib.util.find_spec("metainfer.orchestrator.gpu_broker") is not None
+)
+requires_gpu_broker = pytest.mark.skipif(
+ not _GPU_BROKER_AVAILABLE,
+ reason="metainfer.orchestrator.gpu_broker is not available in this tree",
+)
+
+
+def _assignments(spec):
+ from ..orchestrator.config import WorkerAssignment
+ return [WorkerAssignment(worker_id=w, gpu=g, shape_ids=list(s))
+ for w, g, s in spec]
+
+
+def test_rebalance_merges_workers_when_leases_are_fewer():
+ """2 leases + 4 planned workers -> 2 lanes, every shape kept."""
+ from ..orchestrator.config import rebalance_for_gpus
+ planned = _assignments([
+ ("worker_0", 0, ["a"]), ("worker_1", 1, ["b"]),
+ ("worker_2", 2, ["c"]), ("worker_3", 3, ["d"]),
+ ])
+ lanes = rebalance_for_gpus(planned, [2, 3])
+ assert len(lanes) == 2 # one lane per device
+ assert sorted(l.gpu for l in lanes) == [2, 3]
+ assert {l.worker_id for l in lanes} == {"worker_2", "worker_3"}
+ assert sorted(s for l in lanes for s in l.shape_ids) == ["a", "b", "c", "d"]
+
+
+def test_rebalance_keeps_one_lane_per_worker_when_leases_are_plentiful():
+ from ..orchestrator.config import rebalance_for_gpus
+ planned = _assignments([("worker_0", 0, ["a"]), ("worker_1", 1, ["b"])])
+ lanes = rebalance_for_gpus(planned, [0, 1, 2, 3])
+ assert len(lanes) == 2 # no empty lanes
+ assert sorted(l.gpu for l in lanes) == [0, 1]
+
+
+@requires_gpu_broker
+def test_lease_gpus_for_task_layouts_workers_on_leases(monkeypatch):
+ from ..orchestrator import config as C
+ from metainfer.orchestrator import gpu_broker
+
+ class FakeBroker:
+ def __init__(self, *a, **k):
+ self.calls = []
+
+ def acquire(self, holder, **kwargs):
+ self.calls.append((holder, kwargs))
+ return [{"gpu": 2}, {"gpu": 3}]
+
+ monkeypatch.setattr(gpu_broker, "GpuBroker", FakeBroker)
+ planned = _assignments([
+ ("worker_0", 0, ["a"]), ("worker_1", 1, ["b"]),
+ ("worker_2", 2, ["c"]), ("worker_3", 3, ["d"]),
+ ])
+ lanes, info = C.lease_gpus_for_task(planned, {"task_id": "t-1"})
+ assert info["mode"] == "scheduled" and info["gpus"] == [2, 3]
+ assert info["holder"] == "dkao:t-1"
+ assert info["requested_workers"] == 4
+ assert info["mapping"] == {"worker_2": 2, "worker_3": 3}
+ assert len(lanes) == 2
+
+
+@requires_gpu_broker
+def test_lease_gpus_for_task_fails_loudly_when_nothing_is_free(monkeypatch):
+ from ..orchestrator import config as C
+ from metainfer.orchestrator import gpu_broker
+
+ class EmptyBroker:
+ def __init__(self, *a, **k):
+ pass
+
+ def acquire(self, holder, **kwargs):
+ return []
+
+ monkeypatch.setattr(gpu_broker, "GpuBroker", EmptyBroker)
+ with pytest.raises(RuntimeError, match="no GPU lease"):
+ C.lease_gpus_for_task(_assignments([("worker_0", 0, ["a"])]),
+ {"task_id": "t-2", "gpu_lease_max_waits": 2})
+
+
+def test_scheduled_mode_skips_manual_pin_validation():
+ """manual (worker_N -> GPU N) is only enforced in occupy mode."""
+ from ..orchestrator.config import OptimizerConfig, WorkerAssignment
+ import dataclasses
+ base = dict(operator="Quantized GEMM", dtype="INT8 W8A8", hardware="dcu",
+ kernel_language="HIP C++", claude_model="m",
+ execution_mode="gen_and_opt", target_repo_path=None,
+ shapes={}, assignments=[WorkerAssignment("worker_2", 3, ["a"])],
+ shape_scope="all", mock_iterations=1,
+ minimum_improvement_percent=1.0)
+ occupy = OptimizerConfig(assignment_mode="manual", **base)
+ scheduled = OptimizerConfig(assignment_mode="scheduled", gpu_mode="scheduled",
+ **base)
+ assert occupy.gpu_mode == "occupy"
+ assert scheduled.gpu_mode == "scheduled"
+ # the pipeline only validates worker_N -> GPU N for occupy
+ assert (occupy.assignment_mode == "manual" and occupy.gpu_mode == "occupy")
+ assert not (scheduled.assignment_mode == "manual"
+ and scheduled.gpu_mode == "occupy")
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_preflight.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_preflight.py
new file mode 100644
index 00000000..faa21706
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_preflight.py
@@ -0,0 +1,476 @@
+"""GPU admission pre-check: usage rule + pollution marking."""
+
+from __future__ import annotations
+
+import json
+
+import pytest
+
+from ..orchestrator import gpu_preflight as gp
+
+import importlib.util
+from pathlib import Path
+
+_GPU_BROKER_AVAILABLE = (
+ importlib.util.find_spec("metainfer.orchestrator.gpu_broker") is not None
+)
+requires_gpu_broker = pytest.mark.skipif(
+ not _GPU_BROKER_AVAILABLE,
+ reason="metainfer.orchestrator.gpu_broker is not available in this tree",
+)
+_DTK_HIPPROF_AVAILABLE = Path("/opt/dtk/bin/hipprof").is_file()
+requires_dtk_hipprof = pytest.mark.skipif(
+ not _DTK_HIPPROF_AVAILABLE,
+ reason="/opt/dtk/bin/hipprof is required (DCU toolchain is absent)",
+)
+
+SMI_SAMPLE = """
+================================= System Management Interface ==================================
+================================================================================================
+HCU Temp AvgPwr Perf PwrCap VRAM% HCU% Mode
+0 51.0C 145.0W manual 300.0W 0% 0.0% Normal
+1 49.0C 147.0W manual 300.0W 12% 4.5% Normal
+2 57.0C 251.0W manual 300.0W 40% 97.0% Normal
+3 58.0C 254.0W manual 300.0W 92% 81.7% Normal
+======================================== End of SMI Log =========================================
+"""
+
+#: What hy-smi actually prints inside a container that cannot open the
+#: management interface: the HCU% column (and every other utilisation value)
+#: is a placeholder, NOT a zero.
+SMI_UNAVAILABLE = """
+================================= System Management Interface ==================================
+================================================================================================
+HCU Temp AvgPwr Perf PwrCap VRAM% HCU% Mode
+0 52.0C 146.0W high 300.0W 0% N/A Normal
+1 50.0C 149.0W high 300.0W 0% N/A Normal
+2 50.0C 148.0W high 300.0W 0% N/A Normal
+3 51.0C 148.0W high 300.0W 0% N/A Normal
+======================================== End of SMI Log =========================================
+"""
+
+
+def test_parse_hy_smi_rows():
+ parsed = gp.parse_hy_smi(SMI_SAMPLE)
+ assert set(parsed) == {0, 1, 2, 3}
+ assert parsed[0]["vram_percent"] == 0.0
+ assert parsed[2]["util_percent"] == 97.0
+ assert parsed[3]["power_w"] == 254.0
+ assert parsed[1]["temp_c"] == 49.0
+
+
+def test_na_utilisation_is_unknown_not_zero():
+ """``N/A`` must not become ``0.0``: that is what admitted busy cards."""
+ parsed = gp.parse_hy_smi(SMI_UNAVAILABLE)
+ assert set(parsed) == {0, 1, 2, 3}
+ assert "util_percent" not in parsed[0] # unknown, not idle
+ assert parsed[0]["power_w"] == 146.0 # real fields are kept
+ assert gp.smi_is_incomplete(SMI_UNAVAILABLE) is True
+ assert gp.smi_is_incomplete(SMI_SAMPLE) is False
+
+
+def test_sample_falls_back_to_sysfs_and_gates_busy_cards():
+ """hy-smi cannot see HCU%: sysfs decides, and a busy card stays blocked."""
+ def sysfs():
+ return {0: {"util_percent": 0.0, "vram_percent": 2.0},
+ 1: {"util_percent": 0.0, "vram_percent": 2.0},
+ 2: {"util_percent": 74.0, "vram_percent": 20.0},
+ 3: {"util_percent": 0.0, "vram_percent": 95.0}}
+
+ states = gp.sample_gpu_state([0, 1, 2, 3], samples=1, interval_s=0,
+ smi_reader=lambda: SMI_UNAVAILABLE,
+ sysfs_reader=sysfs)
+ assert states[2]["util_percent"] == 74.0
+ assert states[3]["vram_percent"] == 95.0
+ assert not any(row.get("unavailable") for row in states.values())
+ checks = {g: gp.check_gpu(g, states[g], total_gb=64.0) for g in range(4)}
+ assert [g for g in range(4) if checks[g]["usable"]] == [0, 1]
+ assert checks[2]["usable"] is False and checks[3]["usable"] is False
+
+
+def test_no_source_at_all_means_unavailable(monkeypatch):
+ states = gp.sample_gpu_state([0, 1], samples=1, interval_s=0,
+ smi_reader=lambda: "",
+ sysfs_reader=lambda: {})
+ assert states[0]["unavailable"] == 1.0
+ assert gp.check_gpu(0, states[0])["usable"] is False
+
+
+def test_usability_requires_vram_limit_and_idle_device():
+ """Operator rule: VRAM <= 90% *and* HCU == 0, nothing else counts."""
+ ok = gp.check_gpu(0, {"vram_percent": 89.0, "util_percent": 0.0},
+ total_gb=64.0)
+ assert ok["usable"] is True and ok["measurement_suspect"] is False
+
+ too_full = gp.check_gpu(1, {"vram_percent": 91.0, "util_percent": 0.0},
+ total_gb=64.0)
+ assert too_full["usable"] is False
+ assert any("VRAM 91%" in r for r in too_full["reasons"])
+
+ busy = gp.check_gpu(2, {"vram_percent": 10.0, "util_percent": 3.0},
+ total_gb=64.0)
+ assert busy["usable"] is False # HCU must be 0
+ assert any("HCU 3.0%" in r for r in busy["reasons"])
+
+ assert gp.DEFAULT_VRAM_LIMIT_PERCENT == 90.0
+ assert gp.DEFAULT_UTIL_TOLERANCE == 0.0
+
+
+def test_busy_device_is_unusable_and_suspect():
+ busy = gp.check_gpu(2, {"vram_percent": 40.0, "util_percent": 97.0,
+ "power_w": 251.0}, foreign_pids=[12345])
+ assert busy["usable"] is False # a busy device is never used
+ assert busy["measurement_suspect"] is True
+ assert any("busy" in r for r in busy["reasons"])
+ # power draw is recorded for diagnosis but is not part of the gate
+ assert busy["power_w"] == 251.0
+ assert not any("power" in r for r in busy["reasons"])
+ # foreign pids are informational (idle holders are harmless)
+ assert any("foreign" in n for n in busy["notes"])
+
+
+def test_idle_foreign_holders_are_not_suspicion():
+ """A device at 0% util / idle power is clean even if other containers hold it."""
+ idle = gp.check_gpu(0, {"vram_percent": 0.0, "util_percent": 0.0,
+ "power_w": 144.0}, foreign_pids=[1, 2, 3])
+ assert idle["usable"] is True
+ assert idle["measurement_suspect"] is False
+ assert idle["reasons"] == []
+ assert any("foreign" in n for n in idle["notes"])
+
+
+def test_an_idle_card_passes_even_at_high_board_power():
+ """The gate is exactly "VRAM <= 90% and HCU == 0" — nothing else blocks."""
+ check = gp.check_gpu(1, {"vram_percent": 5.0, "util_percent": 0.0,
+ "power_w": 250.0})
+ assert check["usable"] is True
+ assert check["reasons"] == []
+
+
+def test_preflight_gate_lists_all_and_only_passing_devices(monkeypatch):
+ monkeypatch.setattr(gp, "sample_gpu_state",
+ lambda ids=None, **kw: {
+ 0: {"vram_percent": 1.0, "util_percent": 0.0},
+ 1: {"vram_percent": 5.0, "util_percent": 0.0,
+ "power_w": 250.0},
+ 2: {"vram_percent": 30.0, "util_percent": 4.0},
+ 3: {"vram_percent": 97.0, "util_percent": 0.0},
+ })
+ monkeypatch.setattr(gp, "foreign_kfd_pids", lambda *a, **k: [])
+ monkeypatch.setattr(gp, "total_vram_gb", lambda *a, **k: 64.0)
+ plan = gp.preflight_gpus([0, 1, 2, 3])
+ assert plan["clean_ids"] == [0, 1]
+ assert plan["suspect_ids"] == []
+ assert plan["blocked_ids"] == [2, 3]
+ assert plan["over_limit_ids"] == [3]
+ # only gate-passing devices are ever offered
+ assert plan["preferred"] == [0, 1]
+ assert "HCU 4.0%" in plan["gate_reasons"]["2"][0]
+
+
+def test_blocked_devices_are_never_offered_for_a_measurement(monkeypatch):
+ """Both cards full: the gate offers nothing so the caller waits and re-checks.
+
+ Waiting 30 min and retrying (up to 48 times) is the operator rule; silently
+ measuring on a device that failed the gate is what this prevents.
+ """
+ monkeypatch.setattr(gp, "sample_gpu_state",
+ lambda ids=None, **kw: {
+ 0: {"vram_percent": 97.0, "util_percent": 5.0},
+ 1: {"vram_percent": 99.0, "util_percent": 5.0},
+ })
+ monkeypatch.setattr(gp, "foreign_kfd_pids", lambda *a, **k: [])
+ monkeypatch.setattr(gp, "total_vram_gb", lambda *a, **k: 64.0)
+ plan = gp.preflight_gpus([0, 1])
+ assert plan["preferred"] == []
+ assert plan["clean_ids"] == []
+ assert plan["blocked_ids"] == [0, 1]
+ assert "VRAM 97% > 90% limit" in plan["gate_reasons"]["0"]
+
+
+def test_unreadable_device_state_is_not_usable(monkeypatch):
+ """No HCU/VRAM reading must never be treated as "idle"."""
+ monkeypatch.setattr(gp, "foreign_kfd_pids", lambda *a, **k: [])
+ check = gp.check_gpu(0, {"unavailable": 1.0,
+ "missing_fields": ["vram_percent", "util_percent"]})
+ assert check["usable"] is False
+ assert check["unavailable"] is True
+ assert any("state unavailable" in r for r in check["reasons"])
+
+
+def test_foreign_kfd_pids_detects_invisible_pids(tmp_path):
+ (tmp_path / "1").write_text("", encoding="utf-8") # our own pid
+ (tmp_path / "99999999").write_text("", encoding="utf-8") # foreign
+ pids = gp.foreign_kfd_pids(tmp_path)
+ assert 99999999 in pids
+ assert 1 not in pids
+
+
+def test_preflight_enabled_env_and_answers(monkeypatch):
+ monkeypatch.delenv("METAINFER_GPU_PREFLIGHT", raising=False)
+ assert gp.preflight_enabled({}) is True # default on
+ assert gp.preflight_enabled({"gpu_preflight": "false"}) is False
+ monkeypatch.setenv("METAINFER_GPU_PREFLIGHT", "1")
+ assert gp.preflight_enabled({"gpu_preflight": "false"}) is True
+ monkeypatch.setenv("METAINFER_GPU_PREFLIGHT", "off")
+ assert gp.preflight_enabled({}) is False
+
+
+def test_disabled_preflight_returns_default_order():
+ plan = gp.preflight_gpus([0, 1, 2, 3], enabled=False)
+ assert plan["enabled"] is False
+ assert plan["preferred"] == [0, 1, 2, 3]
+ assert plan["gpus"] == {}
+
+
+# ---------------------------------------------------------------- sub-task gate
+# A DKAO child iterates for hours, so the device can be taken over long after
+# the parent admitted it. These tests pin the *child-side* rule: no timed
+# benchmark and no PMC profile may start on a device that fails the gate.
+
+def test_measurement_gate_waits_then_admits(monkeypatch, tmp_path):
+ from ..orchestrator import w8a8_pipeline as wp
+
+ blocked = {"gpus": {2: {"usable": False, "util_percent": 91.0,
+ "vram_percent": 10.0,
+ "reasons": ["device busy: HCU 91.0% > 0.0%"]}}}
+ ok = {"gpus": {2: {"usable": True, "util_percent": 0.0,
+ "vram_percent": 5.0, "state_source": "hy-smi"}}}
+ plans = [blocked, ok]
+ calls = []
+
+ def fake_preflight(ids, **kw):
+ calls.append((list(ids), kw.get("vram_limit_percent"),
+ kw.get("util_tolerance")))
+ return plans.pop(0)
+
+ sleeps = []
+ monkeypatch.setattr("metainfer.tasks.dcu_kernel_auto_opt.orchestrator"
+ ".gpu_preflight.preflight_gpus", fake_preflight)
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ monkeypatch.setattr(wp.time, "sleep", lambda s: sleeps.append(s))
+
+ check = wp.ensure_measurement_gate(2, state_dir=tmp_path, wait_seconds=1800,
+ max_waits=48, site="benchmark")
+ assert check["usable"] is True
+ assert sleeps == [1800] # exactly one 30-minute re-check
+ assert calls[0][0] == [2] # only the assigned device
+ assert calls[0][1] == 90.0 and calls[0][2] == 0.0 # the operator's gate
+
+ rows = [json.loads(l) for l in
+ (tmp_path / "measurement_gate.jsonl").read_text().splitlines()]
+ assert [r["event"] for r in rows] == ["gate_blocked", "gate_ok"]
+ assert rows[0]["site"] == "benchmark" and rows[0]["gpu"] == 2
+ assert rows[1]["util_percent"] == 0.0
+
+
+def test_measurement_gate_gives_up_after_max_waits(monkeypatch, tmp_path):
+ from ..orchestrator import w8a8_pipeline as wp
+
+ always = {"gpus": {3: {"usable": False, "util_percent": 0.0,
+ "vram_percent": 96.0,
+ "reasons": ["VRAM 96% > 90% limit"]}}}
+ monkeypatch.setattr("metainfer.tasks.dcu_kernel_auto_opt.orchestrator"
+ ".gpu_preflight.preflight_gpus",
+ lambda ids, **kw: always)
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ sleeps = []
+ monkeypatch.setattr(wp.time, "sleep", lambda s: sleeps.append(s))
+
+ with pytest.raises(wp.MeasurementGateBlocked) as excinfo:
+ wp.ensure_measurement_gate(3, state_dir=tmp_path, wait_seconds=1800,
+ max_waits=3)
+ assert excinfo.value.gpu == 3
+ assert excinfo.value.checks == 3 # 3 re-checks, then give up
+ assert "VRAM 96%" in excinfo.value.gate_reason
+ assert sleeps == [1800, 1800, 1800]
+ rows = (tmp_path / "measurement_gate.jsonl").read_text().splitlines()
+ assert len(rows) == 4 # 3 blocked + the final blocked
+
+
+def test_measurement_gate_fails_closed_when_the_probe_breaks(monkeypatch, tmp_path):
+ from ..orchestrator import w8a8_pipeline as wp
+
+ def boom(*_a, **_k):
+ raise RuntimeError("hy-smi exploded")
+
+ monkeypatch.setattr("metainfer.tasks.dcu_kernel_auto_opt.orchestrator"
+ ".gpu_preflight.preflight_gpus", boom)
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ monkeypatch.setattr(wp.time, "sleep", lambda s: None)
+ with pytest.raises(wp.MeasurementGateBlocked):
+ wp.ensure_measurement_gate(0, state_dir=tmp_path, wait_seconds=0,
+ max_waits=1)
+
+
+def test_runner_gates_every_measurement(monkeypatch, tmp_path):
+ """benchmark() and profile_pmc() both pass the gate before launching."""
+ from ..orchestrator import w8a8_pipeline as wp
+
+ gated = []
+ monkeypatch.setattr(wp.W8A8Runner, "_ensure_gate",
+ lambda self, site: gated.append((self.gpu, site)))
+ monkeypatch.setattr(wp.time, "sleep", lambda s: None) # no real gate waits
+
+ runner = object.__new__(wp.W8A8Runner)
+ runner.worker_root = tmp_path
+ runner.source = tmp_path
+ runner.harness = tmp_path / "w8a8_bench.py"
+ runner.env = {"METAINFER_GPU_PREFLIGHT": "1"}
+ runner.gpu = 1
+ monkeypatch.setattr(wp.W8A8Runner, "_prepare_compile_cache",
+ lambda self: {"build_key": "k"})
+ monkeypatch.setattr(wp.W8A8Runner, "_ensure_reference_cache",
+ lambda self, m, n, k: None)
+ monkeypatch.setattr(wp, "_run",
+ lambda *a, **k: type("R", (), {"stdout": "{}"})())
+ monkeypatch.setattr(wp, "_last_json", lambda out: {})
+
+ runner.benchmark({"M": 16, "N": 64, "K": 32})
+ assert gated == [(1, "benchmark")]
+
+ (tmp_path / "profile_pmc.sh").write_text("#!/bin/bash\n", encoding="utf-8")
+ monkeypatch.setattr(wp, "parse_pmc_csv",
+ lambda path: {"available": False, "counters": {}})
+ runner.profile_pmc({"M": 16, "N": 64, "K": 32}, tmp_path / "out")
+ assert gated == [(1, "benchmark"), (1, "profile_pmc")]
+
+
+def test_any_measurement_gate_picks_an_idle_sibling(monkeypatch, tmp_path):
+ """A busy GPU 0 must not park the task while GPU 1 is free."""
+ from ..orchestrator import w8a8_pipeline as wp
+
+ plan = {"gpus": {
+ 0: {"usable": False, "util_percent": 98.0, "vram_percent": 97.0,
+ "reasons": ["device busy: HCU 98.0% > 0.0%"]},
+ 1: {"usable": True, "util_percent": 0.0, "vram_percent": 5.0,
+ "state_source": "sysfs"},
+ }}
+ calls = []
+ monkeypatch.setattr(
+ "metainfer.tasks.dcu_kernel_auto_opt.orchestrator.gpu_preflight"
+ ".preflight_gpus",
+ lambda ids, **kw: (calls.append(list(ids)), plan)[1])
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ sleeps = []
+ monkeypatch.setattr(wp.time, "sleep", lambda s: sleeps.append(s))
+
+ chosen = wp.ensure_any_measurement_gate([0, 1, 2, 3], state_dir=tmp_path,
+ wait_seconds=1800, site="probe")
+ assert chosen == 1 # an idle sibling, not the busy GPU 0
+ assert sleeps == [] # and no waiting at all
+ assert calls == [[0, 1, 2, 3]] # the whole candidate set is checked
+ rows = [json.loads(l) for l in
+ (tmp_path / "measurement_gate.jsonl").read_text().splitlines()]
+ assert rows[0]["event"] == "gate_ok" and rows[0]["passing"] == [1]
+
+
+def test_any_measurement_gate_waits_when_no_candidate_passes(monkeypatch, tmp_path):
+ from ..orchestrator import w8a8_pipeline as wp
+
+ busy = {"gpus": {d: {"usable": False, "util_percent": 99.0,
+ "vram_percent": 96.0,
+ "reasons": [f"gpu{d} VRAM 96% > 90% limit"]}
+ for d in (0, 1)}}
+ monkeypatch.setattr(
+ "metainfer.tasks.dcu_kernel_auto_opt.orchestrator.gpu_preflight"
+ ".preflight_gpus", lambda ids, **kw: busy)
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ sleeps = []
+ monkeypatch.setattr(wp.time, "sleep", lambda s: sleeps.append(s))
+
+ with pytest.raises(wp.MeasurementGateBlocked) as excinfo:
+ wp.ensure_any_measurement_gate([0, 1], state_dir=tmp_path,
+ wait_seconds=1800, max_waits=2)
+ assert sleeps == [1800, 1800] # 30-minute re-checks, then give up
+ assert "gpu0" in excinfo.value.gate_reason
+ assert "gpu1" in excinfo.value.gate_reason
+
+
+def test_any_measurement_gate_prefers_the_tasks_own_device(monkeypatch, tmp_path):
+ from ..orchestrator import w8a8_pipeline as wp
+
+ free = {"gpus": {d: {"usable": True, "util_percent": 0.0,
+ "vram_percent": 1.0} for d in (0, 3)}}
+ monkeypatch.setattr(
+ "metainfer.tasks.dcu_kernel_auto_opt.orchestrator.gpu_preflight"
+ ".preflight_gpus", lambda ids, **kw: free)
+ monkeypatch.setattr(wp, "_env_flag_off", lambda env: False)
+ # 3 is the task's own card and is free: the probe stays on it
+ assert wp.ensure_any_measurement_gate([3, 0], env={"METAINFER_GPU_PREFLIGHT": "1"}) == 3
+
+
+def _fake_scaffold(root):
+ """A scaffold that satisfies the trusted Generate preflight file checks."""
+ from ..orchestrator import gen_and_opt_pipeline as gp
+
+ root.mkdir(parents=True, exist_ok=True)
+ for name in (gp._HARNESS_FILENAME, gp.W8A8_API_FILENAME,
+ gp._GIT_ATTRIBUTES_FILE, "scaffold_manifest.json",
+ *gp._SCAFFOLD_FILES):
+ path = root / name
+ path.parent.mkdir(parents=True, exist_ok=True)
+ path.write_text("", encoding="utf-8")
+ (root / "profile_pmc.sh").write_text(
+ "#!/bin/bash\n/opt/dtk/bin/hipprof --pmc --pmc-type 3 "
+ '"$source_dir" "$output_dir"\n', encoding="utf-8")
+ return root
+
+
+@requires_dtk_hipprof
+def test_generate_probe_runs_on_a_device_that_passed_the_gate(monkeypatch, tmp_path):
+ """The probe is gated, but not pinned to GPU 0 (a busy card must not park it)."""
+ from ..orchestrator import gen_and_opt_pipeline as gp
+
+ source = _fake_scaffold(tmp_path / "scaffold")
+ contract = gp.file_digest(source / gp.W8A8_API_FILENAME)
+ seen = {}
+
+ def fake_run(cmd, **kwargs):
+ if cmd[-1] == "--self-test":
+ return type("R", (), {"stdout": '{"passed": true}'})()
+ if cmd[-1] == "--probe":
+ env = kwargs.get("env") or {}
+ seen["probe_gpu"] = env.get("HIP_VISIBLE_DEVICES")
+ # DTK filters twice when both variables name the same non-zero
+ # index, which hides the device ("No HIP GPUs are available" on
+ # GPU 1 while GPU 0 survived): binding must set HIP only.
+ seen["probe_env"] = dict(env)
+ return type("R", (), {"stdout": json.dumps({
+ "visible_devices": 1, "cudagraph_available": True})})()
+ return type("R", (), {"stdout": ""})()
+
+ monkeypatch.setattr(gp, "_run", fake_run)
+ monkeypatch.setattr(gp, "_gate_preflight_enabled", lambda: True)
+
+ calls = []
+
+ def fake_gate(devices, **kwargs):
+ calls.append((list(devices), kwargs))
+ return 1
+
+ monkeypatch.setattr(gp, "ensure_any_measurement_gate", fake_gate)
+ out = gp._validate_generate_scaffold(
+ source, contract_sha256=contract, shapes={"s1": {"M": 16}},
+ probe_devices=[0, 1, 2, 3], state_dir=tmp_path / "state")
+
+ assert calls[0][0] == [0, 1, 2, 3] # every candidate may answer
+ # the wait is auditable in the task's own state dir, not silently invisible
+ assert calls[0][1]["state_dir"] == tmp_path / "state"
+ assert calls[0][1]["site"] == "generate_probe"
+ assert seen["probe_gpu"] == "1" # and the probe honours the choice
+ assert "ROCR_VISIBLE_DEVICES" not in seen["probe_env"]
+ assert out["gpu_probe"]["visible_devices"] == 1
+
+
+@requires_gpu_broker
+def test_probe_candidates_put_the_tasks_own_devices_first():
+ from ..orchestrator import gen_and_opt_pipeline as gp
+
+ class _Cfg:
+ assignments = [type("A", (), {"gpu": 2})(), type("A", (), {"gpu": 0})()]
+
+ devices = gp._probe_gate_devices(_Cfg())
+ assert devices[:2] == [2, 0] # the task's own cards, in order
+ assert 1 in devices and 3 in devices # then the idle siblings
+ assert len(devices) == len(set(devices))
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_harness_io.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_harness_io.py
new file mode 100644
index 00000000..013375af
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_harness_io.py
@@ -0,0 +1,105 @@
+"""Tests for orchestrator.harness_io (M1 slice 1: harness seed + reader)."""
+
+from __future__ import annotations
+
+import os
+from pathlib import Path
+
+import pytest
+import yaml
+
+from ..orchestrator import w8a8_pipeline
+from ..orchestrator.config import ROUND_ACCEPTANCE_IMPROVEMENT_PERCENT
+from ..orchestrator.harness_io import (
+ default_harness_dir,
+ harness_root,
+ load_gates,
+ load_manifest,
+ plugin_dir,
+ seed_workspace,
+)
+
+
+def test_plugin_dir_and_default_seed_exist():
+ root = plugin_dir()
+ assert root.name == "dcu_kernel_auto_opt"
+ seed = default_harness_dir()
+ assert seed.is_dir()
+ assert (seed / "manifest.yaml").is_file()
+ assert (seed / "gates.yaml").is_file()
+
+
+def test_harness_root_defaults_to_seed(monkeypatch):
+ monkeypatch.delenv("METAINFER_HARNESS_ROOT", raising=False)
+ assert harness_root() == default_harness_dir().resolve()
+
+
+def test_harness_root_env_override(monkeypatch, tmp_path):
+ marker = tmp_path / "marker.txt"
+ marker.write_text("x", encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(tmp_path))
+ assert harness_root() == tmp_path.resolve()
+
+
+def test_manifest_lists_gates_component():
+ manifest = load_manifest()
+ assert manifest.get("schema_version") == 1
+ components = manifest.get("components", {})
+ assert "gates" in components
+ # gates is read at runtime now (orchestrator/gate_policy.py), so it is wired
+ assert components["gates"]["wired"] is True
+ assert "gate_policy" in (components["gates"].get("role") or "") or (
+ "runtime" in (components["gates"].get("role") or ""))
+
+
+def test_gates_seed_matches_builtin_defaults():
+ """Drift guard: the seed gates.yaml must equal the built-in defaults.
+
+ The runtime now reads gates.yaml through orchestrator/gate_policy.py (see
+ tests/test_gate_policy.py); this test keeps the *seed file* honest so a
+ fresh harness starts exactly where the hard-coded behaviour was.
+ """
+ from ..orchestrator import gate_policy as gp
+
+ gates = load_gates().get("gates", {})
+ defaults = gp.builtin_gates()
+ assert gates["round_acceptance_improvement_percent"] == (
+ defaults["round_acceptance_improvement_percent"])
+ assert gates["shadow"]["min_improvement_percent"] == (
+ defaults["shadow"]["min_improvement_percent"])
+ assert gates["shadow"]["max_exclusive_percent"] == (
+ defaults["shadow"]["max_exclusive_percent"])
+ assert gates["plateau"]["max_regression_percent"] == (
+ defaults["plateau"]["max_regression_percent"])
+ assert gates["plateau"]["recent_valid_rounds"] == (
+ defaults["plateau"]["recent_valid_rounds"])
+ assert gates["isa_gate"]["required_valid_isa_guided_rounds"] == (
+ defaults["isa_gate"]["required_valid_isa_guided_rounds"])
+ assert gates["task_budget"]["default_max_iterations"] == (
+ defaults["task_budget"]["default_max_iterations"])
+
+
+def test_seed_manifest_marks_data_components_wired():
+ manifest = load_manifest()
+ components = manifest.get("components") or {}
+ assert components["gates"]["wired"] is True
+ assert components["planner_policy"]["wired"] is True
+ assert components["planner_catalog"]["wired"] is True
+ assert components["manifest"]["wired"] is True
+
+
+def test_seed_workspace_copies_tree(tmp_path):
+ dst = seed_workspace(tmp_path / "ws")
+ assert dst.is_dir()
+ for name in ("manifest.yaml", "gates.yaml", "README.md"):
+ assert (dst / name).is_file()
+ # content equality with the seed
+ with (dst / "gates.yaml").open(encoding="utf-8") as fh:
+ copied = yaml.safe_load(fh)
+ original = load_gates()
+ assert copied == original
+
+
+def test_seed_workspace_missing_root_raises(tmp_path):
+ with pytest.raises(FileNotFoundError):
+ seed_workspace(tmp_path / "out", root=tmp_path / "does-not-exist")
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plan_render.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plan_render.py
new file mode 100644
index 00000000..20475fb3
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plan_render.py
@@ -0,0 +1,102 @@
+"""Tests for planner rendering + controlled-state parity with the legacy menu.
+
+Controlled-state parity (not corpus heuristic): on states where the planner's
+legacy-aligned branches deliberately mirror the current round menus, assert the
+mandate texts agree on key direction tokens.
+"""
+
+from __future__ import annotations
+
+from ..orchestrator.planner import catalog, choose_plan_from_history, render_plan
+from ..orchestrator.prompts import w8a8_round_strategy
+from ..orchestrator.w8a8_pipeline import isa_round_policy
+
+
+def _valid(improvement: float = 0.5) -> dict:
+ return {
+ "build_success": True,
+ "correctness_passed": True,
+ "speedup": 1.0,
+ "metrics": {"graph_capture_passed": True},
+ "acceptance": {"improvement_percent": improvement},
+ }
+
+
+def test_render_all_catalog_plans_nonempty_and_deterministic():
+ cat = catalog()
+ ctx = {"iteration": 3, "max_iterations": 10, "rounds_left": 8}
+ seen = {}
+ for pid in cat:
+ text = render_plan(pid, ctx=ctx, cat=cat)
+ assert pid in text and len(text) > 10
+ seen[pid] = text
+ # deterministic
+ for pid, text in seen.items():
+ assert render_plan(pid, ctx=ctx, cat=cat) == text
+
+
+def test_fresh_lane_m16_matches_menu_on_dumma():
+ # Fresh m16 lane: planner (establish_arch) and legacy menu (round 1 DUMMA
+ # bootstrap instruction) both direct at DUMMA.
+ plan = choose_plan_from_history([], iteration=1, shape={"M": 16})
+ menu = w8a8_round_strategy(
+ {"M": 16, "N": 1536, "K": 4096},
+ 1,
+ [],
+ {},
+ max_iterations=10,
+ isa_policy=isa_round_policy(
+ iteration=1, max_iterations=10, history=[]
+ ),
+ )
+ rendered = render_plan(plan, cat=catalog())
+ assert "DUMMA" in rendered and "DUMMA" in menu
+
+
+def test_plateau_open_matches_menu_on_isa():
+ history = [_valid(0.5) for _ in range(5)] + [
+ _valid(v) for v in (0.5, 0.9, 1.2)
+ ]
+ shape = {"M": 16, "N": 1536, "K": 4096}
+ iteration, max_iterations = 9, 10
+ policy = isa_round_policy(
+ iteration=iteration, max_iterations=max_iterations, history=history
+ )
+ assert policy.get("phase") == "isa_guided_hip"
+
+ plan = choose_plan_from_history(
+ history, iteration=iteration, max_iterations=max_iterations,
+ shape=shape,
+ )
+ menu = w8a8_round_strategy(
+ shape, iteration, history, {},
+ max_iterations=max_iterations, isa_policy=policy,
+ )
+ rendered = render_plan(plan, cat=catalog())
+ assert plan == "isa_guided_hip"
+ assert "ISA" in rendered and "ISA" in menu
+
+
+def test_faster_wrong_matches_menu_on_repair():
+ history = [_valid(0.5), {
+ "build_success": True,
+ "correctness_passed": False,
+ "speedup": 1.9,
+ "metrics": {"graph_capture_passed": True},
+ }]
+ shape = {"M": 16, "N": 1536, "K": 4096}
+ iteration, max_iterations = 4, 10
+ plan = choose_plan_from_history(
+ history, iteration=iteration, max_iterations=max_iterations,
+ shape=shape,
+ )
+ policy = isa_round_policy(
+ iteration=iteration, max_iterations=max_iterations, history=history
+ )
+ menu = w8a8_round_strategy(
+ shape, iteration, history, {},
+ max_iterations=max_iterations, isa_policy=policy,
+ )
+ rendered = render_plan(plan, cat=catalog())
+ assert plan == "repair_faster_wrong"
+ assert "repair" in rendered.lower() and "repair" in menu.lower()
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner.py
new file mode 100644
index 00000000..ffe79f98
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner.py
@@ -0,0 +1,204 @@
+"""Tests for orchestrator.planner (M1 slice 2: state-conditioned selector v0)."""
+
+from __future__ import annotations
+
+from ..orchestrator import planner
+
+
+def _valid_record(improvement: float = 0.5, *, speedup: float = 1.0) -> dict:
+ return {
+ "build_success": True,
+ "correctness_passed": True,
+ "speedup": speedup,
+ "metrics": {"graph_capture_passed": True},
+ "acceptance": {"improvement_percent": improvement},
+ }
+
+
+def _wrong_fast_record() -> dict:
+ return {
+ "build_success": True,
+ "correctness_passed": False,
+ "speedup": 1.8,
+ "metrics": {"graph_capture_passed": True},
+ }
+
+
+def _build_fail_record() -> dict:
+ return {"build_success": False, "failure_reason": "compile error line 42"}
+
+
+def _infra_record() -> dict:
+ return {
+ "build_success": False,
+ "failure_reason": "timed out after 900s",
+ }
+
+
+def test_catalog_loads_from_yaml_seed():
+ cat = planner.catalog()
+ for pid in ("repair_faster_wrong", "isa_guided_hip", "consolidate"):
+ assert pid in cat
+
+
+def test_p0_faster_wrong_beats_everything():
+ history = [_valid_record(), _wrong_fast_record()]
+ assert planner.choose_plan_from_history(history, iteration=3) == (
+ "repair_faster_wrong"
+ )
+
+
+def test_empty_history_is_not_build_failure():
+ # A fresh lane (no previous round) must not be routed to fix_build.
+ assert planner.choose_plan_from_history([]) == "establish_arch"
+
+
+def test_p0_infra_retry_and_build_fix():
+ assert planner.choose_plan_from_history([_valid_record(), _infra_record()]) == (
+ "retry_same"
+ )
+ assert planner.choose_plan_from_history(
+ [_valid_record(), _build_fail_record()]
+ ) == "fix_build"
+
+
+def test_p1_last_round_consolidates():
+ history = [_valid_record() for _ in range(3)]
+ assert planner.choose_plan_from_history(
+ history, iteration=10, max_iterations=10
+ ) == "consolidate"
+
+
+def test_p1_plateau_opens_isa_then_asm():
+ # 8 valid HIP rounds + last three improvements inside [-2, 2) prove
+ # plateau with the ISA gate open (required_hip_rounds = max_iterations - 2).
+ history = [_valid_record(0.5) for _ in range(5)] + [
+ _valid_record(v) for v in (0.5, 0.9, 1.2)
+ ]
+ plan = planner.choose_plan_from_history(
+ history, iteration=9, max_iterations=10
+ )
+ assert plan == "isa_guided_hip"
+ plan2 = planner.choose_plan_from_history(
+ history,
+ iteration=9,
+ max_iterations=10,
+ compiler_limitation_confirmed=True,
+ )
+ assert plan2 == "conditional_inline_asm"
+
+
+def test_p2_occupancy_signature_selects_resource_round():
+ history = [_valid_record() for _ in range(2)]
+ pmc = {"waves_per_cu": 12, "target_waves_per_cu": 16}
+ assert planner.choose_plan_from_history(
+ history, pmc=pmc, iteration=5
+ ) == "occupancy_resource"
+
+
+def test_p2_bank_conflict_signature():
+ history = [_valid_record() for _ in range(2)]
+ pmc = {"lds_instructions": 1000, "lds_bank_conflicts": 2000}
+ assert planner.choose_plan_from_history(
+ history, pmc=pmc, iteration=5
+ ) == "memory_layout"
+
+
+def test_p4_fallback_matches_legacy_menu_intent(monkeypatch):
+ # The no-evidence fallback (uncertainty -> legacy_menu) is on by default;
+ # disable it here so the P4 portfolio table itself is exercised.
+ pol = dict(planner._builtin_policy())
+ pol["uncertainty"] = {"enabled": False}
+ monkeypatch.setattr(planner, "policy", lambda root=None: pol)
+ history = [_valid_record() for _ in range(2)]
+ plan = planner.choose_plan_from_history(
+ history, iteration=2, max_iterations=10, shape={"M": 4096}
+ )
+ assert plan == "memory_layout" # large_m legacy menu, step after bootstrap
+
+
+def test_coverage_guard_breaks_same_plan_repetition():
+ history = [_valid_record() for _ in range(2)]
+ pmc = {
+ "waves_per_cu": 12,
+ "target_waves_per_cu": 16,
+ "l2_hit_rate": 40.0,
+ }
+ # two prior identical occupancy rounds -> P3 should route to the next
+ # candidate (memory_layout from l2_low) instead of repeating occupancy.
+ plan = planner.choose_plan_from_history(
+ history,
+ pmc=pmc,
+ iteration=6,
+ plan_tags=["occupancy_resource", "occupancy_resource"],
+ )
+ assert plan == "memory_layout"
+
+
+def test_fallback_portfolio_cycles_instead_of_pinning_the_tail():
+ """Long runs must not be pinned to the table's last row (consolidate).
+
+ 9-8-8 iteration 3 showed 9 of 11 planner picks were ``consolidate``
+ because the fallback indexed past the table end and clamped. The
+ portfolio now cycles, so a long run keeps revisiting exploration plans.
+ """
+ from ..orchestrator.planner import _legacy_fallback
+
+ picks = [
+ _legacy_fallback({"shape": {"M": 16}, "iteration": it})
+ for it in range(2, 14)
+ ]
+ assert len(set(picks)) > 2, picks
+ tail = picks[-4:]
+ assert len(set(tail)) > 1, tail
+ assert tail.count("consolidate") < len(tail), tail
+
+
+def test_uncertainty_falls_back_to_legacy_menu(monkeypatch):
+ """No bottleneck evidence -> defer to the hand-tuned menu."""
+ from ..orchestrator import planner as P
+
+ pol = dict(P._builtin_policy())
+ pol["uncertainty"] = {"enabled": True, "plan": "legacy_menu",
+ "min_valid_rounds": 1}
+ monkeypatch.setattr(P, "policy", lambda root=None: pol)
+ monkeypatch.setattr(P, "bottleneck_tags", lambda pmc: [])
+ ctx = {"valid_hip_rounds": 2, "pmc": {}, "shape": {"M": 16},
+ "iteration": 3, "max_iterations": 11, "rounds_left": 9,
+ "tried_counts": {}, "consecutive_same": 0,
+ "last_present": False, "faster_wrong": False}
+ assert P.choose_plan(ctx) == "legacy_menu"
+
+
+def test_uncertainty_is_skipped_when_a_bottleneck_is_known(monkeypatch):
+ from ..orchestrator import planner as P
+
+ pol = dict(P._builtin_policy())
+ pol["uncertainty"] = {"enabled": True, "plan": "legacy_menu"}
+ pol["bottleneck_to_plan"] = {"lds_bank_conflicts": "memory_layout"}
+ monkeypatch.setattr(P, "policy", lambda root=None: pol)
+ monkeypatch.setattr(P, "bottleneck_tags",
+ lambda pmc: ["lds_bank_conflicts"])
+ ctx = {"valid_hip_rounds": 2, "pmc": {"available": True},
+ "shape": {"M": 16}, "iteration": 3, "max_iterations": 11,
+ "rounds_left": 9, "tried_counts": {}, "consecutive_same": 0,
+ "last_present": False, "faster_wrong": False}
+ assert P.choose_plan(ctx) == "memory_layout"
+
+
+def test_legacy_menu_plan_renders_the_legacy_menu(tmp_path, monkeypatch):
+ import json as _json
+ from ..orchestrator import w8a8_pipeline as W
+
+ monkeypatch.setenv("METAINFER_PLANNER", "1")
+ monkeypatch.setattr(W, "choose_plan_from_history",
+ lambda *a, **k: "legacy_menu")
+ sink = tmp_path / "planner_plans.jsonl"
+ text = W._round_strategy_text(
+ {"M": 16, "N": 4096, "K": 2048}, 2, [], {"available": False},
+ {"max_iterations": 5, "skill_allowed": True},
+ plan_sink=sink, shape_id="hy3_tp4_o_proj_m16",
+ )
+ assert text and len(text) > 50
+ row = _json.loads(sink.read_text(encoding="utf-8").splitlines()[-1])
+ assert row["plan_id"] == "legacy_menu" and row["source"] == "planner"
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_policy.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_policy.py
new file mode 100644
index 00000000..56e4a4cd
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_policy.py
@@ -0,0 +1,62 @@
+"""Prove planner_policy.yaml is a live, evolvable state->plan component."""
+
+from __future__ import annotations
+
+import shutil
+from pathlib import Path
+
+import yaml
+
+from ..orchestrator import planner
+from ..orchestrator.harness_io import default_harness_dir
+
+
+def _valid():
+ return {"build_success": True, "correctness_passed": True,
+ "metrics": {"graph_capture_passed": True},
+ "acceptance": {"improvement_percent": 5.0}}
+
+
+def test_policy_seed_loads():
+ pol = planner.policy()
+ assert pol["repair_priority"]["faster_wrong"] == "repair_faster_wrong"
+ assert pol["bottleneck_to_plan"]["occupancy_limited"] == "occupancy_resource"
+
+
+def test_workspace_policy_change_changes_selected_plan(monkeypatch, tmp_path):
+ root = tmp_path / "harness"
+ shutil.copytree(default_harness_dir(), root)
+ policy_path = root / "planner_policy.yaml"
+ data = yaml.safe_load(policy_path.read_text(encoding="utf-8"))
+ # AHE-like change: route occupancy pressure to pipeline_tune instead of
+ # occupancy_resource. The selector must obey the evolved workspace file.
+ data["bottleneck_to_plan"]["occupancy_limited"] = "pipeline_tune"
+ policy_path.write_text(yaml.safe_dump(data, sort_keys=False), encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+
+ plan = planner.choose_plan_from_history(
+ [_valid(), _valid()],
+ pmc={"waves_per_cu": 8, "target_waves_per_cu": 16},
+ iteration=4,
+ max_iterations=10,
+ shape={"M": 4096},
+ )
+ assert plan == "pipeline_tune"
+
+
+def test_policy_fallback_order_is_evolvable(monkeypatch, tmp_path):
+ root = tmp_path / "harness"
+ shutil.copytree(default_harness_dir(), root)
+ path = root / "planner_policy.yaml"
+ data = yaml.safe_load(path.read_text(encoding="utf-8"))
+ data["fallback"]["prefill"][0] = "occupancy_resource"
+ # exercise the evolvable portfolio itself, not the no-evidence fallback
+ data.setdefault("uncertainty", {})["enabled"] = False
+ path.write_text(yaml.safe_dump(data, sort_keys=False), encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+
+ plan = planner.choose_plan_from_history(
+ [_valid(), _valid()], pmc={}, iteration=2, max_iterations=10,
+ shape={"M": 4096},
+ )
+ assert plan == "occupancy_resource"
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_wiring.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_wiring.py
new file mode 100644
index 00000000..9261414f
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_planner_wiring.py
@@ -0,0 +1,84 @@
+"""Tests for planner wiring env toggle (METAINFER_PLANNER)."""
+
+from __future__ import annotations
+
+from ..orchestrator.w8a8_pipeline import (
+ _planner_enabled,
+ _round_strategy_text,
+ isa_round_policy,
+)
+from ..orchestrator.prompts import w8a8_round_strategy
+
+
+def _isa_policy(iteration: int, max_iterations: int, history) -> dict:
+ return isa_round_policy(
+ iteration=iteration, max_iterations=max_iterations, history=history
+ )
+
+
+def test_planner_disabled_by_default(monkeypatch):
+ monkeypatch.delenv("METAINFER_PLANNER", raising=False)
+ assert _planner_enabled() is False
+
+
+def test_env_toggle_enables(monkeypatch):
+ monkeypatch.setenv("METAINFER_PLANNER", "1")
+ assert _planner_enabled() is True
+
+
+def test_disabled_returns_exact_legacy_text(monkeypatch):
+ monkeypatch.delenv("METAINFER_PLANNER", raising=False)
+ shape = {"M": 16, "N": 1536, "K": 4096}
+ history = []
+ policy = _isa_policy(1, 10, history)
+ got = _round_strategy_text(shape, 1, history, {}, policy)
+ expected = w8a8_round_strategy(
+ shape, 1, history, {},
+ max_iterations=10, isa_policy=policy,
+ )
+ assert got == expected
+
+
+def test_enabled_renders_planner_mandate(monkeypatch):
+ monkeypatch.setenv("METAINFER_PLANNER", "1")
+ shape = {"M": 16, "N": 1536, "K": 4096}
+ history = []
+ policy = _isa_policy(1, 10, history)
+ got = _round_strategy_text(shape, 1, history, {}, policy)
+ assert got.startswith("Mandatory decision for this round:")
+ assert "establish_arch" in got
+
+
+def test_round_strategy_text_records_planner_decision(tmp_path, monkeypatch):
+ """The planner's own choice is written to planner_plans.jsonl (hard evidence)."""
+ import json as _json
+ from ..orchestrator.w8a8_pipeline import _round_strategy_text
+
+ monkeypatch.setenv("METAINFER_PLANNER", "1")
+ sink = tmp_path / "planner_plans.jsonl"
+ text = _round_strategy_text(
+ {"M": 16, "N": 4096, "K": 2048}, 1, [], {"available": False},
+ {"max_iterations": 3, "skill_allowed": True},
+ plan_sink=sink, shape_id="hy3_tp4_o_proj_m16",
+ )
+ assert text
+ assert sink.is_file()
+ row = _json.loads(sink.read_text(encoding="utf-8").splitlines()[0])
+ assert row["source"] == "planner"
+ assert row["plan_id"]
+ assert row["iteration"] == 1
+ assert row["shape_id"] == "hy3_tp4_o_proj_m16"
+
+
+def test_round_strategy_text_writes_nothing_without_planner(tmp_path, monkeypatch):
+ from ..orchestrator.w8a8_pipeline import _round_strategy_text
+
+ monkeypatch.delenv("METAINFER_PLANNER", raising=False)
+ sink = tmp_path / "planner_plans.jsonl"
+ text = _round_strategy_text(
+ {"M": 16, "N": 4096, "K": 2048}, 1, [], {"available": False},
+ {"max_iterations": 3, "skill_allowed": True},
+ plan_sink=sink, shape_id="hy3_tp4_o_proj_m16",
+ )
+ assert text
+ assert not sink.exists()
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plugin.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plugin.py
index 176eb25d..db897b5b 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plugin.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_plugin.py
@@ -49,7 +49,7 @@ def test_form_schema_is_available():
assert model["default"] == "Opus"
assert model["override_component"] == "agent-model"
assert [option["label"] for option in model["options"]] == [
- "Opus", "Sonnet", "deepseek-v4-flash"
+ "Opus", "Sonnet", "deepseek-flash-4.1", "deepseek-v4-flash"
]
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_predictions.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_predictions.py
new file mode 100644
index 00000000..6b047a5b
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_predictions.py
@@ -0,0 +1,94 @@
+"""Tests for orchestrator.predictions (M1 slice 3: inner decision hook)."""
+
+from __future__ import annotations
+
+from ..orchestrator.predictions import check_prediction, parse_prediction, plan_tag
+
+
+def test_no_structured_prediction_returns_none():
+ assert check_prediction(
+ {"hypothesis": "plain prose"},
+ candidate_us=100.0,
+ best_us=110.0,
+ passed=True,
+ p90_guard_passed=True,
+ ) is None
+ assert parse_prediction({"hypothesis": "x"}) is None
+
+
+def test_infra_failure_is_na():
+ out = check_prediction(
+ {"prediction": {"direction": "improve"}},
+ candidate_us=float("inf"),
+ best_us=100.0,
+ passed=False,
+ p90_guard_passed=False,
+ failure_reason="timed out after 900s",
+ )
+ assert out is not None and out["checked"] == "na"
+
+
+def test_range_hit_and_miss():
+ proposal = {"prediction": {"expected_us_range": [90.0, 105.0]}}
+ hit = check_prediction(
+ proposal,
+ candidate_us=98.0,
+ best_us=110.0,
+ passed=True,
+ p90_guard_passed=True,
+ )
+ assert hit["checked"] == "hit"
+ miss = check_prediction(
+ proposal,
+ candidate_us=120.0,
+ best_us=110.0,
+ passed=True,
+ p90_guard_passed=True,
+ )
+ assert miss["checked"] == "miss"
+ assert "outside declared range" in miss["reason"]
+
+
+def test_direction_improve_requires_band():
+ ok = check_prediction(
+ {"prediction": {"direction": "improve"}},
+ candidate_us=100.0,
+ best_us=110.0, # ~9% better
+ passed=True,
+ p90_guard_passed=True,
+ )
+ assert ok["checked"] == "hit"
+ flat = check_prediction(
+ {"prediction": {"direction": "flat"}},
+ candidate_us=109.0, # ~0.9% delta -> inside flat band
+ best_us=110.0,
+ passed=True,
+ p90_guard_passed=True,
+ )
+ assert flat["checked"] == "hit"
+ wrong = check_prediction(
+ {"prediction": {"direction": "improve"}},
+ candidate_us=109.0, # only ~0.9% better -> below the 2% band
+ best_us=110.0,
+ passed=True,
+ p90_guard_passed=True,
+ )
+ assert wrong["checked"] == "miss"
+
+
+def test_correctness_failure_means_miss():
+ out = check_prediction(
+ {"prediction": {"expected_us_range": [90.0, 105.0]}},
+ candidate_us=98.0,
+ best_us=110.0,
+ passed=False,
+ p90_guard_passed=False,
+ )
+ assert out["checked"] == "miss"
+ assert "correctness failed" in out["reason"]
+
+
+def test_plan_tag_helpers():
+ assert plan_tag({"hypothesis": "x"}) is None
+ assert plan_tag({"plan_id": "occupancy_resource"}) == "occupancy_resource"
+ assert plan_tag({"plan_id": " "}) is None
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_prompt_schema.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_prompt_schema.py
new file mode 100644
index 00000000..a3dd5acd
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_prompt_schema.py
@@ -0,0 +1,47 @@
+"""Prompt-schema tests: worker prompt now carries optional decision fields."""
+
+from __future__ import annotations
+
+from pathlib import Path
+
+from ..orchestrator.config import WorkerAssignment
+from ..orchestrator.w8a8_pipeline import RealW8A8OptimizationPipeline
+
+
+def _round_prompt(**overrides):
+ args = dict(
+ assignment=WorkerAssignment("worker_0", 0, ["m2"]),
+ shape_id="m2",
+ shape={"M": 2, "N": 16, "K": 32},
+ best={"median_us": 10.0},
+ root=Path("/tmp/worker"),
+ iteration=1,
+ guidance=None,
+ history=[],
+ pmc_evidence={},
+ )
+ args.update(overrides)
+ return RealW8A8OptimizationPipeline._worker_prompt(**args)
+
+
+def test_first_turn_prompt_mentions_decision_fields():
+ prompt = _round_prompt()
+ assert '"plan_id"' in prompt
+ assert '"prediction"' in prompt
+ assert '"expected_us_range"' in prompt
+ assert '"direction": "improve"' in prompt
+ assert "checks `prediction` against the measured round" in prompt
+
+
+def test_continuation_prompt_mentions_optional_fields():
+ prompt = _round_prompt(
+ continuation=True,
+ history=[{
+ "iteration": 1,
+ "accepted": True,
+ "build_success": True,
+ "correctness_passed": True,
+ "metrics": {"median_us": 9.0},
+ }],
+ )
+ assert "optional plan_id and prediction" in prompt
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_retest_tool.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_retest_tool.py
new file mode 100644
index 00000000..18ed4186
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_retest_tool.py
@@ -0,0 +1,52 @@
+"""A/B retest tool: variant lookup, header parsing, verdict and environment."""
+
+from __future__ import annotations
+
+from pathlib import Path
+
+from ..tools.retest_variant import (
+ BENCH_ENV, bench_env, decide, find_variant, parse_variant_header_median,
+)
+
+
+def test_find_variant_tolerates_model_label_drift(tmp_path):
+ """glm should find the shipped glm52 directory."""
+ root = tmp_path / "variant" / "int8w8a8-gemm"
+ glm = root / "glm52" / "TP8" / "M4096"
+ glm.mkdir(parents=True)
+ (glm / "shared_down_proj.hip").write_text("// kernel\n", encoding="utf-8")
+ assert find_variant("glm", 8, 4096, "shared_down_proj", root) == (
+ glm / "shared_down_proj.hip")
+ assert find_variant("glm", 8, 16, "shared_down_proj", root) is None
+ assert find_variant("hy3", 4, 4096, "o_proj", root) is None
+
+
+def test_parse_variant_header_median_reads_the_record():
+ text = ("// @@variant shape=hy3_tp4_o_proj_m4096 commit=abc added=2026-09-14\n"
+ "// median_us=453.2 p90_us=454.2 speedup=43.87 baseline_us=1.988e+04\n"
+ "// code follows\n")
+ path = Path("/tmp/_variant_header_test.hip")
+ path.write_text(text, encoding="utf-8")
+ try:
+ assert parse_variant_header_median(path) == 453.2
+ finally:
+ path.unlink(missing_ok=True)
+ assert parse_variant_header_median(Path("/tmp/does-not-exist.hip")) is None
+
+
+def test_bench_env_pins_gfx928_and_the_device():
+ env = bench_env(2, base={})
+ assert env["PYTORCH_ROCM_ARCH"] == BENCH_ENV["PYTORCH_ROCM_ARCH"] == "gfx928"
+ assert env["HIP_VISIBLE_DEVICES"] == "2"
+ assert env["CUDA_VISIBLE_DEVICES"] == "2"
+
+
+def test_decide_requires_margin_and_reproducibility():
+ # a clear, reproducible win
+ assert decide(140.0, 280.0, recorded_us=140.0)["verdict"] == "promote"
+ # too small a margin
+ assert decide(275.0, 280.0, recorded_us=275.0)["verdict"] == "keep"
+ # fast but the re-measurement disagrees with the record -> review
+ assert decide(140.0, 280.0, recorded_us=100.0)["verdict"] == "needs_review"
+ # missing measurement
+ assert decide(None, 280.0)["verdict"] == "measurement_failed"
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_round_strategy.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_round_strategy.py
new file mode 100644
index 00000000..0c3cd43d
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_round_strategy.py
@@ -0,0 +1,490 @@
+"""systemprompt/round_strategy.yaml is wired: the worker-facing mandate is data.
+
+What these tests pin down:
+
+* the built-in texts are byte-identical to the historical Python literals, so an
+ un-evolved harness tells the worker exactly what it always did;
+* the seed file reproduces those defaults (drift guard, like gates.yaml);
+* editing the seed changes the runtime mandate, including inside the pipeline's
+ own call path;
+* ``wired: false`` in the manifest makes the component inert;
+* a malformed / partial harness file degrades to built-in text instead of
+ raising (a broken harness is a performance problem, not a crash).
+"""
+
+from __future__ import annotations
+
+import shutil
+
+import pytest
+import yaml
+
+from ..orchestrator import round_strategy as rs
+from ..orchestrator.harness_io import default_harness_dir
+from ..orchestrator.prompts import w8a8_round_strategy
+
+import importlib.util
+
+import pytest
+_HARNESS_EVOLVE_AVAILABLE = (
+ importlib.util.find_spec(
+ "metainfer.tasks.harness_evolve.orchestrator.pipeline"
+ )
+ is not None
+)
+requires_harness_evolve = pytest.mark.skipif(
+ not _HARNESS_EVOLVE_AVAILABLE,
+ reason="metainfer.tasks.harness_evolve is not available in this tree",
+)
+
+M16 = {"model": "hy3", "M": 16, "N": 2048, "K": 256}
+SMALL = {"model": "hy3", "M": 2, "N": 2048, "K": 256}
+PREFILL = {"model": "hy3", "M": 3072, "N": 2048, "K": 256}
+
+
+@pytest.fixture(autouse=True)
+def _clear_cache():
+ rs.reset_cache()
+ yield
+ rs.reset_cache()
+
+
+def _workspace(tmp_path, *, wired=True):
+ root = tmp_path / "harness"
+ shutil.copytree(default_harness_dir(), root)
+ manifest = yaml.safe_load((root / "manifest.yaml").read_text(encoding="utf-8"))
+ manifest["components"]["systemprompt"]["wired"] = wired
+ (root / "manifest.yaml").write_text(
+ yaml.safe_dump(manifest, sort_keys=False), encoding="utf-8")
+ return root
+
+
+def _strategy_file(root):
+ return root / rs.RELATIVE_PATH
+
+
+def _edit(root, mutate):
+ data = yaml.safe_load(_strategy_file(root).read_text(encoding="utf-8"))
+ mutate(data)
+ _strategy_file(root).write_text(
+ yaml.safe_dump(data, sort_keys=False, allow_unicode=True),
+ encoding="utf-8")
+
+
+# --------------------------------------------------------------------------- #
+# Built-in defaults and the seed file agree
+# --------------------------------------------------------------------------- #
+
+def test_seed_matches_builtin_defaults():
+ """The seed file mirrors the built-in texts (so nothing drifts silently)."""
+ seed = yaml.safe_load(_strategy_file(default_harness_dir()).read_text(
+ encoding="utf-8"))
+ defaults = rs.builtin_strategy()
+ for section in ("repair", "phases", "append"):
+ for key, value in defaults[section].items():
+ assert seed[section][key].strip() == value, f"{section}.{key}"
+ for regime, portfolio in defaults["portfolios"].items():
+ for slot, value in portfolio.items():
+ assert seed["portfolios"][regime][slot].strip() == value, (
+ f"portfolios.{regime}.{slot}")
+
+
+def test_seed_round_trips_through_the_loader(tmp_path, monkeypatch):
+ """Loading the seed yields exactly the built-in texts."""
+ root = _workspace(tmp_path)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.source.endswith("round_strategy.yaml")
+ assert loaded.repair == rs.builtin_strategy()["repair"]
+ assert loaded.phases == rs.builtin_strategy()["phases"]
+ assert loaded.append == rs.builtin_strategy()["append"]
+ assert loaded.portfolios == rs.builtin_strategy()["portfolios"]
+
+
+def test_no_harness_root_uses_the_plugin_seed(monkeypatch):
+ """Without METAINFER_HARNESS_ROOT the plugin's own seed is the workspace.
+
+ Resolution falls back to ``harness_default/`` (not to an in-code sentinel),
+ so the seed file is what production tasks read — and the drift test above
+ proves it carries the built-in texts.
+ """
+ monkeypatch.delenv("METAINFER_HARNESS_ROOT", raising=False)
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.source.endswith("harness_default/systemprompt/round_strategy.yaml")
+ assert loaded.portfolios["m16"]["2"].startswith("Architecture round")
+ assert loaded.portfolios == rs.builtin_strategy()["portfolios"]
+
+
+# --------------------------------------------------------------------------- #
+# Editing the file changes the runtime mandate
+# --------------------------------------------------------------------------- #
+
+def test_editing_the_seed_changes_runtime_mandate(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["portfolios"]["m16"]["2"] = (
+ "EVOLVED: benchmark 3-wave and 4-wave geometries before anything else.")
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+
+ rendered = w8a8_round_strategy(M16, 2, [], {}, 10, None)
+ assert rendered.startswith("EVOLVED: benchmark 3-wave")
+ # Untouched branches keep their historical text.
+ assert w8a8_round_strategy(M16, 1, [], {}, 10, None).startswith(
+ "Establish a minimal 16x16x32 DUMMA tile")
+
+
+def test_edited_text_is_seen_by_the_pipeline_call_path(tmp_path, monkeypatch):
+ """w8a8_pipeline's non-planner path renders the harness text too."""
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["repair"]["build_failure"] = "EVOLVED BUILD REPAIR for {iteration}"
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ monkeypatch.delenv("METAINFER_PLANNER", raising=False)
+ rs.reset_cache()
+
+ from ..orchestrator import w8a8_pipeline as W
+
+ history = [{"iteration": 4, "build_success": False,
+ "failure_reason": "hipcc: undefined template",
+ "artifact_dir": "/tmp/a4"}]
+ text = W.w8a8_round_strategy(M16, 4, history, {}, max_iterations=10,
+ isa_policy={})
+ assert text == "EVOLVED BUILD REPAIR for 4"
+
+
+def test_repeated_scale_can_be_edited(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["phases"]["isa_guided_hip"] = "EVOLVED ISA round {completed} of many"
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ text = w8a8_round_strategy(
+ M16, 3, [], {}, 10,
+ {"phase": "isa_guided_hip", "valid_isa_guided_rounds": 1})
+ assert text == "EVOLVED ISA round 2 of many"
+
+
+def test_append_suffixes_can_be_edited(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["append"]["split_candidates"] = " SUSPECTS={split_candidates}."
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ text = w8a8_round_strategy(
+ M16, 3, [], {"grid_blocks": 500, "device_cu_count": 60}, 10, None)
+ assert "SUSPECTS=" in text
+ assert "Trusted occupancy-probe split candidates" not in text
+
+
+# --------------------------------------------------------------------------- #
+# The manifest decides whether the file applies at all
+# --------------------------------------------------------------------------- #
+
+def test_unwired_component_falls_back_to_builtin(tmp_path, monkeypatch):
+ root = _workspace(tmp_path, wired=False)
+
+ def mutate(data):
+ data["portfolios"]["m16"]["2"] = "EVOLVED: should never be used"
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.source == "builtin"
+ assert loaded.portfolios["m16"]["2"].startswith("Architecture round")
+ assert any("wired:false" in note for note in loaded.notes)
+
+
+def test_missing_component_entry_still_applies(tmp_path, monkeypatch):
+ """A manifest without the component key must not disable the file."""
+ root = _workspace(tmp_path)
+ manifest = yaml.safe_load((root / "manifest.yaml").read_text(encoding="utf-8"))
+ manifest["components"].pop("systemprompt")
+ (root / "manifest.yaml").write_text(
+ yaml.safe_dump(manifest, sort_keys=False), encoding="utf-8")
+
+ def mutate(data):
+ data["portfolios"]["m16"]["2"] = "EVOLVED without a manifest entry"
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ assert w8a8_round_strategy(M16, 2, [], {}, 10, None).startswith("EVOLVED without")
+
+
+# --------------------------------------------------------------------------- #
+# Degradation: a broken harness must not break a round
+# --------------------------------------------------------------------------- #
+
+def test_missing_file_uses_builtin(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ _strategy_file(root).unlink()
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.source == "builtin"
+ assert loaded.repair == rs.builtin_strategy()["repair"]
+
+
+def test_malformed_yaml_is_reported_not_raised(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ _strategy_file(root).write_text("repair: [unclosed\n", encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.source == "builtin"
+ assert loaded.notes and "round_strategy.yaml" in loaded.notes[0]
+ assert w8a8_round_strategy(M16, 2, [], {}, 10, None).startswith("Architecture")
+
+
+def test_partial_file_keeps_unlisted_branches(tmp_path, monkeypatch):
+ """Omitting a key keeps its built-in text instead of blanking the branch."""
+ root = _workspace(tmp_path)
+ # A file that only carries a reformatted large-M round 3.
+ _strategy_file(root).write_text(
+ "portfolios:\n large:\n \"3\": EVOLVED large round three\n",
+ encoding="utf-8")
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.portfolios["large"]["3"] == "EVOLVED large round three"
+ assert loaded.portfolios["large"]["2"] == (
+ rs.builtin_strategy()["portfolios"]["large"]["2"])
+ assert loaded.portfolios["m16"]["2"].startswith("Architecture round")
+
+
+def test_unknown_fields_degrade_to_builtin_text(tmp_path, monkeypatch):
+ """A renamed format field must not raise mid-round."""
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["append"]["grid_warning"] = " warning for {cus} CUs"
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ text = w8a8_round_strategy(
+ M16, 3, [], {"grid_blocks": 4, "device_cu_count": 60}, 10, None)
+ # The suffix falls back to the built-in wording rather than raising.
+ assert "Trusted control-plane warning" in text
+
+
+def test_unknown_regime_is_ignored(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["portfolios"]["medium"] = {"2": "nonsense"}
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert "medium" not in loaded.portfolios
+ assert any("unknown regime" in note for note in loaded.notes)
+
+
+def test_empty_string_is_ignored(tmp_path, monkeypatch):
+ """An empty value keeps the built-in text (a mandate is never blank)."""
+ root = _workspace(tmp_path)
+
+ def mutate(data):
+ data["portfolios"]["m16"]["2"] = " "
+
+ _edit(root, mutate)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ loaded = rs.load_round_strategy()
+ assert loaded.portfolios["m16"]["2"].startswith("Architecture round")
+
+
+# --------------------------------------------------------------------------- #
+# Cache invalidation: a mid-run harness swap is picked up
+# --------------------------------------------------------------------------- #
+
+def test_cache_is_invalidated_by_file_change(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ assert w8a8_round_strategy(M16, 2, [], {}, 10, None).startswith("Architecture")
+
+ _edit(root, lambda data: data["portfolios"]["m16"].__setitem__(
+ "2", "EVOLVED after a mid-run swap"))
+ import os
+ stat = _strategy_file(root).stat()
+ os.utime(_strategy_file(root), (stat.st_atime, stat.st_mtime + 10))
+ assert w8a8_round_strategy(M16, 2, [], {}, 10, None) == "EVOLVED after a mid-run swap"
+
+
+def test_cache_is_invalidated_by_manifest_change(tmp_path, monkeypatch):
+ root = _workspace(tmp_path)
+ monkeypatch.setenv("METAINFER_HARNESS_ROOT", str(root))
+ rs.reset_cache()
+ assert rs.load_round_strategy().source.endswith("round_strategy.yaml")
+
+ manifest = yaml.safe_load((root / "manifest.yaml").read_text(encoding="utf-8"))
+ manifest["components"]["systemprompt"]["wired"] = False
+ (root / "manifest.yaml").write_text(
+ yaml.safe_dump(manifest, sort_keys=False), encoding="utf-8")
+ import os
+ stat = (root / "manifest.yaml").stat()
+ os.utime(root / "manifest.yaml", (stat.st_atime, stat.st_mtime + 10))
+ assert rs.load_round_strategy().source == "builtin"
+
+
+# --------------------------------------------------------------------------- #
+# The small-M late-round crash is fixed
+# --------------------------------------------------------------------------- #
+
+@pytest.mark.parametrize("max_iterations", [9, 10, 11, 20])
+@pytest.mark.parametrize("iteration", [8, 9, 10, 11, 12])
+def test_small_m_late_rounds_never_raise(max_iterations, iteration):
+ """M<16 has no slot 9/10; it used to raise KeyError for max_iterations 9..10."""
+ text = w8a8_round_strategy(SMALL, iteration, [], {}, max_iterations, None)
+ assert text
+ assert isinstance(text, str)
+
+
+def test_small_m_keeps_its_consolidation_round_late():
+ text = w8a8_round_strategy(SMALL, 10, [], {}, 20, None)
+ assert text.startswith("HIP-only consolidation round")
+
+
+def test_normal_regimes_are_unaffected_by_the_fix():
+ # 16..127 and >=128 still reach the late slots.
+ assert w8a8_round_strategy(M16, 9, [], {}, 10, None).startswith(
+ "Late ISA-diagnosis round")
+ assert w8a8_round_strategy(M16, 10, [], {}, 10, None).startswith(
+ "Final conditional inline-asm round")
+ assert w8a8_round_strategy(PREFILL, 20, [], {}, 20, None).startswith(
+ "Final conditional inline-asm round")
+
+
+# --------------------------------------------------------------------------- #
+# End-to-end: seeding, and the AHE write-scope check reaches the new component
+# --------------------------------------------------------------------------- #
+
+def test_seed_workspace_carries_the_systemprompt_component(tmp_path):
+ """HE seeds the experiment workspace from harness_default/, so the file ships."""
+ from ..orchestrator.harness_io import seed_workspace
+
+ dst = tmp_path / "seeded"
+ seed_workspace(dst)
+ assert (dst / rs.RELATIVE_PATH).is_file()
+ manifest = yaml.safe_load((dst / "manifest.yaml").read_text(encoding="utf-8"))
+ entry = manifest["components"]["systemprompt"]
+ assert entry["path"] == rs.RELATIVE_PATH
+ assert entry["wired"] is True
+
+
+#: Rendered outputs of a 5400-case matrix (M regime x iteration x history x PMC
+#: x ISA phase), frozen after the externalisation refactor. Any edit to a
+#: mandate — in the built-in defaults *or* in the seed file — changes this
+#: digest, which makes an accidental behaviour change impossible to merge
+#: unnoticed. Regenerate deliberately when a wording change is intended.
+_RENDERED_DIGEST = "8a3bb533a7779a45e6ad2a7b45128d0cc1b406155f83d04f1fc74fcb15255fea"
+
+_SHAPES = {
+ "m16": {"model": "hy3", "M": 16, "N": 2048, "K": 256},
+ "m64": {"model": "hy3", "M": 64, "N": 4096, "K": 512},
+ "large": {"model": "hy3", "M": 3072, "N": 2048, "K": 256},
+ "prefill": {"model": "hy3", "M": 4096, "N": 1024, "K": 256},
+ "small": {"model": "hy3", "M": 2, "N": 2048, "K": 256},
+ "tiny": {"model": "hy3", "M": 1, "N": 2048, "K": 256},
+}
+_HISTORIES = {
+ "empty": [],
+ "faster_wrong": [
+ {"iteration": 2, "correctness_passed": False, "speedup": 1.4,
+ "artifact_dir": "/tmp/arch/iter2", "failure_reason": "mismatch_count 12"},
+ {"iteration": 4, "correctness_passed": True, "speedup": 0.9,
+ "artifact_dir": "/tmp/arch/iter4", "failure_reason": ""},
+ ],
+ "infra": [{"iteration": 3, "correctness_passed": True, "build_success": True,
+ "failure_reason": "agent timeout after 900s",
+ "artifact_dir": "/tmp/i3"}],
+ "build": [{"iteration": 3, "correctness_passed": True, "build_success": False,
+ "failure_reason": "hipcc: undefined template",
+ "artifact_dir": "/tmp/i3"}],
+ "assert": [{"iteration": 3, "correctness_passed": True, "build_success": True,
+ "failure_reason": "mismatch_count 3", "artifact_dir": "/tmp/i3"}],
+}
+_PMCS = [
+ {},
+ {"grid_blocks": 4, "device_cu_count": 60, "l2_hit_rate": 0.55},
+ {"grid_blocks": 400, "device_cu_count": 60, "l2_hit_rate": 0.9},
+]
+_ISA_POLICIES = [
+ {},
+ {"phase": "hip_only"},
+ {"phase": "isa_guided_hip", "valid_isa_guided_rounds": 0},
+ {"phase": "isa_guided_hip", "valid_isa_guided_rounds": 3},
+ {"phase": "conditional_inline_asm"},
+]
+
+
+def test_rendered_mandates_match_the_frozen_digest():
+ """The worker-facing mandates are frozen: 5400 rendered strings, one hash."""
+ import hashlib
+ import json
+
+ payload = {}
+ for regime, shape in _SHAPES.items():
+ for iteration in range(1, 13):
+ for hlabel, history in _HISTORIES.items():
+ for pidx, pmc in enumerate(_PMCS):
+ for iidx, isa in enumerate(_ISA_POLICIES):
+ payload[
+ f"{regime}-it{iteration:02d}-{hlabel}"
+ f"-pmc{pidx}-isa{iidx}"
+ ] = w8a8_round_strategy(shape, iteration, history, pmc, 10, isa)
+ blob = json.dumps(payload, ensure_ascii=False, sort_keys=True).encode("utf-8")
+ assert len(payload) == 5400
+ assert hashlib.sha256(blob).hexdigest() == _RENDERED_DIGEST, (
+ "a round mandate changed; if that was intended, regenerate the digest "
+ "intentionally and describe the new wording in the change manifest")
+
+
+@requires_harness_evolve
+def test_ahe_wired_scope_check_covers_the_systemprompt_component(tmp_path):
+ """An edit to the component while it is wired:false is reported by AHE.
+
+ This is what makes the wider write scope auditable: harness_evolve compares
+ each component's declared path, so a new component is covered as soon as it
+ is registered in the manifest.
+ """
+ from metainfer.tasks.harness_evolve.orchestrator.pipeline import (
+ _wired_scope_violations,
+ )
+
+ before = tmp_path / "before"
+ after = tmp_path / "after"
+ shutil.copytree(default_harness_dir(), before)
+ shutil.copytree(default_harness_dir(), after)
+ manifest = yaml.safe_load((after / "manifest.yaml").read_text(encoding="utf-8"))
+ manifest["components"]["systemprompt"]["wired"] = False
+ (after / "manifest.yaml").write_text(
+ yaml.safe_dump(manifest, sort_keys=False), encoding="utf-8")
+ # Unchanged file: not a violation.
+ assert _wired_scope_violations(before, after) == []
+ # Edited file while declared unwired: a violation, with the component named.
+ text = (after / rs.RELATIVE_PATH).read_text(encoding="utf-8")
+ (after / rs.RELATIVE_PATH).write_text(
+ text + "\n# evolved by hand\n", encoding="utf-8")
+ violations = _wired_scope_violations(before, after)
+ assert violations == [{"component": "systemprompt",
+ "path": rs.RELATIVE_PATH}]
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_shape_catalog_panel.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_shape_catalog_panel.py
new file mode 100644
index 00000000..10be54fb
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_shape_catalog_panel.py
@@ -0,0 +1,126 @@
+"""Panel <-> contract <-> baseline integration for the model shape catalog.
+
+``static/dkao-shape-input.js`` is what the DKAO "New Task" panel offers the
+operator, so it is the source of truth for which shapes a task can submit.
+Every shape it can submit must
+
+* satisfy the frozen operator contract (``validate_optimization_shape``), and
+* have a fixed Triton Graph baseline (``fixed_triton_graph_baseline``),
+
+otherwise the run dies in the ``baseline`` phase with a ValueError long after
+the task was created. These tests read the *real* frontend catalog (by
+evaluating its self-contained leading section under node, so the assertions
+cannot drift from the shipped UI) and check that the three layers agree.
+"""
+
+from __future__ import annotations
+
+import json
+import shutil
+import subprocess
+from pathlib import Path
+
+import pytest
+
+from ..api.int8w8a8gemm.int8_w8a8_gemm_api import (
+ MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES,
+ validate_optimization_shape,
+)
+from ..orchestrator.w8a8_baselines import fixed_triton_graph_baseline
+
+
+STATIC_JS = Path(__file__).resolve().parent.parent / "static" / "dkao-shape-input.js"
+MODEL_CATALOG_MARKERS = (
+ "var W8A8_M_VALUES",
+ 'var DEFAULT_W8A8_SHAPES = modelCatalog("DeepSeek V4 Flash");',
+)
+# Models whose TP8 catalog carries the M=8 short-decode boundary.
+M8_MODELS = ("Hy3 (Hunyuan 3)", "MiniMax M3", "GLM5.2")
+
+
+def _frontend_catalogs() -> dict[str, list[dict[str, object]]]:
+ """Return {model label: [shape, ...]} from the shipped panel catalog."""
+ node = shutil.which("node")
+ if node is None:
+ pytest.skip("node is required to evaluate the panel shape catalog")
+ text = STATIC_JS.read_text(encoding="utf-8")
+ start, end = MODEL_CATALOG_MARKERS
+ head = text.index(start)
+ tail = text.index(end) + len(end)
+ harness = (
+ text[head:tail]
+ + "\nvar out = {};\n"
+ + "Object.keys(MODEL_WORKLOADS).forEach(function (label) {\n"
+ + " out[label] = modelCatalog(label);\n"
+ + "});\n"
+ + "console.log(JSON.stringify(out));\n"
+ )
+ proc = subprocess.run(
+ [node, "-e", harness],
+ capture_output=True,
+ text=True,
+ check=False,
+ )
+ assert proc.returncode == 0, proc.stderr
+ return json.loads(proc.stdout)
+
+
+@pytest.fixture(scope="module")
+def catalogs() -> dict[str, list[dict[str, object]]]:
+ return _frontend_catalogs()
+
+
+def test_panel_catalog_exposes_m8_for_model_tp8(catalogs):
+ # The M=8 short-decode boundary added on 2026-09-20 must be offered for
+ # Hy3 / MiniMax M3 / GLM5.2 TP8 operators.
+ for model in M8_MODELS:
+ tp8 = [s for s in catalogs[model] if s["tp_size"] == 8]
+ assert tp8, model
+ assert all(s["M"] in (8, 2, 16, 3072, 4096) for s in tp8), model
+ assert any(s["M"] == 8 for s in tp8), model
+ assert sorted({int(s["M"]) for s in tp8}) == [2, 8, 16, 3072, 4096]
+
+
+def test_panel_catalog_keeps_deepseek_tp8_defaults(catalogs):
+ # DeepSeek TP8 intentionally keeps the three original M values, so the
+ # DeepSeek-only default workload and its validation scope are unchanged.
+ tp8 = [s for s in catalogs["DeepSeek V4 Flash"] if s["tp_size"] == 8]
+ assert sorted({int(s["M"]) for s in tp8}) == [2, 16, 3072]
+
+
+def test_panel_m8_shapes_match_the_extra_m_value_contract(catalogs):
+ assert 8 in MODEL_TP8_EXTRA_OPTIMIZATION_M_VALUES
+ for model in M8_MODELS:
+ m8 = [s for s in catalogs[model] if s["tp_size"] == 8 and s["M"] == 8]
+ # One M=8 shape per catalog operator.
+ assert len(m8) == len(
+ {s["operator"] for s in catalogs[model] if s["tp_size"] == 8}
+ )
+ for shape in m8:
+ validate_optimization_shape(shape)
+
+
+def test_every_panel_tp4_tp8_shape_is_submittable(catalogs):
+ # A shape the panel offers must both pass the frozen contract and resolve a
+ # fixed baseline; otherwise the task fails in the baseline phase.
+ checked = 0
+ for model, shapes in catalogs.items():
+ for shape in shapes:
+ if shape["tp_size"] not in (4, 8):
+ continue
+ validate_optimization_shape(shape)
+ record = fixed_triton_graph_baseline(str(shape["id"]), shape)
+ assert record["median_us"] > 0, (model, shape)
+ checked += 1
+ assert checked >= 70
+
+
+def test_panel_m8_shapes_resolve_decode_baselines(catalogs):
+ for model in M8_MODELS:
+ for shape in catalogs[model]:
+ if shape["tp_size"] != 8 or shape["M"] != 8:
+ continue
+ record = fixed_triton_graph_baseline(str(shape["id"]), shape)
+ assert record["timing_scope"] == "decode_graph_replay"
+ assert record["baseline_kind"] == "triton_graph"
+ assert record["case"], shape
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_skill_store.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_skill_store.py
index 54333e68..9eda815c 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_skill_store.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_skill_store.py
@@ -3,19 +3,31 @@
import pytest
from ..orchestrator.config import load_config
+from ..orchestrator import skill_store as skill_store_mod
from ..orchestrator.skill_store import (
_apply_fuse_decision,
_parse_fuse_decision,
+ bundled_skills_root,
dsh_skills_root,
generate_merged_skill,
generate_worker_skill,
list_skill_library,
publish_skill,
rollback_skill,
+ seed_bundled_skills,
sync_skill_libraries,
)
+@pytest.fixture(autouse=True)
+def _isolated_bundled_skills(tmp_path, monkeypatch):
+ """Point the plugin's bundled-skills root at an empty tmp dir by default,
+ so seeding never pollutes library assertions in unrelated tests."""
+ bundled = tmp_path / "bundled-skills"
+ bundled.mkdir(exist_ok=True)
+ monkeypatch.setattr(skill_store_mod, "bundled_skills_root", lambda: bundled)
+
+
def test_skill_library_lists_and_publishes_without_overwrite(
tmp_path, monkeypatch
):
@@ -84,6 +96,28 @@ def test_sync_skill_libraries_mirrors_dsh_to_ccb(tmp_path, monkeypatch):
assert summary2["skipped"] == ["alpha", "beta"]
+def test_seed_bundled_skills_copies_missing_only(tmp_path, monkeypatch):
+ dsh_root = tmp_path / "dsh"
+ bundled = tmp_path / "bundled"
+ (bundled / "zeta").mkdir(parents=True)
+ (bundled / "zeta" / "SKILL.md").write_text("# zeta seed\n", encoding="utf-8")
+ (bundled / "eta").mkdir(parents=True)
+ (bundled / "eta" / "SKILL.md").write_text("# eta seed\n", encoding="utf-8")
+ # dsh already has eta (different content) -> must not be overwritten
+ (dsh_root / "eta").mkdir(parents=True)
+ (dsh_root / "eta" / "SKILL.md").write_text("# eta live\n", encoding="utf-8")
+ monkeypatch.setenv("DSH_SKILLS_DIR", str(dsh_root))
+ monkeypatch.setattr(skill_store_mod, "bundled_skills_root", lambda: bundled)
+
+ seeded = seed_bundled_skills()
+ assert seeded["added"] == ["zeta"]
+ assert (dsh_root / "zeta" / "SKILL.md").read_text() == "# zeta seed\n"
+ # an existing live skill is never clobbered by the bundled seed
+ assert (dsh_root / "eta" / "SKILL.md").read_text() == "# eta live\n"
+ # idempotent
+ assert seed_bundled_skills()["added"] == []
+
+
def test_rollback_skill_restores_backup(tmp_path, monkeypatch):
dsh_root = tmp_path / "dsh"
ccb_root = tmp_path / "ccb"
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_validation_budget.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_validation_budget.py
new file mode 100644
index 00000000..fbd26361
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_validation_budget.py
@@ -0,0 +1,54 @@
+"""Validation-budget knobs (scope + bench sampling) shared by DKAO and AHE."""
+
+from __future__ import annotations
+
+import pytest
+
+from ..orchestrator.gen_and_opt_pipeline import _validation_shape_list
+from ..orchestrator.validation_budget import (
+ ENV_BENCH_SAMPLES, ENV_BENCH_WARMUPS, ENV_VALIDATE_SCOPE, QUICK_BENCH,
+ resolve_bench_kwargs, resolve_validation_scope,
+)
+
+
+@pytest.fixture(autouse=True)
+def _clean_env(monkeypatch):
+ for name in (ENV_VALIDATE_SCOPE, ENV_BENCH_WARMUPS, ENV_BENCH_SAMPLES,
+ "METAINFER_BENCH_REPLAYS"):
+ monkeypatch.delenv(name, raising=False)
+
+
+def test_scope_defaults_to_full_api(monkeypatch):
+ assert resolve_validation_scope({}) == "api"
+ assert resolve_validation_scope(None) == "api"
+
+
+def test_scope_reads_form_labels(monkeypatch):
+ assert resolve_validation_scope({"validation_scope": "Task shapes only"}) == "task"
+ assert resolve_validation_scope({"validation_scope": "All API shapes"}) == "api"
+ assert resolve_validation_scope({"validation_scope": "task"}) == "task"
+
+
+def test_scope_env_wins(monkeypatch):
+ monkeypatch.setenv(ENV_VALIDATE_SCOPE, "task")
+ assert resolve_validation_scope({"validation_scope": "All API shapes"}) == "task"
+
+
+def test_bench_kwargs_default_and_quick(monkeypatch):
+ assert resolve_bench_kwargs({}) == {}
+ assert resolve_bench_kwargs({"bench_profile": "quick"}) == QUICK_BENCH
+ assert resolve_bench_kwargs({"bench_profile": "full"}) == {}
+
+
+def test_bench_kwargs_explicit_and_env_priority(monkeypatch):
+ out = resolve_bench_kwargs({"bench_profile": "quick", "bench_samples": 7})
+ assert out["samples"] == 7 and out["warmups"] == QUICK_BENCH["warmups"]
+ monkeypatch.setenv(ENV_BENCH_SAMPLES, "5")
+ assert resolve_bench_kwargs({"bench_samples": 7})["samples"] == 5
+
+
+def test_validation_shape_list_scopes():
+ optimized = [{"id": "a"}]
+ fallback = [{"id": "b"}, {"id": "c"}]
+ assert [s["id"] for s in _validation_shape_list(optimized, fallback, "task")] == ["a"]
+ assert [s["id"] for s in _validation_shape_list(optimized, fallback, "api")] == ["a", "b", "c"]
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_promote.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_promote.py
new file mode 100644
index 00000000..cd23a824
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_promote.py
@@ -0,0 +1,149 @@
+"""Kernel promotion gates (variant tree write) — no GPU calls."""
+
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import pytest
+
+from ..orchestrator import variant_store as vs
+from ..orchestrator.variant_promote import (
+ accepted_kernel_for, model_label_for, promote_variant,
+)
+
+ANSWERS = {"operator": "Quantized GEMM", "dtype": "INT8 W8A8"}
+SHAPE = "hy3_tp4_o_proj_m16"
+
+
+@pytest.fixture()
+def variant_root(tmp_path, monkeypatch):
+ root = tmp_path / "variant"
+ root.mkdir(parents=True, exist_ok=True)
+ monkeypatch.setattr(vs, "variant_root", lambda: root)
+ return root
+
+
+def _workspace(tmp_path: Path, median_us: float, *, shape: str = SHAPE,
+ commit: str = "deadbeef") -> Path:
+ ws = tmp_path / "ws"
+ leaf = ws / "workers" / "worker_0" / "accepted" / shape
+ leaf.mkdir(parents=True, exist_ok=True)
+ (leaf / "kernel.hip").write_text("// kernel source\n", encoding="utf-8")
+ (leaf / "manifest.json").write_text(json.dumps({
+ "commit": commit,
+ "shape": {"M": 16, "N": 4096, "K": 2048},
+ "metrics": {"median_us": median_us, "p90_us": median_us * 1.02},
+ }), encoding="utf-8")
+ return ws
+
+
+def _seed_variant(variant_root: Path, median_us: float) -> Path:
+ meta = vs.derive_variant_meta(ANSWERS | {"model": "Hy3 (Hunyuan 3)"}, SHAPE)
+ target = vs.variant_path(meta)
+ target.parent.mkdir(parents=True, exist_ok=True)
+ target.write_text(vs.section_header(meta, commit="old", metrics={
+ "median_us": median_us, "p90_us": median_us}) + "// old kernel\n// @@end\n",
+ encoding="utf-8")
+ return target
+
+
+def test_adds_when_no_existing_variant(tmp_path, variant_root):
+ ws = _workspace(tmp_path, 20.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=3.0,
+ source_task="he-test")
+ assert out["action"] == "added"
+ assert Path(out["path"]).is_file()
+ assert "he-test" in Path(out["path"]).read_text(encoding="utf-8")
+
+
+def test_updates_when_faster_than_threshold(tmp_path, variant_root):
+ _seed_variant(variant_root, 100.0)
+ ws = _workspace(tmp_path, 90.0) # 10% faster
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=3.0)
+ assert out["action"] == "updated"
+ assert out["improvement_percent"] == pytest.approx(10.0, abs=1e-6)
+ assert out["backup"] and Path(out["backup"]).is_file()
+ assert "median_us=90" in Path(out["path"]).read_text(encoding="utf-8")
+
+
+def test_skips_improvement_below_threshold(tmp_path, variant_root):
+ _seed_variant(variant_root, 100.0)
+ ws = _workspace(tmp_path, 98.5) # only 1.5% faster
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=3.0)
+ assert out["action"] == "skipped"
+ assert "required 3.00%" in out["reason"]
+ # existing variant untouched (still the old median)
+ assert "median_us=100" in Path(out["path"]).read_text(encoding="utf-8")
+
+
+def test_slower_candidate_is_skipped_by_default_gate(tmp_path, variant_root):
+ """With the default gate (min 0%) any regression is skipped up front."""
+ _seed_variant(variant_root, 100.0)
+ ws = _workspace(tmp_path, 120.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=0.0)
+ assert out["action"] == "skipped"
+ assert "required 0.00%" in out["reason"]
+
+
+def test_add_variant_guard_rejects_slower_when_gate_is_loosened(tmp_path, variant_root):
+ """A loosened threshold still cannot overwrite with a slower kernel."""
+ _seed_variant(variant_root, 100.0)
+ ws = _workspace(tmp_path, 101.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=-10.0)
+ assert out["action"] == "rejected"
+ assert "slower" in out["reason"]
+
+
+def test_skips_when_correctness_failed(tmp_path, variant_root):
+ ws = _workspace(tmp_path, 10.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=False, min_improvement_percent=3.0)
+ assert out["action"] == "skipped"
+ assert "correctness" in out["reason"]
+
+
+def test_no_kernel_found(tmp_path, variant_root):
+ out = promote_variant(workspace_dir=tmp_path / "empty", answers=ANSWERS,
+ shape_id=SHAPE)
+ assert out["action"] == "no_kernel"
+
+
+def test_dry_run_does_not_write(tmp_path, variant_root):
+ ws = _workspace(tmp_path, 10.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True, min_improvement_percent=3.0,
+ dry_run=True)
+ assert out["action"] == "would-add"
+ assert not Path(out["path"]).exists()
+
+
+def test_meta_taxonomy_and_overrides(tmp_path, variant_root):
+ ws = _workspace(tmp_path, 10.0)
+ out = promote_variant(workspace_dir=ws, answers=ANSWERS, shape_id=SHAPE,
+ correctness_ok=True)
+ meta = out["meta"]
+ assert (meta["family"], meta["model"], meta["tp"], meta["m"],
+ meta["operator_name"]) == ("int8w8a8-gemm", "hy3", 4, 16, "o_proj")
+
+ # shapes without tp/M in the id take them from the caller (pool contract)
+ ws2 = _workspace(tmp_path / "two", 10.0, shape="m16_wqkv_a")
+ out2 = promote_variant(workspace_dir=ws2, answers=ANSWERS,
+ shape_id="m16_wqkv_a", correctness_ok=True,
+ tp=8, m=16, model_label="DeepSeek V4 Flash")
+ assert out2["meta"]["tp"] == 8 and out2["meta"]["m"] == 16
+ assert "/TP8/M16/" in out2["path"]
+
+
+def test_model_label_and_accepted_lookup(tmp_path):
+ assert model_label_for("minimax_tp8_o_proj_m16") == "MiniMax M3"
+ assert model_label_for("glm_tp8_o_proj_m4096") == "GLM5.2"
+ ws = _workspace(tmp_path, 5.0)
+ found = accepted_kernel_for(ws, SHAPE)
+ assert found and found["metrics"]["median_us"] == 5.0
+ assert accepted_kernel_for(ws, "nope") is None
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_store.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_store.py
index 6310e5e2..0f5a29c4 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_store.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_variant_store.py
@@ -196,3 +196,31 @@ def test_reject_slower_guard_skips_when_median_unknown(tmp_path, monkeypatch):
)
assert result["action"] == "updated"
assert "NEW" in vs.variant_path(meta).read_text()
+
+
+def test_backfill_variant_baselines(tmp_path, monkeypatch):
+ monkeypatch.setattr(vs, "variant_root", lambda: tmp_path)
+ target = (
+ tmp_path / "int8w8a8-gemm" / "deepseek-v4" / "TP4" / "M16"
+ / "wqkv_a.hip"
+ )
+ target.parent.mkdir(parents=True)
+ target.write_text(
+ "// @@variant shape=tp4_wqkv_a_m16 commit=abc added=2026-08-26\n"
+ "// median_us=20.0 p90_us=21.0\n"
+ "// source=task-1\n"
+ "// @@end\nkernel\n",
+ encoding="utf-8",
+ )
+ r = vs.backfill_variant_baselines()
+ # tp4 wqkv_a m16 fixed Triton baseline is 66.597 us -> speedup ~3.33
+ assert r["updated"] == ["int8w8a8-gemm/deepseek-v4/TP4/M16/wqkv_a.hip"]
+ text = target.read_text()
+ assert "baseline_us=66.597" in text
+ assert "speedup=3.329" in text
+ assert text.index("baseline_us") < text.index("median_us")
+ # idempotent and surfaced by the index
+ assert vs.backfill_variant_baselines()["updated"] == []
+ v = vs.list_variant_index()[0]
+ assert v["speedup"] is not None
+ assert abs(v["baseline_us"] - 66.597) < 1e-3
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_w8a8_baselines.py b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_w8a8_baselines.py
index 0eb9d65a..b3a56910 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_w8a8_baselines.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_w8a8_baselines.py
@@ -86,6 +86,37 @@ def test_model_catalog_tp8_m4096_baselines(shape, expected):
assert record["timing_scope"] == "prefill_graph_replay"
+@pytest.mark.parametrize(
+ ("shape", "expected"),
+ [
+ # Hy3 TP8 M=8 (measured 2026-09-20, CUDA-graph replay).
+ ({"tp_size": 8, "M": 8, "N": 1280, "K": 4096}, 66.432),
+ ({"tp_size": 8, "M": 8, "N": 4096, "K": 1024}, 31.472),
+ ({"tp_size": 8, "M": 8, "N": 384, "K": 4096}, 57.008),
+ ({"tp_size": 8, "M": 8, "N": 4096, "K": 192}, 14.592),
+ # MiniMax M3 TP8 M=8.
+ ({"tp_size": 8, "M": 8, "N": 1280, "K": 6144}, 98.865),
+ ({"tp_size": 8, "M": 8, "N": 1536, "K": 6144}, 99.473),
+ ({"tp_size": 8, "M": 8, "N": 6144, "K": 1024}, 32.960),
+ ({"tp_size": 8, "M": 8, "N": 768, "K": 6144}, 92.928),
+ ({"tp_size": 8, "M": 8, "N": 6144, "K": 384}, 19.984),
+ # GLM5.2 TP8 M=8.
+ ({"tp_size": 8, "M": 8, "N": 2624, "K": 6144}, 87.200),
+ ({"tp_size": 8, "M": 8, "N": 2048, "K": 2048}, 36.704),
+ ({"tp_size": 8, "M": 8, "N": 3584, "K": 512}, 15.808),
+ ({"tp_size": 8, "M": 8, "N": 6144, "K": 2048}, 58.864),
+ ({"tp_size": 8, "M": 8, "N": 512, "K": 6144}, 88.448),
+ ({"tp_size": 8, "M": 8, "N": 6144, "K": 256}, 15.136),
+ ],
+)
+def test_model_catalog_tp8_m8_baselines(shape, expected):
+ # M=8 is the model-catalog short-decode boundary added on 2026-09-20 and
+ # resolves as a decode baseline (M <= 16).
+ record = fixed_triton_graph_baseline("shape", shape)
+ assert record["median_us"] == expected
+ assert record["timing_scope"] == "decode_graph_replay"
+
+
def test_bootstrap_metrics_are_kept_separate():
bootstrap = {"passed": True, "median_us": 123.0}
record = fixed_triton_graph_baseline(
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/__init__.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/__init__.py
new file mode 100644
index 00000000..e69de29b
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp4_m4096.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp4_m4096.py
new file mode 100644
index 00000000..6da463c1
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp4_m4096.py
@@ -0,0 +1,277 @@
+#!/usr/bin/env python3
+"""Fair Triton W8A8 GEMM baseline for the six TP4 logical shapes at M=4096.
+
+Timed region launches only lmslim-style int8_utils.matmul_kernel into a
+preallocated bf16 output (output allocation/clear excluded), matching the
+MetaInfer fair-baseline scope. Uses GPU events, hot cache, and reports
+median/P90/min in microseconds plus effective INT8 TOPS.
+
+The Triton config replicates matmul_int8's built-in default for M > 1024:
+ BLOCK_SIZE_M=256, BLOCK_SIZE_N=256, BLOCK_SIZE_K=64,
+ GROUP_SIZE_M=8, SPLIT_K=1, num_stages=0, num_warps=8
+"""
+
+from __future__ import annotations
+
+import argparse
+import importlib.util
+import json
+import math
+import statistics
+import sys
+from pathlib import Path
+
+import torch
+import triton
+
+
+UTILS_PATH = str(Path(__file__).resolve().parent / "int8_utils.py")
+M = 4096
+
+# (shape_id, operator, K, N) -- the six TP4 logical operators.
+TP4_CASES = [
+ ("tp4_wqkv_a_m4096", "wqkv_a", 4096, 1536),
+ ("tp4_wq_b_m4096", "wq_b", 1024, 8192),
+ ("tp4_indexer_wq_b_m4096", "indexer.wq_b", 1024, 8192),
+ ("tp4_wo_b_m4096", "wo_b", 2048, 4096),
+ ("tp4_shared_gate_up_proj_m4096", "shared_gate_up_proj", 4096, 1024),
+ ("tp4_shared_down_proj_m4096", "shared_down_proj", 512, 4096),
+]
+
+# Exact default config used by matmul_int8 when M > 1024.
+CONFIG = {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8,
+}
+
+
+def load_int8_utils(path: str):
+ spec = importlib.util.spec_from_file_location(
+ "baseline_int8_utils", path
+ )
+ if spec is None or spec.loader is None:
+ raise RuntimeError(f"cannot load {path}")
+ module = importlib.util.module_from_spec(spec)
+ sys.modules[spec.name] = module
+ spec.loader.exec_module(module)
+ return module
+
+
+def launch(out, a, a_scale, b, b_scale, m, n, k, utils) -> None:
+ grid = (
+ triton.cdiv(m, CONFIG["BLOCK_SIZE_M"])
+ * triton.cdiv(n, CONFIG["BLOCK_SIZE_N"]),
+ CONFIG["SPLIT_K"],
+ )
+ utils.matmul_kernel[grid](
+ a,
+ a_scale,
+ b,
+ b_scale,
+ out,
+ m,
+ n,
+ k,
+ a.stride(0),
+ a.stride(1),
+ a_scale.stride(0),
+ b.stride(0),
+ b.stride(1),
+ b_scale.stride(0),
+ out.stride(0),
+ out.stride(1),
+ **CONFIG,
+ )
+
+
+def percentile(values, fraction):
+ ordered = sorted(values)
+ return ordered[max(0, math.ceil(fraction * len(ordered)) - 1)]
+
+
+def measure(fn, warmups, samples, launches_per_sample):
+ for _ in range(warmups):
+ fn()
+ torch.cuda.synchronize()
+
+ start = torch.cuda.Event(enable_timing=True)
+ end = torch.cuda.Event(enable_timing=True)
+ times_us = []
+ for _ in range(samples):
+ start.record()
+ for _ in range(launches_per_sample):
+ fn()
+ end.record()
+ end.synchronize()
+ times_us.append(
+ start.elapsed_time(end) * 1000.0 / launches_per_sample
+ )
+ return {
+ "median_us": statistics.median(times_us),
+ "p90_us": percentile(times_us, 0.90),
+ "min_us": min(times_us),
+ "max_us": max(times_us),
+ "samples_us": [round(v, 3) for v in times_us],
+ }
+
+
+@torch.no_grad()
+def reference_us(a, b, a_scale, b_scale, device):
+ dot = torch.mm(
+ a.to("cpu", dtype=torch.int64),
+ b.to("cpu", dtype=torch.int64),
+ )
+ scaled = (
+ dot.to(torch.float32)
+ * a_scale.to("cpu", dtype=torch.float32)
+ * b_scale.to("cpu", dtype=torch.float32).T
+ )
+ return scaled.to(torch.bfloat16).to(device)
+
+
+def run_case(utils, case_id, operator, k, n, warmups, samples,
+ launches_per_sample, skip_check, mode):
+ device = torch.device("cuda:0")
+ torch.manual_seed(20260724 + M + n + k)
+ a = torch.randint(-127, 128, (M, k), device=device, dtype=torch.int8)
+ b = torch.randint(-127, 128, (k, n), device=device, dtype=torch.int8)
+ a_scale = torch.rand((M, 1), device=device, dtype=torch.float32) + 0.01
+ b_scale = torch.rand((n, 1), device=device, dtype=torch.float32) + 0.01
+ out = torch.empty((M, n), device=device, dtype=torch.bfloat16)
+
+ fn = lambda: launch(out, a, a_scale, b, b_scale, M, n, k, utils)
+
+ # One untimed JIT compile + correctness sanity pass.
+ fn()
+ torch.cuda.synchronize()
+ errors = None
+ if not skip_check:
+ ref = reference_us(a, b, a_scale, b_scale, device)
+ diff = (out.float() - ref.float()).abs()
+ denom = ref.float().abs().clamp_min(1e-6)
+ errors = {
+ "max_abs": float(diff.max().item()),
+ "max_rel": float((diff / denom).max().item()),
+ "mismatch_count": int((out != ref).sum().item()),
+ }
+ del ref
+
+ result = {
+ "shape_id": case_id,
+ "operator": operator,
+ "M": M,
+ "N": n,
+ "K": k,
+ "config": CONFIG,
+ "errors": errors,
+ }
+
+ if mode in ("eager", "both"):
+ result["eager"] = measure(
+ fn, warmups, samples, launches_per_sample
+ )
+ if mode in ("graph", "both"):
+ fn()
+ torch.cuda.synchronize()
+ graph = torch.cuda.CUDAGraph()
+ with torch.cuda.graph(graph):
+ fn()
+ torch.cuda.synchronize()
+ result["graph"] = measure(
+ graph.replay, warmups, samples, launches_per_sample
+ )
+ del graph
+
+ del a, b, a_scale, b_scale, out
+ return result
+
+
+def main() -> int:
+ parser = argparse.ArgumentParser()
+ parser.add_argument("--utils", default=UTILS_PATH)
+ parser.add_argument("--warmups", type=int, default=10)
+ parser.add_argument("--samples", type=int, default=20)
+ parser.add_argument("--launches-per-sample", type=int, default=5)
+ parser.add_argument("--mode", choices=["eager", "graph", "both"],
+ default="both")
+ parser.add_argument("--skip-check", action="store_true")
+ parser.add_argument("--json", type=str, default="")
+ args = parser.parse_args()
+
+ utils = load_int8_utils(args.utils)
+ print(
+ "BENCHMARK triton=int8_utils.matmul_kernel output=preallocated "
+ f"quantization=excluded allocation=excluded epilogue=included "
+ f"M={M} config={json.dumps(CONFIG)} mode={args.mode}",
+ flush=True,
+ )
+ print(
+ f"{'shape_id':28s} {'N':5s} {'K':5s} {'eager_us':>10s} "
+ f"{'graph_us':>10s} {'tops':>9s}",
+ flush=True,
+ )
+ results = []
+ for case_id, operator, k, n in TP4_CASES:
+ res = run_case(
+ utils,
+ case_id,
+ operator,
+ k,
+ n,
+ args.warmups,
+ args.samples,
+ args.launches_per_sample,
+ args.skip_check,
+ args.mode,
+ )
+ results.append(res)
+ eager_us = res.get("eager", {}).get("median_us")
+ graph_us = res.get("graph", {}).get("median_us")
+ tops = (
+ 2.0 * M * n * k / ((graph_us or eager_us) * 1.0e-6) / 1.0e12
+ if (graph_us or eager_us) else float("nan")
+ )
+ print(
+ f"{case_id:28s} {n:5d} {k:5d} "
+ f"{(eager_us if eager_us is not None else float('nan')):10.3f} "
+ f"{(graph_us if graph_us is not None else float('nan')):10.3f} "
+ f"{tops:9.3f}",
+ flush=True,
+ )
+ for key in ("eager", "graph"):
+ if key in res:
+ m_ = res[key]
+ print(
+ f" {key}: median={m_['median_us']:.3f}us "
+ f"p90={m_['p90_us']:.3f}us min={m_['min_us']:.3f}us "
+ f"max={m_['max_us']:.3f}us errors={res['errors']}",
+ flush=True,
+ )
+ if args.json:
+ with open(args.json, "w", encoding="utf-8") as handle:
+ json.dump(
+ {
+ "scope": "gemm_out",
+ "triton_source": args.utils,
+ "M": M,
+ "warmups": args.warmups,
+ "samples": args.samples,
+ "launches_per_sample": args.launches_per_sample,
+ "mode": args.mode,
+ "results": results,
+ },
+ handle,
+ indent=2,
+ ensure_ascii=False,
+ )
+ print(f"wrote {args.json}", flush=True)
+ return 0
+
+
+if __name__ == "__main__":
+ sys.exit(main())
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m4096.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m4096.py
new file mode 100644
index 00000000..3dc94b39
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m4096.py
@@ -0,0 +1,233 @@
+#!/usr/bin/env python3
+"""Fair Triton W8A8 GEMM baselines for the model-catalog TP8 shapes at M=4096.
+
+Same protocol as bench_triton_tp4_m4096.py / the 2026-08-12 catalog entries:
+lmslim-style int8_utils.matmul_kernel into a preallocated bf16 output (output
+allocation/clear excluded), GPU events, hot cache, CUDA-Graph replay,
+warmups=10 / samples=20 / launches_per_sample=5. Reports median/P90/min in
+microseconds.
+
+Triton config replicates matmul_int8's built-in default for M > 1024:
+ BLOCK_SIZE_M=256, BLOCK_SIZE_N=256, BLOCK_SIZE_K=64,
+ GROUP_SIZE_M=8, SPLIT_K=1, num_stages=0, num_warps=8
+"""
+
+from __future__ import annotations
+
+import argparse
+import importlib.util
+import json
+import math
+import statistics
+import sys
+from pathlib import Path
+
+import torch
+import triton
+
+
+UTILS_PATH = str(Path(__file__).resolve().parent / "int8_utils.py")
+M = 4096
+
+# (shape_id, operator, K, N) -- Hy3 / MiniMax M3 / GLM5.2 TP8 at M=4096.
+TP8_CASES = [
+ ("hy3_tp8_qkv_proj_m4096", "qkv_proj", 4096, 1280),
+ ("hy3_tp8_o_proj_m4096", "o_proj", 1024, 4096),
+ ("hy3_tp8_shared_gate_up_proj_m4096", "shared_gate_up_proj", 4096, 384),
+ ("hy3_tp8_shared_down_proj_m4096", "shared_down_proj", 192, 4096),
+ ("minimax_tp8_qkv_proj_m4096", "qkv_proj", 6144, 1280),
+ ("minimax_tp8_qkv_proj_and_indexer_qk_m4096", "qkv_proj_and_indexer_qk", 6144, 1536),
+ ("minimax_tp8_o_proj_m4096", "o_proj", 1024, 6144),
+ ("minimax_tp8_shared_gate_up_proj_m4096", "shared_gate_up_proj", 6144, 768),
+ ("minimax_tp8_shared_down_proj_m4096", "shared_down_proj", 384, 6144),
+ ("glm52_tp8_fused_qkv_a_proj_m4096", "fused_qkv_a_proj", 6144, 2624),
+ ("glm52_tp8_q_b_proj_m4096", "q_b_proj", 2048, 2048),
+ ("glm52_tp8_kv_b_proj_m4096", "kv_b_proj", 512, 3584),
+ ("glm52_tp8_o_proj_m4096", "o_proj", 2048, 6144),
+ ("glm52_tp8_shared_gate_up_proj_m4096", "shared_gate_up_proj", 6144, 512),
+ ("glm52_tp8_shared_down_proj_m4096", "shared_down_proj", 256, 6144),
+]
+
+# Exact default config used by matmul_int8 when M > 1024.
+CONFIG = {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8,
+}
+
+
+def load_int8_utils(path: str):
+ spec = importlib.util.spec_from_file_location(
+ "baseline_int8_utils", path
+ )
+ if spec is None or spec.loader is None:
+ raise RuntimeError(f"cannot load {path}")
+ module = importlib.util.module_from_spec(spec)
+ sys.modules[spec.name] = module
+ spec.loader.exec_module(module)
+ return module
+
+
+def launch(out, a, a_scale, b, b_scale, m, n, k, utils) -> None:
+ grid = (
+ triton.cdiv(m, CONFIG["BLOCK_SIZE_M"])
+ * triton.cdiv(n, CONFIG["BLOCK_SIZE_N"]),
+ CONFIG["SPLIT_K"],
+ )
+ utils.matmul_kernel[grid](
+ a,
+ a_scale,
+ b,
+ b_scale,
+ out,
+ m,
+ n,
+ k,
+ a.stride(0),
+ a.stride(1),
+ a_scale.stride(0),
+ b.stride(0),
+ b.stride(1),
+ b_scale.stride(0),
+ out.stride(0),
+ out.stride(1),
+ **CONFIG,
+ )
+
+
+def percentile(values, fraction):
+ ordered = sorted(values)
+ return ordered[max(0, math.ceil(fraction * len(ordered)) - 1)]
+
+
+def measure(fn, warmups, samples, launches_per_sample):
+ for _ in range(warmups):
+ fn()
+ torch.cuda.synchronize()
+
+ start = torch.cuda.Event(enable_timing=True)
+ end = torch.cuda.Event(enable_timing=True)
+ times_us = []
+ for _ in range(samples):
+ start.record()
+ for _ in range(launches_per_sample):
+ fn()
+ end.record()
+ end.synchronize()
+ times_us.append(
+ start.elapsed_time(end) * 1000.0 / launches_per_sample
+ )
+ return {
+ "median_us": statistics.median(times_us),
+ "p90_us": percentile(times_us, 0.90),
+ "min_us": min(times_us),
+ "max_us": max(times_us),
+ "samples_us": [round(v, 3) for v in times_us],
+ }
+
+
+def run_case(utils, case_id, operator, k, n, warmups, samples,
+ launches_per_sample, mode):
+ device = torch.device("cuda:0")
+ torch.manual_seed(20260724 + M + n + k)
+ a = torch.randint(-127, 128, (M, k), device=device, dtype=torch.int8)
+ b = torch.randint(-127, 128, (k, n), device=device, dtype=torch.int8)
+ a_scale = torch.rand((M, 1), device=device, dtype=torch.float32) + 0.01
+ b_scale = torch.rand((n, 1), device=device, dtype=torch.float32) + 0.01
+ out = torch.empty((M, n), device=device, dtype=torch.bfloat16)
+
+ fn = lambda: launch(out, a, a_scale, b, b_scale, M, n, k, utils)
+
+ # One untimed JIT compile pass (correctness is the production Triton path).
+ fn()
+ torch.cuda.synchronize()
+
+ result = {
+ "shape_id": case_id,
+ "operator": operator,
+ "M": M,
+ "N": n,
+ "K": k,
+ "config": CONFIG,
+ }
+ if mode in ("eager", "both"):
+ result["eager"] = measure(
+ fn, warmups, samples, launches_per_sample
+ )
+ if mode in ("graph", "both"):
+ fn()
+ torch.cuda.synchronize()
+ graph = torch.cuda.CUDAGraph()
+ with torch.cuda.graph(graph):
+ fn()
+ torch.cuda.synchronize()
+ result["graph"] = measure(
+ graph.replay, warmups, samples, launches_per_sample
+ )
+ del graph
+
+ del a, b, a_scale, b_scale, out
+ return result
+
+
+def main() -> int:
+ parser = argparse.ArgumentParser()
+ parser.add_argument("--utils", default=UTILS_PATH)
+ parser.add_argument("--warmups", type=int, default=10)
+ parser.add_argument("--samples", type=int, default=20)
+ parser.add_argument("--launches-per-sample", type=int, default=5)
+ parser.add_argument("--mode", choices=["eager", "graph", "both"],
+ default="graph")
+ parser.add_argument("--json", type=str, default="")
+ args = parser.parse_args()
+
+ utils = load_int8_utils(args.utils)
+ print(
+ "BENCHMARK triton=int8_utils.matmul_kernel output=preallocated "
+ f"quantization=excluded allocation=excluded epilogue=included "
+ f"M={M} config={json.dumps(CONFIG)} mode={args.mode}",
+ flush=True,
+ )
+ print(
+ f"{'shape_id':42s} {'N':5s} {'K':5s} {'graph_us':>10s} "
+ f"{'p90_us':>10s} {'tops':>9s}",
+ flush=True,
+ )
+ results = []
+ for case_id, operator, k, n in TP8_CASES:
+ res = run_case(
+ utils,
+ case_id,
+ operator,
+ k,
+ n,
+ args.warmups,
+ args.samples,
+ args.launches_per_sample,
+ args.mode,
+ )
+ results.append(res)
+ graph_us = (res.get("graph") or res.get("eager") or {}).get("median_us")
+ p90_us = (res.get("graph") or res.get("eager") or {}).get("p90_us")
+ tops = (
+ 2.0 * M * n * k / (graph_us * 1e-6) / 1e12
+ if graph_us else None
+ )
+ print(
+ f"{case_id:42s} {n:5d} {k:5d} "
+ f"{graph_us:>10.3f} {p90_us:>10.3f} "
+ f"{(f'{tops:.3f}' if tops else '—'):>9s}",
+ flush=True,
+ )
+ if args.json:
+ with open(args.json, "w", encoding="utf-8") as fh:
+ json.dump(results, fh, indent=2)
+ return 0
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m8.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m8.py
new file mode 100644
index 00000000..57de33b6
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/bench_triton_tp8_m8.py
@@ -0,0 +1,238 @@
+#!/usr/bin/env python3
+"""Fair Triton W8A8 GEMM baselines for the model-catalog TP8 shapes at M=8.
+
+Same protocol as bench_triton_tp8_m4096.py / the 2026-08-12 catalog entries:
+lmslim-style int8_utils.matmul_kernel into a preallocated bf16 output (output
+allocation/clear excluded), GPU events, hot cache, CUDA-Graph replay,
+warmups=10 / samples=20 / launches_per_sample=5. Reports median/P90/min in
+microseconds.
+
+Triton config replicates matmul_int8's built-in default for M <= 32:
+ BLOCK_SIZE_M=16, BLOCK_SIZE_N=32, BLOCK_SIZE_K=256,
+ GROUP_SIZE_M=4, SPLIT_K=1, num_stages=0, num_warps=4
+
+``--m`` defaults to 8 (the new decode boundary). Passing ``--m 16`` re-measures
+the already-published M=16 catalog entries as a protocol cross-check.
+"""
+
+from __future__ import annotations
+
+import argparse
+import importlib.util
+import json
+import math
+import statistics
+import sys
+from pathlib import Path
+
+import torch
+import triton
+
+
+UTILS_PATH = str(Path(__file__).resolve().parent / "int8_utils.py")
+
+# (shape_id stem, operator, K, N) -- Hy3 / MiniMax M3 / GLM5.2 TP8 catalog.
+TP8_CASES = [
+ ("hy3_tp8_qkv_proj", "qkv_proj", 4096, 1280),
+ ("hy3_tp8_o_proj", "o_proj", 1024, 4096),
+ ("hy3_tp8_shared_gate_up_proj", "shared_gate_up_proj", 4096, 384),
+ ("hy3_tp8_shared_down_proj", "shared_down_proj", 192, 4096),
+ ("minimax_tp8_qkv_proj", "qkv_proj", 6144, 1280),
+ ("minimax_tp8_qkv_proj_and_indexer_qk", "qkv_proj_and_indexer_qk", 6144, 1536),
+ ("minimax_tp8_o_proj", "o_proj", 1024, 6144),
+ ("minimax_tp8_shared_gate_up_proj", "shared_gate_up_proj", 6144, 768),
+ ("minimax_tp8_shared_down_proj", "shared_down_proj", 384, 6144),
+ ("glm52_tp8_fused_qkv_a_proj", "fused_qkv_a_proj", 6144, 2624),
+ ("glm52_tp8_q_b_proj", "q_b_proj", 2048, 2048),
+ ("glm52_tp8_kv_b_proj", "kv_b_proj", 512, 3584),
+ ("glm52_tp8_o_proj", "o_proj", 2048, 6144),
+ ("glm52_tp8_shared_gate_up_proj", "shared_gate_up_proj", 6144, 512),
+ ("glm52_tp8_shared_down_proj", "shared_down_proj", 256, 6144),
+]
+
+# Exact default config used by matmul_int8 when M <= 32.
+CONFIG = {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4,
+}
+
+
+def load_int8_utils(path: str):
+ spec = importlib.util.spec_from_file_location("baseline_int8_utils", path)
+ if spec is None or spec.loader is None:
+ raise RuntimeError(f"cannot load {path}")
+ module = importlib.util.module_from_spec(spec)
+ sys.modules[spec.name] = module
+ spec.loader.exec_module(module)
+ return module
+
+
+def launch(out, a, a_scale, b, b_scale, m, n, k, utils) -> None:
+ grid = (
+ triton.cdiv(m, CONFIG["BLOCK_SIZE_M"])
+ * triton.cdiv(n, CONFIG["BLOCK_SIZE_N"]),
+ CONFIG["SPLIT_K"],
+ )
+ utils.matmul_kernel[grid](
+ a,
+ a_scale,
+ b,
+ b_scale,
+ out,
+ m,
+ n,
+ k,
+ a.stride(0),
+ a.stride(1),
+ a_scale.stride(0),
+ b.stride(0),
+ b.stride(1),
+ b_scale.stride(0),
+ out.stride(0),
+ out.stride(1),
+ **CONFIG,
+ )
+
+
+def percentile(values, fraction):
+ ordered = sorted(values)
+ return ordered[max(0, math.ceil(fraction * len(ordered)) - 1)]
+
+
+def measure(fn, warmups, samples, launches_per_sample):
+ for _ in range(warmups):
+ fn()
+ torch.cuda.synchronize()
+
+ start = torch.cuda.Event(enable_timing=True)
+ end = torch.cuda.Event(enable_timing=True)
+ times_us = []
+ for _ in range(samples):
+ start.record()
+ for _ in range(launches_per_sample):
+ fn()
+ end.record()
+ end.synchronize()
+ times_us.append(start.elapsed_time(end) * 1000.0 / launches_per_sample)
+ return {
+ "median_us": statistics.median(times_us),
+ "p90_us": percentile(times_us, 0.90),
+ "min_us": min(times_us),
+ "max_us": max(times_us),
+ "samples_us": [round(v, 3) for v in times_us],
+ }
+
+
+def run_case(utils, case_stem, operator, k, n, m, warmups, samples,
+ launches_per_sample, mode):
+ device = torch.device("cuda:0")
+ torch.manual_seed(20260724 + m + n + k)
+ a = torch.randint(-127, 128, (m, k), device=device, dtype=torch.int8)
+ b = torch.randint(-127, 128, (k, n), device=device, dtype=torch.int8)
+ a_scale = torch.rand((m, 1), device=device, dtype=torch.float32) + 0.01
+ b_scale = torch.rand((n, 1), device=device, dtype=torch.float32) + 0.01
+ out = torch.empty((m, n), device=device, dtype=torch.bfloat16)
+
+ def fn():
+ launch(out, a, a_scale, b, b_scale, m, n, k, utils)
+
+ # One untimed JIT compile pass (correctness is the production Triton path).
+ fn()
+ torch.cuda.synchronize()
+
+ result = {
+ "shape_id": f"{case_stem}_m{m}",
+ "operator": operator,
+ "M": m,
+ "N": n,
+ "K": k,
+ "config": CONFIG,
+ }
+ if mode in ("eager", "both"):
+ result["eager"] = measure(fn, warmups, samples, launches_per_sample)
+ if mode in ("graph", "both"):
+ fn()
+ torch.cuda.synchronize()
+ graph = torch.cuda.CUDAGraph()
+ with torch.cuda.graph(graph):
+ fn()
+ torch.cuda.synchronize()
+ result["graph"] = measure(
+ graph.replay, warmups, samples, launches_per_sample
+ )
+ del graph
+
+ del a, b, a_scale, b_scale, out
+ return result
+
+
+def main() -> int:
+ parser = argparse.ArgumentParser()
+ parser.add_argument("--utils", default=UTILS_PATH)
+ parser.add_argument("--m", type=int, default=8)
+ parser.add_argument("--warmups", type=int, default=10)
+ parser.add_argument("--samples", type=int, default=20)
+ parser.add_argument("--launches-per-sample", type=int, default=5)
+ parser.add_argument("--mode", choices=["eager", "graph", "both"],
+ default="graph")
+ parser.add_argument("--json", type=str, default="")
+ parser.add_argument("--only", type=str, default="",
+ help="comma-separated case-stem substrings to keep")
+ args = parser.parse_args()
+ m = args.m
+
+ utils = load_int8_utils(args.utils)
+ cases = TP8_CASES
+ if args.only:
+ needles = [s.strip() for s in args.only.split(",") if s.strip()]
+ cases = [c for c in cases if any(x in c[0] for x in needles)]
+
+ print(
+ "BENCHMARK triton=int8_utils.matmul_kernel output=preallocated "
+ f"quantization=excluded allocation=excluded epilogue=included "
+ f"M={m} config={json.dumps(CONFIG)} mode={args.mode}",
+ flush=True,
+ )
+ print(
+ f"{'shape_id':44s} {'N':5s} {'K':5s} {'graph_us':>10s} "
+ f"{'p90_us':>10s} {'eager_us':>10s}",
+ flush=True,
+ )
+ results = []
+ for case_stem, operator, k, n in cases:
+ res = run_case(
+ utils,
+ case_stem,
+ operator,
+ k,
+ n,
+ m,
+ args.warmups,
+ args.samples,
+ args.launches_per_sample,
+ args.mode,
+ )
+ results.append(res)
+ graph_us = (res.get("graph") or {}).get("median_us")
+ p90_us = (res.get("graph") or {}).get("p90_us")
+ eager_us = (res.get("eager") or {}).get("median_us")
+ print(
+ f"{res['shape_id']:44s} {n:5d} {k:5d} "
+ f"{(f'{graph_us:.3f}' if graph_us else '—'):>10s} "
+ f"{(f'{p90_us:.3f}' if p90_us else '—'):>10s} "
+ f"{(f'{eager_us:.3f}' if eager_us else '—'):>10s}",
+ flush=True,
+ )
+ if args.json:
+ with open(args.json, "w", encoding="utf-8") as fh:
+ json.dump(results, fh, indent=2)
+ return 0
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/int8_utils.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/int8_utils.py
new file mode 100644
index 00000000..11886719
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/int8_utils.py
@@ -0,0 +1,602 @@
+import functools
+import json
+import logging
+import os
+from typing import Any, Dict, List, Optional, Tuple
+import json
+
+import torch
+import triton
+import triton.language as tl
+from triton.language.extra import libdevice
+import lmslim.envs as lsenvs
+
+use_lightop = lsenvs.LMSLIM_USE_LIGHTOP
+device_name = lsenvs.LMSLIM_GPU_NAME
+
+@triton.jit
+def _per_token_quant_int8(
+ x_ptr,
+ xq_ptr,
+ scale_ptr,
+ stride_x,
+ stride_xq,
+ N,
+ BLOCK: tl.constexpr,
+):
+ # Adapted from https://github.com/InternLM/lmdeploy/blob/086481ed84b59bee3b8e4274e5fc69620040c048/lmdeploy/pytorch/kernels/cuda/w8a8_triton_kernels.py#L282
+ row_id = tl.program_id(0)
+
+ cols = tl.arange(0, BLOCK)
+ mask = cols < N
+
+ x = tl.load(x_ptr + row_id * stride_x + cols, mask=mask,
+ other=0.0).to(tl.float32)
+ absmax = tl.maximum(tl.max(tl.abs(x)), 1e-10)
+ scale_x = absmax / 127
+ x_q = x * (127 / absmax)
+ x_q = libdevice.nearbyint(x_q).to(tl.int8)
+
+ tl.store(xq_ptr + row_id * stride_xq + cols, x_q, mask=mask)
+ tl.store(scale_ptr + row_id, scale_x)
+
+
+def per_token_quant_int8(x):
+ M = x.numel() // x.shape[-1]
+ N = x.shape[-1]
+ x_q = torch.empty_like(x, device=x.device, dtype=torch.int8)
+ scales = torch.empty(x.shape[:-1] + (1, ),
+ device=x.device,
+ dtype=torch.float32)
+ BLOCK = triton.next_power_of_2(N)
+ # heuristics for number of warps
+ num_warps = min(max(BLOCK // 256, 1), 8)
+
+ #assert x.is_contiguous()
+ _per_token_quant_int8[(M, )](
+ x,
+ x_q,
+ scales,
+ stride_x=x.stride(-2),
+ stride_xq=x_q.stride(-2),
+ N=N,
+ BLOCK=BLOCK,
+ num_warps=num_warps,
+ num_stages=1,
+ )
+
+ return x_q, scales
+
+
+@triton.jit
+def _per_token_group_quant_int8(
+ # Pointers to inputs and output
+ y_ptr,
+ y_q_ptr,
+ y_s_ptr,
+ # Stride of input
+ group_size,
+ # M,
+ # K,
+ # # Collums of input
+ # N,
+ SIZE,
+ # Avoid to divide zero
+ eps,
+ # Information for int8
+ int8_min,
+ int8_max,
+ # Meta-parameters
+ BLOCK: tl.constexpr,
+ s_num : tl.constexpr,
+):
+ """A Triton-accelerated function to perform per-token-group
+ quantization on a tensor.
+
+ This function converts the tensor values into int8 values.
+ """
+ # Map the program id to the row of X and Y it should compute.
+ g_id = tl.program_id(0)
+ y_ptr += g_id * BLOCK
+ y_q_ptr += g_id * BLOCK
+ y_s_ptr += g_id * s_num
+
+ cols = tl.arange(0, BLOCK) # N <= BLOCK
+ s_cols = tl.arange(0, s_num)
+ mask = g_id * BLOCK + cols < SIZE
+
+ y = tl.load(y_ptr + cols, mask=mask, other=0.0).to(tl.float32)
+ y = tl.reshape(y, (s_num, 128))
+ # Quant
+ _absmax = tl.maximum(tl.max(tl.abs(y), axis=1), eps)
+ y_s = (_absmax / int8_max).reshape(s_num, 1)
+ y_q = tl.clamp(y / y_s, int8_min, int8_max).to(y_q_ptr.dtype.element_ty)
+
+ y_q = tl.reshape(y_q, (s_num*128))
+ y_s = tl.reshape(y_s, (s_num))
+
+ tl.store(y_q_ptr + cols, y_q, mask=mask)
+ tl.store(y_s_ptr + s_cols, y_s.to(y_s_ptr.dtype.element_ty))
+
+
+def per_token_group_quant_int8(
+ x: torch.Tensor,
+ group_size: int,
+ eps: float = 1e-10,
+ dtype: torch.dtype = torch.int8,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+ """Function to perform per-token-group quantization on an input tensor `x`.
+
+ It converts the tensor values into signed int8 values and returns the
+ quantized tensor along with the scaling factor used for quantization.
+
+ Args:
+ x: The input tenosr with ndim >= 2.
+ group_size: The group size used for quantization.
+ eps: The minimum to avoid dividing zero.
+ dtype: The dype of output tensor. Note that only `torch.int8`
+ is supported for now.
+
+ Returns:
+ Tuple[torch.Tensor, torch.Tensor]: The quantized tensor and the
+ scaling factor for quantization.
+ """
+ assert (x.shape[-1] % group_size == 0
+ ), "the last dimension of `x` cannot be divisible by `group_size`"
+ #assert x.is_contiguous(), "`x` is not contiguous"
+
+ iinfo = torch.iinfo(dtype)
+ int8_max = iinfo.max
+ int8_min = iinfo.min
+
+ x_q = torch.empty_like(x, device=x.device, dtype=dtype)
+ N = group_size
+
+ m = x.shape[0]
+ if m<=16:
+ config={"BLOCK":128,"s_num":1,"num_warps":1,"num_stages":1}
+ elif m<=256:
+ config={"BLOCK":1024,"s_num":8,"num_warps":4,"num_stages":1}
+ else:
+ config={"BLOCK":2048,"s_num":16,"num_warps":4,"num_stages":2}
+
+ grid = lambda META: (
+ triton.cdiv(x.numel(), META['BLOCK']),
+ )
+
+ x_s = torch.empty(
+ x.shape[:-1] + (x.shape[-1] // group_size, ),
+ device=x.device,
+ dtype=torch.float32,
+ )
+
+ BLOCK = triton.next_power_of_2(N)
+ # heuristics for number of warps
+ num_warps = min(max(BLOCK // 256, 1), 8)
+ num_stages = 1
+ _per_token_group_quant_int8[grid](
+ x,
+ x_q,
+ x_s,
+ group_size,
+ # M,
+ # K,
+ # N,
+ x.numel(),
+ eps,
+ int8_min=int8_min,
+ int8_max=int8_max,
+ **config
+ )
+
+ return x_q, x_s
+
+
+@triton.jit
+def _w8a8_block_int8_matmul(
+ # Pointers to inputs and output
+ A,
+ B,
+ C,
+ As,
+ Bs,
+ # Shape for matmul
+ M,
+ N,
+ K,
+ # Block size for block-wise quantization
+ group_n,
+ group_k,
+ # Stride for inputs and output
+ stride_am,
+ stride_ak,
+ stride_bk,
+ stride_bn,
+ stride_cm,
+ stride_cn,
+ stride_As_m,
+ stride_As_k,
+ stride_Bs_k,
+ stride_Bs_n,
+ # Meta-parameters
+ BLOCK_SIZE_M: tl.constexpr,
+ BLOCK_SIZE_N: tl.constexpr,
+ BLOCK_SIZE_K: tl.constexpr,
+ GROUP_SIZE_M: tl.constexpr,
+):
+ """Triton-accelerated function used to perform linear operations (dot
+ product) on input tensors `A` and `B` with block-wise quantization,
+ and store the result in output tensor `C`.
+ """
+
+ pid = tl.program_id(axis=0)
+ num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
+ num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
+ num_pid_in_group = GROUP_SIZE_M * num_pid_n
+ group_id = pid // num_pid_in_group
+ first_pid_m = group_id * GROUP_SIZE_M
+ group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
+ pid_m = first_pid_m + (pid % group_size_m)
+ pid_n = (pid % num_pid_in_group) // group_size_m
+
+ offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
+ offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
+
+ # offs_bsn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
+ offs_bsn = pid_n * BLOCK_SIZE_N // group_n
+
+ offs_k = tl.arange(0, BLOCK_SIZE_K)
+ a_ptrs = A + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
+ b_ptrs = B + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
+
+ # a_ptrs = A + (offs_am[:, None] * stride_am)
+ # b_ptrs = B + (offs_bn[None, :] * stride_bn)
+
+ As_ptrs = As + offs_am * stride_As_m
+ # offs_bsn = offs_bn // group_n
+ Bs_ptrs = Bs + offs_bsn * stride_Bs_n
+
+ accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+ for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
+
+ k_start = k * BLOCK_SIZE_K
+ offs_ks = k_start // group_k
+ a_s = tl.load(As_ptrs + offs_ks * stride_As_k)
+ b_s = tl.load(Bs_ptrs + offs_ks * stride_Bs_k)
+
+ a = tl.load(a_ptrs,
+ mask=offs_k[None, :] < K - k * BLOCK_SIZE_K,
+ other=0.0)
+ b = tl.load(b_ptrs,
+ mask=offs_k[:, None] < K - k * BLOCK_SIZE_K,
+ other=0.0)
+
+
+ accumulator += tl.dot(a, b).to(tl.float32) * a_s[:, None] * b_s[None, :]
+
+ a_ptrs += BLOCK_SIZE_K * stride_ak
+ b_ptrs += BLOCK_SIZE_K * stride_bk
+
+
+ if C.dtype.element_ty == tl.bfloat16:
+ c = accumulator.to(tl.bfloat16)
+ elif C.dtype.element_ty == tl.float16:
+ c = accumulator.to(tl.float16)
+ else:
+ c = accumulator.to(tl.float32)
+
+ offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
+ offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+ c_ptrs = C + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
+ c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
+ tl.store(c_ptrs, c, mask=c_mask)
+
+
+def matmul_block_int8(
+ A: torch.Tensor,
+ B: torch.Tensor,
+ As: torch.Tensor,
+ Bs: torch.Tensor,
+ block_size: List[int],
+ output_dtype: torch.dtype = torch.float16,
+ bias:Optional[torch.Tensor]=None,
+ config:Optional[dict] = None
+) -> torch.Tensor:
+ """This function performs matrix multiplication with block-wise
+ quantization.
+
+ It takes two input tensors `A` and `B` with scales `As` and `Bs`.
+ The output is returned in the specified `output_dtype`.
+
+ Args:
+ A: The input tensor, e.g., activation.
+ B: The input tensor, e.g., weight.
+ As: The per-token-group quantization scale for `A`.
+ Bs: The per-block quantization scale for `B`.
+ block_size: The block size for per-block quantization. It should be
+ 2-dim, e.g., [128, 128].
+ output_dytpe: The dtype of the returned tensor.
+
+ Returns:
+ torch.Tensor: The result of matmul.
+ """
+
+ block_n, block_k = block_size[0], block_size[1]
+ M = A.numel() // A.shape[-1]
+ N, K = B.shape
+
+ C_shape = A.shape[:-1] + (N, )
+ C = A.new_empty(C_shape, dtype=output_dtype)
+
+ if config is None:
+ if M<=64:
+ config = {
+ "BLOCK_SIZE_M": 16, #64
+ "BLOCK_SIZE_N": 64,
+ "BLOCK_SIZE_K": 128,
+ "GROUP_SIZE_M": 2,
+ "num_warps": 4,
+ "num_stages": 0,
+ }
+ elif M<128:
+ config = {
+ "BLOCK_SIZE_M": 32, #64
+ "BLOCK_SIZE_N": 64,
+ "BLOCK_SIZE_K": 128,
+ "GROUP_SIZE_M": 2,
+ "num_warps": 4,
+ "num_stages": 0,
+ }
+ elif M<=256:
+ config = {
+ "BLOCK_SIZE_M": 64, #64
+ "BLOCK_SIZE_N": 64,
+ "BLOCK_SIZE_K": 128,
+ "GROUP_SIZE_M": 2,
+ "num_warps": 4,
+ "num_stages": 0,
+ }
+ else :
+ config = {
+ "BLOCK_SIZE_M": 64, #64
+ "BLOCK_SIZE_N": 128,
+ "BLOCK_SIZE_K": 128,
+ "GROUP_SIZE_M": 8,
+ "num_warps": 8,
+ "num_stages": 0,
+ }
+
+ def grid(META):
+ return (triton.cdiv(M, META["BLOCK_SIZE_M"]) *
+ triton.cdiv(N, META["BLOCK_SIZE_N"]), )
+
+
+ _w8a8_block_int8_matmul[grid](
+ A,
+ B,
+ C,
+ As,
+ Bs,
+ M,
+ N,
+ K,
+ block_n,
+ block_k,
+ A.stride(-2),
+ A.stride(-1),
+ B.stride(1),
+ B.stride(0),
+ C.stride(-2),
+ C.stride(-1),
+ As.stride(-2),
+ As.stride(-1),
+ Bs.stride(1),
+ Bs.stride(0),
+ **config,
+ )
+
+ if bias is not None:
+ C = C + bias
+
+ return C
+
+def apply_w8a8_block_int8_linear(
+ input: torch.Tensor,
+ weight: torch.Tensor,
+ block_size: List[int],
+ weight_scale: torch.Tensor,
+ input_scale: Optional[torch.Tensor] = None,
+ bias: Optional[torch.Tensor] = None,
+ config: Optional[dict] = None,
+) -> torch.Tensor:
+ assert input_scale is None
+ # View input as 2D matrix for fp8 methods
+ input_2d = input.view(-1, input.shape[-1])
+ output_shape = [*input.shape[:-1], weight.shape[0]]
+
+ q_input, x_scale = per_token_group_quant_int8(input_2d, block_size[1])
+
+ output = matmul_block_int8(
+ q_input, weight, x_scale, weight_scale, block_size,
+ output_dtype=input.dtype,config=config,
+ )
+
+ if bias is not None:
+ output = output + bias
+ return output.to(dtype=input.dtype).view(*output_shape)
+
+
+def input_to_int8(
+ x: torch.Tensor, dtype: torch.dtype = torch.int8
+) -> Tuple[torch.Tensor, torch.Tensor]:
+ """This function quantizes input values to
+ int8 values with tensor-wise quantization.
+ """
+ iinfo = torch.iinfo(dtype)
+ min_val, max_val = x.aminmax()
+ amax = torch.maximum(min_val.abs(), max_val.abs()).clamp(min=1e-12)
+ int8_min, int8_max = iinfo.min, iinfo.max
+ scale = int8_max / amax
+ x_scl_sat = (x * scale).clamp(min=int8_min, max=int8_max)
+ return x_scl_sat.to(dtype).contiguous(), scale.float().reciprocal()
+
+
+def block_dequant(
+ x_q_block: torch.Tensor,
+ x_s: torch.Tensor,
+ block_size: List[int],
+) -> torch.Tensor:
+ """This function conducts block-wise dequantization.
+ The inputs are block-wise quantization tensor `x_q_block`,
+ block-wise quantization scale and the block size.
+ The outputs are dequantized tensor.
+ """
+ block_n, block_k = block_size[0], block_size[1]
+ n, k = x_q_block.shape
+ n_tiles = (n + block_n - 1) // block_n
+ k_tiles = (k + block_k - 1) // block_k
+ assert n_tiles == x_s.shape[0]
+ assert k_tiles == x_s.shape[1]
+
+ x_dq_block = x_q_block.to(torch.float32)
+
+ for i in range(k_tiles):
+ for j in range(n_tiles):
+ x_dq_block[
+ j * block_n : min((j + 1) * block_n, n),
+ i * block_k : min((i + 1) * block_k, k),
+ ] *= x_s[j][i]
+
+ return x_dq_block
+
+@triton.jit
+def matmul_kernel(
+ # Pointers to matrices
+ a_ptr, as_ptr, b_ptr, bs_ptr, c_ptr,
+ # Matrix dimensions
+ M, N, K,
+ # The stride variables represent how much to increase the ptr by when moving by 1
+ # element in a particular dimension. E.g. `stride_am` is how much to increase `a_ptr`
+ # by to get the element one row down (A has M rows).
+ stride_am, stride_ak,
+ stride_asm,
+ stride_bk, stride_bn,
+ stride_bsn,
+ stride_cm, stride_cn,
+ # Meta-parameters
+ BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
+ GROUP_SIZE_M: tl.constexpr, SPLIT_K: tl.constexpr,
+):
+ """Kernel for computing the matmul C = A x B.
+ A has shape (M, K), B has shape (K, N) and C has shape (M, N)
+ """
+ # -----------------------------------------------------------
+ # Map program ids `pid` to the block of C it should compute.
+ # This is done in a grouped ordering to promote L2 data reuse.
+ # See above `L2 Cache Optimizations` section for details.
+ pid = tl.program_id(axis=0)
+ pid_sp_k = tl.program_id(axis=1)
+ num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
+ num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
+ num_pid_in_group = GROUP_SIZE_M * num_pid_n
+ group_id = pid // num_pid_in_group
+ first_pid_m = group_id * GROUP_SIZE_M
+ group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
+ pid_m = first_pid_m + (pid % group_size_m)
+ pid_n = (pid % num_pid_in_group) // group_size_m
+
+ # ----------------------------------------------------------
+ # Create pointers for the first blocks of A and B.
+ # We will advance this pointer as we move in the K direction
+ # and accumulate
+ # `a_ptrs` is a block of [BLOCK_SIZE_M, BLOCK_SIZE_K] pointers
+ # `b_ptrs` is a block of [BLOCK_SIZE_K, BLOCK_SIZE_N] pointers
+ # See above `Pointer Arithmetics` section for details
+ offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
+ offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
+ offs_k = pid_sp_k * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
+ a_ptrs = a_ptr + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
+ b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
+ as_ptrs = as_ptr + offs_am * stride_asm
+ bs_ptrs = bs_ptr + offs_bn * stride_bsn
+ a_scale = tl.load(as_ptrs, mask=offs_am < M, other=0.0)
+ b_scale = tl.load(bs_ptrs, mask=offs_bn < N, other=0.0)
+ # -----------------------------------------------------------
+ # Iterate to compute a block of the C matrix.
+ # We accumulate into a `[BLOCK_SIZE_M, BLOCK_SIZE_N]` block
+ # of fp32 values for higher accuracy.
+ # `accumulator` will be converted back to fp16 after the loop.
+ accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.int32)
+ for k in range(0, tl.cdiv(K, BLOCK_SIZE_K * SPLIT_K)):
+ # Load the next block of A and B, generate a mask by checking the K dimension.
+ # If it is out of bounds, set it to 0.
+ a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K * SPLIT_K, other=0.0)
+ b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K * SPLIT_K, other=0.0)
+ # We accumulate along the K dimension.
+ accumulator += tl.dot(a, b)
+ # Advance the ptrs to the next K block.
+ a_ptrs += BLOCK_SIZE_K * SPLIT_K * stride_ak
+ b_ptrs += BLOCK_SIZE_K * SPLIT_K * stride_bk
+ # You can fuse arbitrary activation functions here
+ # while the accumulator is still in FP32!
+ if SPLIT_K == 1:
+ c = (accumulator.to(tl.float32) * a_scale[:, None] * b_scale[None, :]).to(c_ptr.dtype.element_ty)
+ else:
+ c = (accumulator.to(tl.float32) * a_scale[:, None] * b_scale[None, :])#.to(c_ptr.dtype.element_ty)
+ # -----------------------------------------------------------
+ # Write back the block of the output matrix C with masks.
+ offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
+ offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+ c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
+ c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
+ if SPLIT_K == 1:
+ tl.store(c_ptrs, c, mask=c_mask)
+ else:
+ tl.atomic_add(c_ptrs, c, mask=c_mask)
+
+def matmul_int8(a, a_scale, b, b_scale,out_dtype,config):
+ # Check constraints.
+ assert a.shape[1] == b.shape[0], "Incompatible dimensions"
+ M, K = a.shape
+ K, N = b.shape
+
+ if config is None:
+ #print("no config")
+ if M<=32:
+ config= {'BLOCK_SIZE_M': 16, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 256, 'GROUP_SIZE_M': 4,'SPLIT_K': 1,'num_stages':0, 'num_warps':4}
+ elif M<=64:
+ config= {'BLOCK_SIZE_M': 32, 'BLOCK_SIZE_N': 32, 'BLOCK_SIZE_K': 256, 'GROUP_SIZE_M': 4,'SPLIT_K': 1,'num_stages':0, 'num_warps':4}
+ elif M<=128:
+ config= {'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_N': 64, 'BLOCK_SIZE_K': 256, 'GROUP_SIZE_M': 4,'SPLIT_K': 1,'num_stages':0, 'num_warps':4}
+ elif M<=1024:
+ config= {'BLOCK_SIZE_M': 128, 'BLOCK_SIZE_N': 128, 'BLOCK_SIZE_K': 256, 'GROUP_SIZE_M': 4,'SPLIT_K': 1,'num_stages':0, 'num_warps':8}
+ else:
+ config= {'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_N': 256, 'BLOCK_SIZE_K': 64, 'GROUP_SIZE_M': 8,'SPLIT_K': 1,'num_stages':0, 'num_warps':8}
+ if config['SPLIT_K']==1:
+ c = torch.zeros((M, N), device=a.device, dtype=out_dtype)
+ else:
+ c = torch.zeros((M, N), device=a.device, dtype=torch.float32)
+
+ grid = lambda META: (
+ triton.cdiv(M, META['BLOCK_SIZE_M']) * triton.cdiv(N, META['BLOCK_SIZE_N']),
+ META['SPLIT_K'],
+ )
+ assert a.shape[1] == b.shape[0], "Incompatible dimensions"
+
+
+ matmul_kernel[grid](
+ a, a_scale, b, b_scale, c,
+ M, N, K,
+ a.stride(0), a.stride(1),
+ a_scale.stride(0),
+ b.stride(0), b.stride(1),
+ b_scale.stride(0),
+ c.stride(0), c.stride(1),
+ **config
+ )
+
+ if config['SPLIT_K']!=1:
+ c=c.to(out_dtype)
+ return c
+
+
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m4096_graph.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m4096_graph.json
new file mode 100644
index 00000000..ccbf2f8d
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m4096_graph.json
@@ -0,0 +1,662 @@
+[
+ {
+ "shape_id": "hy3_tp8_qkv_proj_m4096",
+ "operator": "qkv_proj",
+ "M": 4096,
+ "N": 1280,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 19381.382751464844,
+ "p90_us": 20706.951904296875,
+ "min_us": 15527.166748046875,
+ "max_us": 33025.71716308594,
+ "samples_us": [
+ 33025.717,
+ 19930.211,
+ 15664.891,
+ 17588.31,
+ 18619.633,
+ 18823.849,
+ 21526.091,
+ 15859.666,
+ 19238.075,
+ 20260.184,
+ 15527.167,
+ 19733.803,
+ 19524.69,
+ 19786.154,
+ 20706.952,
+ 17930.505,
+ 20039.809,
+ 19157.535,
+ 16103.629,
+ 19555.281
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_o_proj_m4096",
+ "operator": "o_proj",
+ "M": 4096,
+ "N": 4096,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 14151.79214477539,
+ "p90_us": 15978.352355957031,
+ "min_us": 10805.253601074219,
+ "max_us": 17007.850646972656,
+ "samples_us": [
+ 14512.738,
+ 17007.851,
+ 11326.995,
+ 10805.254,
+ 13926.711,
+ 14670.845,
+ 14319.626,
+ 14161.777,
+ 13924.182,
+ 12676.867,
+ 13472.328,
+ 11711.237,
+ 14111.505,
+ 14375.624,
+ 14141.808,
+ 16150.25,
+ 14195.662,
+ 14111.92,
+ 15978.352,
+ 14732.059
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_gate_up_proj_m4096",
+ "operator": "shared_gate_up_proj",
+ "M": 4096,
+ "N": 384,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 16712.230682373047,
+ "p90_us": 17514.361572265625,
+ "min_us": 15547.166442871094,
+ "max_us": 18294.686889648438,
+ "samples_us": [
+ 16706.294,
+ 18294.687,
+ 16718.167,
+ 17544.826,
+ 17514.362,
+ 16727.061,
+ 15547.166,
+ 16738.069,
+ 16576.123,
+ 16353.539,
+ 16856.978,
+ 16953.551,
+ 16529.98,
+ 16421.12,
+ 16027.951,
+ 16203.271,
+ 16812.18,
+ 16739.862,
+ 16485.535,
+ 16286.533
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_down_proj_m4096",
+ "operator": "shared_down_proj",
+ "M": 4096,
+ "N": 4096,
+ "K": 192,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 10762.615585327148,
+ "p90_us": 12987.09716796875,
+ "min_us": 9852.87094116211,
+ "max_us": 13697.439575195312,
+ "samples_us": [
+ 13056.47,
+ 10556.271,
+ 11678.119,
+ 9955.428,
+ 10320.246,
+ 10068.512,
+ 13697.44,
+ 11907.071,
+ 10484.849,
+ 10555.374,
+ 10817.062,
+ 12246.387,
+ 10072.864,
+ 10708.17,
+ 11191.416,
+ 12355.088,
+ 12987.097,
+ 10983.04,
+ 9852.871,
+ 9935.653
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_m4096",
+ "operator": "qkv_proj",
+ "M": 4096,
+ "N": 1280,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 41000.61798095703,
+ "p90_us": 41881.70166015625,
+ "min_us": 40260.833740234375,
+ "max_us": 42669.35729980469,
+ "samples_us": [
+ 40708.112,
+ 41134.851,
+ 41381.177,
+ 41881.702,
+ 41801.389,
+ 42043.939,
+ 42669.357,
+ 40753.168,
+ 40260.834,
+ 41577.524,
+ 40839.377,
+ 40447.772,
+ 40937.708,
+ 41451.099,
+ 40329.794,
+ 40462.268,
+ 41187.875,
+ 40845.648,
+ 40568.027,
+ 41063.528
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_and_indexer_qk_m4096",
+ "operator": "qkv_proj_and_indexer_qk",
+ "M": 4096,
+ "N": 1536,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 35604.698181152344,
+ "p90_us": 45884.57946777344,
+ "min_us": 27749.08447265625,
+ "max_us": 46322.772216796875,
+ "samples_us": [
+ 46322.772,
+ 45884.579,
+ 45077.631,
+ 45431.604,
+ 45714.41,
+ 44870.535,
+ 45646.893,
+ 46203.641,
+ 36398.19,
+ 27957.782,
+ 32282.748,
+ 31533.688,
+ 35115.964,
+ 29992.844,
+ 32868.619,
+ 27749.084,
+ 32278.049,
+ 31587.637,
+ 36093.433,
+ 28834.07
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_o_proj_m4096",
+ "operator": "o_proj",
+ "M": 4096,
+ "N": 6144,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 19897.51739501953,
+ "p90_us": 20350.270080566406,
+ "min_us": 16113.107299804688,
+ "max_us": 22143.16864013672,
+ "samples_us": [
+ 18625.305,
+ 16587.906,
+ 19008.813,
+ 19965.579,
+ 19805.585,
+ 20406.717,
+ 20196.259,
+ 20116.774,
+ 22143.169,
+ 19389.856,
+ 17160.941,
+ 20140.356,
+ 19829.456,
+ 20350.27,
+ 20321.152,
+ 19153.064,
+ 16113.107,
+ 18237.416,
+ 20135.303,
+ 20016.586
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_gate_up_proj_m4096",
+ "operator": "shared_gate_up_proj",
+ "M": 4096,
+ "N": 768,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 31882.257080078125,
+ "p90_us": 32810.479736328125,
+ "min_us": 28815.545654296875,
+ "max_us": 33757.51953125,
+ "samples_us": [
+ 28815.546,
+ 32810.48,
+ 33757.52,
+ 33088.87,
+ 31372.415,
+ 32189.896,
+ 31768.692,
+ 31452.863,
+ 31608.185,
+ 32180.679,
+ 31995.822,
+ 31029.648,
+ 32477.213,
+ 32050.54,
+ 31519.229,
+ 31712.759,
+ 32021.133,
+ 31480.319,
+ 31442.175,
+ 32122.824
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_down_proj_m4096",
+ "operator": "shared_down_proj",
+ "M": 4096,
+ "N": 6144,
+ "K": 384,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 22389.16244506836,
+ "p90_us": 23217.950439453125,
+ "min_us": 19888.59405517578,
+ "max_us": 26559.368896484375,
+ "samples_us": [
+ 26559.369,
+ 23217.95,
+ 23532.243,
+ 22366.364,
+ 22783.244,
+ 22659.154,
+ 22893.546,
+ 20970.54,
+ 22374.107,
+ 21844.141,
+ 21321.919,
+ 21402.269,
+ 22404.218,
+ 19888.594,
+ 23014.03,
+ 22565.503,
+ 20996.107,
+ 22041.0,
+ 22622.835,
+ 20145.161
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_fused_qkv_a_proj_m4096",
+ "operator": "fused_qkv_a_proj",
+ "M": 4096,
+ "N": 2624,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 68661.41662597656,
+ "p90_us": 70210.54077148438,
+ "min_us": 47223.052978515625,
+ "max_us": 70451.49536132812,
+ "samples_us": [
+ 69346.32,
+ 68621.863,
+ 68886.328,
+ 70219.086,
+ 68263.092,
+ 70210.541,
+ 70451.495,
+ 68671.973,
+ 69810.4,
+ 69454.895,
+ 69113.562,
+ 68650.861,
+ 69821.857,
+ 58982.715,
+ 47223.053,
+ 49626.874,
+ 51567.096,
+ 52752.594,
+ 53318.494,
+ 49815.668
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_q_b_proj_m4096",
+ "operator": "q_b_proj",
+ "M": 4096,
+ "N": 2048,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 14385.5712890625,
+ "p90_us": 15009.501647949219,
+ "min_us": 11206.52847290039,
+ "max_us": 16653.09295654297,
+ "samples_us": [
+ 14075.645,
+ 12226.59,
+ 11618.449,
+ 14454.031,
+ 14636.586,
+ 14437.842,
+ 14322.932,
+ 14592.075,
+ 16653.093,
+ 15187.801,
+ 11206.528,
+ 14112.669,
+ 14517.839,
+ 15009.502,
+ 14548.204,
+ 12725.101,
+ 11365.851,
+ 13823.718,
+ 14589.165,
+ 14333.301
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_kv_b_proj_m4096",
+ "operator": "kv_b_proj",
+ "M": 4096,
+ "N": 3584,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 6939.938735961914,
+ "p90_us": 7771.638488769531,
+ "min_us": 5538.338470458984,
+ "max_us": 8975.7568359375,
+ "samples_us": [
+ 6990.929,
+ 8975.757,
+ 8890.351,
+ 5578.849,
+ 5567.842,
+ 5553.794,
+ 7100.013,
+ 7378.083,
+ 7270.152,
+ 7218.505,
+ 6882.964,
+ 7771.638,
+ 6567.775,
+ 6888.949,
+ 5586.017,
+ 5567.042,
+ 5538.338,
+ 6640.636,
+ 7402.626,
+ 7460.033
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_o_proj_m4096",
+ "operator": "o_proj",
+ "M": 4096,
+ "N": 6144,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 54775.79650878906,
+ "p90_us": 56661.138916015625,
+ "min_us": 32334.09423828125,
+ "max_us": 57705.450439453125,
+ "samples_us": [
+ 36803.061,
+ 38816.177,
+ 32334.094,
+ 36348.23,
+ 36142.572,
+ 57260.956,
+ 56661.139,
+ 57705.45,
+ 54939.215,
+ 56308.258,
+ 55439.417,
+ 54385.504,
+ 55639.667,
+ 55714.325,
+ 55518.042,
+ 53696.216,
+ 54612.378,
+ 54527.069,
+ 54463.104,
+ 55575.513
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_gate_up_proj_m4096",
+ "operator": "shared_gate_up_proj",
+ "M": 4096,
+ "N": 512,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 21254.589080810547,
+ "p90_us": 21705.372619628906,
+ "min_us": 20394.827270507812,
+ "max_us": 23713.912963867188,
+ "samples_us": [
+ 20852.121,
+ 20782.909,
+ 20957.654,
+ 21016.469,
+ 21695.837,
+ 20957.654,
+ 20914.935,
+ 21597.601,
+ 22041.522,
+ 23713.913,
+ 21585.057,
+ 21181.616,
+ 21223.471,
+ 21285.707,
+ 21705.373,
+ 21578.625,
+ 21093.587,
+ 20394.827,
+ 21346.121,
+ 21464.134
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_down_proj_m4096",
+ "operator": "shared_down_proj",
+ "M": 4096,
+ "N": 6144,
+ "K": 256,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "graph": {
+ "median_us": 17700.262451171875,
+ "p90_us": 19782.623291015625,
+ "min_us": 15626.028442382812,
+ "max_us": 22184.974670410156,
+ "samples_us": [
+ 17591.882,
+ 19158.612,
+ 19004.921,
+ 16231.48,
+ 17128.282,
+ 19256.914,
+ 19557.191,
+ 19782.623,
+ 22184.975,
+ 19914.011,
+ 17808.643,
+ 18663.495,
+ 15845.126,
+ 17278.868,
+ 17352.115,
+ 17899.68,
+ 16849.059,
+ 15626.028,
+ 17149.53,
+ 17040.669
+ ]
+ }
+ }
+]
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_merged.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_merged.json
new file mode 100644
index 00000000..e5104ff8
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_merged.json
@@ -0,0 +1,197 @@
+[
+ {
+ "shape_id": "hy3_tp8_qkv_proj_m8",
+ "M": 8,
+ "N": 1280,
+ "K": 4096,
+ "operator": "qkv_proj",
+ "pass_medians_us": [
+ 66.432,
+ 66.272,
+ 67.216
+ ],
+ "published_median_us": 66.432
+ },
+ {
+ "shape_id": "hy3_tp8_o_proj_m8",
+ "M": 8,
+ "N": 4096,
+ "K": 1024,
+ "operator": "o_proj",
+ "pass_medians_us": [
+ 31.472,
+ 31.488,
+ 31.424
+ ],
+ "published_median_us": 31.472
+ },
+ {
+ "shape_id": "hy3_tp8_shared_gate_up_proj_m8",
+ "M": 8,
+ "N": 384,
+ "K": 4096,
+ "operator": "shared_gate_up_proj",
+ "pass_medians_us": [
+ 56.896,
+ 57.008,
+ 57.792
+ ],
+ "published_median_us": 57.008
+ },
+ {
+ "shape_id": "hy3_tp8_shared_down_proj_m8",
+ "M": 8,
+ "N": 4096,
+ "K": 192,
+ "operator": "shared_down_proj",
+ "pass_medians_us": [
+ 14.592,
+ 14.592,
+ 14.592
+ ],
+ "published_median_us": 14.592
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_m8",
+ "M": 8,
+ "N": 1280,
+ "K": 6144,
+ "operator": "qkv_proj",
+ "pass_medians_us": [
+ 98.401,
+ 98.865,
+ 99.36
+ ],
+ "published_median_us": 98.865
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_and_indexer_qk_m8",
+ "M": 8,
+ "N": 1536,
+ "K": 6144,
+ "operator": "qkv_proj_and_indexer_qk",
+ "pass_medians_us": [
+ 99.456,
+ 99.473,
+ 99.968
+ ],
+ "published_median_us": 99.473
+ },
+ {
+ "shape_id": "minimax_tp8_o_proj_m8",
+ "M": 8,
+ "N": 6144,
+ "K": 1024,
+ "operator": "o_proj",
+ "pass_medians_us": [
+ 32.96,
+ 33.008,
+ 32.768
+ ],
+ "published_median_us": 32.96
+ },
+ {
+ "shape_id": "minimax_tp8_shared_gate_up_proj_m8",
+ "M": 8,
+ "N": 768,
+ "K": 6144,
+ "operator": "shared_gate_up_proj",
+ "pass_medians_us": [
+ 92.544,
+ 93.936,
+ 92.928
+ ],
+ "published_median_us": 92.928
+ },
+ {
+ "shape_id": "minimax_tp8_shared_down_proj_m8",
+ "M": 8,
+ "N": 6144,
+ "K": 384,
+ "operator": "shared_down_proj",
+ "pass_medians_us": [
+ 19.968,
+ 33.76,
+ 19.984
+ ],
+ "published_median_us": 19.984
+ },
+ {
+ "shape_id": "glm52_tp8_fused_qkv_a_proj_m8",
+ "M": 8,
+ "N": 2624,
+ "K": 6144,
+ "operator": "fused_qkv_a_proj",
+ "pass_medians_us": [
+ 87.2,
+ 88.032,
+ 86.976
+ ],
+ "published_median_us": 87.2
+ },
+ {
+ "shape_id": "glm52_tp8_q_b_proj_m8",
+ "M": 8,
+ "N": 2048,
+ "K": 2048,
+ "operator": "q_b_proj",
+ "pass_medians_us": [
+ 37.92,
+ 36.704,
+ 35.792
+ ],
+ "published_median_us": 36.704
+ },
+ {
+ "shape_id": "glm52_tp8_kv_b_proj_m8",
+ "M": 8,
+ "N": 3584,
+ "K": 512,
+ "operator": "kv_b_proj",
+ "pass_medians_us": [
+ 16.0,
+ 15.808,
+ 15.744
+ ],
+ "published_median_us": 15.808
+ },
+ {
+ "shape_id": "glm52_tp8_o_proj_m8",
+ "M": 8,
+ "N": 6144,
+ "K": 2048,
+ "operator": "o_proj",
+ "pass_medians_us": [
+ 58.864,
+ 58.864,
+ 58.064
+ ],
+ "published_median_us": 58.864
+ },
+ {
+ "shape_id": "glm52_tp8_shared_gate_up_proj_m8",
+ "M": 8,
+ "N": 512,
+ "K": 6144,
+ "operator": "shared_gate_up_proj",
+ "pass_medians_us": [
+ 89.216,
+ 88.336,
+ 88.448
+ ],
+ "published_median_us": 88.448
+ },
+ {
+ "shape_id": "glm52_tp8_shared_down_proj_m8",
+ "M": 8,
+ "N": 6144,
+ "K": 256,
+ "operator": "shared_down_proj",
+ "pass_medians_us": [
+ 16.096,
+ 15.136,
+ 15.072
+ ],
+ "published_median_us": 15.136
+ }
+]
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run1.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run1.json
new file mode 100644
index 00000000..a5ec1aa7
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run1.json
@@ -0,0 +1,662 @@
+[
+ {
+ "shape_id": "hy3_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 66.4322018623352,
+ "p90_us": 67.2324001789093,
+ "min_us": 66.20839834213257,
+ "max_us": 67.4884021282196,
+ "samples_us": [
+ 67.488,
+ 67.232,
+ 67.2,
+ 66.56,
+ 66.432,
+ 66.368,
+ 66.432,
+ 66.304,
+ 66.368,
+ 67.456,
+ 66.304,
+ 66.208,
+ 66.464,
+ 66.272,
+ 66.272,
+ 66.432,
+ 66.336,
+ 66.368,
+ 66.464,
+ 67.232
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 31.47220015525818,
+ "p90_us": 31.584200263023376,
+ "min_us": 31.296199560165405,
+ "max_us": 32.41620063781738,
+ "samples_us": [
+ 32.416,
+ 31.36,
+ 31.52,
+ 31.552,
+ 31.456,
+ 31.584,
+ 31.616,
+ 31.392,
+ 31.456,
+ 31.52,
+ 31.392,
+ 31.36,
+ 31.552,
+ 31.296,
+ 31.488,
+ 31.584,
+ 31.424,
+ 31.584,
+ 31.392,
+ 31.328
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 384,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 56.89629912376404,
+ "p90_us": 57.69619941711426,
+ "min_us": 56.64039850234985,
+ "max_us": 57.95239806175232,
+ "samples_us": [
+ 56.672,
+ 57.696,
+ 56.864,
+ 56.96,
+ 56.8,
+ 57.184,
+ 56.704,
+ 56.928,
+ 57.728,
+ 57.056,
+ 56.8,
+ 56.992,
+ 56.704,
+ 57.12,
+ 56.64,
+ 56.96,
+ 56.704,
+ 56.864,
+ 56.768,
+ 57.952
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 192,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 14.592200517654419,
+ "p90_us": 14.688000082969666,
+ "min_us": 14.528200030326843,
+ "max_us": 17.952199280261993,
+ "samples_us": [
+ 14.656,
+ 14.528,
+ 14.624,
+ 14.592,
+ 14.592,
+ 14.56,
+ 14.624,
+ 14.592,
+ 14.624,
+ 14.688,
+ 14.528,
+ 14.592,
+ 14.656,
+ 14.528,
+ 14.56,
+ 14.592,
+ 14.592,
+ 14.688,
+ 17.952,
+ 14.624
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 98.40050041675568,
+ "p90_us": 98.68860244750977,
+ "min_us": 98.27240109443665,
+ "max_us": 99.52059984207153,
+ "samples_us": [
+ 98.497,
+ 98.336,
+ 99.361,
+ 98.337,
+ 98.337,
+ 98.401,
+ 98.401,
+ 98.337,
+ 98.465,
+ 98.272,
+ 98.432,
+ 98.689,
+ 98.273,
+ 98.273,
+ 98.561,
+ 98.529,
+ 98.305,
+ 99.521,
+ 98.336,
+ 98.4
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_and_indexer_qk_m8",
+ "operator": "qkv_proj_and_indexer_qk",
+ "M": 8,
+ "N": 1536,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 99.45639967918396,
+ "p90_us": 100.32060146331787,
+ "min_us": 99.2326021194458,
+ "max_us": 100.41660070419312,
+ "samples_us": [
+ 100.321,
+ 100.353,
+ 99.233,
+ 99.265,
+ 99.553,
+ 99.233,
+ 99.456,
+ 99.328,
+ 99.617,
+ 99.233,
+ 99.361,
+ 99.361,
+ 99.361,
+ 100.224,
+ 99.456,
+ 100.417,
+ 99.553,
+ 99.457,
+ 99.457,
+ 99.297
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 32.96020030975342,
+ "p90_us": 33.024200797080994,
+ "min_us": 32.864201068878174,
+ "max_us": 33.05619955062866,
+ "samples_us": [
+ 32.96,
+ 32.992,
+ 32.96,
+ 33.056,
+ 32.96,
+ 33.056,
+ 32.96,
+ 32.864,
+ 32.96,
+ 32.992,
+ 32.864,
+ 32.928,
+ 33.024,
+ 32.992,
+ 32.96,
+ 33.024,
+ 32.992,
+ 32.96,
+ 32.96,
+ 32.928
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 768,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 92.54440069198608,
+ "p90_us": 93.34459900856018,
+ "min_us": 92.06439852714539,
+ "max_us": 93.4723973274231,
+ "samples_us": [
+ 93.345,
+ 92.545,
+ 92.129,
+ 92.577,
+ 92.385,
+ 92.481,
+ 92.193,
+ 92.417,
+ 92.449,
+ 93.377,
+ 92.288,
+ 92.544,
+ 93.152,
+ 92.576,
+ 92.064,
+ 93.472,
+ 92.256,
+ 92.544,
+ 92.896,
+ 92.609
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 384,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 19.968199729919434,
+ "p90_us": 20.03200054168701,
+ "min_us": 19.904200732707977,
+ "max_us": 20.064200460910797,
+ "samples_us": [
+ 19.904,
+ 20.032,
+ 20.0,
+ 19.968,
+ 19.968,
+ 19.968,
+ 19.968,
+ 19.968,
+ 20.0,
+ 20.0,
+ 20.064,
+ 19.968,
+ 19.968,
+ 19.904,
+ 19.968,
+ 19.936,
+ 20.0,
+ 19.968,
+ 20.064,
+ 20.0
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_fused_qkv_a_proj_m8",
+ "operator": "fused_qkv_a_proj",
+ "M": 8,
+ "N": 2624,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 87.20049858093262,
+ "p90_us": 89.47259783744812,
+ "min_us": 86.17640137672424,
+ "max_us": 97.60059714317322,
+ "samples_us": [
+ 87.329,
+ 87.265,
+ 86.208,
+ 86.24,
+ 94.241,
+ 87.136,
+ 86.177,
+ 86.241,
+ 86.336,
+ 86.176,
+ 86.369,
+ 86.177,
+ 87.136,
+ 97.601,
+ 89.473,
+ 89.408,
+ 89.152,
+ 89.248,
+ 89.152,
+ 89.313
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_q_b_proj_m8",
+ "operator": "q_b_proj",
+ "M": 8,
+ "N": 2048,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 37.92020082473755,
+ "p90_us": 38.24020028114319,
+ "min_us": 36.80019974708557,
+ "max_us": 38.624200224876404,
+ "samples_us": [
+ 38.624,
+ 38.336,
+ 37.888,
+ 37.824,
+ 38.048,
+ 37.92,
+ 37.856,
+ 37.856,
+ 37.92,
+ 37.952,
+ 37.792,
+ 37.984,
+ 37.888,
+ 38.08,
+ 38.24,
+ 37.92,
+ 38.112,
+ 38.08,
+ 37.92,
+ 36.8
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_kv_b_proj_m8",
+ "operator": "kv_b_proj",
+ "M": 8,
+ "N": 3584,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 16.00010022521019,
+ "p90_us": 24.86419975757599,
+ "min_us": 15.71200042963028,
+ "max_us": 26.464200019836426,
+ "samples_us": [
+ 26.464,
+ 23.808,
+ 24.448,
+ 23.84,
+ 24.864,
+ 24.896,
+ 24.832,
+ 17.024,
+ 15.776,
+ 15.744,
+ 15.84,
+ 15.872,
+ 16.032,
+ 15.872,
+ 15.808,
+ 15.872,
+ 15.968,
+ 15.936,
+ 15.712,
+ 16.736
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 58.864399790763855,
+ "p90_us": 59.29639935493469,
+ "min_us": 57.79240131378174,
+ "max_us": 59.52039957046509,
+ "samples_us": [
+ 59.136,
+ 59.104,
+ 58.88,
+ 59.2,
+ 59.52,
+ 58.432,
+ 58.848,
+ 58.432,
+ 59.136,
+ 58.848,
+ 59.296,
+ 57.792,
+ 59.04,
+ 57.984,
+ 58.848,
+ 59.456,
+ 58.016,
+ 58.304,
+ 58.912,
+ 57.92
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 512,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 89.21639919281006,
+ "p90_us": 89.60059881210327,
+ "min_us": 88.22460174560547,
+ "max_us": 89.60059881210327,
+ "samples_us": [
+ 89.249,
+ 89.601,
+ 88.225,
+ 89.28,
+ 88.512,
+ 89.152,
+ 89.345,
+ 88.609,
+ 88.225,
+ 89.536,
+ 88.576,
+ 89.12,
+ 89.313,
+ 89.601,
+ 89.184,
+ 89.472,
+ 88.608,
+ 89.249,
+ 89.057,
+ 89.601
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 256,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 16.096200048923492,
+ "p90_us": 16.223999857902527,
+ "min_us": 15.072199702262878,
+ "max_us": 18.592000007629395,
+ "samples_us": [
+ 16.096,
+ 16.096,
+ 16.096,
+ 15.104,
+ 16.096,
+ 16.192,
+ 16.096,
+ 16.064,
+ 16.128,
+ 15.072,
+ 16.224,
+ 15.2,
+ 16.192,
+ 18.592,
+ 16.128,
+ 16.064,
+ 15.136,
+ 16.224,
+ 16.192,
+ 16.224
+ ]
+ }
+ }
+]
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run2.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run2.json
new file mode 100644
index 00000000..2970f303
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run2.json
@@ -0,0 +1,662 @@
+[
+ {
+ "shape_id": "hy3_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 66.27230048179626,
+ "p90_us": 67.26440191268921,
+ "min_us": 66.11220240592957,
+ "max_us": 74.72040057182312,
+ "samples_us": [
+ 67.264,
+ 66.24,
+ 66.272,
+ 66.304,
+ 66.24,
+ 74.72,
+ 66.336,
+ 67.328,
+ 66.112,
+ 66.368,
+ 66.208,
+ 66.24,
+ 66.208,
+ 66.112,
+ 66.304,
+ 67.168,
+ 66.112,
+ 66.304,
+ 66.272,
+ 66.24
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 31.488101184368134,
+ "p90_us": 31.58400058746338,
+ "min_us": 31.391999125480652,
+ "max_us": 32.41620063781738,
+ "samples_us": [
+ 32.416,
+ 31.456,
+ 31.424,
+ 31.488,
+ 31.584,
+ 31.392,
+ 31.552,
+ 31.424,
+ 31.488,
+ 31.424,
+ 31.616,
+ 31.52,
+ 31.552,
+ 31.424,
+ 31.488,
+ 31.52,
+ 31.552,
+ 31.456,
+ 31.456,
+ 31.456
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 384,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 57.00819790363312,
+ "p90_us": 57.856398820877075,
+ "min_us": 56.83220028877258,
+ "max_us": 57.98439979553223,
+ "samples_us": [
+ 57.984,
+ 57.856,
+ 56.832,
+ 57.152,
+ 57.024,
+ 57.92,
+ 56.992,
+ 56.896,
+ 56.928,
+ 56.992,
+ 56.832,
+ 57.088,
+ 56.96,
+ 57.088,
+ 56.928,
+ 57.12,
+ 56.96,
+ 57.056,
+ 56.896,
+ 57.088
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 192,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 14.591999351978302,
+ "p90_us": 14.655999839305878,
+ "min_us": 14.496199786663055,
+ "max_us": 15.96820056438446,
+ "samples_us": [
+ 14.624,
+ 14.56,
+ 14.592,
+ 14.592,
+ 14.56,
+ 14.592,
+ 14.56,
+ 15.968,
+ 14.56,
+ 14.496,
+ 14.592,
+ 14.56,
+ 14.592,
+ 14.592,
+ 14.56,
+ 14.656,
+ 14.592,
+ 14.592,
+ 14.592,
+ 14.656
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 98.8646000623703,
+ "p90_us": 99.64860081672668,
+ "min_us": 98.56060147285461,
+ "max_us": 100.03260374069214,
+ "samples_us": [
+ 99.617,
+ 99.745,
+ 98.561,
+ 99.649,
+ 98.913,
+ 98.656,
+ 98.624,
+ 98.976,
+ 98.72,
+ 98.593,
+ 100.033,
+ 98.625,
+ 98.721,
+ 98.817,
+ 99.52,
+ 98.72,
+ 98.752,
+ 99.552,
+ 99.616,
+ 99.041
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_and_indexer_qk_m8",
+ "operator": "qkv_proj_and_indexer_qk",
+ "M": 8,
+ "N": 1536,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 99.47250187397003,
+ "p90_us": 100.35239458084106,
+ "min_us": 99.29659962654114,
+ "max_us": 100.41639804840088,
+ "samples_us": [
+ 100.352,
+ 100.416,
+ 99.393,
+ 99.521,
+ 100.321,
+ 99.393,
+ 99.457,
+ 99.297,
+ 100.416,
+ 99.328,
+ 99.648,
+ 99.328,
+ 99.488,
+ 99.361,
+ 99.553,
+ 99.489,
+ 99.457,
+ 99.329,
+ 99.361,
+ 99.52
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 33.00819993019104,
+ "p90_us": 33.05619955062866,
+ "min_us": 32.92819857597351,
+ "max_us": 34.04819965362549,
+ "samples_us": [
+ 34.048,
+ 32.992,
+ 32.928,
+ 33.024,
+ 33.024,
+ 33.056,
+ 32.928,
+ 33.056,
+ 32.992,
+ 32.992,
+ 32.992,
+ 32.96,
+ 32.992,
+ 32.992,
+ 33.056,
+ 32.992,
+ 33.024,
+ 33.088,
+ 33.024,
+ 33.024
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 768,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 93.93649995326996,
+ "p90_us": 95.87259888648987,
+ "min_us": 92.83260107040405,
+ "max_us": 99.07259941101074,
+ "samples_us": [
+ 93.792,
+ 93.952,
+ 93.024,
+ 92.833,
+ 93.921,
+ 93.888,
+ 93.024,
+ 92.833,
+ 93.025,
+ 93.024,
+ 93.024,
+ 99.073,
+ 95.873,
+ 95.585,
+ 96.001,
+ 95.873,
+ 95.648,
+ 95.776,
+ 95.84,
+ 95.744
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 384,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 33.76020044088364,
+ "p90_us": 37.216201424598694,
+ "min_us": 21.312199532985687,
+ "max_us": 39.61620032787323,
+ "samples_us": [
+ 39.04,
+ 34.944,
+ 33.824,
+ 34.112,
+ 33.792,
+ 33.6,
+ 37.216,
+ 34.016,
+ 32.992,
+ 33.248,
+ 33.088,
+ 33.728,
+ 39.616,
+ 32.96,
+ 35.424,
+ 33.216,
+ 33.792,
+ 22.08,
+ 21.6,
+ 21.312
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_fused_qkv_a_proj_m8",
+ "operator": "fused_qkv_a_proj",
+ "M": 8,
+ "N": 2624,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 88.03249895572662,
+ "p90_us": 88.64039778709412,
+ "min_us": 87.07240223884583,
+ "max_us": 89.47259783744812,
+ "samples_us": [
+ 89.473,
+ 87.841,
+ 88.32,
+ 88.064,
+ 88.096,
+ 88.64,
+ 87.52,
+ 87.84,
+ 88.225,
+ 88.832,
+ 87.36,
+ 87.808,
+ 88.032,
+ 88.256,
+ 87.232,
+ 88.384,
+ 88.033,
+ 88.0,
+ 87.648,
+ 87.072
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_q_b_proj_m8",
+ "operator": "q_b_proj",
+ "M": 8,
+ "N": 2048,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 36.70420050621033,
+ "p90_us": 37.88819909095764,
+ "min_us": 35.84020137786865,
+ "max_us": 37.95219957828522,
+ "samples_us": [
+ 37.568,
+ 37.888,
+ 35.904,
+ 36.576,
+ 37.952,
+ 35.84,
+ 36.608,
+ 37.888,
+ 35.968,
+ 36.576,
+ 37.952,
+ 35.968,
+ 36.768,
+ 37.088,
+ 35.904,
+ 36.64,
+ 37.056,
+ 36.032,
+ 37.344,
+ 37.056
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_kv_b_proj_m8",
+ "operator": "kv_b_proj",
+ "M": 8,
+ "N": 3584,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 15.808099508285522,
+ "p90_us": 17.728200554847717,
+ "min_us": 15.71200042963028,
+ "max_us": 17.85600036382675,
+ "samples_us": [
+ 16.48,
+ 17.728,
+ 16.96,
+ 15.808,
+ 15.712,
+ 16.448,
+ 17.728,
+ 17.856,
+ 15.776,
+ 15.744,
+ 16.576,
+ 17.728,
+ 16.64,
+ 15.744,
+ 15.744,
+ 15.744,
+ 15.712,
+ 15.712,
+ 15.808,
+ 15.808
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 58.86429846286774,
+ "p90_us": 59.456199407577515,
+ "min_us": 57.888197898864746,
+ "max_us": 64.06440138816833,
+ "samples_us": [
+ 59.232,
+ 58.592,
+ 58.752,
+ 58.592,
+ 58.464,
+ 58.464,
+ 58.432,
+ 57.888,
+ 58.144,
+ 58.144,
+ 58.016,
+ 64.064,
+ 59.456,
+ 59.2,
+ 59.392,
+ 59.264,
+ 58.976,
+ 59.136,
+ 59.456,
+ 59.04
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 512,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 88.33639919757843,
+ "p90_us": 303.5536050796509,
+ "min_us": 87.90439963340759,
+ "max_us": 472.0344066619873,
+ "samples_us": [
+ 89.088,
+ 88.256,
+ 88.448,
+ 88.16,
+ 88.128,
+ 88.48,
+ 88.129,
+ 88.128,
+ 88.352,
+ 88.128,
+ 88.128,
+ 88.416,
+ 87.904,
+ 88.289,
+ 88.32,
+ 162.305,
+ 303.554,
+ 450.85,
+ 295.938,
+ 472.034
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 256,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 15.136100351810455,
+ "p90_us": 15.168200433254242,
+ "min_us": 15.07200002670288,
+ "max_us": 16.128000617027283,
+ "samples_us": [
+ 16.0,
+ 15.136,
+ 15.104,
+ 15.136,
+ 15.072,
+ 15.104,
+ 15.168,
+ 15.136,
+ 15.136,
+ 15.104,
+ 15.136,
+ 15.136,
+ 15.136,
+ 15.104,
+ 16.128,
+ 15.168,
+ 15.168,
+ 15.168,
+ 15.104,
+ 15.168
+ ]
+ }
+ }
+]
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run3.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run3.json
new file mode 100644
index 00000000..c7586f86
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/tp8_m8_graph_run3.json
@@ -0,0 +1,662 @@
+[
+ {
+ "shape_id": "hy3_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 67.21619963645935,
+ "p90_us": 67.36019849777222,
+ "min_us": 66.27219915390015,
+ "max_us": 67.42439866065979,
+ "samples_us": [
+ 67.072,
+ 67.104,
+ 67.072,
+ 67.232,
+ 67.36,
+ 67.168,
+ 67.264,
+ 67.232,
+ 67.424,
+ 67.04,
+ 67.296,
+ 66.272,
+ 67.136,
+ 67.232,
+ 67.168,
+ 67.392,
+ 66.272,
+ 67.264,
+ 67.2,
+ 67.296
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 31.424200534820557,
+ "p90_us": 31.488201022148132,
+ "min_us": 31.327998638153076,
+ "max_us": 32.28819966316223,
+ "samples_us": [
+ 32.288,
+ 31.424,
+ 31.488,
+ 31.424,
+ 31.584,
+ 31.424,
+ 31.424,
+ 31.488,
+ 31.392,
+ 31.36,
+ 31.456,
+ 31.424,
+ 31.456,
+ 31.424,
+ 31.392,
+ 31.328,
+ 31.392,
+ 31.456,
+ 31.328,
+ 31.424
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 384,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 57.79229998588562,
+ "p90_us": 57.95220136642456,
+ "min_us": 56.83220028877258,
+ "max_us": 57.95239806175232,
+ "samples_us": [
+ 57.792,
+ 57.856,
+ 56.832,
+ 56.928,
+ 57.792,
+ 57.888,
+ 57.792,
+ 57.952,
+ 57.824,
+ 57.952,
+ 57.632,
+ 57.888,
+ 57.728,
+ 57.792,
+ 57.76,
+ 57.888,
+ 57.728,
+ 57.952,
+ 57.792,
+ 56.96
+ ]
+ }
+ },
+ {
+ "shape_id": "hy3_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 4096,
+ "K": 192,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 14.591999351978302,
+ "p90_us": 14.62399959564209,
+ "min_us": 14.496000111103058,
+ "max_us": 14.624199271202087,
+ "samples_us": [
+ 14.56,
+ 14.528,
+ 14.624,
+ 14.496,
+ 14.56,
+ 14.592,
+ 14.56,
+ 14.592,
+ 14.592,
+ 14.56,
+ 14.592,
+ 14.528,
+ 14.528,
+ 14.592,
+ 14.624,
+ 14.592,
+ 14.624,
+ 14.624,
+ 14.56,
+ 14.592
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_m8",
+ "operator": "qkv_proj",
+ "M": 8,
+ "N": 1280,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 99.36040043830872,
+ "p90_us": 99.5523989200592,
+ "min_us": 98.43239784240723,
+ "max_us": 99.61640238761902,
+ "samples_us": [
+ 99.392,
+ 99.552,
+ 99.52,
+ 99.424,
+ 98.561,
+ 99.488,
+ 98.432,
+ 99.489,
+ 98.592,
+ 99.425,
+ 98.784,
+ 99.552,
+ 99.297,
+ 98.688,
+ 98.528,
+ 99.328,
+ 99.616,
+ 99.457,
+ 99.232,
+ 98.56
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_qkv_proj_and_indexer_qk_m8",
+ "operator": "qkv_proj_and_indexer_qk",
+ "M": 8,
+ "N": 1536,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 99.9684989452362,
+ "p90_us": 100.28860569000244,
+ "min_us": 99.16840195655823,
+ "max_us": 100.67239999771118,
+ "samples_us": [
+ 99.872,
+ 100.385,
+ 99.456,
+ 100.224,
+ 99.553,
+ 100.672,
+ 100.289,
+ 100.16,
+ 100.256,
+ 99.425,
+ 100.192,
+ 99.328,
+ 100.065,
+ 99.328,
+ 100.192,
+ 99.52,
+ 100.256,
+ 99.489,
+ 99.168,
+ 99.328
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 32.76819884777069,
+ "p90_us": 32.96020030975342,
+ "min_us": 32.70399868488312,
+ "max_us": 33.82419943809509,
+ "samples_us": [
+ 33.824,
+ 32.864,
+ 32.896,
+ 32.704,
+ 32.768,
+ 32.896,
+ 32.736,
+ 32.8,
+ 32.832,
+ 32.736,
+ 32.704,
+ 33.024,
+ 32.736,
+ 32.736,
+ 32.896,
+ 32.736,
+ 32.736,
+ 32.768,
+ 32.96,
+ 32.736
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 768,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 92.92849898338318,
+ "p90_us": 93.15239787101746,
+ "min_us": 92.00040102005005,
+ "max_us": 93.34459900856018,
+ "samples_us": [
+ 93.216,
+ 92.992,
+ 92.192,
+ 93.12,
+ 92.225,
+ 92.16,
+ 92.897,
+ 92.224,
+ 93.345,
+ 93.056,
+ 92.0,
+ 93.088,
+ 93.152,
+ 92.96,
+ 92.128,
+ 93.024,
+ 92.224,
+ 92.192,
+ 93.056,
+ 92.192
+ ]
+ }
+ },
+ {
+ "shape_id": "minimax_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 384,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 19.98410001397133,
+ "p90_us": 20.03220021724701,
+ "min_us": 19.904200732707977,
+ "max_us": 20.064200460910797,
+ "samples_us": [
+ 20.0,
+ 19.968,
+ 19.968,
+ 19.936,
+ 20.032,
+ 19.968,
+ 19.904,
+ 20.032,
+ 20.064,
+ 19.968,
+ 20.0,
+ 19.968,
+ 20.032,
+ 19.968,
+ 20.032,
+ 20.032,
+ 20.064,
+ 19.936,
+ 19.968,
+ 20.032
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_fused_qkv_a_proj_m8",
+ "operator": "fused_qkv_a_proj",
+ "M": 8,
+ "N": 2624,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 86.9763970375061,
+ "p90_us": 87.13639974594116,
+ "min_us": 86.01620197296143,
+ "max_us": 87.13639974594116,
+ "samples_us": [
+ 87.04,
+ 86.944,
+ 86.144,
+ 87.136,
+ 86.176,
+ 87.136,
+ 87.04,
+ 86.912,
+ 86.976,
+ 87.04,
+ 86.912,
+ 87.072,
+ 86.016,
+ 87.136,
+ 87.008,
+ 86.24,
+ 86.976,
+ 87.136,
+ 86.112,
+ 86.976
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_q_b_proj_m8",
+ "operator": "q_b_proj",
+ "M": 8,
+ "N": 2048,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 35.79210042953491,
+ "p90_us": 35.90399920940399,
+ "min_us": 35.74399948120117,
+ "max_us": 36.7682009935379,
+ "samples_us": [
+ 36.768,
+ 35.808,
+ 35.808,
+ 35.808,
+ 35.776,
+ 35.744,
+ 35.776,
+ 35.808,
+ 35.808,
+ 35.904,
+ 35.744,
+ 35.872,
+ 36.704,
+ 35.776,
+ 35.776,
+ 35.744,
+ 35.776,
+ 35.744,
+ 35.84,
+ 35.776
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_kv_b_proj_m8",
+ "operator": "kv_b_proj",
+ "M": 8,
+ "N": 3584,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 15.744000673294067,
+ "p90_us": 15.776200592517853,
+ "min_us": 15.680000185966492,
+ "max_us": 21.376200020313263,
+ "samples_us": [
+ 15.712,
+ 15.744,
+ 15.744,
+ 15.744,
+ 15.68,
+ 21.376,
+ 15.712,
+ 15.744,
+ 15.744,
+ 15.712,
+ 15.776,
+ 15.776,
+ 15.68,
+ 15.744,
+ 15.808,
+ 15.68,
+ 15.712,
+ 15.712,
+ 15.744,
+ 15.776
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_o_proj_m8",
+ "operator": "o_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 58.06429982185364,
+ "p90_us": 58.56040120124817,
+ "min_us": 57.53620266914368,
+ "max_us": 59.296202659606934,
+ "samples_us": [
+ 58.56,
+ 59.296,
+ 57.824,
+ 58.112,
+ 58.208,
+ 57.6,
+ 57.792,
+ 58.528,
+ 57.76,
+ 57.536,
+ 58.208,
+ 58.24,
+ 58.848,
+ 57.888,
+ 58.24,
+ 57.696,
+ 58.208,
+ 58.016,
+ 57.856,
+ 57.824
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_gate_up_proj_m8",
+ "operator": "shared_gate_up_proj",
+ "M": 8,
+ "N": 512,
+ "K": 6144,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 88.44839930534363,
+ "p90_us": 88.70440125465393,
+ "min_us": 87.52040266990662,
+ "max_us": 88.73639702796936,
+ "samples_us": [
+ 88.736,
+ 88.448,
+ 88.608,
+ 88.64,
+ 88.384,
+ 88.416,
+ 87.712,
+ 87.744,
+ 87.584,
+ 88.64,
+ 87.52,
+ 88.224,
+ 87.68,
+ 88.704,
+ 88.448,
+ 88.609,
+ 87.52,
+ 88.576,
+ 88.672,
+ 88.704
+ ]
+ }
+ },
+ {
+ "shape_id": "glm52_tp8_shared_down_proj_m8",
+ "operator": "shared_down_proj",
+ "M": 8,
+ "N": 6144,
+ "K": 256,
+ "config": {
+ "BLOCK_SIZE_M": 16,
+ "BLOCK_SIZE_N": 32,
+ "BLOCK_SIZE_K": 256,
+ "GROUP_SIZE_M": 4,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 4
+ },
+ "graph": {
+ "median_us": 15.07200002670288,
+ "p90_us": 15.136000514030457,
+ "min_us": 15.007999539375305,
+ "max_us": 15.136200189590454,
+ "samples_us": [
+ 15.136,
+ 15.072,
+ 15.04,
+ 15.072,
+ 15.008,
+ 15.008,
+ 15.072,
+ 15.072,
+ 15.072,
+ 15.04,
+ 15.136,
+ 15.104,
+ 15.136,
+ 15.104,
+ 15.072,
+ 15.072,
+ 15.104,
+ 15.104,
+ 15.104,
+ 15.072
+ ]
+ }
+ }
+]
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_eager.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_eager.json
new file mode 100644
index 00000000..48a063d7
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_eager.json
@@ -0,0 +1,281 @@
+{
+ "scope": "gemm_out",
+ "triton_source": "/tmp/baseline_int8_utils.py",
+ "M": 4096,
+ "warmups": 10,
+ "samples": 20,
+ "launches_per_sample": 5,
+ "mode": "eager",
+ "results": [
+ {
+ "shape_id": "tp4_wqkv_a_m4096",
+ "operator": "wqkv_a",
+ "M": 4096,
+ "N": 1536,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 13287.435150146484,
+ "p90_us": 13315.707397460938,
+ "min_us": 13170.518493652344,
+ "max_us": 13707.110595703125,
+ "samples_us": [
+ 13226.106,
+ 13308.22,
+ 13230.266,
+ 13292.188,
+ 13315.707,
+ 13190.199,
+ 13218.904,
+ 13296.187,
+ 13313.756,
+ 13294.81,
+ 13282.683,
+ 13296.956,
+ 13353.532,
+ 13170.518,
+ 13203.448,
+ 13707.111,
+ 13215.833,
+ 13229.721,
+ 13219.096,
+ 13292.923
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_wq_b_m4096",
+ "operator": "wq_b",
+ "M": 4096,
+ "N": 8192,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 20623.723602294922,
+ "p90_us": 20699.789428710938,
+ "min_us": 20575.30517578125,
+ "max_us": 21157.14569091797,
+ "samples_us": [
+ 20627.628,
+ 20606.154,
+ 20615.851,
+ 20721.645,
+ 20621.803,
+ 20619.339,
+ 20625.644,
+ 21157.146,
+ 20654.027,
+ 20667.628,
+ 20578.441,
+ 20637.163,
+ 20699.789,
+ 20575.305,
+ 20658.06,
+ 20611.241,
+ 20650.603,
+ 20599.562,
+ 20597.418,
+ 20594.986
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_indexer_wq_b_m4096",
+ "operator": "indexer.wq_b",
+ "M": 4096,
+ "N": 8192,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 20654.634857177734,
+ "p90_us": 20685.16387939453,
+ "min_us": 20580.072021484375,
+ "max_us": 21134.840393066406,
+ "samples_us": [
+ 20639.98,
+ 20622.218,
+ 20637.993,
+ 20662.732,
+ 20598.09,
+ 20663.66,
+ 20658.06,
+ 20632.042,
+ 20633.29,
+ 20651.21,
+ 20665.291,
+ 21134.84,
+ 20685.164,
+ 20580.072,
+ 20664.265,
+ 20590.665,
+ 20696.204,
+ 20669.418,
+ 20670.186,
+ 20642.058
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_wo_b_m4096",
+ "operator": "wo_b",
+ "M": 4096,
+ "N": 4096,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 19915.618896484375,
+ "p90_us": 19963.380432128906,
+ "min_us": 19849.97100830078,
+ "max_us": 20379.873657226562,
+ "samples_us": [
+ 19945.428,
+ 19849.971,
+ 20379.874,
+ 19900.435,
+ 19941.397,
+ 19923.093,
+ 19873.17,
+ 19899.539,
+ 19913.683,
+ 19905.396,
+ 19963.38,
+ 19887.508,
+ 19876.372,
+ 19917.555,
+ 19924.883,
+ 20352.672,
+ 19895.027,
+ 19882.227,
+ 19918.611,
+ 19960.02
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_shared_gate_up_proj_m4096",
+ "operator": "shared_gate_up_proj",
+ "M": 4096,
+ "N": 1024,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 8774.919891357422,
+ "p90_us": 8838.265228271484,
+ "min_us": 8717.046356201172,
+ "max_us": 9269.158172607422,
+ "samples_us": [
+ 8721.815,
+ 8792.313,
+ 8838.265,
+ 8745.751,
+ 8775.352,
+ 8793.529,
+ 8774.488,
+ 8747.095,
+ 8717.046,
+ 8845.21,
+ 8749.047,
+ 8780.152,
+ 8805.529,
+ 8770.008,
+ 8790.744,
+ 8806.841,
+ 8736.246,
+ 8773.912,
+ 9269.158,
+ 8771.128
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_shared_down_proj_m4096",
+ "operator": "shared_down_proj",
+ "M": 4096,
+ "N": 4096,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "eager": {
+ "median_us": 5582.3015213012695,
+ "p90_us": 5592.509841918945,
+ "min_us": 5555.356979370117,
+ "max_us": 5619.230270385742,
+ "samples_us": [
+ 5585.566,
+ 5592.51,
+ 5555.357,
+ 5586.398,
+ 5571.326,
+ 5581.853,
+ 5582.75,
+ 5586.045,
+ 5586.43,
+ 5582.973,
+ 5576.606,
+ 5571.358,
+ 5567.805,
+ 5563.549,
+ 5576.189,
+ 5619.23,
+ 5587.838,
+ 5566.749,
+ 5568.093,
+ 5595.134
+ ]
+ }
+ }
+ ]
+}
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_graph.json b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_graph.json
new file mode 100644
index 00000000..a6464631
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/baseline/triton_tp4_m4096_graph.json
@@ -0,0 +1,281 @@
+{
+ "scope": "gemm_out",
+ "triton_source": "/tmp/baseline_int8_utils.py",
+ "M": 4096,
+ "warmups": 10,
+ "samples": 20,
+ "launches_per_sample": 5,
+ "mode": "graph",
+ "results": [
+ {
+ "shape_id": "tp4_wqkv_a_m4096",
+ "operator": "wqkv_a",
+ "M": 4096,
+ "N": 1536,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 13247.252655029297,
+ "p90_us": 13320.054626464844,
+ "min_us": 13144.914245605469,
+ "max_us": 13643.904113769531,
+ "samples_us": [
+ 13336.472,
+ 13262.518,
+ 13314.262,
+ 13197.523,
+ 13320.055,
+ 13643.904,
+ 13300.246,
+ 13269.365,
+ 13284.662,
+ 13253.653,
+ 13234.421,
+ 13224.02,
+ 13218.292,
+ 13228.885,
+ 13147.346,
+ 13240.852,
+ 13214.676,
+ 13229.396,
+ 13144.914,
+ 13285.077
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_wq_b_m4096",
+ "operator": "wq_b",
+ "M": 4096,
+ "N": 8192,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 20590.225219726562,
+ "p90_us": 20648.818969726562,
+ "min_us": 20520.016479492188,
+ "max_us": 21081.182861328125,
+ "samples_us": [
+ 20596.434,
+ 21081.183,
+ 20520.016,
+ 20586.673,
+ 20577.713,
+ 20597.105,
+ 20556.367,
+ 20648.819,
+ 20589.33,
+ 20599.345,
+ 20590.61,
+ 20578.609,
+ 20576.752,
+ 20544.048,
+ 21078.076,
+ 20597.136,
+ 20589.841,
+ 20524.046,
+ 20621.681,
+ 20624.913
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_indexer_wq_b_m4096",
+ "operator": "indexer.wq_b",
+ "M": 4096,
+ "N": 8192,
+ "K": 1024,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 20602.671813964844,
+ "p90_us": 20697.555541992188,
+ "min_us": 20566.477966308594,
+ "max_us": 21020.53985595703,
+ "samples_us": [
+ 20609.616,
+ 20639.314,
+ 20574.065,
+ 20697.556,
+ 20587.505,
+ 21020.54,
+ 20601.36,
+ 20644.786,
+ 20570.096,
+ 20603.824,
+ 20603.375,
+ 20586.96,
+ 20610.159,
+ 20601.968,
+ 20566.478,
+ 20597.328,
+ 20613.904,
+ 20600.912,
+ 20981.018,
+ 20566.574
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_wo_b_m4096",
+ "operator": "wo_b",
+ "M": 4096,
+ "N": 4096,
+ "K": 2048,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 19881.948852539062,
+ "p90_us": 19916.25213623047,
+ "min_us": 19814.553833007812,
+ "max_us": 20356.455993652344,
+ "samples_us": [
+ 19870.557,
+ 19886.205,
+ 19901.534,
+ 19921.021,
+ 19859.677,
+ 19874.237,
+ 19911.804,
+ 19833.531,
+ 19829.051,
+ 20356.456,
+ 19863.837,
+ 19898.973,
+ 19877.693,
+ 19857.532,
+ 19838.747,
+ 19896.861,
+ 19889.693,
+ 19896.861,
+ 19814.554,
+ 19916.252
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_shared_gate_up_proj_m4096",
+ "operator": "shared_gate_up_proj",
+ "M": 4096,
+ "N": 1024,
+ "K": 4096,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 8790.192031860352,
+ "p90_us": 8815.552520751953,
+ "min_us": 8717.342376708984,
+ "max_us": 9243.083953857422,
+ "samples_us": [
+ 8815.553,
+ 8757.567,
+ 8770.112,
+ 9243.084,
+ 8726.078,
+ 8750.815,
+ 8791.904,
+ 8795.233,
+ 8717.342,
+ 8759.872,
+ 8832.194,
+ 8806.433,
+ 8799.168,
+ 8779.68,
+ 8723.199,
+ 8802.465,
+ 8798.4,
+ 8788.48,
+ 8806.753,
+ 8740.222
+ ]
+ }
+ },
+ {
+ "shape_id": "tp4_shared_down_proj_m4096",
+ "operator": "shared_down_proj",
+ "M": 4096,
+ "N": 4096,
+ "K": 512,
+ "config": {
+ "BLOCK_SIZE_M": 256,
+ "BLOCK_SIZE_N": 256,
+ "BLOCK_SIZE_K": 64,
+ "GROUP_SIZE_M": 8,
+ "SPLIT_K": 1,
+ "num_stages": 0,
+ "num_warps": 8
+ },
+ "errors": null,
+ "graph": {
+ "median_us": 5545.821189880371,
+ "p90_us": 5561.069488525391,
+ "min_us": 5528.492736816406,
+ "max_us": 6035.929489135742,
+ "samples_us": [
+ 5570.702,
+ 5552.142,
+ 5551.213,
+ 5546.317,
+ 5534.701,
+ 5537.453,
+ 5546.029,
+ 5541.453,
+ 5545.197,
+ 5557.133,
+ 5537.933,
+ 5528.493,
+ 5561.069,
+ 5544.845,
+ 5552.333,
+ 5551.566,
+ 5545.613,
+ 5533.453,
+ 5544.685,
+ 6035.929
+ ]
+ }
+ }
+ ]
+}
\ No newline at end of file
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/planner_parity.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/planner_parity.py
new file mode 100644
index 00000000..046255fd
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/planner_parity.py
@@ -0,0 +1,257 @@
+#!/usr/bin/env python3
+"""Offline parity analysis: planner (v0) vs the legacy round menu.
+
+Read-only tool over historical task repos under ``kernel-repos`` (and optionally
+workspace ``iterations``). For each recorded round it reconstructs the state
+(prior same-worker/same-shape records + pmc evidence), reproduces what the
+legacy menu (``w8a8_round_strategy``) would mandate, asks what the
+state-conditioned ``planner`` would pick, maps both to coarse plan families, and
+reports agreement/divergence by M regime and iteration bucket, plus samples.
+
+Caveats (read the output with these in mind):
+- State reconstruction is approximate: pmc evidence is frequently absent in
+ records, so the planner often falls back to its P4 legacy approximation.
+- The menu text -> family mapping is a keyword heuristic.
+- This tool never writes into the repos; it only prints / dumps a report.
+
+Usage:
+ python3 tools/planner_parity.py [--kernel-repos ROOT] [--limit-repos N]
+ [--out /tmp/planner_parity_report.json]
+"""
+
+from __future__ import annotations
+
+import argparse
+import collections
+import json
+import re
+import sys
+from pathlib import Path
+from typing import Any, Dict, List, Optional, Tuple
+
+# Make the plugin importable when run as a script from the plugin tree.
+sys.path.insert(0, str(Path(__file__).resolve().parents[2])) # .../dcu_kernel_auto_opt
+sys.path.insert(0, str(Path(__file__).resolve().parents[4])) # MetaInfer/
+
+from metainfer.tasks.dcu_kernel_auto_opt.orchestrator import planner as planner_mod # noqa: E402
+from metainfer.tasks.dcu_kernel_auto_opt.orchestrator.planner import choose_plan_from_history # noqa: E402
+from metainfer.tasks.dcu_kernel_auto_opt.orchestrator.prompts import w8a8_round_strategy # noqa: E402
+from metainfer.tasks.dcu_kernel_auto_opt.orchestrator.w8a8_pipeline import isa_round_policy # noqa: E402
+
+_MENU_FAMILY_RULES: List[Tuple[str, str]] = [
+ # Explicit markers first so generic words inside round prose do not win.
+ ("faster but incorrect", "repair"),
+ ("faster_wrong", "repair"),
+ ("repair", "repair"),
+ ("conditional inline-asm", "inline_asm"),
+ ("inline asm", "inline_asm"),
+ ("ISA-guided", "isa_guided"),
+ ("ISA", "isa"),
+ ("split-K", "grid_splitk"),
+ ("split_k", "grid_splitk"),
+ ("consolidat", "consolidate"),
+ ("epilogue", "epilogue"),
+ ("Architecture round", "architecture"),
+ ("Architecture/pipeline round", "architecture"),
+ ("macro-tile", "architecture"),
+ ("tile-shape", "architecture"),
+ ("geometry", "architecture"),
+ ("tile", "architecture"),
+ ("architecture", "architecture"),
+ ("packing", "memory_layout"),
+ ("packed", "memory_layout"),
+ ("swizzle", "memory_layout"),
+ ("bank", "memory_layout"),
+ ("layout", "memory_layout"),
+ ("double buff", "pipeline"),
+ ("prefetch", "pipeline"),
+ ("staging", "pipeline"),
+ ("barrier", "pipeline"),
+ ("pipeline", "pipeline"),
+ ("occupancy", "occupancy"),
+ ("waves per block", "occupancy"),
+ ("VGPR", "occupancy"),
+ ("register", "occupancy"),
+ ("bootstrap", "bootstrap"),
+ ("scalar", "bootstrap"),
+]
+
+_PLAN_TO_FAMILY = {
+ "repair_faster_wrong": "repair",
+ "fix_build": "repair",
+ "retry_same": "repair",
+ "bootstrap_correctness": "bootstrap",
+ "establish_arch": "architecture",
+ "architecture_explore": "architecture",
+ "grid_splitk": "grid_splitk",
+ "pipeline_tune": "pipeline",
+ "memory_layout": "memory_layout",
+ "occupancy_resource": "occupancy",
+ "epilogue_fusion": "epilogue",
+ "isa_guided_hip": "isa_guided",
+ "conditional_inline_asm": "inline_asm",
+ "consolidate": "consolidate",
+}
+
+
+def menu_family(text: str) -> str:
+ lowered = text.lower()
+ for needle, family in _MENU_FAMILY_RULES:
+ if needle in lowered:
+ return family
+ return "unknown"
+
+
+def iter_number_from_path(path: Path) -> int:
+ m = re.search(r"iteration(\d+)", path.name)
+ if m:
+ return int(m.group(1))
+ # parent dirs like .../iteration13/iteration.json
+ m = re.search(r"iteration(\d+)", path.parent.name)
+ return int(m.group(1)) if m else -1
+
+
+def load_records(root: Path, limit_repos: Optional[int]) -> List[Tuple[str, Dict[str, Any]]]:
+ """Return (repo_dir_name, record) for every candidate iteration.json."""
+ records: List[Tuple[str, Dict[str, Any]]] = []
+ repos = sorted(p for p in root.iterdir() if p.is_dir())
+ if limit_repos:
+ repos = repos[:limit_repos]
+ for repo in repos:
+ for path in sorted(
+ p for p in repo.rglob("iteration.json")
+ if "candidates" in p.parts
+ ):
+ try:
+ data = json.loads(path.read_text(encoding="utf-8"))
+ except (OSError, ValueError):
+ continue
+ if isinstance(data, dict):
+ records.append((repo.name, data))
+ return records
+
+
+def analyze(root: Path, limit_repos: Optional[int]) -> Dict[str, Any]:
+ records = load_records(root, limit_repos)
+ chains: Dict[Tuple[str, str, str], List[Dict[str, Any]]] = collections.defaultdict(list)
+ for repo, rec in records:
+ worker = str(rec.get("worker_id", "?"))
+ shape = rec.get("shape_id", "?")
+ iteration = iter_number_from_path(Path(str(rec.get("record_path", "iteration0")))) if False else int(rec.get("iteration", 0))
+ chains[(repo, worker, shape)].append(
+ {"record": rec, "iteration": iteration}
+ )
+
+ stats = collections.Counter()
+ buckets: Dict[str, collections.Counter] = collections.defaultdict(collections.Counter)
+ samples: List[Dict[str, Any]] = []
+
+ for (repo, worker, shape_id), entries in chains.items():
+ entries.sort(key=lambda e: e["iteration"])
+ history: List[Dict[str, Any]] = []
+ for entry in entries:
+ rec = entry["record"]
+ iteration = entry["iteration"]
+ shape = rec.get("shape", {}) or {}
+ m = int(shape.get("M", 0) or 0)
+ regime = "prefill(M>=128)" if m >= 128 else ("m16" if m >= 16 else "small(M<16)")
+ pmc = rec.get("pmc_evidence") or {}
+
+ try:
+ policy = isa_round_policy(
+ iteration=iteration,
+ max_iterations=10,
+ history=history,
+ )
+ menu_text = w8a8_round_strategy(
+ shape, iteration, history, pmc,
+ max_iterations=10, isa_policy=policy,
+ )
+ except Exception as exc: # noqa: BLE001
+ stats["menu_error"] += 1
+ history.append(rec)
+ continue
+
+ menu_fam = menu_family(menu_text)
+ try:
+ plan_id = choose_plan_from_history(
+ history, pmc, iteration=iteration, max_iterations=10,
+ shape=shape,
+ compiler_limitation_confirmed=(
+ policy.get("phase") == "conditional_inline_asm"
+ ),
+ )
+ except Exception as exc: # noqa: BLE001
+ stats["planner_error"] += 1
+ history.append(rec)
+ continue
+
+ plan_fam = _PLAN_TO_FAMILY.get(plan_id, "unknown")
+ agree = (menu_fam == plan_fam)
+ stats["total"] += 1
+ stats["agree" if agree else "disagree"] += 1
+ bucket_key = f"{regime}|iter{min(iteration, 12)}"
+ buckets[bucket_key]["agree" if agree else "disagree"] += 1
+
+ if not agree and len(samples) < 40:
+ samples.append({
+ "repo": repo,
+ "shape_id": shape_id,
+ "iteration": iteration,
+ "regime": regime,
+ "menu_family": menu_fam,
+ "plan_family": plan_fam,
+ "plan_id": plan_id,
+ "hypothesis": str(rec.get("hypothesis") or "")[:200],
+ })
+ history.append(rec)
+
+ total = stats.get("total", 0)
+ return {
+ "corpus": {"repos_scanned": len(set(r for r, _ in records)) if records else 0,
+ "records": len(records)},
+ "stats": dict(stats),
+ "agree_rate": (stats.get("agree", 0) / total) if total else None,
+ "buckets": {
+ k: dict(v) for k, v in sorted(buckets.items())
+ },
+ "samples": samples,
+ }
+
+
+def main() -> int:
+ parser = argparse.ArgumentParser(description=__doc__)
+ parser.add_argument("--kernel-repos", type=Path,
+ default=Path(__file__).resolve().parents[2].parent.parent.parent / "kernel-repos")
+ parser.add_argument("--limit-repos", type=int, default=None)
+ parser.add_argument("--out", type=Path, default=None)
+ args = parser.parse_args()
+
+ if not args.kernel_repos.is_dir():
+ print(f"kernel-repos not found at {args.kernel_repos}", file=sys.stderr)
+ return 2
+
+ report = analyze(args.kernel_repos, args.limit_repos)
+ total = report["stats"].get("total", 0)
+ print(f"corpus: {report['corpus']['records']} records / {report['corpus']['repos_scanned']} repos")
+ print(f"agree={report['stats'].get('agree',0)} disagree={report['stats'].get('disagree',0)} "
+ f"agree_rate={report['agree_rate']:.1%}" if report["agree_rate"] is not None
+ else f"agree={report['stats'].get('agree',0)} disagree={report['stats'].get('disagree',0)}")
+ print("--- buckets (regime|iteration): agree/disagree ---")
+ for k, v in list(report["buckets"].items())[:40]:
+ print(f" {k}: {v}")
+ print("--- divergence samples ---")
+ for s in report["samples"][:10]:
+ print(f" [{s['repo']} {s['shape_id']} it{s['iteration']} {s['regime']}] "
+ f"menu={s['menu_family']} plan={s['plan_family']} ({s['plan_id']})")
+ print(f" hypo: {s['hypothesis'][:150]}")
+
+ if args.out:
+ args.out.write_text(json.dumps(report, ensure_ascii=False, indent=2),
+ encoding="utf-8")
+ print(f"report -> {args.out}")
+ return 0
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/tools/retest_variant.py b/metainfer/tasks/dcu_kernel_auto_opt/tools/retest_variant.py
new file mode 100644
index 00000000..744f0207
--- /dev/null
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tools/retest_variant.py
@@ -0,0 +1,331 @@
+"""A/B retest tool: compare a candidate kernel against the shipped variant.
+
+Why this exists
+---------------
+DKAO records each candidate's median in its ``accepted//manifest.json``
+and each shipped variant's median in its header comment. Those are *records*,
+not measurements: if the round that produced them ran while another workload
+shared the device, the recorded number can be suppressed (we saw 884us -> 2043
+-> 3047us for one unchanged shape). Before promoting a kernel — especially one
+copied out of an old workspace — re-measure both sides under identical
+conditions.
+
+What "A/B" means here
+---------------------
+* **A** = the *candidate* kernel (what we might promote).
+* **B** = the *variant* currently shipped in DKAO (the incumbent).
+* Both are compiled from the same repository skeleton with only the kernel
+ source swapped, benchmarked on the **same idle device**, with the **same
+ sampling parameters**, and measured **alternately** (A, B, A, B...).
+
+Alternating matters: GPU clock/thermal state and any background load drift over
+time, so "measure A fully, then B fully" would mix that drift into the
+difference. Interleaving makes both sides absorb the same conditions, so the
+remaining difference is attributable to the kernel itself. Two rounds are run
+so a single fluke cannot decide the outcome; medians are compared and a
+candidate must win by ``--threshold`` percent (default 3, matching DKAO's
+promotion rule) *and* reproduce its own recorded median (``--recorded-tol``)
+before it is written.
+
+Usage
+-----
+ python -m metainfer.tasks.dcu_kernel_auto_opt.tools.retest_variant \
+ --candidate-workspace \
+ --shape hy3_tp4_o_proj_m4096 --rounds 2 [--promote]
+
+ # or straight from an accepted kernel file
+ python ...retest_variant --shape ... --candidate-kernel path/kernel.hip \
+ --kernel-root
+"""
+
+from __future__ import annotations
+
+import argparse
+import json
+import os
+import shutil
+import statistics
+import subprocess
+import sys
+from pathlib import Path
+from typing import Any, Dict, List, Optional, Tuple
+
+PLUGIN_DIR = Path(__file__).resolve().parents[1]
+REPO_ROOT = Path(__file__).resolve().parents[4]
+BENCH = PLUGIN_DIR / "assets" / "w8a8_bench.py"
+VARIANT_ROOT = PLUGIN_DIR / "variant" / "int8w8a8-gemm"
+
+#: DKAO compiles for gfx928 only. The int8 DUMMA fragment specialisations in
+#: /opt/dtk/include/du_mma.h sit behind ``#if defined(__gfx928__)``, so a
+#: default multi-arch torch build fails with "undefined template DUFragment
+#: <... signed char ...>". w8a8_pipeline sets this for the same reason.
+BENCH_ENV = {"PYTORCH_ROCM_ARCH": "gfx928"}
+
+
+# --------------------------------------------------------------- helpers ----
+
+def find_variant(model: str, tp: int, m: int, operator: str,
+ root: Optional[Path] = None) -> Optional[Path]:
+ """Shipped variant for a shape, tolerating model-label drift (glm -> glm52)."""
+ root = Path(root or VARIANT_ROOT)
+ exact = root / model / f"TP{tp}" / f"M{m}" / f"{operator}.hip"
+ if exact.is_file():
+ return exact
+ for model_dir in sorted(root.iterdir()) if root.is_dir() else []:
+ if not model_dir.name.startswith(model):
+ continue
+ candidate = model_dir / f"TP{tp}" / f"M{m}" / f"{operator}.hip"
+ if candidate.is_file():
+ return candidate
+ hits = sorted(root.glob(f"*/TP{tp}/M{m}/{operator}.hip")) if root.is_dir() else []
+ return hits[0] if hits else None
+
+
+def parse_variant_header_median(path: Path) -> Optional[float]:
+ """Median recorded in a variant's header comment (recorded, not measured)."""
+ try:
+ for line in path.read_text(encoding="utf-8",
+ errors="replace").splitlines()[:8]:
+ if "median_us=" in line:
+ return float(line.split("median_us=")[1].split()[0])
+ except (OSError, ValueError, IndexError):
+ return None
+ return None
+
+
+def bench_env(device: int, base: Optional[Dict[str, str]] = None) -> Dict[str, str]:
+ """Environment for an independent measurement on one device."""
+ env = dict(base if base is not None else os.environ)
+ env["HIP_VISIBLE_DEVICES"] = str(device)
+ env["CUDA_VISIBLE_DEVICES"] = str(device)
+ env.update(BENCH_ENV)
+ return env
+
+
+def decide(candidate_us: Optional[float], variant_us: Optional[float], *,
+ threshold_percent: float = 3.0,
+ recorded_us: Optional[float] = None,
+ recorded_tol_percent: float = 5.0) -> Dict[str, Any]:
+ """Promotion verdict from measured medians (+ reproducibility check)."""
+ if not candidate_us or not variant_us:
+ return {"verdict": "measurement_failed",
+ "reason": "one side produced no median"}
+ improvement = (variant_us - candidate_us) / variant_us * 100.0
+ recorded_delta = None
+ if recorded_us:
+ recorded_delta = (candidate_us - float(recorded_us)) / float(recorded_us) * 100.0
+ if improvement < float(threshold_percent):
+ return {"verdict": "keep", "improvement_percent": improvement,
+ "recorded_delta_percent": recorded_delta,
+ "reason": (f"improvement {improvement:.2f}% < required "
+ f"{float(threshold_percent):.2f}%")}
+ if recorded_delta is not None and abs(recorded_delta) > float(recorded_tol_percent):
+ return {"verdict": "needs_review", "improvement_percent": improvement,
+ "recorded_delta_percent": recorded_delta,
+ "reason": (f"win {improvement:.2f}% but the re-measurement "
+ f"differs from the record by {recorded_delta:.2f}%")}
+ return {"verdict": "promote", "improvement_percent": improvement,
+ "recorded_delta_percent": recorded_delta,
+ "reason": f"candidate faster by {improvement:.2f}%"}
+
+
+def accepted_kernel_for(workspace: Path, shape_id: str) -> Optional[Path]:
+ hits = sorted(Path(workspace).glob(
+ f"workers/*/accepted/{shape_id}/kernel.hip"))
+ return hits[0] if hits else None
+
+
+def recorded_median_for(kernel: Path) -> Optional[float]:
+ try:
+ manifest = json.loads((kernel.parent / "manifest.json")
+ .read_text(encoding="utf-8"))
+ except (OSError, ValueError):
+ return None
+ value = (manifest.get("metrics") or {}).get("median_us")
+ try:
+ return float(value)
+ except (TypeError, ValueError):
+ return None
+
+
+def idle_device(gpu_ids: Optional[List[int]] = None) -> Optional[int]:
+ """First device satisfying the operating gate (VRAM<=90 and HCU==0)."""
+ try:
+ sys.path.insert(0, str(REPO_ROOT))
+ from metainfer.tasks.dcu_kernel_auto_opt.orchestrator.gpu_preflight import ( # noqa: E501
+ preflight_gpus,
+ )
+ except Exception: # noqa: BLE001
+ return None
+ plan = preflight_gpus(gpu_ids or [0, 1, 2, 3], samples=3, interval_s=2)
+ clean = plan.get("clean_ids") or []
+ return int(clean[0]) if clean else None
+
+
+def prepare_repo(kernel_root: Path, kernel: Path, target: Path) -> Path:
+ """Copy a kernel repo skeleton and swap in one kernel source."""
+ if target.exists():
+ shutil.rmtree(target)
+ shutil.copytree(kernel_root, target)
+ (target / "csrc" / "w8a8_gemm_hip.hip").write_text(
+ Path(kernel).read_text(encoding="utf-8", errors="replace"),
+ encoding="utf-8")
+ return target
+
+
+def measure(repo: Path, shape: Tuple[int, int, int], device: int, *,
+ warmups: int = 30, samples: int = 30, replays: int = 50,
+ reference_cache: Optional[Path] = None,
+ timeout_s: int = 3600) -> Optional[float]:
+ m, n, k = shape
+ cmd = [sys.executable, str(BENCH), "--source", str(repo),
+ "--m", str(m), "--n", str(n), "--k", str(k),
+ "--warmups", str(warmups), "--samples", str(samples),
+ "--replays-per-sample", str(replays), "--skip-correctness"]
+ if reference_cache is not None and Path(reference_cache).is_dir():
+ cmd += ["--reference-cache-dir", str(reference_cache)]
+ proc = subprocess.run(cmd, capture_output=True, text=True, cwd=str(repo),
+ env=bench_env(device), timeout=timeout_s)
+ for line in reversed((proc.stdout or "").splitlines()):
+ line = line.strip()
+ if line.startswith("{") and line.endswith("}"):
+ try:
+ payload = json.loads(line)
+ except ValueError:
+ continue
+ value = payload.get("median_us")
+ return float(value) if value else None
+ return None
+
+
+def ab_retest(*, shape_id: str, candidate_kernel: Path, kernel_root: Path,
+ model: str, tp: int, m: int, operator: str,
+ shape: Tuple[int, int, int], rounds: int = 2,
+ device: Optional[int] = None, warmups: int = 30, samples: int = 30,
+ replays: int = 50, threshold: float = 3.0,
+ recorded_tol: float = 5.0,
+ reference_cache: Optional[Path] = None,
+ build_root: Path = Path("/tmp/retest_build"),
+ variant_path: Optional[Path] = None) -> Dict[str, Any]:
+ """Interleaved A/B measurement of one candidate against the variant."""
+ device = device if device is not None else idle_device()
+ if device is None:
+ return {"shape": shape_id, "verdict": "no_idle_device"}
+ variant = variant_path or find_variant(model, tp, m, operator)
+ if variant is None:
+ return {"shape": shape_id, "verdict": "no_variant",
+ "reason": f"no variant for {model}/TP{tp}/M{m}/{operator}"}
+ cand_repo = prepare_repo(kernel_root, candidate_kernel,
+ Path(build_root) / shape_id / "cand")
+ var_repo = prepare_repo(kernel_root, variant,
+ Path(build_root) / shape_id / "var")
+ measured: Dict[str, List[float]] = {"cand": [], "var": []}
+ for _ in range(max(1, int(rounds))):
+ for label, repo in (("cand", cand_repo), ("var", var_repo)):
+ value = measure(repo, shape, device, warmups=warmups,
+ samples=samples, replays=replays,
+ reference_cache=reference_cache)
+ if value:
+ measured[label].append(value)
+ cand_median = statistics.median(measured["cand"]) if measured["cand"] else None
+ var_median = statistics.median(measured["var"]) if measured["var"] else None
+ verdict = decide(cand_median, var_median, threshold_percent=threshold,
+ recorded_us=recorded_median_for(candidate_kernel),
+ recorded_tol_percent=recorded_tol)
+ return {"shape": shape_id, "device": device,
+ "candidate_medians_us": measured["cand"],
+ "variant_medians_us": measured["var"],
+ "candidate_median_us": cand_median, "variant_median_us": var_median,
+ "variant_path": str(variant),
+ "recorded_median_us": recorded_median_for(candidate_kernel),
+ **verdict}
+
+
+# ------------------------------------------------------------------- CLI ----
+
+def main(argv: Optional[List[str]] = None) -> int:
+ ap = argparse.ArgumentParser(
+ prog="retest-variant",
+ description="A/B re-measure a candidate kernel against the shipped variant")
+ ap.add_argument("--shape", required=True, help="shape id, e.g. hy3_tp4_o_proj_m4096")
+ ap.add_argument("--candidate-workspace", type=Path,
+ help="AHE child workspace (accepted kernel is looked up)")
+ ap.add_argument("--candidate-kernel", type=Path, help="explicit kernel.hip")
+ ap.add_argument("--kernel-root", type=Path,
+ help="repo skeleton containing csrc/ (defaults to the "
+ "candidate's worker source)")
+ ap.add_argument("--model", help="model label (default: parsed from the shape id)")
+ ap.add_argument("--tp", type=int, help="tensor-parallel size (default: parsed)")
+ ap.add_argument("--m", type=int, help="M (default: from the accepted manifest)")
+ ap.add_argument("--operator", help="operator name (default: parsed from the shape id)")
+ ap.add_argument("--shape-mnk", help="M,N,K of the problem, e.g. 4096,4096,2048")
+ ap.add_argument("--rounds", type=int, default=2)
+ ap.add_argument("--warmups", type=int, default=30)
+ ap.add_argument("--samples", type=int, default=30)
+ ap.add_argument("--replays", type=int, default=50)
+ ap.add_argument("--threshold", type=float, default=3.0)
+ ap.add_argument("--recorded-tol", type=float, default=5.0)
+ ap.add_argument("--device", type=int, help="GPU id (default: pick an idle one)")
+ ap.add_argument("--promote", action="store_true",
+ help="write the variant when the verdict is 'promote'")
+ ap.add_argument("--json", type=Path, help="write the report to this path")
+ args = ap.parse_args(argv)
+
+ from metainfer.tasks.harness_evolve.orchestrator.pool import Pool
+ pool_file = Path("/root/zth_agent/ahe-kernel-repos/registered_pool.yaml")
+ pool = Pool.from_yaml(pool_file) if pool_file.is_file() else None
+ inst = pool.get(args.shape) if pool is not None else None
+
+ candidate = args.candidate_kernel
+ if candidate is None:
+ if not args.candidate_workspace:
+ ap.error("--candidate-workspace or --candidate-kernel is required")
+ candidate = accepted_kernel_for(args.candidate_workspace, args.shape)
+ if candidate is None or not Path(candidate).is_file():
+ ap.error(f"candidate kernel not found for {args.shape}")
+
+ kernel_root = args.kernel_root or Path(candidate).parents[2] / "source"
+ model = args.model or (inst.model if inst else args.shape.split("_tp")[0])
+ tp = args.tp or (inst.tp_size if inst else 4)
+ m = args.m or (inst.M if inst else 4096)
+ operator = args.operator or args.shape.split("_m")[0].split("_", 2)[-1]
+ if args.shape_mnk:
+ mnk = tuple(int(x) for x in args.shape_mnk.split(","))
+ elif inst is not None:
+ mnk = (inst.M, inst.N, inst.K)
+ else:
+ ap.error("--shape-mnk is required when the shape is not in the pool")
+
+ report = ab_retest(
+ shape_id=args.shape, candidate_kernel=Path(candidate),
+ kernel_root=Path(kernel_root), model=model, tp=int(tp), m=int(m),
+ operator=operator, shape=mnk, rounds=args.rounds, device=args.device,
+ warmups=args.warmups, samples=args.samples, replays=args.replays,
+ threshold=args.threshold, recorded_tol=args.recorded_tol,
+ )
+ print(json.dumps(report, ensure_ascii=False, indent=2))
+ if args.json:
+ Path(args.json).write_text(json.dumps(report, ensure_ascii=False, indent=2),
+ encoding="utf-8")
+
+ if args.promote and report.get("verdict") == "promote":
+ from metainfer.tasks.dcu_kernel_auto_opt.orchestrator.variant_promote import ( # noqa: E501
+ model_label_for, promote_variant,
+ )
+ workspace = args.candidate_workspace or Path(candidate).parents[4]
+ outcome = promote_variant(
+ workspace_dir=Path(workspace),
+ answers={"operator": "Quantized GEMM", "dtype": "INT8 W8A8",
+ "model": model, "tp_size": tp},
+ shape_id=args.shape,
+ source_task=f"{args.shape} (A/B retested)",
+ correctness_ok=True, min_improvement_percent=args.threshold,
+ tp=int(tp), m=int(m),
+ model_label=(model_label_for(args.shape, str(model)) or None),
+ )
+ print(json.dumps({"promotion": outcome}, ensure_ascii=False, indent=2))
+ return 0
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/indexer_wq_b.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/indexer_wq_b.hip
index 143caa0b..c6a575d2 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/indexer_wq_b.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/indexer_wq_b.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_indexer_wq_b_m16 commit=9b5b710d0a2c4cd7ac22e2037dfe8735059eec0b added=2026-08-26
+// baseline_us=48.274 speedup=2.70443
// median_us=17.85 p90_us=17.88
// source=hy3-dsh-tp8-m16-1-adf021ab
// W8A8 INT8 GEMM — gfx928 implementation (worker_1, physical GPU 1)
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_down_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_down_proj.hip
index 487c6b3a..2cce19a8 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_down_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_down_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_shared_down_proj_m16 commit=119c281feefbb681f88dba89897484424bb07937 added=2026-08-26
+// baseline_us=14.849 speedup=1.75665
// median_us=8.453 p90_us=9.13
// source=hy3-dsh-tp8-m16-1-adf021ab
// W8A8 INT8 GEMM bootstrap kernel for Hygon K500SM_AI / gfx928.
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_gate_up_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_gate_up_proj.hip
index 15e74f81..4c63671f 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_gate_up_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/shared_gate_up_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_shared_gate_up_proj_m16 commit=ea0f87b2e586e93141907abbdfc541b1585f4c50 added=2026-08-26
+// baseline_us=59.21 speedup=5.75973
// median_us=10.28 p90_us=10.33
// source=hy3-dsh-tp8-m16-1-adf021ab
// W8A8 INT8 GEMM bootstrap kernel for Hygon K500SM_AI / gfx928.
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wo_b.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wo_b.hip
index f420f456..c67c8445 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wo_b.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wo_b.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_wo_b_m16 commit=eb0fbbad60c58b1cf09e6ceff6cb9aaa68a2c1cb added=2026-08-26
+// baseline_us=32.137 speedup=3.27594
// median_us=9.81 p90_us=9.84
// source=hy3-dsh-tp8-m16-1-adf021ab
// MetaInfer INT8 W8A8 GEMM for gfx928 (worker_2 / tp8_wo_b_m16).
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wq_b.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wq_b.hip
index ff0598e7..0b19eb88 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wq_b.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wq_b.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_wq_b_m16 commit=bfaef2fc76d0e4a095b2ffae2b466f2863511d43 added=2026-08-26
+// baseline_us=32.137 speedup=2.19066
// median_us=14.67 p90_us=14.7
// source=hy3-dsh-tp8-m16-1-adf021ab
// W8A8 INT8 GEMM — gfx928 implementation (worker_1, physical GPU 1)
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wqkv_a.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wqkv_a.hip
index bfa665ac..a383ad61 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wqkv_a.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/deepseek-v4/TP8/M16/wqkv_a.hip
@@ -1,4 +1,5 @@
// @@variant shape=tp8_wqkv_a_m16 commit=b67e227c20b3349d46020f9f52cd2484b2142ff3 added=2026-08-26
+// baseline_us=66.591 speedup=3.8052
// median_us=17.5 p90_us=17.52
// source=hy3-dsh-tp8-m16-1-adf021ab
// MetaInfer W8A8 GEMM - worker_0 HIP implementation (gfx928).
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/fused_qkv_a_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/fused_qkv_a_proj.hip
index a7e83daf..7748f78d 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/fused_qkv_a_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/fused_qkv_a_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_fused_qkv_a_proj_m16 commit=bf306c40c48713a81efbd7ace430c9cf21cbb429 added=2026-08-31
+// baseline_us=97.99 speedup=2.33476
// median_us=41.97 p90_us=42.09
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// INT8 W8A8 GEMM kernel for Hygon K500SM_AI / gfx928 (worker_0).
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/kv_b_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/kv_b_proj.hip
index 33985f73..93be3411 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/kv_b_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/kv_b_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_kv_b_proj_m16 commit=a150485d1a63d39ca9251b13d3da304013431efd added=2026-08-31
+// baseline_us=18.833 speedup=2.1881
// median_us=8.607 p90_us=9.766
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// csrc/w8a8_gemm_hip.hip
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/o_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/o_proj.hip
index 3ec28b39..caab0383 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/o_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/o_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_o_proj_m16 commit=c1cfe1975c2fcb8d4bc23c86e8026afae225a4c4 added=2026-08-31
+// baseline_us=66.98 speedup=2.79083
// median_us=24 p90_us=24.02
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// MetaInfer INT8 W8A8 GEMM for Hygon K500SM_AI / gfx928.
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/q_b_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/q_b_proj.hip
index 8d888111..ad036b39 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/q_b_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/q_b_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_q_b_proj_m16 commit=8a9a28e409277a6f266a59e0c053a3d2bb516f12 added=2026-08-31
+// baseline_us=42.035 speedup=2.97698
// median_us=14.12 p90_us=14.14
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// csrc/w8a8_gemm_hip.hip
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_down_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_down_proj.hip
index fdf64fd2..b841e219 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_down_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_down_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_shared_down_proj_m16 commit=42d49afdb24af21a5480786c88ca11a84d57f44a added=2026-08-31
+// baseline_us=18.177 speedup=2.20621
// median_us=8.239 p90_us=9.122
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// W8A8 INT8 GEMM for Hygon K500SM_AI / gfx928 (worker_3).
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_gate_up_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_gate_up_proj.hip
index 3455553e..b33439e2 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_gate_up_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M16/shared_gate_up_proj.hip
@@ -1,4 +1,5 @@
// @@variant shape=glm_tp8_shared_gate_up_proj_m16 commit=fa6ea50630e35036b22141edc48f72975fca041d added=2026-08-31
+// baseline_us=99.318 speedup=9.17911
// median_us=10.82 p90_us=10.87
// source=glm5-2-dsh-tp8-m16-test1-cb4d262c
// W8A8 INT8 GEMM for Hygon K500SM_AI / gfx928 (worker_3).
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/fused_qkv_a_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/fused_qkv_a_proj.hip
index aaa4a2a8..b8032fe2 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/fused_qkv_a_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/fused_qkv_a_proj.hip
@@ -1,266 +1,78 @@
-// @@variant shape=glm_tp8_fused_qkv_a_proj_m4096 commit=935c285f78b3c36ce8f49ef13a68b0cf4e8ff4d5 added=2026-08-31
-// median_us=937.6 p90_us=939.5 speedup=73.23 baseline_us=6.866e+04
-// source=glm5-2-dsh-tp8-m4096-1-e6a280a2
-// @@variant shape=glm_tp8_fused_qkv_a_proj_m4096
-// MetaInfer W8A8 INT8 GEMM for Hygon K500SM_AI / gfx928 (worker_0, GPU 0).
+// @@variant shape=glm_tp8_fused_qkv_a_proj_m4096 commit=7cf8d67fe9e0c38f2b4c3c67aded1bf8cc894342 added=2026-09-14
+// median_us=693.8 p90_us=695.3 speedup=98.96 baseline_us=6.866e+04
+// source=test-9-11-1/iteration_002/glm_tp8_fused_qkv_a_proj_m4096
+// MetaInfer W8A8 INT8 GEMM -- round-1 architecture establishment for
+// Hygon K500SM_AI / gfx928 (worker_0, physical GPU 0).
//
// Assigned shape:
// glm_tp8_fused_qkv_a_proj_m4096 : M=4096, N=2624, K=6144
+// (GLM5.2 TP=8 fused_qkv_a_proj, logical (K, N) = (6144, 2624))
//
-// Operator contract (fixed by the control plane):
+// Operator contract (fixed by csrc/bindings.cpp and int8_w8a8_gemm_api.py):
// out[m, n] = bf16( int32_dot(x_q[m, :], raw_weight[:, n])
-// * x_scale[m] * weight_scale[n] )
+// * x_scale[m, 0] * weight_scale[n, 0] )
//
-// Iteration 1 (correctness-first DUMMA bootstrap):
-// * Packing bootstrap: launch_pack_w8a8_weight is the mandated identity
-// device-to-device copy for every (K, N) (raw [K, N] row-major weight is
-// copied unchanged and the fp32 scales are copied unchanged). The GEMM
-// therefore interprets packed_weight as the logical [K, N] row-major
-// layout everywhere. Later Parallel explore rounds may change the pack
-// kernel and the matching GEMM interpretation together.
-// * Iteration 1 dispatched tile <64, 64, 128> (41 x 64 = 2624 blocks,
-// 16,384 B LDS/block, 4 blocks/CU, official median 6391.92 us = 20.66
-// logical TOPS). Its exact-source ISA shows the stage loop is
-// LDS-issue-bound, not MMA-bound: per stage per wavefront 136 ds_read
-// (mostly du_load_matrix_sync byte-reassembly ds_read_u8) + 107 waitcnt
-// vs only 32 v_mmac and 2 barriers; 2624 blocks x 96 barriers =
-// 251,904 barriers and 2.06 GB of staged global traffic per replay.
-// * The epilogue is the coalesced fragment store: the verified gfx928
-// accumulator mapping (row = lane & 15, col_mod4 = lane >> 4,
-// frag.x[i] -> column col_mod4 + 4*i) is transposed in registers (8
-// shfl_xor + 8 v_cndmask) so each lane owns four contiguous bf16 columns
-// and writes ONE 8-byte store (100% store sector efficiency). The
-// per-element float multiply order (float(dot) * x_scale[m] *
-// weight_scale[n]) and the bf16 rounding are the reference order, so the
-// stored bits are identical to the scalar path.
-// * Generic arms: the <64,64,128> tiled kernel also serves every
-// large-M shape with N % 64 == 0 and K % 128 == 0 (identity layout); a
-// scalar int8/int32 grid-stride kernel is the fallback for every other
-// (m, n, k), including the paired M=2 and M=16 API shapes with the same
-// (N, K), which must never reach the tiled path (m >= 128 guard).
-// * The timed operator (launch_w8a8_gemm) performs no allocation,
-// compilation, autotuning, packing, host/device synchronization, or
-// default-stream launch: it only dispatches kernels on the
-// caller-provided HIP stream and ignores the workspace (split-K capacity
-// is zero for this shape: m*n*4 = 42,990,080 B > 16 MiB budget).
+// Mandatory decision this round: the tile-shape (M-tile vs N-tile aspect ratio)
+// round. The packed_weight architecture, the exact-shape packed-B swizzle, the
+// double-buffered 64-K stage pipeline, the bank-safe A LDS stride, the
+// conflict-free swizzled B fragment reads and the fused coalesced bf16 epilogue
+// are all RETAINED unchanged; only the macro-tile aspect ratio and its derived
+// wave split are retuned for M=4096, N=2624, K=6144. No split-K, no workspace
+// use, no allocation/compilation/packing/synchronization inside the timed API.
//
-// Iteration 1 (mandated 2-D macro-tile baseline: benchmark 64x64 / 64x128 /
-// 128x64, pick with code-object resources + occupancy):
-// * Tile comparison (all three instantiations already compiled by the
-// iteration-1 build; VGPR/SGPR/LDS/spills are exact code-object facts
-// from profiles/.../iteration1/current-best-isa/metadata.txt; grid and
-// staged traffic are arithmetic for M=4096, N=2624, K=6144):
-// tile VGPR SGPR LDS B blocks/CU waves/CU grid staged GB
-// <64, 64,128> 57 36 16384 4 16 2624 2.06
-// <64,128,128> 92 36 24576 2 8 1344 1.56*
-// <128,64,128> 88 34 24576 2 8 1312 1.55
-// (*) 64x128 needs an N tail: N=2624 = 41*64 is not a multiple of
-// 128, so the n0=2560 block would zero-fill B columns >= 2624 and an
-// epilogue column guard, wasting ~2.4% of all MMAs. No split-K in any
-// arm: 1312 blocks dwarf the 120 CUs.
-// * Iteration 1 dispatched the exact-fit 2-D macro-tile <128, 64, 128>:
-// M=4096 = 32*128 and N=2624 = 41*64 are both exact, so no tail
-// predication is introduced anywhere (the generic A row guard and the
-// epilogue row guard are retained for the generic path). Grid = (N/64) x
-// (M/128) = 41 x 32 = 1312 blocks; 256 threads = 4 wavefronts, each wave
-// owns a 64x32 quadrant = eight m16n16k32 int32 accumulator fragments;
-// single-buffered 128-K LDS stage (A[128,128] + B[128,64] = 24,576
-// B/block) with TWO __syncthreads per stage, K=6144 = 48 stages. Versus
-// the iteration-1 64x64 tile this halves B re-reads (B staged by 32
-// M-blocks instead of 64; staged global traffic 2.06 -> 1.55 GB per
-// replay) and cuts total barriers 251,904 -> 125,952, at the cost of
-// occupancy 4 -> 2 blocks/CU (16 -> 8 resident waves/CU). The int32
-// accumulation order (k0-outer over 128-K stages, kk-inner ascending,
-// same m16n16k32 element-to-slot mapping) is bit-identical to the
-// iteration-1 kernel and the scalar reference.
-// * The <64,128,128> instantiation stays compiled as the next benchmark
-// candidate (needs the N-tail handling described above before dispatch);
-// the generic <64,64,128> arm and the scalar fallback are untouched.
+// (1) launch_pack_w8a8_weight packs the exact (k, n) == (6144, 2624) weight
+// ONCE, outside the timed region and out of Graph capture, into a
+// fragment-interleaved swizzle of 512-byte 32-k x 16-n DUMMA B fragment
+// blocks laid out [k32][n16][k8][n][8]:
+// packed[((kk>>5)*n16g + (n>>4))*512 +
+// (((kk>>3)&3)*16 + (n&15))*8 + (kk&7)] = raw[kk*n + n],
+// n16g = n>>4 = 164. Byte-exact bijection 192 x 164 x 512 = k*n =
+// 16,121,856: same byte count and same buffer allocation, so every
+// captured address is graph-stable. Every other (K, N) keeps the identity
+// device-to-device copy.
+// (2) the exact-shape guard dispatches the new
+// w8a8_dumma_prefill_packedb_db_kernel<256, 128, 64, 512, true>:
+// grid = ((N + 127)/128) x (M/256) = 21 x 16 = 336 blocks,
+// block = 512 threads = 8 gfx928 wavefronts,
+// a 64 x 64 quadrant of sixteen m16n16k32 int32 accumulator fragments
+// per wave (4 wave rows x 2 wave columns, the balanced square panel),
+// i.e. a 256 x 128 M-major operand-reuse macro-tile, 64-K
+// DOUBLE-buffered stages with ONE __syncthreads per stage and the
+// stage-(s+1) global loads issued into staging VGPRs before the stage-s
+// consume, k0-outer over 96 ascending 64-K stages with kk-inner
+// ascending m16n16k32 MMAs (exact int32 accumulation order of the
+// reference), and a fused branch-free per-row/column scale + bf16
+// epilogue emitting one coalesced 8-byte store per lane.
+// The A LDS tile uses the bank-safe 72-byte row stride (StageK + 8: an
+// odd 8-byte count, so the 16 rows of each 16-lane ds_read_b64 phase land
+// on 16 distinct bank pairs); the swizzled B tile is a flat BN x 64-byte
+// copy of the packed stage, so every col_major B fragment read is one
+// conflict-free ds_read_b64 (load_frag8_swz).
+// LDS = 2 x (A[256, 72] + B[128, 64]) = 2 x 18,432 + 2 x 8,192 =
+// 53,248 B per block, so ONE resident block/CU = 8 waves/CU (the
+// two-operand global traffic falls from M*K*21 + N*K*32 = 1,044 MB to
+// M*K*21 + N*K*16 = 786 MB, -24.7%, at unchanged wave occupancy).
+// (3) the scalar int8/int32 fallback decodes the swizzled layout element-wise
+// in ascending kk for (k, n) == (6144, 2624), so the paired M=2 / M=16
+// API shapes with the same (N, K) stay byte-exact; the generic
+// single-buffered <64, 64, 128> identity arm covers every other large-M
+// shape with n % 64 == 0 and k % 128 == 0. The fallback is never replaced
+// by an unconditional shape-specialized launch.
//
-// Iteration 2 (mandated operand-reuse round: cooperative A+B LDS staging
-// vs direct loads, quantified A/B reuse, vectorized coalesced global loads,
-// bank-safe LDS layout):
-// * Operand-reuse accounting for the <128,64,128> macro-tile (M=4096,
-// N=2624, K=6144): per block A = 128 x 6144 = 786,432 B and B = 6144 x
-// 64 = 393,216 B of unique bytes; the 41 N-blocks re-read every A
-// row-slab (A global traffic M*K*41 = 1.03 GB/replay, 41x amplification
-// of the 25.2 MB unique A) and the 32 M-blocks re-read every B
-// column-slab (K*N*32 = 0.52 GB/replay, 32x of the 16.1 MB unique B);
-// total staged traffic 1.55 GB/replay. Cooperative staging additionally
-// buys 4x wave-level reuse inside each block (one 16-B vectorized global
-// load per byte, four waves re-read it from LDS), so it is retained: a
-// direct per-wave global fragment path would issue 4x more vmem (4 waves
-// x 48 x 8-B fragment loads per stage = ~9,216 vs ~288 staged
-// 16-B loads per wave) and every 8-B fragment load would be a
-// row-strided 16-sector access (vs 4 sectors for the coalesced 16-B
-// staging) -- both directions strictly worse.
-// * The exact-source ISA of the iteration-1 <128,64,128> kernel shows the
-// per-stage per-wavefront consume issues 80 ds_read (64 ds_read_u8 +
-// 16 ds_read2_b32) + 99 waitcnt + ~96 byte-reassembly VALU
-// (v_and/v_or3/v_lshlrev) per stage against only 32 v_mmac: the
-// row_major matrix_b loader of du_load_matrix_sync lowers each B
-// fragment to 8 ds_read_u8 at 16-B spacing plus a byte-reassembly chain,
-// and the LDS row strides 128 (A) / 64 (B) alias every row onto one
-// bank phase (32r and 16r mod 32 == 0), matching the PMC
-// lds_bank_conflicts = 274,071,552 (~12.5 per LDS instruction). The
-// kernel is LDS-issue/conflict-bound, not MMA- or HBM-bound.
-// * Focused change (one mechanism, all inside csrc/w8a8_gemm_hip.hip):
-// (1) launch_pack_w8a8_weight now packs the exact (k,n) == (6144,2624)
-// weight once, outside the timed region and out of Graph capture, into
-// the n-major layout packed[n*K + kk] = raw[kk*N + n] (same byte count
-// and buffer, captured addresses unchanged; every other (K, N) keeps
-// the identity copy);
-// (2) the exact-shape guard dispatches a new
-// w8a8_dumma_prefill_packedb_kernel<128,64,128> (grid 41x32 = 1312
-// blocks, 256 threads = 4 waves, same single-buffered 128-K stage with
-// two __syncthreads, same k0-outer/kk-inner int32 order, same epilogue)
-// that stages A row-major and B n-major into bank-safe LDS rows of 136 B
-// (StageK + 8: 8-byte-aligned so every fragment read is one ds_read_b64,
-// 34 dwords so the 16 rows of each 16-lane ds_read_b64 phase land on 16
-// distinct bank pairs -> 4-phase conflict-free floor) and replaces both
-// library loaders with the lineage-validated explicit load_frag8 (one
-// aligned 8-B load per lane per fragment: x[0..7] = 8 consecutive bytes
-// of row (lane&15) at k-offset ((lane>>4)*8) for row_major A and the
-// same 8 k-bytes of column n0+(lane&15) for col_major B over the
-// packed [N, K] tile -- identical fragment values, so the m16n16k32
-// v_mmac inputs and the exact int32 accumulation order are bit-identical
-// to the iteration-1 kernel and the scalar reference);
-// (3) the scalar fallback decodes the n-major layout when
-// (k,n) == (6144,2624) so the paired M=2/M=16 API shapes with the same
-// (N, K) stay byte-exact; the generic <64,64,128> identity arm and the
-// <64,128,128>/<128,64,128> sibling instantiations are untouched.
-// * Expected ISA/PMC deltas: ds_read per stage per wavefront 80 -> 24
-// (16 A + 8 B ds_read_b64), waitcnt 99 -> ~24-32, reassembly VALU ~0,
-// lds_bank_conflicts 274M -> conflict-free floor (~1-3M), LDS/block
-// 24,576 -> 26,112 B (2 blocks/CU unchanged; no spills).
+// Correctness: the same weight elements reach the same fragment slots in the
+// same ascending k-chunk sequence (k0-outer over 96 ascending 64-K stages,
+// kk-inner ascending over two m16n16k32 steps, i.e. the identical (k0 + kk)
+// k-chunk sequence 0, 32, 64, 96, ...), and the epilogue applies the exact
+// reference multiply order float(dot) * x_scale[row] * weight_scale[col], so
+// every bf16 output byte is bit-identical to the scalar reference. Widening BM
+// only changes WHICH wave computes which output element; it never reorders the
+// K stages or the int32 accumulation.
//
-// Iteration 3 (mandated pipeline round: single vs double buffering across K
-// tiles; retain double buffering only on ISA/PMC evidence of reduced VMEM
-// stalls without harmful LDS or occupancy growth):
-// * Bottleneck evidence from the accepted iteration-2 source (fresh PMC
-// digest 20f3bc9a..., profiles/.../iteration3/pmc.json: profiled
-// 2041.095 us, vmem_read 1,574,400, lds_instructions 7,892,992 with
-// lds_wait 4,571,574, lds_bank_conflicts 0, 26,112 B LDS, 2 blocks/CU):
-// the exact-source ISA of w8a8_dumma_prefill_packedb_kernel<128,64,128>
-// shows the stage loop is [global_load_dwordx4 ... -> s_waitcnt
-// vmcnt(0) -> ds_write2_b64 -> s_barrier -> 24x ds_read_b64 -> 32x
-// v_mmac -> s_barrier] -- the FULL global-load round trip (issue, vmcnt
-// wait, LDS flush) is serialized in front of every consume with ZERO
-// compute overlap, and the 14 vmcnt + 57 lgkmcnt waits in the symbol sit
-// on that critical path. The K=6144 grid (1312 blocks) and 2-block
-// co-residency do not hide this per-stage latency.
-// * Focused change (one mechanism, all inside csrc/w8a8_gemm_hip.hip):
-// the exact-shape guard now dispatches a new
-// w8a8_dumma_prefill_packedb_db_kernel<128,64,64> -- SAME macro-tile
-// (grid 41x32 = 1312 blocks, 256 threads = 4 waves, 64x32 quadrant per
-// wave, eight m16n16k32 int32 accumulators), SAME load_frag8 fragment
-// reads, SAME coalesced epilogue -- with the K stage changed from a
-// single-buffered 128-K stage (2 __syncthreads, full VMEM latency
-// exposed) to a DOUBLE-buffered 64-K stage with ONE __syncthreads per
-// stage: the stage-(s+1) global loads are issued into 3 int4 staging
-// VGPR per thread BEFORE the stage-s consume, and the prefetched vectors
-// are flushed into the idle LDS buffer after the consume (the validated
-// M=3072 prefetch order global_load_dwordx4 -> v_mmac -> vmcnt(0) ->
-// ds_write2_b64, here in pure HIP: the vmcnt wait lands on the data
-// dependency before the ds_write, i.e. AFTER the MMAC burst).
-// * Resources: LDS = 2 x (A[128,72] + B[64,72]) = 2 x 13,824 = 27,648
-// B/block (26,112 -> 27,648 B, +5.9%); 2 x 27,648 = 55,296 <= 65,536 B
-// keeps TWO resident blocks/CU = 8 waves/CU, occupancy UNCHANGED (the
-// StageK=128 double-buffer alternative needs 52,224 B/block -> 1
-// block/CU = 4 waves/CU and is rejected up front as harmful occupancy
-// loss). kLdsStride = 72 = 64 + 8: 8-aligned with an odd 8-byte count
-// (72/8 = 9) so the 16 rows of each 16-lane ds_read_b64 phase land on 16
-// distinct bank pairs (18r mod 32, r = 0..15) -- the same conflict-free
-// 4-phase floor as iteration 2 -- and 72 % 16 == 8 keeps the two-half
-// ds_write2_b64 staging form. VGPR grows by the 12 staging registers
-// (~72 -> ~84), far below the 128/thread budget for 2 blocks/CU; 0
-// spills expected. Barriers per block: 1 (prologue) + 95 (loop) = 96,
-// the same as iteration 2's 2 x 48 (the final stage is consumed without
-// a trailing barrier).
-// * Correctness: k0-outer over 96 ASCENDING 64-K stages with kk-inner
-// ascending over two m16n16k32 steps gives the identical (k0+kk)
-// k-chunk sequence 0, 32, 64, 96, ... as iteration 2's 48 x 128-K
-// stages, and the element-to-slot fragment mapping is unchanged, so
-// every int32 accumulator is bit-identical (mismatch 0 / max_abs_error
-// 0.0 expected).
-// * Falsifiable retention: double buffering is retained only if the exact
-// shape shows median_us < 2276.930923461914 AND p90_us <
-// 2540.756378173828 with mismatch_count == 0, max_abs_error == 0.0,
-// graph_capture_passed == true, the exact code object shows 0 spills /
-// LDS == 27,648 B (2 blocks/CU still fits: 55,296 <= 65,536) and the
-// stage ISA shows the prefetch loads issued BEFORE the ds_read_b64 /
-// v_mmac burst with the vmcnt wait after the MMACs (load -> mmac ->
-// vmcnt(0) -> ds_write2_b64), i.e. reduced exposed VMEM stalls; any
-// median/p90 >= current best, any mismatch, any spill, any occupancy
-// loss, or any ISA that still serializes the vmcnt wait in front of the
-// consume falsifies the round and the control plane restores the
-// iteration-2 source (single buffering retained).
-//
-// Iteration 5 (mandated packing round: test one weight packing/swizzle that
-// makes each DUMMA B tile vector-loadable and LDS-bank-safe):
-// * Baseline: the accepted iteration-3 kernel
-// w8a8_dumma_prefill_packedb_db_kernel<128,64,64> (official median
-// 1070.23 us = 123.4 logical TOPS; fresh PMC digest 47542b38...,
-// profiles/.../iteration4/pmc.json of the current source: profiled
-// 1101.96 us, grid 1312, 96 arch_vgpr / 27,648 B LDS, lds_instructions
-// 4,901,632 = 9.7 per wave-stage (6 merged ds_read2_b64 + 3
-// ds_write2_b64), lds_wait 3,701,906, lds_bank_conflicts 12,091,392 =
-// 6 x 4 x 4 x 96 x 1312 EXACTLY -- one conflict per 16-lane phase of
-// every merged ds_read2_b64, the inherent 2-cycle alias of the +16-row
-// (16 x 72 = 1152 B = 9 x 128 B) second-fragment offset, NOT a
-// layout defect (the iteration-2 unmerged ds_read_b64 control had 0
-// conflicts at the same 72/136-B strides), vmem_read 1,574,400, L2 hit
-// 79.3% with 5,155,404 misses, TCC read 269.3 MB/replay). The kernel is
-// LDS-latency-bound, not HBM-bound; the global B operand is the
-// remaining layout lever.
-// * Mechanism: replace the iteration-2 n-major pack
-// packed[n*K + kk] (each 64-K B tile = 64 column chunks of 64 B at
-// 6144-B stride -> 64 x 128-B lines touched per tile, 50% line
-// utilization, scattered L2 footprint) with a B-PANEL pack
-// packed[(n>>6)*(k*64) + (kk>>4)*1024 + (n&63)*16 + (kk&15)] =
-// raw[kk*n + n] (41 panels of [K, 64] int8, 16-byte k-runs per column,
-// byte-exact bijection verified by simulation): every [64,64] B tile of
-// a 64-K stage is then FOUR CONTIGUOUS 1-KiB slabs (one per 16-k-run),
-// each staging vector is a 16-byte-aligned contiguous k-run, and a wave
-// reads 4 x 256 B of one panel region per stage -- the tile is
-// vector-loadable as one coalesced stream with 8x fewer 128-B lines
-// touched per tile (32 vs 256) and full line utilization, so L2 tag
-// pressure and B HBM fetch traffic drop without any change in LDS
-// layout or instruction count.
-// * Kernel-side changes are ONLY the two B global-load sites (prologue +
-// steady-state prefetch) of w8a8_dumma_prefill_packedb_db_kernel
-// (dispatched solely by the exact-shape guard; static_assert BN == 64):
-// vector (n_row, k16) is loaded from byte
-// (n0>>6)*(k*64) + ((k0>>4)+k16)*1024 + n_row*16 instead of
-// (n0+n_row)*k + k0 + k16*16. LDS staging layout (b_tile[64][72],
-// two-half ds_write2_b64), the col_major load_frag8 consumes, the
-// k0-outer/kk-inner int32 accumulation order, the coalesced epilogue,
-// the grid (41x32 = 1312 blocks), threads (256 = 4 waves), occupancy
-// (2 blocks/CU, LDS 27,648 B) and VGPR budget are ALL unchanged, so the
-// round is a pure packing/swizzle A/B.
-// * Correctness: the same weight elements reach the same fragment slots
-// in the same ascending k-chunk sequence, so every int32 accumulator is
-// bit-identical (mismatch 0 / max_abs_error 0.0 expected); the scalar
-// fallback decodes the panel layout when (k,n) == (6144,2624) (ascending
-// 16-byte k-runs keep the exact int32 order for the paired M=2/M=16 API
-// shapes); the generic identity arms and sibling kernels are untouched.
-// * Packing stays outside the timed region and out of Graph capture
-// (launch_pack_w8a8_weight, same byte count k*n = 16,121,856 and same
-// buffer -> graph-stable captured addresses unchanged; pack relation
-// verified by a byte-exact index simulation before dispatch).
-// * Predicted deltas: vmem_read_instructions unchanged (~1,574,400), LDS
-// counters unchanged (conflicts stay at the 12,091,392 merged-read alias
-// floor), L2 misses < 5,155,404 with hit rate UP, TCC read bytes <
-// 269,329,728 (B line fetches halve), median ~1000-1070 us. Falsifiable:
-// retained only if median_us < 1070.2328491210938 AND p90_us <
-// 1081.9308471679688 with mismatch_count == 0, max_abs_error == 0.0,
-// graph_capture_passed == true, the exact code object shows 0 spills /
-// LDS == 27,648 B / VGPR <= 128 (2 blocks/CU), and PMC confirms L2
-// misses and TCC read bytes strictly below the iteration-4 values with
-// vmem_read/lds/conflicts not growing; any median/p90 >= current best,
-// any mismatch, any spill, any occupancy loss, or any L2/HBM regression
-// falsifies the pack and the control plane restores the iteration-3
-// source.
+// Reported launch geometry for this round (iteration 4, tile shape):
+// grid_blocks = 21 x 16 = 336
+// waves_per_block = 512 / 64 = 8
+// estimated_active_cus = 120 (K500SM_AI has 120 CUs; 336 blocks with
+// 1 resident block/CU = 2.8 blocks/CU occupy every CU)
//
// Include order is fixed by the control plane: hip_runtime, hip_bfloat16,
// then du_mma (this DTK's du_mma.h is not self-contained otherwise).
@@ -273,12 +85,18 @@
namespace {
-constexpr int kWaveSize = 64; // gfx928 native wavefront
+constexpr int kWaveSize = 64; // gfx928 native wavefront size
constexpr int kTileM = 16;
constexpr int kTileN = 16;
constexpr int kTileK = 32; // INT8 DUMMA unit: m16n16k32, int32 accumulation
// 4 wavefronts; must remain a multiple of the gfx928 wavefront size (64).
constexpr int kThreadsPerBlock = 256;
+// Exact-shape packed-B prefill workgroup: 8 wavefronts = 512 threads, with the
+// M-major operand-reuse macro-tile (BM=256 x BN=128) selected by the
+// iteration-4 tile-shape round.
+constexpr int kPackedThreads = 512;
+constexpr int kPackedBM = 256;
+constexpr int kPackedBN = 128;
// Exact assigned shape (K, N) after TP=8 partitioning.
constexpr int kExactN = 2624;
@@ -286,21 +104,39 @@ constexpr int kExactK = 6144;
using namespace du::dumma;
-// Coalesced accumulator epilogue for gfx928 INT8 m16n16k32 (verified against
-// du_store_matrix_sync in the accepted worker-29 lineage): the accumulator
-// lane mapping is row = lane & 15, col_mod4 = lane >> 4, frag.x[i] maps to
-// columns col_mod4 + 4*i. A register 4x4 transpose (two 2x2 steps with
-// shfl_xor 16 then 32, one v_cndmask per element per step) re-routes the
-// int32 values so lane (r, c4) owns the four CONTIGUOUS columns
-// base_col + 4*c4 .. +3; they are scaled (float(dot) * x_scale[row] *
-// weight_scale[col], the exact reference multiply order), converted to bf16,
-// packed and stored with ONE 8-byte store per lane. The row >= m guard is
-// wavefront-uniform (lane & 15 cycles the same 16 rows in every 16-lane
-// group), so the shuffles never mix active and inactive lanes. base_col is a
-// multiple of 64 for the dispatched tiles, so the float4 weight_scale load
-// (col0 % 4 == 0) and the 8-byte store (n even -> out offset a multiple of
-// 8) are aligned.
-template
+// Branch-free round-to-nearest-even fp32 -> bf16 high-16 payload:
+// bf16 = (b + 0x7fff + ((b >> 16) & 1)) >> 16, b = __float_as_uint(v)
+// For every FINITE fp32 this is exactly the normal path of __float2bfloat16,
+// while dropping the ~13-instruction EXEC-bracketed Inf/NaN sequence this DTK
+// emits for the library call (adding 0x7fff + the retained bit-16 LSB cannot
+// carry out of the sign; the +((b>>16)&1) term rounds exact ties to an even
+// LSB, max-float still rounds to +/-Inf and +/-0.0/subnormal payloads are
+// preserved). Everything converted here is finite by construction: |int32 dot|
+// < 2^31 and the dequantization scales are finite.
+__device__ __forceinline__ unsigned int bf16_rne_high16(float v) {
+ const unsigned int b = __float_as_uint(v);
+ return (b + 0x7fffu + ((b >> 16) & 1u)) >> 16;
+}
+
+// Fused per-row/per-column dequantization + bf16 epilogue for the m16n16k32
+// int32 accumulator of gfx928 INT8 DUMMA. Accumulator lane mapping:
+// row = lane & 15, col_mod4 = lane >> 4, frag.x[i] -> column col_mod4 + 4*i.
+// A register 4x4 transpose (two 2x2 steps via shfl_xor 16 then 32) re-routes
+// the values so lane (r, c4) owns the four CONTIGUOUS columns base_col +
+// 4*c4 .. +3, which are then scaled in the exact reference order
+// float(dot) * x_scale[row] * weight_scale[col], rounded to bf16 and stored
+// with ONE 8-byte store per lane. base_col is a multiple of 64 for the
+// dispatched tiles, so the float4 weight_scale load (col0 % 4 == 0) and the
+// 8-byte store are both aligned. The row >= m guard is wavefront-uniform
+// (lane & 15 cycles the same 16 rows in each 16-lane group), so the shuffles
+// never mix active and inactive lanes.
+//
+// kGuarded == false is the exact-shape full-tile specialization (m % BM == 0):
+// every row of every block tile is strictly inside the output, so the dead
+// `row >= m` EXEC bracket around each of the eight unrolled stores is removed
+// at compile time. The default kGuarded == true keeps the exact guarded
+// semantics for every other m.
+template
__device__ __forceinline__ void store_prefill_fragment_coalesced(
const AccFragment& frag,
const float* __restrict__ x_scale,
@@ -312,8 +148,10 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced(
int base_col,
int lane) {
const int row = base_row + (lane & 15);
- if (row >= m) {
- return;
+ if (kGuarded) {
+ if (row >= m) {
+ return;
+ }
}
const int c4 = lane >> 4; // 0..3
const int x0 = frag.x[0];
@@ -341,8 +179,8 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced(
const int f2 = hi ? t0 : a2;
const int f3 = hi ? t1 : a3;
- // Lane (r, c4) now owns columns base_col + 4*c4 .. +3 (8 B, 8-byte
- // aligned: base_col is a multiple of 64, n*2 is a multiple of 8).
+ // Lane (r, c4) now owns columns base_col + 4*c4 .. +3 (8 B, 8-byte aligned:
+ // base_col is a multiple of 64 and n*2 is a multiple of 8).
const int col0 = base_col + 4 * c4;
const float xs = x_scale[row];
const float4 ws = *reinterpret_cast(weight_scale + col0);
@@ -352,53 +190,126 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced(
const float v3 = static_cast(f3) * xs * ws.w;
const uint64_t packed =
static_cast(
- static_cast(__float2bfloat16(v0))) |
- (static_cast(
- static_cast(__float2bfloat16(v1)))
- << 16) |
+ bf16_rne_high16(v0) |
+ (bf16_rne_high16(v1) << 16)) |
(static_cast(
- static_cast(__float2bfloat16(v2)))
- << 32) |
- (static_cast(
- static_cast(__float2bfloat16(v3)))
- << 48);
+ bf16_rne_high16(v2) |
+ (bf16_rne_high16(v3) << 16))
+ << 32);
*reinterpret_cast(out + static_cast(row) * n + col0) =
packed;
}
-// Simple native INT8 DUMMA m16n16k32 tiled prefill kernel (identity [K, N]
-// weight layout; launch_pack_w8a8_weight is the identity bootstrap, so
-// packed_weight == raw logical weight). Template parameters: BM x BN output
-// tile per block, StageK K rows per single-buffered LDS stage. 256 threads =
-// 4 wavefronts; each wave owns a (BM/2) x (BN/2) quadrant of m16n16k32
-// int32 accumulator fragments. Grid = (N/BN) x (ceil(M/BM)).
+// Explicit 8-byte int8 fragment loader for BOTH the row_major matrix_a and the
+// col_major matrix_b m16n16k32 fragments: lane l -> row (l & 15), k-quarter
+// (l >> 4) holds the EIGHT CONTIGUOUS bytes at p[row * ldm + (l >> 4) * 8 .. +7].
+// Loading the 8 bytes as one int64 forces the ds_read_b64 form (16-lane phases,
+// distinct bank pairs per group -> the 4-phase conflict-free floor with an odd
+// 8-byte LDS row stride), identical fragment values to the library loader, so
+// the v_mmac inputs and the exact int32 accumulation order are unchanged.
+// `lane` is threadIdx.x % kWaveSize, so no wave-id builtin is required.
+template
+__device__ __forceinline__ void load_frag8(
+ Frag& f,
+ const int8_t* __restrict__ p,
+ unsigned ldm,
+ int lane) {
+ const unsigned row = static_cast(lane) & 0xfu;
+ const unsigned kq = static_cast(lane) >> 4;
+ const int64_t v = *reinterpret_cast(p + row * ldm + (kq << 3));
+ reinterpret_cast(f.x)[0] = v;
+}
+
+// Fragment-interleaved B swizzle helpers.
//
-// K is staged cooperatively in a SINGLE LDS buffer at StageK granularity
-// with TWO __syncthreads per stage (one after the staging stores, one after
-// all waves consumed the buffer):
-// * A[BM, StageK] is staged row-major with flat 16-byte vectors (each
-// thread strides over kAVectors = BM*StageK/16 vectors); rows past M are
-// zero-filled and masked in the epilogue.
-// * B[StageK, BN] is staged row-major from the logical [K, N] weight with
-// flat 16-byte vectors (each thread strides over kBVectors =
-// StageK*BN/16 vectors). The dispatch guarantees n % BN == 0 and
-// k % StageK == 0, so every global vector is aligned and in-bounds.
-// * Fragment loads use the library du_load_matrix_sync (row_major for both
-// operands) and accumulate with du_mma_sync in ascending k order --
-// k0-outer over stages, kk-inner over kTileK steps -- preserving the
-// exact int32 accumulation order of the scalar reference.
-// * The epilogue is the coalesced fragment store.
+// Each 32-k x 16-n DUMMA B fragment block is 512 B = [k8 0..3][n 0..15][8 B]
+// (n-interleaved k8 groups). The 16-byte chunk (k32, n16, k8, n_pair) holds the
+// two 8-byte k-runs of n-rows 2*n_pair and 2*n_pair+1 of one k8 group, and sits
+// at byte (k32 * n16g + n16) * 512 + k8 * 128 + n_pair * 16 of the packed
+// [n>>4][k>>5][512 B] array (n16g = n>>4). Global elements satisfy
+// packed element (col, kk) at ((kk>>5)*n16g + (col>>4))*512 +
+// (((kk>>3)&3)*16 + (col&15))*8 + (kk&7) -- the exact inverse of this chunk
+// map, so staging is a flat 16-byte copy and the fragment load reads the same
+// bytes the library col_major loader consumes.
+__device__ __forceinline__ int64_t packed_b_swz_offset(
+ int k32,
+ int n16,
+ int n16g,
+ int k8,
+ int n_pair) {
+ return (static_cast(k32) * n16g + n16) * 512 +
+ static_cast(k8) * 128 + (n_pair << 4);
+}
+
+// col_major B fragment loader on the swizzled tile: lane (n = lane&15,
+// k8 = lane>>4) reads p + k8*128 + n*8. One ds_read_b64 per fragment and each
+// 16-lane phase spans exactly one 128-byte k8 block (each bank exactly once),
+// so the read is conflict-free by construction. Byte placement is identical to
+// load_frag8, so the v_mmac operands and the int32 accumulation stay bit-exact.
+template
+__device__ __forceinline__ void load_frag8_swz(
+ Frag& f,
+ const int8_t* __restrict__ p,
+ int lane) {
+ const unsigned n = static_cast(lane) & 0xfu;
+ const unsigned k8 = static_cast(lane) >> 4;
+ const int64_t v =
+ *reinterpret_cast(p + (k8 << 7) + (n << 3));
+ reinterpret_cast(f.x)[0] = v;
+}
+
+// Tile-shape round: balanced per-wave quadrant. Pick kWaveRows (hence
+// kWaveCols = kWaves / kWaveRows) so that the wave's m16n16k32 panel
+// (kWaveM16 x kWaveN16 fragments) is as square as possible. A square panel
+// minimises the LDS fragment reads per v_mmac, (kWaveM16 + kWaveN16) reads for
+// kWaveM16 * kWaveN16 MMAs, and keeps the live operand fragment set -- and
+// therefore the VGPR bill of the 16-accumulator M-major macro-tile -- small.
+// The search is restricted to exact divisor splits of both the tile fragment
+// counts and the wave count so the kernel's quadrant arithmetic stays valid;
+// for the retained 128 x 128 / 8-wave geometry it reproduces the historical
+// 2 wave rows x 4 wave columns split exactly.
+__host__ __device__ constexpr int balanced_wave_rows(int bm_frags,
+ int bn_frags,
+ int waves) {
+ int best_rows = 1;
+ int best_diff = 1 << 30;
+ for (int rows = 1; rows <= waves; ++rows) {
+ if (waves % rows != 0) continue;
+ if (bm_frags % rows != 0) continue;
+ const int cols = waves / rows;
+ if (bn_frags % cols != 0) continue;
+ const int m16 = bm_frags / rows;
+ const int n16 = bn_frags / cols;
+ const int diff = (m16 > n16) ? (m16 - n16) : (n16 - m16);
+ if (diff < best_diff) {
+ best_diff = diff;
+ best_rows = rows;
+ }
+ }
+ return best_rows;
+}
+
+// Generic native INT8 DUMMA m16n16k32 tiled prefill kernel with the identity
+// [K, N] weight layout (packed_weight == raw logical weight for every (K, N)
+// other than the exact assigned pair). Template parameters: BM x BN output tile
+// per block, StageK K rows per single-buffered LDS stage. 256 threads =
+// 4 wavefronts; each wave owns a (BM/2) x (BN/2) quadrant of m16n16k32 int32
+// accumulator fragments. Grid = (N/BN) x (ceil(M/BM)).
//
-// Bootstrap keeps this direct/single-buffered by design (no double buffering,
-// no split-K, no raw asm); later rounds tune tile, stage, and packing.
+// K is staged cooperatively in a SINGLE LDS buffer at StageK granularity with
+// TWO __syncthreads per stage. Fragment loads use the library
+// du_load_matrix_sync (row_major for both operands) and accumulate with
+// du_mma_sync in ascending k order -- k0-outer over stages, kk-inner over
+// kTileK steps -- preserving the exact int32 accumulation order of the scalar
+// reference. The epilogue is the coalesced fragment store above.
template
__global__ __launch_bounds__(kThreadsPerBlock) void
w8a8_dumma_prefill_tiled_kernel(
- const int8_t* __restrict__ x_q, // [M, K] row-major
- const int8_t* __restrict__ weight, // [K, N] row-major (identity pack)
- const float* __restrict__ x_scale, // [M]
+ const int8_t* __restrict__ x_q, // [M, K] row-major
+ const int8_t* __restrict__ weight, // [K, N] row-major (identity pack)
+ const float* __restrict__ x_scale, // [M]
const float* __restrict__ weight_scale, // [N]
- hip_bfloat16* __restrict__ out, // [M, N]
+ hip_bfloat16* __restrict__ out, // [M, N]
int m,
int n,
int k) {
@@ -509,269 +420,147 @@ w8a8_dumma_prefill_tiled_kernel(
}
}
-// Explicit 8-byte int8 fragment loader (lineage-validated load_frag8): for
-// both the row_major matrix_a and the col_major matrix_b m16n16k32
-// fragments, lane l -> row (l & 15), k-quarter (l >> 4) holds the EIGHT
-// CONTIGUOUS bytes at p[row*ldm + (l>>4)*8 .. +7]. Loading the 8 bytes as
-// one int64 forces the ds_read_b64 form (16-lane phases, distinct bank
-// pairs per group -> 4-phase conflict-free floor with an odd 8-byte LDS row
-// stride), identical fragment values to the library loader, so the v_mmac
-// inputs and the exact int32 accumulation order are unchanged.
-template
-__device__ __forceinline__ void load_frag8(Frag& f, const int8_t* p,
- unsigned ldm) {
- const unsigned row = __lane_id() & 0xf;
- const unsigned kq = __lane_id() >> 4;
- const int64_t v =
- *reinterpret_cast(p + row * ldm + (kq << 3));
- reinterpret_cast(f.x)[0] = v;
-}
-
-// Iteration-2 exact-shape packed-B prefill kernel (dispatched only for
-// (m >= 128, n == 2624, k == 6144), where launch_pack_w8a8_weight produced
-// the n-major packed_weight[n*K + kk] = raw[kk*N + n] layout). Same
-// geometry as w8a8_dumma_prefill_tiled_kernel<128,64,128> (grid (N/64) x
-// (M/128), 256 threads = 4 wavefronts, 64x32 quadrant per wave, eight
-// m16n16k32 int32 accumulators, single-buffered 128-K LDS stage with TWO
-// __syncthreads per stage, same k0-outer/kk-inner int32 accumulation order,
-// same coalesced fragment epilogue). The two differences, both inside the
-// operand path:
-// * B is staged n-major (b_tile[n][kk] from packed [N, K]) so col_major
-// B fragments read 8 contiguous K bytes per lane;
-// * A and B LDS row strides are 136 B (StageK + 8): 8-byte aligned for
-// the ds_read_b64 fragment loads (odd 8-byte stride count -> the 16
-// rows of each 16-lane phase land on 16 distinct bank pairs, 34r mod 32
-// = 2r mod 32, r = 0..15 -> all 16 even dword starts, conflict-free
-// 4-phase floor) and 136 % 16 == 8 so staging stores are the two-half
-// int64 ds_write2_b64 form (single ds_write_b128 would be misaligned on
-// odd rows).
-// Both fragment loaders are the explicit load_frag8 above; the staged bytes
-// are the same matrix elements in the same fragment slots as the library
-// loaders produced from the identity layout, so every v_mmac input and the
-// exact int32 accumulation (k0-outer over 48 stages, kk-inner ascending
-// over four m16n16k32 steps) are bit-identical to the accepted iteration-1
-// kernel and to the scalar reference.
-template
-__global__ __launch_bounds__(kThreadsPerBlock) void
-w8a8_dumma_prefill_packedb_kernel(
- const int8_t* __restrict__ x_q, // [M, K] row-major
- const int8_t* __restrict__ packed_w, // [N, K] n-major (exact shape)
- const float* __restrict__ x_scale, // [M]
- const float* __restrict__ weight_scale, // [N]
- hip_bfloat16* __restrict__ out, // [M, N]
+// Exact-shape A-staging load. For the dispatched M=4096 every 256-row block
+// tile lies fully inside the output, so the `(g_row < m)` row guard is dead:
+// the exact code object would otherwise lower it to unconditional zero-fills
+// plus EXEC save/restore uniform-branch pairs per K stage, i.e. dead
+// instructions and an EXEC serialization in front of the very
+// global_load_dwordx4s that prefetch the next tile. The kFullTile
+// instantiation removes the select (one coalesced load per staging vector, no
+// zero-fill), while kFullTile == false keeps the zero-fill + row mask
+// semantics identical for every m % BM != 0.
+template
+__device__ __forceinline__ int4 load_a_stage_vec(
+ const int8_t* __restrict__ p,
int m,
- int n,
- int k) {
- constexpr int kWaveM16 = BM / 32;
- constexpr int kWaveN16 = BN / 32;
- static_assert(BM % 32 == 0 && BN % 32 == 0,
- "block tile must be a multiple of the wave quadrant");
- static_assert(kThreadsPerBlock % kWaveSize == 0,
- "blockDim must be a multiple of the gfx928 wavefront size");
- // Bank-safe LDS row stride: StageK + 8 bytes = 34 dwords. 8-aligned with
- // an odd 8-byte count -> every lane's 8-byte fragment is ds_read_b64 at
- // the 4-phase conflict-free floor (see the kernel comment above).
- constexpr int kLdsStride = StageK + 8;
- static_assert(kLdsStride % 8 == 0 && (kLdsStride / 8) % 2 == 1,
- "LDS row stride must be 8-aligned with an odd 8-byte count");
- static_assert(kLdsStride % 16 == 8,
- "LDS row stride must be 8 mod 16 for the two-half staging");
-
- const int tid = static_cast(threadIdx.x);
- const int wave = tid / kWaveSize;
- const int lane = tid % kWaveSize;
- const int wave_row = wave >> 1; // 0 or 1
- const int wave_col = wave & 1; // 0 or 1
- const int m0 = static_cast(blockIdx.y) * BM;
- const int n0 = static_cast(blockIdx.x) * BN;
-
- __shared__ __align__(16) int8_t a_tile[BM * kLdsStride];
- __shared__ __align__(16) int8_t b_tile[BN * kLdsStride];
-
- DUFragment
- a_frag[kWaveM16];
- DUFragment
- b_frag[kWaveN16];
- DUFragment
- acc[kWaveM16][kWaveN16];
-#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
-#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- du_fill_fragment(acc[j][i], 0);
- }
- }
-
- // Flat 16-byte staging vectors per stage (dispatch guarantees
- // k % StageK == 0 and n % BN == 0). 8 vectors per 128-K row.
- constexpr int kAVectors = BM * StageK / static_cast(sizeof(int4));
- constexpr int kBVectors = StageK * BN / static_cast(sizeof(int4));
- constexpr int kVecsPerRow = StageK / static_cast(sizeof(int4));
-
- for (int k0 = 0; k0 < k; k0 += StageK) {
- // Stage A[BM, StageK] row-major with row stride kLdsStride: thread
- // `vec` owns the 16-byte chunk at (local_row, k16); rows past M are
- // zero-filled and masked in the epilogue.
- for (int vec = tid; vec < kAVectors; vec += kThreadsPerBlock) {
- const int local_row = vec / kVecsPerRow;
- const int k16 = vec - local_row * kVecsPerRow;
- const int g_row = m0 + local_row;
- int4 v{0, 0, 0, 0};
- if (g_row < m) {
- v = *reinterpret_cast(
- x_q + static_cast(g_row) * k + k0 + k16 * 16);
- }
- // Two 8-byte halves (stride % 16 == 8 -> odd rows are not 16-byte
- // aligned, so a single int4 store is illegal; the int64 pair lowers
- // to one ds_write2_b64 per thread).
- const int64_t* src64 = reinterpret_cast(&v);
- int64_t* dst64 = reinterpret_cast(
- a_tile + local_row * kLdsStride + k16 * 16);
- dst64[0] = src64[0];
- dst64[1] = src64[1];
- }
- // Stage B[BN, StageK] n-major from the packed [N, K] weight: thread
- // `vec` owns the 16-byte chunk at (n_row, k16); the K axis is always
- // in range.
- for (int vec = tid; vec < kBVectors; vec += kThreadsPerBlock) {
- const int n_row = vec / kVecsPerRow;
- const int k16 = vec - n_row * kVecsPerRow;
- const int4 v = *reinterpret_cast(
- packed_w + static_cast(n0 + n_row) * k + k0 + k16 * 16);
- const int64_t* src64 = reinterpret_cast(&v);
- int64_t* dst64 = reinterpret_cast(
- b_tile + n_row * kLdsStride + k16 * 16);
- dst64[0] = src64[0];
- dst64[1] = src64[1];
- }
- // Make the staged tiles visible to every wavefront.
- __syncthreads();
-
- // Consume the stage: each wave accumulates its (BM/2) x (BN/2) quadrant
- // with kk-ascending m16n16k32 MMAs; one ds_read_b64 per fragment.
- const int local_row = wave_row * (BM / 2);
- const int local_col = wave_col * (BN / 2);
-#pragma unroll
- for (int kk = 0; kk < StageK; kk += kTileK) {
-#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
- load_frag8(
- a_frag[j],
- a_tile + (local_row + j * kTileM) * kLdsStride + kk,
- kLdsStride);
- }
-#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- load_frag8(
- b_frag[i],
- b_tile + (local_col + i * kTileN) * kLdsStride + kk,
- kLdsStride);
- }
-#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
-#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- du_mma_sync(acc[j][i], a_frag[j], b_frag[i], acc[j][i]);
- }
- }
- }
- // Single buffer: every wavefront must finish reading LDS before the
- // next stage's stores overwrite it.
- __syncthreads();
+ int g_row) {
+ if (kFullTile) {
+ return *reinterpret_cast(p);
}
+ return (g_row < m) ? *reinterpret_cast(p) : int4{0, 0, 0, 0};
+}
- const int base_row = m0 + wave_row * (BM / 2);
- const int base_col = n0 + wave_col * (BN / 2);
-#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
-#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- store_prefill_fragment_coalesced(
- acc[j][i], x_scale, weight_scale, out, m, n,
- base_row + j * kTileM, base_col + i * kTileN, lane);
- }
- }
+// Exact-shape A-staging load through the GLOBAL (non-generic) address form:
+// the caller passes the kernel's `x_q` base plus the per-stage byte offset, so
+// the load visits exactly the same bytes as the pointer form but through the
+// global address space (global_load_dwordx4) instead of the generic
+// flat_load_dwordx4 path the pointer form lowers to. Used only by the
+// exact-shape full-tile instantiation.
+__device__ __forceinline__ int4 load_a_stage_vec_global(
+ const int8_t* __restrict__ x_q,
+ int64_t offset) {
+ return *reinterpret_cast(x_q + offset);
}
-// Iteration-3 double-buffered packed-B prefill kernel (dispatched only for
-// (m >= 128, n == 2624, k == 6144)): SAME macro-tile as the accepted
-// iteration-2 w8a8_dumma_prefill_packedb_kernel<128,64,128> (grid (N/64) x
-// (M/128) = 41x32 = 1312 blocks, 256 threads = 4 wavefronts, 64x32 quadrant
-// per wave, eight m16n16k32 int32 accumulators, same load_frag8 fragment
-// reads, same bank-safe LDS row strides, same coalesced epilogue). The ONLY
-// change is the K pipeline (see the file-header iteration-3 section);
-// iteration 5 additionally reads the packed B operand from the B-panel
-// layout (see the file-header iteration-5 section and
-// w8a8_pack_panel_i8_kernel) -- the LDS staging layout, fragment consumes,
-// and accumulation order are untouched:
-// * The single-buffered 128-K stage (two __syncthreads; the exact-source
-// ISA serializes global_load_dwordx4 -> s_waitcnt vmcnt(0) ->
-// ds_write2_b64 -> s_barrier -> consume with no compute overlapping the
-// VMEM round trip) becomes a 64-K DOUBLE-buffered stage with ONE
-// __syncthreads per stage: the stage-(s+1) global loads are issued into
-// 3 int4 staging VGPR per thread BEFORE the stage-s consume, and the
-// prefetched vectors are flushed into the idle buffer after the consume,
-// so the VMEM latency of stage s+1 overlaps the LDS/MMAC work of stage s
-// (validated M=3072 prefetch order, in pure HIP: the vmcnt wait lands on
-// the register data dependency right before the ds_write2_b64, i.e.
-// after the v_mmac burst).
-// * LDS = 2 x (A[128,72] + B[64,72]) = 2 x 13,824 = 27,648 B/block
-// (26,112 -> 27,648 B, +5.9%); 2 x 27,648 = 55,296 <= 65,536 B keeps
-// TWO resident blocks/CU = 8 waves/CU, occupancy unchanged. kLdsStride =
-// 72 = StageK + 8 (8-aligned, odd 8-byte count 9 -> the 16 rows of each
-// 16-lane ds_read_b64 phase land on 16 distinct bank pairs, 18r mod 32,
-// r = 0..15 -> the same 4-phase conflict-free floor; 72 % 16 == 8 keeps
-// the two-half ds_write2_b64 staging form). 12 extra VGPR for the
-// staging registers (~72 -> ~84, still <= 128 for 2 blocks/CU).
-// * k0-outer over 96 ASCENDING 64-K stages with kk-inner ascending over
-// two m16n16k32 steps: the (k0+kk) k-chunk sequence 0, 32, 64, 96, ...
-// is identical to iteration 2 (48 x 128-K stages), so every int32
-// accumulator is bit-identical (mismatch 0 / max_abs_error 0.0).
-// * Barriers per block: 1 (prologue) + 95 (loop) = 96, the same as
-// iteration 2's 2 x 48; the final stage is consumed without a trailing
-// barrier.
-template
-__global__ __launch_bounds__(kThreadsPerBlock) void
+// Double-buffered packed-B prefill kernel, dispatched ONLY for
+// (m >= 128, n == 2624, k == 6144). Geometry: grid ((N+127)/128) x (M/256) =
+// 21 x 16 = 336 blocks, 512 threads = 8 wavefronts, and each wave owns a
+// kRowsPerWave x kColsPerWave = 64 x 64 quadrant of sixteen m16n16k32 int32
+// accumulator fragments (wave_row = wave / kWaveCols, wave_col = wave %
+// kWaveCols, with the balanced 4 x 2 row/column split chosen by
+// balanced_wave_rows), so the block macro-tile is BM x BN = 256 x 128.
+//
+// OPERAND REUSE (the mandated comparison). Per macro-tile every staged A
+// element is consumed by the TWO wave columns (A reuse 2) and every staged B
+// element by the FOUR wave rows (B reuse 4). Direct per-wave global fragment
+// loads cannot replace the LDS stage: a row_major matrix_a fragment phase is
+// lane (row = lane & 15, kq = lane >> 4), so 16 lanes touch 16 rows that are
+// K = 6,144 B apart -- 16 x 32 B sectors for 128 useful bytes, a 4x fetch
+// amplification -- and the swizzled col_major matrix_b fragments would lose
+// their on-chip reuse and multiply their fetch. Global REUSE is the lever, so
+// the M-major macro-tile is widened to BM = 256: A is read once per 128-column
+// tile (21 re-reads) and B once per 256-row tile (16 re-reads), for a total
+// global traffic of M*K*21 + N*K*16 = 786 MB against the 128 x 128 tile's
+// 1,044 MB (-24.7%), at the same 8 waves/CU. The staging loads stay
+// vectorized/coalesced 16-byte-per-thread and the A LDS stride stays the
+// bank-safe 72 B with the flat swizzled B tile. The odd N = 20*128 + 64 is
+// covered by clamping the column tile origin to n - BN, so the last block
+// recomputes 64 overlapping columns (identical int32 accumulation and identical
+// bf16 bytes) instead of predicating a partial tile; every staged B chunk
+// therefore addresses a valid n16 group of the packed weight.
+//
+// The K pipeline is a 64-K DOUBLE-buffered stage: the stage-(s+1) global loads
+// are issued into 2 int4 staging registers per thread BEFORE the stage-s consume
+// and flushed into the idle LDS buffer after it, so the VMEM latency of stage
+// s+1 overlaps the LDS/v_mmac work of stage s (one __syncthreads per stage
+// instead of two). k0-outer runs over 96 ASCENDING 64-K stages with kk-inner
+// ascending over two m16n16k32 steps, i.e. the same (k0 + kk) k-chunk sequence
+// 0, 32, 64, ... as any other staging, so every int32 accumulator is
+// bit-identical.
+//
+// kLdsStride = StageK + 8 = 72 B: 8-aligned with an odd 8-byte count (9) so the
+// 16 rows of each 16-lane ds_read_b64 phase land on 16 distinct bank pairs
+// (18r mod 32, r = 0..15) -- the conflict-free 4-phase floor -- and
+// 72 % 16 == 8 keeps the two-half ds_write2_b64 staging form.
+//
+// The packed B operand is the fragment-interleaved swizzle (see
+// w8a8_pack_bswz_i8_kernel): the LDS b_tile is the SAME [k32][n16][k8][n][8]
+// layout as the pack (2 x (BN/16) 512-byte fragment blocks per 64-K stage), so
+// staging is one flat 16-byte copy per thread and each col_major B fragment is
+// one conflict-free ds_read_b64 (load_frag8_swz). The int32 accumulation order
+// is untouched.
+//
+// kFullTile is true exactly when the caller's m is a multiple of BM, so every
+// block tile is fully inside the output; in that instantiation the A staging
+// loads drop the row guard and the epilogue drops the per-fragment row mask.
+template
+__global__ __launch_bounds__(kBlockThreads) void
w8a8_dumma_prefill_packedb_db_kernel(
- const int8_t* __restrict__ x_q, // [M, K] row-major
- const int8_t* __restrict__ packed_w, // [N/64][K][64] B panels (exact shape, iter 5)
- const float* __restrict__ x_scale, // [M]
+ const int8_t* __restrict__ x_q, // [M, K] row-major
+ const int8_t* __restrict__ packed_w, // [n>>4][k>>5][512 B] swizzled B
+ // fragment blocks
+ const float* __restrict__ x_scale, // [M]
const float* __restrict__ weight_scale, // [N]
- hip_bfloat16* __restrict__ out, // [M, N]
+ hip_bfloat16* __restrict__ out, // [M, N]
int m,
int n,
int k) {
- constexpr int kWaveM16 = BM / 32;
- constexpr int kWaveN16 = BN / 32;
- static_assert(BM % 32 == 0 && BN % 32 == 0,
+ constexpr int kWaves = kBlockThreads / kWaveSize; // 8
+ // Balanced, tile-aware wave split: 4 wave rows x 2 wave columns for the
+ // 256 x 128 macro-tile, i.e. a square 64 x 64 per-wave quadrant.
+ constexpr int kWaveRows =
+ balanced_wave_rows(BM / kTileM, BN / kTileN, kWaves); // 4
+ constexpr int kWaveCols = kWaves / kWaveRows; // 2
+ constexpr int kRowsPerWave = BM / kWaveRows; // 64
+ constexpr int kColsPerWave = BN / kWaveCols; // 64
+ constexpr int kWaveM16 = kRowsPerWave / kTileM; // 4
+ constexpr int kWaveN16 = kColsPerWave / kTileN; // 4
+ static_assert(kWaveRows * kWaveCols == kWaves,
+ "wave split must tile the block exactly");
+ static_assert(BM % kWaveRows == 0 && BN % kWaveCols == 0,
"block tile must be a multiple of the wave quadrant");
- static_assert(kThreadsPerBlock % kWaveSize == 0,
+ static_assert(kBlockThreads % kWaveSize == 0,
"blockDim must be a multiple of the gfx928 wavefront size");
- // Bank-safe LDS row stride: StageK + 8 bytes = 18 dwords for StageK = 64.
constexpr int kLdsStride = StageK + 8;
static_assert(kLdsStride % 8 == 0 && (kLdsStride / 8) % 2 == 1,
"LDS row stride must be 8-aligned with an odd 8-byte count");
static_assert(kLdsStride % 16 == 8,
"LDS row stride must be 8 mod 16 for the two-half staging");
- // Iteration 5: the packed B weight is the panel layout (see
- // w8a8_pack_panel_i8_kernel), which is defined for BN == 64 columns per
- // panel and 16-byte k-runs (StageK % 16 == 0). Only <128,64,64> is ever
- // dispatched by the exact-shape guard, so the panel index math below is
- // exact for every dispatch.
- static_assert(BN == 64 && StageK % 16 == 0,
- "panel B pack requires BN == 64 and StageK % 16 == 0");
+ // The packed B weight is the fragment-interleaved swizzle: one 32-k x 16-n
+ // fragment block per 512 B; a 64-K stage holds 2 k32 blocks x (BN/16) n16
+ // groups x 4 k8 groups x 8 n-pairs = BN*4 16-byte chunks. With 512 threads
+ // and BN = 128 that is exactly one chunk per thread.
+ static_assert(StageK == 64, "swizzled B pack requires StageK == 64");
+ static_assert(BN % 16 == 0, "swizzled B pack requires BN % 16 == 0");
const int tid = static_cast(threadIdx.x);
const int wave = tid / kWaveSize;
const int lane = tid % kWaveSize;
- const int wave_row = wave >> 1; // 0 or 1
- const int wave_col = wave & 1; // 0 or 1
+ const int wave_row = wave / kWaveCols; // 0 .. kWaveRows-1
+ const int wave_col = wave % kWaveCols; // 0 .. kWaveCols-1
const int m0 = static_cast(blockIdx.y) * BM;
- const int n0 = static_cast(blockIdx.x) * BN;
+ // Clamp the column origin so the odd (n % BN) columns are covered by an
+ // overlapping tile: every staged B chunk then addresses a valid n16 group of
+ // the packed weight, and the duplicated columns recompute byte-identically.
+ const int n0i = static_cast(blockIdx.x) * BN;
+ const int n0 = (n0i + BN <= n) ? n0i : (n - BN);
const int kStages = k / StageK; // dispatch guarantees k % StageK == 0
+ const int n16g = n >> 4; // 16-n row groups per 32-k block of the pack
__shared__ __align__(16) int8_t a_tile[2][BM * kLdsStride];
- __shared__ __align__(16) int8_t b_tile[2][BN * kLdsStride];
+ // Swizzled B tile: flat [k32][n16][k8][n_pair][16 B], byte-identical to the
+ // packed global layout of that stage: 2 x (BN/16) x 512 = BN x 64 B.
+ __shared__ __align__(16) int8_t b_tile[2][BN * StageK];
DUFragment
a_frag[kWaveM16];
@@ -787,55 +576,87 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
}
- // Flat 16-byte staging vectors per stage. For the dispatched <128,64,64>
- // instantiation the 512 A vectors and 256 B vectors divide evenly over the
- // 256 threads (2 + 1 per thread), so the prefetch payload is 3 int4
- // staging registers held live across the consume.
+ // Flat 16-byte staging vectors per stage. For the dispatched
+ // <256, 128, 64, 512> instantiation the 1024 A vectors (2 per thread) and the
+ // 512 B vectors (1 per thread) map exactly onto the 512 threads, so the
+ // prefetch payload is 3 int4 staging registers held live across the consume.
constexpr int kVecsPerRow = StageK / static_cast(sizeof(int4));
constexpr int kAVectors = BM * StageK / static_cast(sizeof(int4));
constexpr int kBVectors = StageK * BN / static_cast(sizeof(int4));
- constexpr int kAVecPerThread = kAVectors / kThreadsPerBlock;
- constexpr int kBVecPerThread = kBVectors / kThreadsPerBlock;
- static_assert(kAVectors % kThreadsPerBlock == 0 &&
- kBVectors % kThreadsPerBlock == 0,
+ constexpr int kAVecPerThread = kAVectors / kBlockThreads;
+ constexpr int kBVecPerThread = kBVectors / kBlockThreads;
+ static_assert(kAVectors % kBlockThreads == 0 &&
+ kBVectors % kBlockThreads == 0,
"staging vectors must divide evenly across the block");
+ // The swizzled B stage is exactly one 16-byte chunk per thread, so thread
+ // `vec` = tid is also the flat [k32][n16][k8][n_pair] chunk index.
+ static_assert(kBVecPerThread == 1 && kBVectors == kBlockThreads,
+ "swizzled B staging must be one 16-B chunk per thread");
int4 a_pre[kAVecPerThread];
int4 b_pre[kBVecPerThread];
- // ---- Prologue: issue the stage-0 global loads, flush them into buffer 0
- // and make the tile visible to every wavefront.
+ // Per-thread GLOBAL staging byte pointers, held as loop-carried induction
+ // variables: the prologue loads stage 0 from them, they are then advanced to
+ // stage 1, and every steady-state iteration issues its prefetch from them and
+ // walks them by the compile-time stage stride -- StageK bytes for A and one
+ // 64-K swizzle band `(StageK>>5)*n16g*512` for B.
+ const int8_t* a_src[kAVecPerThread];
+ int a_grow[kAVecPerThread];
+#pragma unroll
+ for (int v = 0; v < kAVecPerThread; ++v) {
+ const int vec = tid + v * kBlockThreads;
+ const int local_row = vec / kVecsPerRow;
+ const int k16 = vec - local_row * kVecsPerRow;
+ a_grow[v] = m0 + local_row;
+ a_src[v] = x_q + static_cast(a_grow[v]) * k + k16 * 16;
+ }
+ const int8_t* b_src[kBVecPerThread];
+#pragma unroll
+ for (int v = 0; v < kBVecPerThread; ++v) {
+ // Chunk (k32, n16, k8, n_pair) of the stage: BN*2 chunks per k32 block
+ // (4 k8 groups x 8 n-pairs per n16), 32 chunks per n16 group.
+ const int vec = tid + v * kBlockThreads;
+ const int k32 = vec / (BN * 2);
+ const int rem = vec - k32 * (BN * 2);
+ b_src[v] = packed_w +
+ packed_b_swz_offset(
+ k32, (n0 >> 4) + ((rem >> 5) & (BN / 16 - 1)), n16g,
+ (rem >> 3) & 3, rem & 7);
+ }
+
+ // ---- Prologue: issue the stage-0 global loads, flush them into buffer 0 and
+ // make the tile visible to every wavefront.
{
- const int k0 = 0;
#pragma unroll
for (int v = 0; v < kAVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
+ const int vec = tid + v * kBlockThreads;
const int local_row = vec / kVecsPerRow;
const int k16 = vec - local_row * kVecsPerRow;
const int g_row = m0 + local_row;
- a_pre[v] = (g_row < m)
- ? *reinterpret_cast(
- x_q + static_cast(g_row) * k + k0 +
- k16 * 16)
- : int4{0, 0, 0, 0};
+ if (kFullTile) {
+ // Exact-shape full-tile A staging through the global (non-generic)
+ // address form -- the same bytes as the pointer form, but through the
+ // load path that emits global_load_dwordx4.
+ a_pre[v] = load_a_stage_vec_global(
+ x_q, static_cast(g_row) * k + k16 * 16);
+ } else {
+ a_pre[v] = load_a_stage_vec(a_src[v], m, a_grow[v]);
+ }
}
#pragma unroll
for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int n_row = vec / kVecsPerRow;
- const int k16 = vec - n_row * kVecsPerRow;
- // Iteration 5: B is read from the panel-packed layout; the [64, 64]
- // tile is 4 contiguous 1-KiB slabs (one per 16-k-run) and vector
- // (n_row, k16) sits at byte (n0>>6)*(k*64) + ((k0>>4)+k16)*1024 +
- // n_row*16 (16-byte aligned, contiguous k-run of column n0+n_row).
- const int64_t panel_off =
- static_cast(n0 >> 6) * (static_cast(k) * 64) +
- (static_cast(k0 >> 4) + k16) * 1024 + n_row * 16;
- b_pre[v] = *reinterpret_cast(packed_w + panel_off);
+ // The 16-byte chunk (k32 = vec/(BN*2), n16 = (vec>>5)&(BN/16-1),
+ // k8 = (vec>>3)&3, n_pair = vec&7) of the swizzled pack sits at byte
+ // (k0>>5 + k32)*n16g*512 + ((n0>>4)+n16)*512 + k8*128 + n_pair*16; the
+ // per-thread pointer already carries the k32/n16/k8/n_pair part for
+ // k0 == 0. Consecutive lanes read consecutive 16-byte chunks, so a warp's
+ // staging loads touch 8 fully-used 128-byte lines.
+ b_pre[v] = *reinterpret_cast(b_src[v]);
}
#pragma unroll
for (int v = 0; v < kAVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
+ const int vec = tid + v * kBlockThreads;
const int local_row = vec / kVecsPerRow;
const int k16 = vec - local_row * kVecsPerRow;
const int64_t* src64 = reinterpret_cast(&a_pre[v]);
@@ -846,118 +667,157 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
#pragma unroll
for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int n_row = vec / kVecsPerRow;
- const int k16 = vec - n_row * kVecsPerRow;
- const int64_t* src64 = reinterpret_cast(&b_pre[v]);
- int64_t* dst64 = reinterpret_cast(
- b_tile[0] + n_row * kLdsStride + k16 * 16);
- dst64[0] = src64[0];
- dst64[1] = src64[1];
+ const int vec = tid + v * kBlockThreads;
+ // The swizzled tile is a flat copy of the packed stage, so the commit is
+ // ONE 16-byte-aligned int4 (ds_write_b128) per thread.
+ *reinterpret_cast(b_tile[0] + vec * 16) = b_pre[v];
}
}
__syncthreads();
- // ---- Steady state: for s = 0..kStages-2 issue the stage-(s+1) global
- // loads into the staging registers BEFORE consuming stage s from buffer
- // s&1, then flush the prefetched vectors into the idle buffer (s+1)&1 and
- // close the stage with ONE __syncthreads (the barrier's vmcnt wait lands
- // after the consume, so the VMEM latency of stage s+1 overlaps the
- // LDS/MMAC work of stage s). The last stage is consumed after the loop
- // without a trailing barrier.
- for (int s = 0; s < kStages - 1; ++s) {
- // Prefetch stage s+1: issue the global loads (kept in VGPR while the
- // stage-s consume runs).
- const int k0 = (s + 1) * StageK;
+ // Advance the staging pointers to stage 1. From here on every steady-state
+ // half prefetches stage s+1 from these pointers and then walks them by the
+ // compile-time stage stride, so both global staging addresses are pure
+ // pointer increments instead of a per-stage address expression.
+ const int64_t b_step = static_cast(StageK >> 5) * n16g * 512;
#pragma unroll
- for (int v = 0; v < kAVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int local_row = vec / kVecsPerRow;
- const int k16 = vec - local_row * kVecsPerRow;
- const int g_row = m0 + local_row;
- a_pre[v] = (g_row < m)
- ? *reinterpret_cast(
- x_q + static_cast(g_row) * k + k0 +
- k16 * 16)
- : int4{0, 0, 0, 0};
- }
+ for (int v = 0; v < kAVecPerThread; ++v) {
+ a_src[v] += StageK;
+ }
#pragma unroll
- for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int n_row = vec / kVecsPerRow;
- const int k16 = vec - n_row * kVecsPerRow;
- // Iteration 5: panel-packed B tile read (see the prologue comment).
- const int64_t panel_off =
- static_cast(n0 >> 6) * (static_cast(k) * 64) +
- (static_cast(k0 >> 4) + k16) * 1024 + n_row * 16;
- b_pre[v] = *reinterpret_cast(packed_w + panel_off);
- }
+ for (int v = 0; v < kBVecPerThread; ++v) {
+ b_src[v] += b_step;
+ }
- // Consume stage s from buffer s&1: each wave accumulates its (BM/2) x
- // (BN/2) quadrant with kk-ascending m16n16k32 MMAs; one ds_read_b64 per
- // fragment.
- const int local_row = wave_row * (BM / 2);
- const int local_col = wave_col * (BN / 2);
+ // ---- Steady state (double-buffer parity specialization). For stage s, issue
+ // the stage-(s+1) global loads into the staging registers BEFORE consuming
+ // stage s from buffer s&1, then flush the prefetched vectors into the idle
+ // buffer (s+1)&1 and close the stage with ONE __syncthreads (so the vmcnt wait
+ // for stage s+1 lands after the MMAC burst of stage s, overlapping the VMEM
+ // latency of s+1 with the stage-s LDS/v_mmac work). The last stage is consumed
+ // after the loop without a trailing barrier.
+ //
+ // Walking the stages in compile-time PAIRS makes each half's parity a
+ // constant (s = 2*sp + h with h unrolled, so s & 1 == h and (s+1) & 1 == h^1),
+ // which folds the double-buffer parity offsets into the per-half LDS base
+ // registers / instruction immediates. The odd leftover for the dispatched
+ // k = 6144 (kStages = 96 -> kSteady = 95, one unpaired stage) is guarded by ONE
+ // uniform compare in the second half of the last pair, so correctness does not
+ // depend on the unroll succeeding.
+ const int kSteady = kStages - 1;
+ const int kPairIters = (kSteady + 1) >> 1;
+ for (int sp = 0; sp < kPairIters; ++sp) {
#pragma unroll
- for (int kk = 0; kk < StageK; kk += kTileK) {
+ for (int h = 0; h < 2; ++h) {
+ const int s = 2 * sp + h;
+ if (h == 0 || s < kSteady) {
+ // Compile-time buffer parity: cur == s & 1, nxt == (s + 1) & 1.
+ const int cur = h;
+ const int nxt = h ^ 1;
+ // Prefetch stage s+1: issue the global loads (held in VGPRs while the
+ // stage-s consume runs) from the loop-carried per-thread stage
+ // pointers, which already sit at the stage-(s+1) addresses. The
+ // exact-shape full-tile instantiation forms the two A staging addresses
+ // in the `x_q + ` global form (so the backend
+ // keeps the non-generic global_load path), while B keeps the
+ // loop-carried staging pointer.
+ const int k0 = (s + 1) * StageK;
#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
- load_frag8(
- a_frag[j],
- a_tile[s & 1] + (local_row + j * kTileM) * kLdsStride + kk,
- kLdsStride);
- }
+ for (int v = 0; v < kAVecPerThread; ++v) {
+ const int vec = tid + v * kBlockThreads;
+ const int local_row = vec / kVecsPerRow;
+ const int k16 = vec - local_row * kVecsPerRow;
+ const int g_row = m0 + local_row;
+ if (kFullTile) {
+ a_pre[v] = load_a_stage_vec_global(
+ x_q, static_cast(g_row) * k + k0 + k16 * 16);
+ } else {
+ a_pre[v] = load_a_stage_vec(a_src[v], m, a_grow[v]);
+ }
+ }
#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- load_frag8(
- b_frag[i],
- b_tile[s & 1] + (local_col + i * kTileN) * kLdsStride + kk,
- kLdsStride);
- }
+ for (int v = 0; v < kBVecPerThread; ++v) {
+ // One plain 16-byte staging load per thread from the stage-(s+1)
+ // swizzle chunk; consecutive lanes read consecutive 16-byte chunks,
+ // so a warp's staging loads touch 8 fully-used 128-byte lines.
+ b_pre[v] = *reinterpret_cast(b_src[v]);
+ }
+
+ // Consume stage s from buffer cur: each wave accumulates its
+ // kRowsPerWave x kColsPerWave quadrant with kk-ascending m16n16k32 MMAs;
+ // one ds_read_b64 per fragment.
+ const int local_row = wave_row * kRowsPerWave;
#pragma unroll
- for (int j = 0; j < kWaveM16; ++j) {
+ for (int kk = 0; kk < StageK; kk += kTileK) {
#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- du_mma_sync(acc[j][i], a_frag[j], b_frag[i], acc[j][i]);
+ for (int j = 0; j < kWaveM16; ++j) {
+ load_frag8(
+ a_frag[j],
+ a_tile[cur] + (local_row + j * kTileM) * kLdsStride + kk,
+ static_cast(kLdsStride),
+ lane);
+ }
+#pragma unroll
+ for (int i = 0; i < kWaveN16; ++i) {
+ // Fragment block (k32 = kk>>5, n16 = wave_col*kWaveN16 + i) of the
+ // swizzled tile -- 512 B per 16-n x 32-k DUMMA B fragment, one
+ // conflict-free ds_read_b64 per lane.
+ load_frag8_swz(
+ b_frag[i],
+ b_tile[cur] +
+ ((kk >> 5) * (BN / 16) + wave_col * kWaveN16 + i) * 512,
+ lane);
+ }
+#pragma unroll
+ for (int j = 0; j < kWaveM16; ++j) {
+#pragma unroll
+ for (int i = 0; i < kWaveN16; ++i) {
+ du_mma_sync(acc[j][i], a_frag[j], b_frag[i], acc[j][i]);
+ }
+ }
}
- }
- }
- // Flush the prefetched stage-(s+1) vectors into the idle buffer.
- {
- const int buf = (s + 1) & 1;
+ // Flush the prefetched stage-(s+1) vectors into the idle buffer nxt.
+ {
#pragma unroll
- for (int v = 0; v < kAVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int local_row = vec / kVecsPerRow;
- const int k16 = vec - local_row * kVecsPerRow;
- const int64_t* src64 = reinterpret_cast(&a_pre[v]);
- int64_t* dst64 = reinterpret_cast(
- a_tile[buf] + local_row * kLdsStride + k16 * 16);
- dst64[0] = src64[0];
- dst64[1] = src64[1];
- }
+ for (int v = 0; v < kAVecPerThread; ++v) {
+ const int vec = tid + v * kBlockThreads;
+ const int local_row = vec / kVecsPerRow;
+ const int k16 = vec - local_row * kVecsPerRow;
+ const int64_t* src64 = reinterpret_cast(&a_pre[v]);
+ int64_t* dst64 = reinterpret_cast(
+ a_tile[nxt] + local_row * kLdsStride + k16 * 16);
+ dst64[0] = src64[0];
+ dst64[1] = src64[1];
+ }
+#pragma unroll
+ for (int v = 0; v < kBVecPerThread; ++v) {
+ const int vec = tid + v * kBlockThreads;
+ // Flat 16-byte swizzled-tile commit (ds_write_b128).
+ *reinterpret_cast(b_tile[nxt] + vec * 16) = b_pre[v];
+ }
+ }
+ // Walk the loop-carried staging pointers to the next stage. StageK is a
+ // compile-time constant, so each A pointer advances with a single 64-bit
+ // add pair; b_step is the loop-invariant swizzle band.
#pragma unroll
- for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- const int n_row = vec / kVecsPerRow;
- const int k16 = vec - n_row * kVecsPerRow;
- const int64_t* src64 = reinterpret_cast(&b_pre[v]);
- int64_t* dst64 = reinterpret_cast(
- b_tile[buf] + n_row * kLdsStride + k16 * 16);
- dst64[0] = src64[0];
- dst64[1] = src64[1];
+ for (int v = 0; v < kAVecPerThread; ++v) {
+ a_src[v] += StageK;
+ }
+#pragma unroll
+ for (int v = 0; v < kBVecPerThread; ++v) {
+ b_src[v] += b_step;
+ }
+ __syncthreads();
}
}
- __syncthreads();
}
// ---- Tail: consume the last stage (its data was flushed before the last
// barrier; no trailing barrier needed).
{
const int s = kStages - 1;
- const int local_row = wave_row * (BM / 2);
- const int local_col = wave_col * (BN / 2);
+ const int local_row = wave_row * kRowsPerWave;
#pragma unroll
for (int kk = 0; kk < StageK; kk += kTileK) {
#pragma unroll
@@ -965,14 +825,16 @@ w8a8_dumma_prefill_packedb_db_kernel(
load_frag8(
a_frag[j],
a_tile[s & 1] + (local_row + j * kTileM) * kLdsStride + kk,
- kLdsStride);
+ static_cast(kLdsStride),
+ lane);
}
#pragma unroll
for (int i = 0; i < kWaveN16; ++i) {
- load_frag8(
+ load_frag8_swz(
b_frag[i],
- b_tile[s & 1] + (local_col + i * kTileN) * kLdsStride + kk,
- kLdsStride);
+ b_tile[s & 1] +
+ ((kk >> 5) * (BN / 16) + wave_col * kWaveN16 + i) * 512,
+ lane);
}
#pragma unroll
for (int j = 0; j < kWaveM16; ++j) {
@@ -984,13 +846,15 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
}
- const int base_row = m0 + wave_row * (BM / 2);
- const int base_col = n0 + wave_col * (BN / 2);
+ const int base_row = m0 + wave_row * kRowsPerWave;
+ const int base_col = n0 + wave_col * kColsPerWave;
#pragma unroll
for (int j = 0; j < kWaveM16; ++j) {
#pragma unroll
for (int i = 0; i < kWaveN16; ++i) {
- store_prefill_fragment_coalesced(
+ // kGuarded == !kFullTile -- the full-tile instantiation skips the dead
+ // per-fragment `row >= m` EXEC bracket.
+ store_prefill_fragment_coalesced(
acc[j][i], x_scale, weight_scale, out, m, n,
base_row + j * kTileM, base_col + i * kTileN, lane);
}
@@ -998,24 +862,33 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
// Generic scalar int8/int32 fallback for unmatched shapes and small-M API
-// cases (including the paired M=2 and M=16 shapes with the same (N, K) --
-// they never reach the m >= 128 tiled path). One output element per
-// grid-stride step; exact int32 accumulation (max assigned K = 6144 keeps
-// the int8 dot well within int32 range: |dot| <= 6144*127*127 < 2^31), then
-// the fused float scale (dot * x_scale[m] * weight_scale[n]) and bf16 store.
-// The bootstrap pack is the identity copy, so packed_weight is always the
-// logical [K, N] row-major layout here.
+// cases (including the paired M=2 and M=16 shapes with the same (N, K) -- they
+// never reach the m >= 128 tiled/packed paths). One output element per
+// grid-stride step; exact int32 accumulation (max assigned K = 6144 keeps the
+// int8 dot well within int32 range: |dot| <= 6144*127*127 < 2^31), then the
+// fused float scale (dot * x_scale[m] * weight_scale[n]) and a manual uint16
+// bf16 store.
+//
+// For the exact (k, n) == (6144, 2624) the packed_weight buffer holds the
+// fragment-interleaved swizzle (produced once, outside the timed region), so
+// the fallback decodes that layout there with ascending kk (every other (K, N)
+// keeps the identity [K, N] row-major copy):
+// packed[((kk>>5)*n16g + (col>>4))*512 +
+// (((kk>>3)&3)*16 + (col&15))*8 + (kk&7)], n16g = n>>4,
+// element-by-element in ascending kk -- identical elements, identical int32
+// accumulators, so the paired M=2 / M=16 API shapes stay byte-exact.
__global__ __launch_bounds__(256) void w8a8_scalar_gemm_kernel(
- const int8_t* __restrict__ a, // [M, K] row-major
- const int8_t* __restrict__ b, // [K, N] row-major (identity pack)
- const float* __restrict__ x_scale, // [M]
+ const int8_t* __restrict__ a, // [M, K] row-major
+ const int8_t* __restrict__ b, // packed weight (swizzled/identity)
+ const float* __restrict__ x_scale, // [M]
const float* __restrict__ weight_scale, // [N]
- hip_bfloat16* __restrict__ out, // [M, N]
+ hip_bfloat16* __restrict__ out, // [M, N]
int m,
int n,
int k) {
const int64_t total = static_cast(m) * n;
const int64_t stride = static_cast(gridDim.x) * blockDim.x;
+ const bool packed_swz = (k == kExactK && n == kExactN);
for (int64_t idx = static_cast(blockIdx.x) * blockDim.x +
threadIdx.x;
idx < total;
@@ -1023,38 +896,29 @@ __global__ __launch_bounds__(256) void w8a8_scalar_gemm_kernel(
const int row = static_cast(idx / n);
const int col = static_cast(idx - static_cast(row) * n);
const int8_t* a_row = a + static_cast(row) * k;
- // Iteration 5: for the exact (k, n) == (6144, 2624) the packed_weight
- // buffer holds the B-panel layout packed[(col>>6)*(k*64) + (kk>>4)*1024
- // + (col&63)*16 + (kk&15)] = raw[kk*n + col] (41 panels of [K, 64] with
- // 16-byte k-runs per column, produced once out of the timed region), so
- // the fallback decodes that layout there in ascending-kk 16-byte runs
- // (keeps the paired M=2/M=16 API shapes with the same (N, K)
- // byte-exact); every other (k, n) keeps the identity [K, N] row-major
- // copy. The int32 accumulation stays kk-ascending 0..k-1 in both arms.
- const bool packed_panel = (k == kExactK && n == kExactN);
int32_t acc = 0;
- if (packed_panel) {
- const int8_t* panel =
- b + static_cast(col >> 6) * (static_cast(k) * 64) +
- (col & 63) * 16;
- for (int kkb = 0; kkb < k; kkb += 16) {
- const int8_t* p = panel + (kkb >> 4) * 1024;
- for (int bb = 0; bb < 16; ++bb) {
- acc += static_cast(a_row[kkb + bb]) *
- static_cast(p[bb]);
- }
+ if (packed_swz) {
+ // Swizzle: one 8-byte k-run per (n, k8) inside a 512-byte 32-k x 16-n
+ // fragment block; walk it in ascending kk.
+ const int64_t n16g = static_cast(n) >> 4;
+ for (int kk = 0; kk < k; ++kk) {
+ const int64_t off =
+ ((static_cast(kk >> 5) * n16g + (col >> 4)) << 9) +
+ ((((kk >> 3) & 3) * 16 + (col & 15)) << 3) + (kk & 7);
+ acc += static_cast(a_row[kk]) *
+ static_cast(b[off]);
}
} else {
const int8_t* b_ptr = b + col;
for (int kk = 0; kk < k; ++kk) {
- acc += static_cast(a_row[kk]) *
- static_cast(b_ptr[0]);
+ acc += static_cast(a_row[kk]) * static_cast(b_ptr[0]);
b_ptr += n;
}
}
- const float scaled = static_cast(acc) *
- x_scale[row] * weight_scale[col];
- out[idx] = __float2bfloat16(scaled);
+ const float scaled =
+ static_cast(acc) * x_scale[row] * weight_scale[col];
+ reinterpret_cast(out)[idx] =
+ static_cast(__float2bfloat16(scaled));
}
}
@@ -1073,61 +937,49 @@ __global__ __launch_bounds__(256) void w8a8_pack_identity_i8_kernel(
}
}
-// Exact-shape n-major pack (iteration 2): for (k, n) == (6144, 2624) the
-// weight is transposed once, outside the timed region and out of Graph
-// capture, into packed[n * k + kk] = raw[kk * n + n_row] so every col_major
-// B fragment is 8 contiguous K bytes in the packed buffer. Same byte count
-// and buffer as the identity pack, so captured addresses are unchanged. One
-// thread per output byte; runs once during weight prep. (Kept compiled as
-// the iteration-2/3 layout control; iteration 5 dispatches the panel pack
-// below for the exact shape.)
-__global__ __launch_bounds__(256) void w8a8_pack_nmajor_i8_kernel(
- const int8_t* __restrict__ raw,
- int8_t* __restrict__ packed,
- int k,
- int n) {
- const int64_t total = static_cast(k) * n;
- const int64_t stride = static_cast(gridDim.x) * blockDim.x;
- for (int64_t idx = static_cast(blockIdx.x) * blockDim.x +
- threadIdx.x;
- idx < total;
- idx += stride) {
- const int col = static_cast(idx / k);
- const int kk = static_cast(idx - static_cast(col) * k);
- packed[idx] = raw[static_cast(kk) * n + col];
- }
-}
-
-// Exact-shape B-panel pack (iteration 5): for (k, n) == (6144, 2624) the
+// Exact-shape fragment-interleaved B swizzle: for (k, n) == (6144, 2624) the
// weight is packed once, outside the timed region and out of Graph capture,
-// into 41 panels of [K, 64] int8 with 16-byte k-runs per column:
-// packed[(n >> 6) * (k * 64) + (kk >> 4) * 1024 + (n & 63) * 16 +
-// (kk & 15)] = raw[kk * n + n]
-// (the byte-exact bijection 41 x 384 x 64 x 16 = k*n is verified by
-// simulation; same byte count and buffer, so captured addresses are
-// unchanged). Every [64, 64] B tile of a 64-K stage is then 4 CONTIGUOUS
-// 1-KiB slabs (one per 16-k-run) and every 16-byte staging vector is a
-// 16-byte-aligned contiguous k-run of one column, so each DUMMA B tile is
-// vector-loadable as one coalesced stream with full 128-B line utilization
-// (vs 64 scattered 64-B chunks at 6144-B stride in the n-major layout).
-// One thread per output byte; runs once during weight prep.
-__global__ __launch_bounds__(256) void w8a8_pack_panel_i8_kernel(
+// into 512-byte 32-k x 16-n DUMMA B fragment blocks laid out [k32][n16][k8][n][8]:
+// packed[((kk >> 5) * n16g + (n >> 4)) * 512 +
+// (((kk >> 3) & 3) * 16 + (n & 15)) * 8 + (kk & 7)] = raw[kk * n + n],
+// where n16g = n >> 4 (164 for the exact shape). This is a byte-exact bijection
+// 192 x 164 x 512 = k*n = 16,121,856: same byte count and buffer as the identity
+// pack, so captured addresses are unchanged. Each thread packs one 16-byte chunk
+// = the 8-byte k-runs of two n-rows of one k8 group (a strided 2x8 gather on the
+// read side; packing is outside timing, so the gather is off the critical
+// path). Consecutive chunks are consecutive packed bytes, so the kernel's B
+// staging is a flat 16-byte copy and every col_major B fragment read is one
+// conflict-free ds_read_b64 (see load_frag8_swz).
+__global__ __launch_bounds__(256) void w8a8_pack_bswz_i8_kernel(
const int8_t* __restrict__ raw,
int8_t* __restrict__ packed,
int k,
int n) {
- const int64_t total = static_cast(k) * n;
+ const int64_t chunks = (static_cast(k) * n) >> 4;
const int64_t stride = static_cast(gridDim.x) * blockDim.x;
- for (int64_t idx = static_cast(blockIdx.x) * blockDim.x +
- threadIdx.x;
- idx < total;
- idx += stride) {
- const int col = static_cast(idx / k);
- const int kk = static_cast(idx - static_cast(col) * k);
- // (col >> 6) < n/64 <= 41, (kk >> 4) < k/16 <= 384, (col & 63) < 64,
- // (kk & 15) < 16 -> the mixed-radix index stays < k*n in int.
- packed[(col >> 6) * (k * 64) + (kk >> 4) * 1024 + (col & 63) * 16 +
- (kk & 15)] = raw[static_cast(kk) * n + col];
+ const int n16g = n >> 4;
+ for (int64_t c = static_cast(blockIdx.x) * blockDim.x +
+ threadIdx.x;
+ c < chunks;
+ c += stride) {
+ const int k32 =
+ static_cast(c / (static_cast(n16g) * 32));
+ const int rem =
+ static_cast(c - static_cast(k32) * n16g * 32);
+ const int n16 = rem >> 5;
+ const int w = rem & 31;
+ const int k8 = w >> 3;
+ const int np = w & 7;
+ const int n0 = (n16 << 4) + (np << 1);
+ const int k0 = (k32 << 5) + (k8 << 3);
+ int8_t* dst16 = packed + (c << 4);
+#pragma unroll
+ for (int i = 0; i < 8; ++i) {
+ // 16-byte chunk = [n0 row: k0..k0+7][n0+1 row: k0..k0+7] (n-major within
+ // the k8 group, matching the [k8][n][8] tile layout).
+ dst16[i] = raw[static_cast(k0 + i) * n + n0];
+ dst16[8 + i] = raw[static_cast(k0 + i) * n + n0 + 1];
+ }
}
}
@@ -1148,8 +1000,10 @@ __global__ __launch_bounds__(256) void w8a8_pack_identity_f32_kernel(
} // namespace
// Stable host launch symbol consumed by csrc/bindings.cpp. Dispatches on the
-// caller-provided stream only; never allocates, synchronizes, packs, or
-// touches the default stream.
+// caller-provided stream only; never allocates, synchronizes, packs, or touches
+// the default stream. The workspace is unused for this shape (no split-K):
+// m*n*4 = 42,990,080 B exceeds the 16 MiB workspace budget, so
+// workspace_split_k_capacity(m, n, k) is zero.
extern "C" void launch_w8a8_gemm(
const int8_t* a,
const int8_t* b,
@@ -1170,96 +1024,60 @@ extern "C" void launch_w8a8_gemm(
auto* out_ptr = static_cast(out);
// Exact assigned shape guard: glm_tp8_fused_qkv_a_proj_m4096
- // (M=4096, N=2624, K=6144), m >= 128 so M=2/M=16 paired API shapes with
- // the same (N, K) keep reaching the scalar fallback below. Iteration 5:
- // dispatch the exact-fit 2-D macro-tile <128, 64, 128> geometry with the
- // B-PANEL packed weight (launch_pack_w8a8_weight packs (6144, 2624) once,
- // out of the timed region, into packed[(n>>6)*(k*64) + (kk>>4)*1024 +
- // (n&63)*16 + (kk&15)] = raw[kk*N + n] so every 64-K B tile is 4
- // contiguous 1-KiB slabs), the bank-safe LDS row strides, the explicit
- // 8-byte load_frag8 fragment loads, and the double-buffered 64-K K
- // pipeline (one __syncthreads per stage, stage-(s+1) global loads
- // prefetched into staging VGPR before the stage-s consume) -- the
- // mandated packing/swizzle round at unchanged tile/occupancy (2
- // blocks/CU, 27,648 B LDS/block). The single-buffered iteration-2
- // control kernel stays compiled (see the m < 0 forcing block below) for
- // the A/B and for rollback.
+ // (M=4096, N=2624, K=6144). The m >= 128 guard keeps the paired M=2/M=16 API
+ // shapes with the same (N, K) on the scalar fallback below. M = 16*256 is
+ // exact and N = 2624 = 20*128 + 64 is covered by clamping the last column
+ // tile origin to n - 128, so every block still stages a full 128-column B
+ // macro-tile and no B chunk ever leaves the packed weight. The compile-time
+ // kFullTile specialization is selected when m is a multiple of BM = 256, and
+ // every other m with this (n, k) keeps the guarded instantiation.
if (m >= 128 && n == kExactN && k == kExactK) {
const dim3 grid(
- static_cast(n / 64),
- static_cast((m + 127) / 128));
- const dim3 block(static_cast(kThreadsPerBlock));
- hipLaunchKernelGGL(
- HIP_KERNEL_NAME(w8a8_dumma_prefill_packedb_db_kernel<128, 64, 64>),
- grid,
- block,
- 0,
- stream,
- a,
- b,
- x_scale,
- weight_scale,
- out_ptr,
- m,
- n,
- k);
- return;
- }
-
- // Force the sibling tile instantiations to compile (<64,128,128> is the
- // next benchmark candidate; this branch can never run because m <= 0
- // already returned above). The single-buffered iteration-2 packed-B kernel
- // is kept here so the control arm of the iteration-3 A/B stays compiled
- // (and rollback only needs the guard repointed).
- if (m < 0) {
- hipLaunchKernelGGL(
- HIP_KERNEL_NAME(w8a8_dumma_prefill_tiled_kernel<64, 128, 128>),
- dim3(1),
- dim3(static_cast(kThreadsPerBlock)),
- 0,
- stream,
- a,
- b,
- x_scale,
- weight_scale,
- out_ptr,
- m,
- n,
- k);
- hipLaunchKernelGGL(
- HIP_KERNEL_NAME(w8a8_dumma_prefill_tiled_kernel<128, 64, 128>),
- dim3(1),
- dim3(static_cast(kThreadsPerBlock)),
- 0,
- stream,
- a,
- b,
- x_scale,
- weight_scale,
- out_ptr,
- m,
- n,
- k);
- hipLaunchKernelGGL(
- HIP_KERNEL_NAME(w8a8_dumma_prefill_packedb_kernel<128, 64, 128>),
- dim3(1),
- dim3(static_cast(kThreadsPerBlock)),
- 0,
- stream,
- a,
- b,
- x_scale,
- weight_scale,
- out_ptr,
- m,
- n,
- k);
+ static_cast((n + kPackedBN - 1) / kPackedBN),
+ static_cast((m + kPackedBM - 1) / kPackedBM));
+ const dim3 block(static_cast(kPackedThreads));
+ if ((m % kPackedBM) == 0) {
+ hipLaunchKernelGGL(
+ HIP_KERNEL_NAME(
+ w8a8_dumma_prefill_packedb_db_kernel<
+ kPackedBM, kPackedBN, 64, kPackedThreads, true>),
+ grid,
+ block,
+ 0,
+ stream,
+ a,
+ b,
+ x_scale,
+ weight_scale,
+ out_ptr,
+ m,
+ n,
+ k);
+ } else {
+ hipLaunchKernelGGL(
+ HIP_KERNEL_NAME(
+ w8a8_dumma_prefill_packedb_db_kernel<
+ kPackedBM, kPackedBN, 64, kPackedThreads, false>),
+ grid,
+ block,
+ 0,
+ stream,
+ a,
+ b,
+ x_scale,
+ weight_scale,
+ out_ptr,
+ m,
+ n,
+ k);
+ }
return;
}
- // Generic native INT8 DUMMA prefill path for every large-M shape with
- // N % 64 == 0 and K % 128 == 0 (identity [K, N] layout). The exact shape
- // above already returned; this covers compatible sibling shapes.
+ // Generic native INT8 DUMMA prefill path for every other large-M shape with
+ // N % 64 == 0 and K % 128 == 0 (identity [K, N] layout). The A row guard
+ // zero-fills rows past M and the epilogue masks them, so only the exact shape
+ // above avoids predication.
if (m >= 128 && (n % 64) == 0 && (k % 128) == 0) {
const dim3 grid(
static_cast(n / 64),
@@ -1282,8 +1100,8 @@ extern "C" void launch_w8a8_gemm(
return;
}
- // Scalar int8/int32 fallback for every unmatched (m, n, k) and all
- // small-M API cases.
+ // Scalar int8/int32 fallback for every unmatched (m, n, k) and all small-M
+ // API cases.
const int64_t total = static_cast(m) * n;
constexpr int kFallbackThreads = 256;
int64_t blocks = (total + kFallbackThreads - 1) / kFallbackThreads;
@@ -1311,15 +1129,15 @@ extern "C" void launch_w8a8_gemm(
k);
}
-// Stable host launch symbol consumed by csrc/bindings.cpp. Iteration 5: the
-// exact (k, n) == (6144, 2624) int8 weight is packed once (outside the
-// timed region and out of Graph capture) into the B-panel layout
-// packed[(n>>6)*(k*64) + (kk>>4)*1024 + (n&63)*16 + (kk&15)] = raw[kk*N + n]
-// consumed by the packed-B prefill kernel's panel tile reads and by the
-// scalar fallback's panel decode; every other (K, N) keeps the identity copy
-// (the logical [K, N] row-major weight). The [N] fp32 scales are always
-// copied unchanged. Same byte count (k*n and n), same allocation,
-// graph-stable addresses.
+// Stable host launch symbol consumed by csrc/bindings.cpp. The exact
+// (k, n) == (6144, 2624) int8 weight is packed once (outside the timed region
+// and out of Graph capture) into the fragment-interleaved swizzle
+// packed[((kk>>5)*n16g + (n>>4))*512 + (((kk>>3)&3)*16 + (n&15))*8 + (kk&7)]
+// = raw[kk*n + n] with n16g = n>>4, consumed by the packed-weight prefill
+// kernel's flat B stage copy and by the scalar fallback's element decode; every
+// other (K, N) keeps the identity copy of the logical [K, N] row-major weight.
+// The [N] fp32 scales are always copied unchanged. Same byte count (k*n and n),
+// same allocations, so the captured addresses stay graph-stable.
extern "C" void launch_pack_w8a8_weight(
const int8_t* raw_weight,
const float* weight_scale,
@@ -1330,14 +1148,10 @@ extern "C" void launch_pack_w8a8_weight(
hipStream_t stream) {
constexpr int kPackThreads = 256;
const int64_t weight_count = static_cast(k) * n;
- // Iteration 5: the exact (k, n) == (6144, 2624) weight is packed once,
- // outside the timed region and out of Graph capture, into the B-panel
- // layout (same byte count and buffer addresses, so captured pointers stay
- // valid; the packed-B prefill kernel and the scalar fallback for this
- // (k, n) both decode it). Every other (K, N) keeps the identity copy
- // (raw row-major [K, N]).
if (k == kExactK && n == kExactN) {
- int64_t blocks = (weight_count + kPackThreads - 1) / kPackThreads;
+ // One thread packs one 16-byte chunk (16 bytes = 2 n-rows x 8 B).
+ const int64_t chunks = weight_count >> 4;
+ int64_t blocks = (chunks + kPackThreads - 1) / kPackThreads;
if (blocks > 4096) {
blocks = 4096;
}
@@ -1345,7 +1159,7 @@ extern "C" void launch_pack_w8a8_weight(
blocks = 1;
}
hipLaunchKernelGGL(
- w8a8_pack_panel_i8_kernel,
+ w8a8_pack_bswz_i8_kernel,
dim3(static_cast(blocks)),
dim3(kPackThreads),
0,
@@ -1393,4 +1207,3 @@ extern "C" void launch_pack_w8a8_weight(
}
}
// @@end
-// @@end
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip
index 7de05137..af1b458c 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip
@@ -1,242 +1,320 @@
-// @@variant shape=glm_tp8_kv_b_proj_m4096 commit=bb4aec7f2d04a07b6c3c5cf4c19e1b9fc97fc8a3 added=2026-08-31
-// median_us=137.4 p90_us=137.9 speedup=50.51 baseline_us=6940
-// source=glm5-2-dsh-tp8-m4096-1-e6a280a2
+// @@variant shape=glm_tp8_kv_b_proj_m4096 commit=547eec13f53e3cee7da02bbdd17763f46e71761e added=2026-09-11
+// median_us=120.1 p90_us=120.9 speedup=57.78 baseline_us=6940
+// source=test-9-11-1-b01955b3/iteration_001/glm_tp8_kv_b_proj_m4096
// INT8 W8A8 GEMM HIP implementation for Hygon K500SM_AI / gfx928.
//
-// Worker: worker_1 (physical GPU 1), assigned shapes:
-// glm_tp8_q_b_proj_m4096 : M=4096, N=2048, K=2048
+// Worker: worker_2 (physical GPU 2). Assigned shape:
// glm_tp8_kv_b_proj_m4096 : M=4096, N=3584, K=512
//
// Logical operation (exact contract):
// out[m, n] = bf16( int32_dot(x_q[m, :], raw_weight[:, n])
// * x_scale[m] * weight_scale[n] )
//
-// Iteration 3 (first valid HIP experiment; iterations 1-2 were killed in the
-// agent infrastructure before any proposal). Accepted-best source digest
-// a64b6250... with one bounded mechanism: n-major packed B + col_major B
-// fragments for the exact glm_tp8_q_b_proj_m4096 (k,n) == (2048, 2048).
-// * Large-prefill path (m >= 128 with exact 128x64x64 geometry): native
-// INT8 DUMMA m16n16k32 with int32 accumulation; one 128x64 output tile
-// per block; four wavefronts (256 threads); each wave owns a 64x32
-// quadrant built from eight m16n16k32 int32 accumulator fragments; the
-// block cooperatively vector-loads A[128,64] and B[64,64] into one
-// single-buffered LDS stage (15,360 B total: 128*80 + 64*80 with 16 B
-// padding per row for bank skew); two __syncthreads per stage; fused
-// dot * x_scale[m] * weight_scale[n] epilogue stored directly as bf16
-// from the accumulator fragments using the verified gfx928 int8
-// m16n16k32 lane mapping (row = lane & 15, col = (lane >> 4) + 4*i).
-// A is read with the library du_load_matrix_sync row-major loader.
-// B is layout-templated:
-// - kNMajorB == true (exact (k,n) == (2048, 2048), q_b_proj): weight
-// is packed once to n-major packed[n*K + k] == raw[k*N + n] outside
-// the timed region; the stage stores B n-major (n row, k contiguous,
-// 80 B row stride); each B fragment is one contiguous 8-byte LDS run
-// per lane, loaded with the lineage-validated load_fragment8 (one
-// ds_read2_b64) - kills the 32 ds_read_u8 byte-gathers + mask/OR
-// reassembly VALU seen in the accepted-best ISA.
-// - kNMajorB == false (any other shape, e.g. kv_b_proj (512, 3584)):
-// raw [K, N] row-major identity layout, row-major B loader (the
-// accepted-best path, byte-identical).
-// Grid dim3(N/64, M/128) = 1024 blocks (q_b_proj) / 1792 blocks
-// (kv_b_proj) dwarfs the 120 CUs, so no split-K is needed.
-// * Generic scalar int8/int32 fallback for every unmatched (m, n, k),
-// including all small-M API cases (M=2, M=16), M tails, and M in
-// (0, 128) with the same (K, N); it decodes the n-major packed layout
-// for the exact (2048, 2048) and the identity [K, N] layout otherwise.
-// * launch_pack_w8a8_weight: n-major device-to-device permutation for
-// (k, n) == (2048, 2048), identity copy for every other (k, n) (both the
-// int8 weight and the fp32 scale). Packing never happens inside the
-// timed GEMM and keeps the same byte count (graph-stable addresses).
+// Iteration 9 (consolidate round): the exact-shape arm keeps its winning
+// 128x128 macro-tile, its 512-thread / eight-wavefront geometry, its
+// 16,384-B single-buffered swizzled-A + plane-B staging, its direct 8-byte
+// fragment loads, its k0-outer/kk-inner int32 order and its coalesced
+// 8-byte epilogue. The one consolidated defect is the bf16 CONVERSION
+// lowering: the exact gfx928 code object of the current best
+// (compile_cache_key 0ece2e0cdbddfe2f512228e2bde5e23189f88fc22d01ae5ae49daf4a97a80f77,
+// source_hip_digest 95dab478d6509cacbb6ca0888d90635197ed82d9202aead00ed49f50c4c7c9ac)
+// shows hip's __float2bfloat16 emitting, for EVERY element each lane
+// converts, an exec-masked inf/NaN fixup - v_and 0x7f800000, v_cmp_ne,
+// s_and_saveexec, v_bfe, v_add3, v_or 0x10000, v_cndmask, s_andn2_saveexec
+// (code-object addresses 0x5EEC..0x77E0 of
+// w8a8_dumma_prefill_128x128_nmajor_kernel) - i.e. ~9 vector instructions
+// plus two scalar EXEC manipulations per element against the 3-instruction
+// rounding the value actually needs. With M*N = 14.68M converted elements
+// and 6,931,456 measured valu_instructions, the block-tail epilogue is the
+// single largest VALU consumer in the kernel, and the
+// s_and_saveexec/s_andn2_saveexec pairs serialize the wavefront there where
+// the only other work left is the second co-resident block. This round
+// replaces ONLY that conversion, ONLY in the exact-shape epilogue, with a
+// plain RNE form (bf16_rne_u16: u += 0x7fff + ((u >> 16) & 1); u >> 16) -
+// the same lineage-validated fix the same-shape sibling reference
+// references/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip applies on this
+// exact DTK 26.04 / gfx928. The formula is bit-identical to hip's own
+// non-masked branch for every finite float, and int32 accumulators times
+// finite scales can only produce finite floats, so mismatch_count 0 /
+// max_abs_error 0 is expected on the unchanged pre-seeded reference data.
+// Everything else - the dense swizzle, the tile-contiguous pack and its
+// scalar decode, the fragment placement, the k0/kk accumulation order, the
+// exact (k, n) == (512, 3584) guard, the identity-layout 128x64 arm, the
+// pack kernels and the generic scalar fallback - is byte-identical.
//
-// Iteration 4 (tile-aspect round): the exact (k, n) == (2048, 2048) q_b_proj
-// arm switches from the 128x64 tile to a 128x128 tile
-// (w8a8_dumma_prefill_128x128_kernel). With the n-major packed B + 8-byte
-// load_fragment8 B reads in place (iteration 3), the remaining LDS-read cost
-// is dominated by the A-side row-major fragment loads (4 x ds_read2_b32 per
-// kk per wave). A 128x128 tile with four 64x64 quadrants keeps the same
-// 4-wave x 8-accumulator-fragment count per wave (8 -> 16 m16n16k32 MMACs
-// per kk from 8 LDS fragment reads, i.e. per-MMAC LDS reads drop 6/8 -> 8/16
-// = -33%), and the A-side global re-read halves (A reuse 32 -> 16, A traffic
-// 268 -> 134 MiB) while B reuse stays 32 (134 MiB), balancing the per-byte
-// A:B traffic at 1:1 and cutting total tile traffic 402 -> 268 MiB (-33%).
-// LDS grows 15,360 -> 20,480 B/block (still 2 blocks/CU = 40,960 B <= 64
-// KiB); the grid becomes (N/128, M/128) = 512 blocks (~4.3/CU). The 16
-// accumulator fragments cost ~+32 VGPR; arch VGPR is expected <= 128 so the
-// accepted kernel's 2-blocks/CU co-residency is preserved (falsified if the
-// exact code object shows > 128 VGPR -> 1 block/CU). kv_b_proj keeps the
-// byte-identical 128x64 kNMajorB=false arm and the scalar fallback is
-// untouched.
+// Iteration 10 (consolidate round): final HIP-only polish on the same winning
+// architecture. The one consolidated defect is the dead per-fragment tail-M
+// masking in the exact-shape epilogue. The exact gfx928 code object of the
+// current best (compile_cache_key 0f7ddb1a6c31e6f3cdd7772a8760e5b20dcea6da9ffca5d556e35790d7e25113,
+// source_hip_digest 8c2468b4dfc603de4dd4a957ebf58d891830cbac940a07c58aed7c52030a8def,
+// arch_vgpr 64, 0 scratch, 16,384 B LDS, 896 blocks) shows the block tail
+// split into eight exec-masked basic blocks
+// (s_and_saveexec / s_cbranch_execz / s_or_b64 around each of the eight
+// fragment epilogues) even though the dispatch guard for this arm is
+// m >= 128 with m % 128 == 0, so every row of every 128-row tile is in bounds
+// and neither the (row >= m) early return in the store helper nor the two
+// (rs < m) x_scale ternaries in the kernel can ever fire. Because each
+// fragment epilogue sits in its own basic block, the compiler also repeats the
+// fragment-independent shuffle scaffolding per fragment (v_mbcnt_lo /
+// v_mbcnt_hi lane-id, the two __shfl_xor partner indices, the two comparison
+// and cndmask pairs, ~13 VALU per fragment x 8). This round removes ONLY that
+// dead masking from the exact-shape arm: the store helper drops its
+// (row >= m) early return and its now-unused m argument, and the caller loads
+// x_scale[rs0] / x_scale[rs1] directly. Nothing else changes: the 128x128
+// macro-tile, the 512-thread / eight-wavefront block, the 16,384-B
+// single-buffered swizzled-A + plane-B staging, the direct 8-byte fragment
+// loads, the k0-outer/kk-inner int32 order, the 4x4 shfl_xor transpose and
+// the single 8-byte bf16 store per lane are all untouched, so the written
+// bytes are bit-identical (removing a mask that is never asserted cannot
+// change any output). The 128x64 identity-layout arm keeps its own
+// store_fragment_scaled bounds check and the generic scalar fallback keeps its
+// linear bounds check, both byte-identical. Falsifiable signature: the eight
+// exec-masked guard blocks disappear from the 128x128 code object, the
+// per-fragment duplicate lane-id / shuffle-index VALU is hoisted, and the
+// dynamic valu_instructions fall from 5,512,192 without changing
+// lds_instructions 1,089,536, lds_bank_conflicts 458,752,
+// vmem_write_instructions 57,344 or the 16 v_mmac per stage.
//
-// Iteration 5 (packing round): the exact (k, n) == (2048, 2048) q_b_proj
-// weight is re-packed once (outside timing/Graph, same byte count, same
-// graph-stable buffer) into a swizzled 64-k-stage-major layout
-// packed[((k0*8 + kc)*n + col)*8 + b] == raw[kk*n + col],
-// kk = k0*64 + kc*8 + b (k0 = 64-k stage, kc = 8-byte sub-chunk, b = byte),
-// and the 128x128 kernel stages each B tile as 16-byte plane chunks into an
-// LDS plane layout b_tile[kc][n][8] (plane stride 1024 B, n stride 8 B).
-// The exact code object of the accepted kernel shows every B fragment read
-// is an 8-byte LDS access at (lane&15)*80 + (lane>>4)*8 (n-major stride-80
-// rows), and PMC counts 6,291,456 bank conflicts over 1,048,576 LDS slots
-// (6.0/slot) with 4,078,494 LDS waits (3.9/slot): the loop is LDS-latency
-// bound and the stride-80 rows alias bank phases every 8 rows (80*8 == 640
-// == 0 mod 128), so every fragment read conflicts ~4-6-way. In the plane
-// layout a lane's 8 bytes sit at (lane>>4)*1024 + (lane&15)*8 relative to
-// the fragment origin: 16 lanes per 128-B phase cover all 32 banks exactly
-// once, i.e. B fragment reads become lane-linear and zero-conflict (4-cycle
-// minimum for 512 B). Staging reads stay coalesced (each 16-B chunk is two
-// n rows of one sub-chunk -> one aligned int4) and LDS writes are 2-way
-// (the 16-B optimum). Same fragment operand bytes, same k0-outer/kk-inner
-// int32 accumulation -> bit-identical results. The scalar fallback decodes
-// the new pack for (2048,2048) only; kv_b_proj keeps the byte-identical
-// 128x64 kNMajorB=false identity path.
+// Iteration 13 (consolidate round, 13/13): final HIP-only polish on the
+// unchanged winning exact-shape architecture, one bounded mechanism. The
+// exact current-best code object (source_hip_digest
+// 8010314fdb6ba5ec0da98d662fc3911763b088245fa7c316bf0986572499207d,
+// compile_cache_key 3fbcc466cfb5087e6f7057b3f26684e598046f28d0f71e0d794962688901c1c7,
+// arch_vgpr 64, scratch 0, 16,384 B static LDS, 896 blocks x 512 threads) puts
+// the whole 64-K stage wall on the EXPOSED cooperative-staging round trip: the
+// two global_load_dwordx4 A/B chunks are issued only after the previous stage's
+// trailing s_barrier and are then drained within ~8 instructions by
+// s_waitcnt vmcnt(1) / vmcnt(0) in front of the three ds_write_b64/b128
+// commits (0x5D1C..0x5D5C), so no part of the A+B L2 latency overlaps any
+// arithmetic or the barrier wait. Iteration 7 raised the staging depth to 2 to
+// overlap exactly that round trip, but it doubled LDS to 32,768 B/block and
+// gave up the 2-blocks/CU residency point, so the prefetch mechanism itself was
+// never measured. This round re-introduces ONLY the single-buffered register
+// prefetch: the A row base and the B per-stage stride are hoisted out of the K
+// loop, the next stage's A and B int4 chunks are loaded BEFORE the trailing
+// barrier and consumed by the next stage's ds_write commits, while LDS stays
+// 16,384 B, __launch_bounds__(kNMBlockThreads, 2), the 64-VGPR residency point
+// and the two barriers per stage are all untouched, and the last stage re-reads
+// its own chunk (k1 == k0) instead of reading past the last K block. The
+// swizzled A tile, the tile-contiguous plane-B layout, the direct 8-byte
+// fragment loads, the k0-outer/kk-inner ascending int32 order, the coalesced
+// RNE epilogue, the exact (k, n) == (512, 3584) dispatch guard, the 128x64
+// identity-layout arm, the pack kernels and the generic scalar fallback are
+// byte-identical, so every stored bit is unchanged. Falsifiable signature: the
+// global_load_dwordx4 pair for stage s+1 is scheduled before the trailing
+// s_barrier and its vmcnt drain moves to the next stage's top, while arch_vgpr
+// stays 64 with vgpr_spill_count 0 and scratch_bytes 0; if the machine
+// scheduler keeps the loads where they already are the object is unchanged and
+// the round is flat.
//
-// Iteration 6 (epilogue round): the fused dot * x_scale[m] * weight_scale[n]
-// -> bf16 epilogue has been in-kernel since iteration 1 and the workspace is
-// unused ((void)workspace; no split-K, no combine pass anywhere in the call
-// chain), so the remaining epilogue inefficiency of the exact (2048, 2048)
-// 128x128 arm is the STORE pattern: store_prefill_fragment issues one
-// 2-byte bf16 store per lane per element - 64 scattered stores per wave per
-// fragment set, each wavefront store touching 16 rows x 4 columns so every
-// 32-B sector is only 25% utilized (PMC: 131,072 vmem_write_instructions =
-// 512 blocks x 4 waves x 64). The new epilogue transposes each 16x16
-// fragment's 4-element groups inside the 4-lane column group (lanes r,
-// r+16, r+32, r+48; two 2x2 steps with __shfl_xor 16 then 32 - the
-// lineage-validated 4x4 register transpose accepted on the sibling TP8
-// workers 0/2, same DTK 26.04/gfx928), so lane (r, c4) owns the four
-// CONTIGUOUS columns 4*c4 .. +3 and writes ONE 8-byte store per lane (100%
-// store sector efficiency; vmem_write 131,072 -> 32,768). Only int32
-// values move between lanes: the per-element multiply order
-// (float(dot) * x_scale[row] * weight_scale[col]) and the bf16 rounding are
-// unchanged, so stored bits are identical. The per-row x_scale (4 rows per
-// lane) and per-column weight_scale float4s (4 per lane) are additionally
-// register-batched on the last K stage (before the final protective
-// __syncthreads), so the epilogue is pure compute + 16 coalesced stores with
-// no interleaved vmem loads (vmem_read ~413,696 -> ~278k). kv_b_proj keeps
-// the byte-identical 128x64 kNMajorB=false arm (old per-element epilogue)
-// and the generic scalar fallback is untouched.
+// Iteration 5 (packing round): the exact (k, n) == (512, 3584) weight is
+// re-packed once, outside the timed/Graph region, from the n-major [N, K]
+// transpose to a TILE-CONTIGUOUS 64-K-stage layout
+// packed[(((k0*(n/128) + nt)*8 + kc)*128 + row)*8 + b] = raw[kk*n + col]
+// kk = k0*64 + kc*8 + b, nt = col/128, row = col%128,
+// i.e. every (64-K stage, 128-N tile) DUMMA B tile is ONE contiguous 8192-B
+// global region laid out as [kc = 8-byte k sub-chunk][128 N rows][8 bytes],
+// which is exactly the plane order the B fragments are read from in LDS. The
+// GEMM's cooperative B staging therefore becomes a plain 2 x 16-B contiguous
+// vector copy per thread (global_load_dwordx4 -> ds_write_b128) with
+// consecutive lanes on consecutive 16-B slots: the global reads are fully
+// coalesced (the n-major [N, K] layout made a wave touch 16 rows 512 B apart)
+// and the 16-B LDS stores have no bank conflicts. The B fragment load becomes
+// lane-linear (kc*(128*8) + row*8, no swizzle arithmetic) and each 16-lane
+// phase covers all 32 banks exactly once, so B is now both vector-loadable and
+// LDS-bank-safe; the swizzled 8-byte-group map is retained for A only (x_q is
+// a runtime activation and cannot be pre-packed). The int32 accumulation is
+// untouched: identical per-lane fragment placement (row = row0 + (lane & 15),
+// k group kc = (lane >> 4) + (kk >> 3), 8 consecutive natural-order k bytes)
+// and identical k0-outer/kk-inner ascending order, so every output byte is
+// bit-identical. Falsifiable signature: the B half of the 917,504 measured
+// lds_bank_conflicts disappears (32 -> 16 per stage per wave), the B staging's
+// 4 x ds_write_b64 become 2 x ds_write_b128, lds_instructions fall
+// 688,128 -> 630,784, and the median improves only if the global/LDS staging
+// latency chain, not the 16-MMAC burst, is the limiter.
//
-// Iteration 7 (compute-pipeline round): the exact (k, n) == (2048, 2048)
-// 128x128 arm loads its four A fragments per kk with the lineage-validated
-// direct load_fragment8 (one ds_read2_b64 straight into the v_mmac operand)
-// instead of the library du_load_matrix_sync row_major loader. The exact
-// code object shows the library loader lowers to 8 x ds_read2_b32 plus a
-// redundant byte-reassembly chain per fragment (~7 VALU: v_and 0xff00 /
-// 0xff0000 / 0xff000000 + v_or_b32_sdwa + v_or3, ~50-56 VALU per stage per
-// wave) between the LDS read and the first MMAC, and PMC counts 3,009,148
-// LDS waits (2.3/slot) against a latency-bound loop running at 1 block/CU
-// (160 VGPR). du_mma.hpp defines matrix_a row_major int8 as x[i] =
-// p[(lane&15)*ldm + (lane>>4)*8 + i] (8 consecutive bytes, memory order)
-// and du_mma_sync passes reinterpret(x) unchanged to v_mmac, so
-// load_fragment8 produces byte-identical operand values and the int32
-// accumulation is bit-identical; only the redundant VALU reassembly (and its
-// lgkmcnt wait states) is removed from the LDS->MMAC critical path, letting
-// the 32-MMAC burst issue back-to-back. Expected PMC: valu_instructions
-// 8.66M -> ~4.8-5.6M (-35..-45%), lds_instructions 1.31M -> ~1.05M (-20%,
-// 8 A ds_read2_b32 + 4 B ds_read2_b64 -> 8 ds_read2_b64 per stage),
-// lds_wait_instructions 3.01M -> ~2.2-2.5M, lds_bank_conflicts
-// approximately unchanged (A rows keep the 2-way stride-80 aliasing),
-// vmem_read/vmem_write and the 32 v_mmac/stage unchanged. kv_b_proj keeps
-// the byte-identical 128x64 kNMajorB=false arm (unchanged lowering) and the
-// generic scalar fallback is untouched.
+// Iteration 6 (mandated epilogue round): the fused scale/bf16 epilogue has
+// been inside the compute kernel since iteration 1 and the workspace is
+// already ignored ((void)workspace in launch_w8a8_gemm), so the remaining
+// epilogue inefficiency of the exact (512, 3584) arm is the STORE pattern:
+// store_fragment_scaled issues one 2-byte bf16 store per lane per element -
+// 64 scattered stores per wave, each wavefront store touching 16 rows x 4
+// columns so every 32-B store sector is only 25% utilized (the bootstrapped
+// PMC counts vmem_write_instructions == 229,376 == 896 x 4 x 16 x 4 exactly).
+// The new epilogue is register-batched: the per-row x_scale is loaded once per
+// 16-row row-block and the four per-column scales of each fragment's owned
+// columns are one aligned float4 load, so the epilogue issues 4 scalar + 16
+// vector vmem loads per lane instead of 80 scalar loads. It then transposes
+// the four-element groups within each 4-lane column group in registers (two
+// 2x2 shfl_xor steps, 16 then 32 - the lineage-validated pattern ported from
+// references/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip, which lowers to
+// ds_bpermute at the block tail where the LDS pipe is idle) so lane (r, c4)
+// ends up holding the four CONTIGUOUS columns 4*c4 .. +3, converts them to
+// bf16 and writes ONE 8-byte store per lane per fragment: 16 coalesced stores
+// per wave instead of 64 scattered ones at 100% sector efficiency, i.e.
+// vmem_write 229,376 -> 57,344. The scale multiply order
+// (float(dot) * x_scale[row] * weight_scale[col]) and the __float2bfloat16
+// rounding are unchanged, so the stored bits are identical; no LDS, no
+// barrier, and no workspace/combine pass is added. Everything else - the A
+// swizzle, the tile-contiguous plane-B staging, the k0-outer/kk-inner
+// ascending int32 order, the exact (k, n) == (512, 3584) guard, the
+// identity-layout 128x64 arm, the pack kernels and the generic scalar
+// fallback - is byte-identical.
//
-// Iteration 1 (kv_b_proj tile-baseline round): the exact (k, n) == (512,
-// 3584) kv_b_proj arm is still the iteration-2 bootstrap pipeline (library
-// fragment loaders, identity [K, N] B, per-element scattered epilogue) and
-// runs at 65 TOPS vs 156 TOPS on the fully-modernized q_b_proj arm (same
-// GPU, same DTK, measured 220.14 us on K=2048/N=2048). This round ports the
-// proven pipeline (swizzled plane-pack B, direct 8-byte A/B fragment loads,
-// register-batched coalesced epilogue, fully unrolled kk) to kv_b_proj as
-// one templated 2-D macro-tile family (w8a8_dumma_prefill_tile_kernel) over
-// the three mandated block tiles plus the iteration-4 aspect flip, all
-// compiled, active selected by kActiveKvTile (0 = 64x64, 1 = 64x128,
-// 2 = 128x64, 3 = 64x128 with 32x32 quadrants):
-// config tile waves threads LDS B grid (blocks) est. VGPR
-// 0 64x64 1 64 9,216 (56, 64) = 3584 ~128-140
-// 1 64x128 2 128 13,312 (28, 64) = 1792 ~128-140
-// 2 128x64 4 256 14,336 (56, 32) = 1792 ~88-100
-// 3 64x128 8 512 13,312 (28, 64) = 1792 ~55-60
-// Per-wave LDS fragment-read bytes per MMAC: 256 B for configs 0/1 (4 A + 4
-// B fragments feed 16 MMACs per kk) vs 384 B for config 2 (4 A + 2 B feed 8
-// MMACs) and vs 384 B for the old 128x64 arm; B reads are zero-conflict in
-// the plane layout for every plane stride, and the A stride-80 2-way
-// aliasing matches the accepted q_b_proj arm. Occupancy: 64x64 -> 7
-// blocks/CU = 7 waves/CU (LDS- and VGPR-limited), 64x128 -> 4 blocks/CU = 8
-// waves/CU (LDS-limited; VGPR <= 128 keeps 4 blocks), 128x64 -> 2 blocks/CU
-// = 8 waves/CU (VGPR-limited) - vs 2 blocks/CU = 8 waves/CU today but with
-// ~2.9x fewer LDS instructions (direct 8-byte reads vs library byte-gathers)
-// and 4x fewer store instructions (coalesced epilogue, 229,376 -> 57,344
-// vmem_write). Active config for the measurement: 64x128 (config 1):
-// deepest per-kk MMAC burst of the three alongside the lowest per-MMAC LDS
-// read volume, half the A-side global re-read of 128x64, and 8 waves/CU at
-// 4 blocks/CU.
-// * Iteration-1 measurement: config 1 landed at arch_vgpr 136 (code object
-// 135) -> 3 blocks/CU = 6 waves/CU (1.5/SIMD), median 209.4 us = 71.8
-// TOPS. The kernel is latency-bound (per-stage wall ~2,454 cycles vs
-// ~250 cycles of pure issue), so 6 waves/CU leaves the barrier + LDS +
-// global-latency chain exposed on the SIMDs that hold only one wave.
-// * Iteration-2 (operand-reuse) round selects config 2 (128x64): same
-// cooperative A+B staging machinery and swizzled-pack B, same
-// bit-identical int32 order; its +50% A-side LDS fragment-read bytes per
-// MMAC (384 vs 256 B; A rows are re-read by the two n-half waves in both
-// configs, but config 2's narrower 64x32 quadrant halves the MMACs per
-// A-frag load) is absorbed by the LDS pipe (~10% of per-stage cycles)
-// while the compiled 78 VGPR / 14,336 B gives 3 blocks x 4 waves = 12
-// waves/CU (3/SIMD) - a 2x occupancy increase over the measured config-1
-// residency.
-// The swizzled pack is extended to (512, 3584) in
-// launch_pack_w8a8_weight (same byte count, outside the timed region) and
-// the scalar fallback decodes it for both swizzled pairs, so the paired
-// M=2/(3584,512) API case stays exact. All other arms (q_b_proj 128x128,
-// generic 128x64, identity packs) are byte-identical.
+// Iteration 8 (resource round): the exact-shape 128x128 n-major arm keeps its
+// winning tile family, its tile-contiguous plane-B pack, its A swizzle, its
+// direct 8-byte fragment loads, its k0-outer/kk-inner int32 order and its
+// coalesced epilogue, and changes only the occupancy resources: waves per
+// block, the VGPR live range and the per-thread LDS staging payload.
+// * Measured occupancy of the current best (iteration-6 source): arch_vgpr
+// 120, static LDS 16,384 B/block, __launch_bounds__(256, 2) -> 2 blocks/CU
+// x 4 waves = 8 waves/CU = 2 waves/SIMD with 16 int32 accumulator VGPR
+// fragments + 8 operand fragments + 4 live int4 staging registers per
+// thread. The PMC says the kernel is issue-stalled, not LDS-bandwidth- or
+// MMAC-bound: 195,964 cycles for 7,168 block-stages, 974,848
+// lds_instructions (4.97 per GPU cycle over 120 CUs) and 1,404,478
+// lds_wait_instructions against a theoretical LDS issue rate of 1/cycle/CU
+// - i.e. the SIMDs are starved of co-resident waves, not of LDS slots.
+// * Focused change: the same 128x128 macro-tile is covered by EIGHT
+// wavefronts of 64 lanes (512 threads, kNMBlockThreads = 8*64) instead of
+// four; each wave now owns a 32x64 quadrant = kAccM 2 A fragments x
+// kAccN 4 B fragments = 8 m16n16k32 int32 accumulators, i.e. 32
+// accumulator VGPR per thread instead of 64. The staging payload per
+// thread halves to exactly one int4 of A (row = tid >> 2, 16-byte K group
+// (tid & 3) * 16) and one int4 of B (offset tid * 16), and the arm is
+// pinned by __launch_bounds__(kNMBlockThreads, 2) to the 64-VGPR /
+// 2-blocks/CU point (65536 VGPR per CU / (2 x 512 lanes) = 64), i.e.
+// 2 blocks x 8 waves = 16 waves/CU = 4 waves/SIMD, twice the currently
+// measured wave count. LDS stays 16,384 B/block, so 2 blocks use 32,768 B
+// of the 65,536-B per-CU LDS (no exact-fit cliff like iteration 7's
+// 32,768 B/block double buffer).
+// * Why exactly this point: the trusted occupancy-probe candidate list for
+// the measured 120-CU device is [2, 3, 4, 5, 6, 7, 8] waves per block; 8
+// is the largest value that still tiles the 128x128 macro-tile with
+// rectangular m16n16k32 quadrants (4 M row groups x 2 N column groups) and
+// it gives the top of the trusted residency range (8 waves/CU -> 16
+// waves/CU = 4/SIMD). The direct sibling evidence is the lineage reference
+// references/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip (same exact
+// M=4096/N=3584/K=512 shape, same DTK 26.04/gfx928): its 8-wave/512-thread
+// 16-waves/CU recipe (config 3, arch_vgpr <= 64, 2 blocks/CU) is the
+// recorded 137.4 us winner of that family versus 169.3 us for its
+// 12-waves/CU 128x64 form, which is the same occupancy axis this round
+// probes while keeping our better 1:1 macro-tile and 0.75-fragment/MMAC
+// quadrant instead of that reference's 1.0-fragment/MMAC 32x32 quadrants.
+// * Falsifiers: the exact code object must show vgpr_spill_count == 0,
+// scratch_bytes == 0 and arch_vgpr <= 64 (else residency falls to 1
+// block/CU = 8 waves/CU, cancelling the whole point); the LDS-read volume
+// per MMAC rises 0.5 -> 0.75 fragments, so the median improves only if the
+// doubled wave count hides more LDS/barrier latency than the +50% fragment
+// reads cost. The int32 accumulation is untouched (exact integer
+// additions of the same products), every fragment placement is the same
+// 8-byte per-lane operand, and the exact (k, n) == (512, 3584) guard, the
+// identity-layout 128x64 arm, the pack kernels and the generic scalar
+// fallback are byte-identical.
//
-// Iteration 5 (kv_b_proj packing round): the exact (k, n) == (512, 3584)
-// weight is re-packed once (outside timing/Graph, same byte count, same
-// graph-stable buffer) from the k-stage-major swizzle to a TILE-CONTIGUOUS
-// stage-major layout
-// packed[(((k0*(n/128) + nt)*8 + kc)*128 + row)*8 + b] == raw[kk*n + col]
-// kk = k0*64 + kc*8 + b, nt = col/128, row = col%128,
-// so every (64-k stage, 128-n tile) B tile is ONE contiguous 8192-B global
-// region. The tile-family kernel (w8a8_dumma_prefill_tile_kernel, all four
-// configs) stages it as 8 consecutive 1024-B wavefront streams instead of 8
-// streams 28,672 B apart in the k-stage-major layout; the LDS plane tile
-// [kc][n][8], the zero-conflict lane-linear 8-byte B fragment reads
-// (load_fragment8_plane_stride), the 16-B staging writes and the
-// k0-outer/kk-inner int32 accumulation are byte-identical, so mismatch 0 /
-// max_abs_error 0 is expected without tolerance debate. q_b_proj (2048,
-// 2048) keeps the old k-stage-major swizzle byte-identical (its 128x128
-// kernel, pack kernel and fallback decode are untouched); the scalar
-// fallback decodes the new tile-contiguous layout for (512,3584) and the
-// old swizzle for (2048,2048).
+// Iteration 4 (tile-shape round): the exact-shape n-major arm moves from the
+// 128x64 macro-tile (M:N = 2:1) to a 128x128 macro-tile (M:N = 1:1) for the
+// exact (M, N, K) = (4096, 3584, 512) shape; the four-wavefront/256-thread
+// block and the direct 8-byte swizzled fragment path are unchanged. Each
+// wave now owns a 64x64 quadrant built from sixteen m16n16k32 int32
+// accumulators (kAccM = 4 A fragments x kAccN = 4 B fragments = 16 MMACs per
+// kk) instead of a 64x32 quadrant with eight. Why the aspect/area change is
+// the lever for this exact M/N/K:
+// * Per-block-stage cooperative-load bytes are fixed by the tile shape:
+// 128*64 (A) + 128*64 (B) = 16 KiB for 128x128 versus 8 KiB + 4 KiB =
+// 12 KiB for 128x64, but the 128x128 tile holds 2x the outputs, so the
+// global/L2 read traffic for the whole operator drops from
+// M*N*K*(1/128 + 1/64) = 176.2 MB to M*N*K*(1/128 + 1/128) = 117.4 MB
+// (-33%). Concretely: A is re-read once per N-tile column (56 -> 28) so
+// A traffic halves 117.4 -> 58.7 MB, while B keeps its 32 M-tile re-reads
+// at 58.7 MB. The 128x64 arm's 2:1 A:B read imbalance becomes 1:1, which
+// is the aspect balance this round is for; the smaller A (2 MiB) and B
+// (1.75 MiB) both stay L2-resident either way.
+// * LDS fragment reads per MMAC drop 6/8 -> 8/16 (-33%): the 64x32 quadrant
+// needed 4 A + 2 B = 6 direct 8-byte reads for 8 MMACs, the 64x64
+// quadrant needs 4 A + 4 B = 8 reads for 16 MMACs.
+// * The measured per-stage wall is dominated by the exposed A+B L2 round
+// trip in front of the MMAC burst (exact code object: 3 x
+// global_load_dwordx4 -> s_waitcnt vmcnt(2)/(1)/(0) -> 6 x ds_write_b64,
+// per 64-K stage). Halving the macro-tile count (grid 1792 -> 896) halves
+// the number of exposed stage walls for the same total work.
+// * 896 blocks is still ~3.7 waves of the 240 co-resident blocks
+// (2 blocks/CU x 120 CU), i.e. enough blocks to occupy every CU; no
+// split-K and the workspace stays unused.
+// Residency: LDS 16,384 B/block (2 x 16,384 = 32,768 <= 64 KiB) and the
+// sixteen accumulators cost 64 VGPR, so the arm is pinned to 2 blocks/CU =
+// 8 waves/CU by an explicit __launch_bounds__(256, 2) register cap at
+// 128 VGPR (floor(16384 VGPR/SIMD / (128 x 64 lanes)) = 2 waves/SIMD);
+// falsified if the exact code object shows > 128 VGPR or any spill -> 1
+// block/CU = 4 waves/CU.
+// Everything else is byte-identical to iteration 2/3: the dense 8-byte-group
+// swizzle, load_frag8_swizzled placement, the one-time n-major packed B, the
+// k0-outer/kk-inner ascending int32 order, the exact (k, n) == (512, 3584)
+// guard, the identity-layout 128x64 arm, the pack kernels and the generic
+// scalar fallback.
//
-// Iteration 6 (kv_b_proj epilogue round): the mandate (per-row x_scale,
-// per-column weight_scale, bf16 conversion and the final coalesced store
-// fused into the compute kernel; no workspace/combine pass) has been in
-// place since iteration 1 - the launcher is a single kernel with
-// (void)workspace and the epilogue is register-batched (xs scalar + ws
-// float4 prefetched on the last K stage) with one 8-byte bf16 store per lane
-// per fragment (57,344 vmem_write). The remaining epilogue inefficiency of
-// the exact (512, 3584) tile family is the bf16 CONVERSION lowering: the
-// exact code object shows hip's __float2bfloat16 emits, per element, an
-// exec-masked inf/NaN fixup (v_and exp + v_cmp_ne + s_and_saveexec + s_xor +
-// v_mov 0 + v_or 0x10000 + v_cmp_eq_u32_sdwa + v_cndmask + s_or_b64) around
-// the 3-instruction RNE rounding (v_bfe + v_add3 + shift-in-pack) - ~11
-// instructions and three scalar exec manipulations that serialize the
-// block-tail epilogue. This round replaces that path, ONLY in the
-// w8a8_dumma_prefill_tile_kernel epilogue, with a plain RNE conversion
-// (bf16_rne_u16: u += 0x7fff + ((u >> 16) & 1); u >> 16) behind a new
-// kRneOnly template flag on store_prefill_fragment_coalesced_scaled
-// (default false, so the q_b_proj 128x128 / 128x64 arms and the scalar
-// fallback stay byte-identical). The RNE formula is bit-identical to hip's
-// own non-masked branch for every finite float - the only values this GEMM
-// can produce (int32 accumulators x finite scales) - so mismatch 0 /
-// max_abs_error 0 is expected on the unchanged pre-seeded reference data.
-// PMC prediction: valu_instructions 7.31M -> ~5.9-6.1M (-15..-20%),
-// lds_instructions 1,089,536 unchanged, lds_bank_conflicts 2,293,760
-// unchanged, vmem_read 229,376 / vmem_write 57,344 unchanged, arch_vgpr <=
-// 64 / 0 scratch / 13,312 B LDS / 2 blocks/CU = 16 waves/CU residency
-// unchanged; the exact code object must show no s_and_saveexec/s_xor/
-// s_andn2_saveexec around the conversion (v_bfe + v_add3 straight into the
-// 8-byte pack).
+// Iteration 2 (operand-reuse round): keep the cooperative A+B LDS staging
+// (the reuse argument below favours it over per-fragment global loads) but
+// replace the two structural costs the bootstrap PMC exposes on the assigned
+// (M, N, K) = (4096, 3584, 512) shape:
+// * lds_instructions 2,465,792 and lds_wait_instructions 7,165,860:
+// the installed du_load_matrix_sync int8 loaders lower to a per-byte
+// ds_read_u8 gather plus mask/OR reassembly VALU for EVERY fragment.
+// * lds_bank_conflicts 6,881,280: the 80-byte padded row stride has bank
+// period 8 over a fragment's 16 rows, so lanes r and r+8 collide.
+// Reuse accounting for one 128x64 macro-tile / 64-K stage (the reason direct
+// global operand loads are rejected):
+// A[128,64] = 8 KiB staged; each A fragment row-block feeds two B fragments
+// (kAccN = 2) and two kk sub-steps, i.e. every staged A byte is consumed
+// by 4 or more MMAC operand reads; a direct global fragment load would
+// re-fetch it per fragment (4x redundant global traffic).
+// B[64,64] = 4 KiB staged; each B fragment feeds all four A row-blocks
+// (kAccM = 4), i.e. 4x reuse. So cooperative staging is retained and the
+// real lever is the *read* path and the LDS conflict structure.
+// Change (one mechanism, both tiles):
+// * Both LDS tiles become dense 64-byte-row swizzled 8-byte-group layouts
+// (no pad). group(row, c) = (((row & 15) + 4*(c & 3)) & 15) + 16*c +
+// 128*(row >> 4); the writers emit two 8-byte LDS stores per 16-byte
+// global chunk at group(row, 2j) / group(row, 2j+1) and the fragment
+// loader reads the swizzle back. Every 16-lane fragment load then touches
+// 16 distinct bank pairs (zero conflicts) instead of the period-8
+// collision.
+// * Every A and B fragment read becomes ONE 8-byte LDS load per lane
+// (load_frag8_swizzled), reproducing du_load_matrix_sync's per-lane
+// (row = row0 + (lane & 15), c = (lane >> 4) + (kk >> 3), 8 consecutive
+// k bytes in natural order) placement, so the int32 accumulation is
+// bit-identical.
+// * B is read col_major from a one-time n-major transpose
+// packed[n*k + kk] = raw[kk*n + n] (produced outside the timed/Graph
+// region by launch_pack_w8a8_weight for exactly (k, n) == (512, 3584));
+// this makes each lane's 8 k bytes contiguous. The identity [K, N] copy
+// is kept for every other (k, n).
+// Iteration-2 geometry (superseded for the n-major arm by the iteration-4
+// 128x128 tile above; the identity-layout arm still uses it): one 128x64
+// tile/block, four gfx928 wavefronts of 64 lanes (256 threads), each wave
+// owning a 64x32 quadrant (kAccM = 4, kAccN = 2, eight m16n16k32 int32
+// accumulators), 2 barriers per 64-K stage, no split-K, grid
+// dim3(N/64, M/128) = (56, 32) = 1792 blocks, LDS = 8 KiB + 4 KiB = 12,288
+// B/block.
+// * The fused int32 * x_scale * weight_scale -> bf16 epilogue keeps the
+// verified gfx928 int8 m16n16k32 accumulator ownership (row = lane & 15,
+// col = (lane >> 4) + 4*i), multiply order (float(dot) * x_scale[row] *
+// weight_scale[col]) and bf16 rounding of the harness reference.
+// * The identity-layout 128x64 DUMMA arm is retained for every non-packed
+// large shape, and the generic scalar int8/int32 fallback covers every
+// unmatched (m, n, k), including all small-M API cases (M=2, M=16) and
+// M < 128 with the same (K, N): one thread per output element, exact
+// int32 dot over K, then float(acc) * x_scale[row] * weight_scale[col] ->
+// bf16. For exactly (512, 3584) the fallback decodes the n-major pack.
+// * launch_pack_w8a8_weight runs outside the timed/Graph region, allocates
+// nothing, keeps the same byte count and graph-stable addresses: the
+// n-major transpose for exactly (k, n) == (512, 3584), identity for every
+// other (k, n), fp32 scale always copied identity.
+//
+// The timed operator performs no allocation, compilation, autotuning, weight
+// packing, host/device synchronization, or default-stream launch; it uses only
+// the caller-provided out and workspace and launches on PyTorch's current HIP
+// stream via hipLaunchKernelGGL.
#include
#include
@@ -251,71 +329,55 @@ using namespace du::dumma;
constexpr int kTileM = 16;
constexpr int kTileN = 16;
constexpr int kTileK = 32; // gfx928 INT8 DUMMA primitive: m16n16k32
-constexpr int kWaveSize = 64; // gfx928 native wavefront
+constexpr int kWaveSize = 64; // gfx928 native wavefront, not 32
constexpr int kBlockM = 128;
constexpr int kBlockN = 64;
constexpr int kStageK = 64;
-constexpr int kBPad = 16; // 64 -> 80-byte LDS row stride (five bank phases)
-constexpr int kBStride = kBlockN + kBPad;
-constexpr int kAStride = kStageK + kBPad; // 64 -> 80-byte A row stride
-constexpr int kBlockThreads = 4 * kWaveSize;
+constexpr int kPad = 16; // 16 B LDS row padding (64 -> 80 B rows)
+constexpr int kAStride = kStageK + kPad;
+constexpr int kBStride = kBlockN + kPad;
+constexpr int kBlockThreads = 4 * kWaveSize; // 256 = four gfx928 wavefronts
-// Exact (k, n) pairs whose weight buffer is packed with the iteration-5
-// swizzled 64-k-stage layout (packed[((k0*8+kc)*n+col)*8+b] == raw[kk*n+col]):
-// the q_b_proj pair (2048, 2048) and - since this kv_b_proj tile round - the
-// kv_b_proj pair (512, 3584). Every other (k, n) keeps the identity layout.
-constexpr int kPackSwizzleK1 = 2048;
-constexpr int kPackSwizzleN1 = 2048;
-constexpr int kPackSwizzleK2 = 512;
-constexpr int kPackSwizzleN2 = 3584;
+// Iteration-4 tile-shape round: the exact-shape n-major arm uses a 128x128
+// macro-tile (M:N = 1:1) instead of the 128x64 identity-arm geometry.
+// Iteration-8 resource round: the same 128x128 macro-tile is covered by EIGHT
+// wavefronts of 64 lanes (512 threads), each owning a 32x64 quadrant built
+// from kAccM = 2 A fragments x kAccN = 4 B fragments = eight m16n16k32 int32
+// accumulators (32 VGPR of accumulator per thread instead of 64). The block
+// is pinned to 2 blocks/CU by __launch_bounds__(512, 2), i.e. a 64-VGPR cap:
+// 2 blocks x 8 waves = 16 waves/CU = 4 waves/SIMD instead of the
+// measured 2 blocks x 4 waves = 8 waves/CU = 2 waves/SIMD.
+constexpr int kNMBlockM = 128;
+constexpr int kNMBlockN = 128;
+constexpr int kNMQuadM = 32; // per-wave M extent (4 row groups)
+constexpr int kNMQuadN = 64; // per-wave N extent (2 col groups)
+constexpr int kNMBlockThreads = 8 * kWaveSize; // 512 = eight gfx928 wavefronts
-inline bool is_swizzled_pack(int k, int n) {
- return (k == kPackSwizzleK1 && n == kPackSwizzleN1) ||
- (k == kPackSwizzleK2 && n == kPackSwizzleN2);
-}
+// Exact (k, n) pair whose weight buffer uses the iteration-5 tile-contiguous
+// 64-K-stage pack packed[(((k0*(n/128) + nt)*8 + kc)*128 + row)*8 + b] =
+// raw[kk*n + col] (identity [K, N] for every other pair).
+constexpr int kPackK = 512;
+constexpr int kPackN = 3584;
+constexpr int kPackTileN = 128; // N rows per tile-contiguous pack region
+constexpr int kPackStageK = 64; // k bytes per pack stage
+
+using bf16_t = hip_bfloat16;
-// Iteration 5 (kv_b_proj packing round): the (512, 3584) pair switched from
-// the k-stage-major swizzle to the tile-contiguous layout
-// packed[(((k0*(n/128) + nt)*8 + kc)*128 + row)*8 + b]; (2048, 2048) keeps
-// the k-stage-major swizzle byte-identical.
inline bool is_tile_major_pack(int k, int n) {
- return k == kPackSwizzleK2 && n == kPackSwizzleN2;
+ return k == kPackK && n == kPackN;
}
-// Active 2-D macro-tile for the exact glm_tp8_kv_b_proj_m4096 shape among the
-// three mandated candidates plus the iteration-4 aspect flip (0 = 64x64,
-// 1 = 64x128, 2 = 128x64, 3 = 64x128 with 32x32 quadrants). All four
-// instantiations are compiled and dispatched behind the exact (k, n) guard;
-// flipping this constant re-routes the kv_b_proj arm without source surgery.
-constexpr int kActiveKvTile = 3; // 3 = 64x128 with 32x32 quadrants
- // (iteration-4 tile-shape round): the
- // down_proj-validated 16-waves/CU recipe
- // (512 threads, 8 wavefronts/block, 4
- // accs/wave, arch_vgpr <= 64 -> 2 blocks/CU
- // = 16 waves/CU = 4/SIMD) applied as the
- // M-tile 128 -> 64 / N-tile 64 -> 128 aspect
- // flip; config 2 (128x64, 3 blocks x 4 waves
- // = 12 waves/CU, accepted best 169.257 us)
- // stays compiled and is the retention
- // fallback if config 3 regresses (arch_vgpr
- // > 64 -> 1 block/CU = 8 waves -> reject).
-constexpr int kKvTileM = 64;
-constexpr int kKvTileN = 128;
-constexpr int kKvQuadN = 64;
-
-using bf16_t = hip_bfloat16;
-
// ---------------------------------------------------------------------------
// Direct fragment epilogue for one m16n16k32 accumulator fragment.
// Verified gfx928 int8 m16n16k32 accumulator ownership (matches
// du_store_matrix_sync): lane & 15 selects the row, lane >> 4 selects
// col % 4, and x[i] maps to columns col%4 + 4*i. The scale multiply order
-// (float(dot) * x_scale[row] * weight_scale[col]) and the bf16 rounding are
-// identical to the harness reference ((dot.float() * x_scale) * ws.T then
-// .to(bfloat16)), so stored bf16 bits match exactly.
+// (float(dot) * x_scale[row] * weight_scale[col]) and the bf16 rounding match
+// the harness reference ((dot.float() * x_scale) * weight_scale.T -> bf16), so
+// the stored bf16 bits are identical.
// ---------------------------------------------------------------------------
template
-__device__ __forceinline__ void store_prefill_fragment(
+__device__ __forceinline__ void store_fragment_scaled(
const AccFragment& frag,
const float* __restrict__ x_scale,
const float* __restrict__ weight_scale,
@@ -340,45 +402,75 @@ __device__ __forceinline__ void store_prefill_fragment(
}
// ---------------------------------------------------------------------------
-// Iteration 6 (epilogue round): register-batched COALESCED direct-fragment
-// epilogue for the exact (2048, 2048) 128x128 arm. The per-element variant
-// (store_prefill_fragment) above issues one 2-byte bf16 store per lane per
-// element: 64 scattered stores per wave per fragment set, each touching 16
-// rows x 4 columns, i.e. every 32-B store sector only 25% utilized. This
-// epilogue transposes the 4-element groups within each 4-lane column group
-// (lanes r, r+16, r+32, r+48 - a 4x4 transpose, two 2x2 steps with
-// shfl_xor 16 then 32; the lineage-validated pattern accepted on the sibling
-// TP8 workers 0/2 on this DTK, which lowers __shfl_xor to ds_bpermute at the
-// block tail where the LDS pipe is idle), so lane (r, c4) ends up holding
-// the four CONTIGUOUS columns 4*c4 .. 4*c4+3, converts them to bf16, packs
-// 4 bf16 (8 B) and writes ONE 8-byte store per lane (100% store sector
-// efficiency). Only the int32 values are re-routed between lanes; the
-// per-element scale multiply order (float(dot) * x_scale[row] *
-// weight_scale[col]) and the __float2bfloat16 rounding are unchanged, so
-// the stored bits are identical to the per-element store. xs/ws come from
-// caller registers (batched per wave on the last K stage), so the epilogue
-// issues no vmem loads. The row >= m guard is wavefront-uniform (all 64
-// lanes of a wave share the same 16-row window), so the shuffles never mix
-// active and inactive lanes; base_col is a multiple of 64 and n*2 a
-// multiple of 8, so the 8-byte store is aligned.
+// Bank-safe dense 8-byte-group LDS map for the exact-shape swizzled arm.
+// group(row, c) is the 8-byte group index of tile position (row, k 8-byte
+// group c). For a fragment load the 16 lanes of a group share row0 and kk, so
+// row & 15 ranges over 0..15 and c is fixed; (row & 15) -> ((row & 15) +
+// 4*(c & 3)) & 15 is a permutation, giving 16 distinct bank pairs and zero
+// bank conflicts over the whole fragment read.
+// ---------------------------------------------------------------------------
+__device__ __forceinline__ int swizzle_group(int row, int c) {
+ return (((row & 15) + 4 * (c & 3)) & 15) + 16 * c + 128 * (row >> 4);
+}
+
+// One 8-byte LDS store of each half of a 16-byte global chunk.
+__device__ __forceinline__ void store_chunk16_swizzled(
+ int8_t* __restrict__ p, int row, int j, const int4& v) {
+ *reinterpret_cast(p + swizzle_group(row, 2 * j) * 8) =
+ *reinterpret_cast(&v);
+ *reinterpret_cast(p + swizzle_group(row, 2 * j + 1) * 8) =
+ *reinterpret_cast(reinterpret_cast(&v) + 8);
+}
+
+// Direct 8-byte fragment load that reproduces du_load_matrix_sync's per-lane
+// placement exactly: row = row0 + (lane & 15), k group c = (lane >> 4) +
+// (kk >> 3), eight consecutive k bytes in natural order. Writing them
+// straight into the fragment storage lets the compiler emit one 8-byte LDS
+// read per fragment with none of the library's per-byte gather/reassembly
+// VALU, so the int32 accumulation is bit-identical.
+__device__ __forceinline__ void load_frag8_swizzled(
+ DUFragmentBase& frag,
+ const signed char* __restrict__ p,
+ int lane,
+ int row0,
+ int kk) {
+ const int row = row0 + (lane & 15);
+ const int c = (lane >> 4) + (kk >> 3);
+ *reinterpret_cast(&frag.x[0]) =
+ *reinterpret_cast(p + swizzle_group(row, c) * 8);
+}
+
+// Iteration-5 (packing round) B fragment loader for the tile-contiguous plane
+// LDS layout b_tile[kc][n][8] (kc = 8-byte k sub-chunk; plane stride
+// kNMBlockN*8 = 1024 B). Lane l reads the same 8 consecutive k bytes as
+// load_frag8_swizzled's placement (n row = row0 + (lane & 15), k group
+// kc = (lane >> 4) + (kk >> 3)) but from the lane-linear address
+// kc * (kNMBlockN * 8) + row * 8,
+// i.e. ((lane >> 4) * 1024) + ((lane & 15) * 8) relative to row0's plane
+// origin, so each 16-lane group of a 128-B LDS phase spans all 32 banks
+// exactly once (zero bank conflicts) and the compiler emits one ds_read_b64.
+__device__ __forceinline__ void load_frag8_plane(
+ DUFragmentBase& frag,
+ const signed char* __restrict__ p,
+ int lane,
+ int row0,
+ int kk) {
+ const int row = row0 + (lane & 15);
+ const int kc = (lane >> 4) + (kk >> 3);
+ *reinterpret_cast(&frag.x[0]) =
+ *reinterpret_cast(
+ p + kc * (kNMBlockN * 8) + row * 8);
+}
+
// ---------------------------------------------------------------------------
-// Iteration 6 (kv_b_proj epilogue round): plain round-to-nearest-even bf16
-// conversion without the inf/NaN fixup path. hip's __float2bfloat16 lowers
-// (verified on this exact DTK 26.04 / gfx928 toolchain with a micro compile)
-// to an exec-masked sequence per element: v_and (exp mask) + v_cmp_ne +
-// s_and_saveexec + s_xor + v_bfe + v_add3 + v_mov 0 + v_or 0x10000 +
-// v_cmp_eq_u32_sdwa + v_cndmask + s_or_b64 (~11 instructions including three
-// scalar exec manipulations that serialize the wavefront). The masking only
-// changes the result for inf/NaN inputs; for every finite float - the only
-// values this GEMM can produce from int32 accumulators and finite scales -
-// the RNE rounding is exactly u += 0x7fff + ((u >> 16) & 1); return u >> 16,
-// bit-identical to hip's own non-masked branch (same v_bfe + v_add3 +
-// shift). The kernel epilogue is the only place that converts all M*N =
-// 14.68M elements; dropping the mask removes ~6 VALU + 5 SALU per element
-// (~1.38M VALU + 1.15M SALU per launch, ~19% of the measured 7.31M VALU
-// stream) and - the falsifiable axis - removes the
-// s_and_saveexec / s_xor / s_andn2_saveexec / s_or_b64 pairs that serialize
-// the block-tail epilogue behind scalar exec manipulation.
+// Iteration 9 (consolidate round): plain round-to-nearest-even fp32 -> bf16
+// conversion, bit-identical to hip's __float2bfloat16 for every finite input
+// (the only values this GEMM can produce: int32 accumulators times finite
+// scales), but without the exec-masked inf/NaN fixup sequence the DTK 26.04
+// gfx928 lowering emits around it (see the header note). The canonical RNE
+// bit trick: add 0x7fff plus the retained LSB so ties round to even, then
+// truncate. Same form as the sibling reference
+// references/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip on this DTK.
// ---------------------------------------------------------------------------
__device__ __forceinline__ unsigned short bf16_rne_u16(float x) {
union {
@@ -390,21 +482,52 @@ __device__ __forceinline__ unsigned short bf16_rne_u16(float x) {
return static_cast(u >> 16);
}
-template
-__device__ __forceinline__ void store_prefill_fragment_coalesced_scaled(
+// ---------------------------------------------------------------------------
+// Iteration-6 (epilogue round): register-batched COALESCED direct-fragment
+// epilogue for the exact-shape arm, ported from the lineage-validated sibling
+// reference references/int8w8a8-gemm/glm52/TP8/M4096/kv_b_proj.hip. The
+// per-element variant (store_fragment_scaled) issues one 2-byte bf16 store per
+// lane per element: four scattered stores per fragment, 64 per wave, each
+// wavefront store touching 16 rows x 4 columns so every 32-B store sector is
+// only 25% utilized (the bootstrap PMC's 229,376 vmem_write instructions =
+// 896 blocks x 4 waves x 16 fragments x 4). This epilogue transposes the
+// four-element groups within each 4-lane column group (lanes r, r+16, r+32,
+// r+48 - a 4x4 transpose done as two 2x2 steps with shfl_xor 16 then 32, the
+// lineage-validated pattern that lowers to ds_bpermute at the block tail
+// where the LDS pipe is idle), so lane (r, c4) ends up holding the four
+// CONTIGUOUS columns base_col + 4*c4 .. +3, converts them to bf16 and writes
+// ONE 8-byte store (100% store sector efficiency; 16 stores per wave instead
+// of 64). Only the int32 values are re-routed between lanes; the per-element
+// scale multiply order (float(dot) * x_scale[row] * weight_scale[col]) and the
+// __float2bfloat16 rounding are unchanged, so the stored bits are identical.
+// xs is batched by the caller (one load per 16-row row-block) and ws is the
+// float4 of the four per-column scales at col0, 16-B aligned because base_col
+// is a multiple of 16, so the epilogue issues 2 scalar + 8 vector vmem loads
+// per lane instead of 80 scalar loads. Only the exact-shape n-major arm calls
+// this helper, and its dispatch guard is m >= 128 with m % 128 == 0, so every
+// row of every block is in bounds and no tail-M mask is needed: the shuffles
+// never mix active and inactive lanes, and n*2 is a multiple of 8, so the
+// 8-byte store is aligned.
+//
+// Iteration 9 (consolidate round): the bf16 conversion below is a plain RNE
+// (bf16_rne_u16) instead of __float2bfloat16. The exact code object of the
+// current best showed the DTK 26.04 lowering wrapping every element in an
+// exec-masked inf/NaN fixup (~9 vector instructions + two s_and_saveexec /
+// s_andn2_saveexec scalar manipulations) whose only effect vs RNE is on
+// inf/NaN inputs; int32 accumulators times finite scales are always finite,
+// and for finite floats the two forms produce the identical 16-bit pattern.
+// ---------------------------------------------------------------------------
+template
+__device__ __forceinline__ void store_fragment_coalesced_scaled(
const AccFragment& frag,
float xs,
- const float4& ws,
+ const float* __restrict__ weight_scale,
bf16_t* __restrict__ out,
int base_row,
int base_col,
- int m,
int n,
int lane) {
const int row = base_row + (lane & 15);
- if (row >= m) {
- return; // tail-M masking: padded rows never write (wavefront-uniform)
- }
const int c4 = lane >> 4; // 0..3
const int x0 = frag.x[0];
const int x1 = frag.x[1];
@@ -431,165 +554,62 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced_scaled(
const int f2 = hi ? t0 : a2;
const int f3 = hi ? t1 : a3;
- // Lane (r, c4) now owns columns base_col + 4*c4 .. +3 (8 B, 8-byte
- // aligned: base_col is a multiple of 16, n*2 is a multiple of 8); ws is
- // the float4 at weight_scale + base_col + 4*c4 preloaded by the caller.
+ // Lane (r, c4) now owns columns base_col + 4*c4 .. +3 (8 B, 8-byte aligned:
+ // base_col is a multiple of 16 and n is a multiple of 4); ws is the float4
+ // of per-column scales at col0.
const int col0 = base_col + 4 * c4;
+ const float4 ws = *reinterpret_cast(weight_scale + col0);
const float v0 = static_cast(f0) * xs * ws.x;
const float v1 = static_cast(f1) * xs * ws.y;
const float v2 = static_cast(f2) * xs * ws.z;
const float v3 = static_cast(f3) * xs * ws.w;
- uint64_t packed;
- if constexpr (kRneOnly) {
- // Iteration 6 (kv_b_proj epilogue round): plain RNE (see bf16_rne_u16) -
- // bit-identical to __float2bfloat16 for every finite input, without the
- // exec-masked inf/NaN fixup path.
- packed = static_cast(bf16_rne_u16(v0)) |
- (static_cast(bf16_rne_u16(v1)) << 16) |
- (static_cast(bf16_rne_u16(v2)) << 32) |
- (static_cast(bf16_rne_u16(v3)) << 48);
- } else {
- packed =
- static_cast(
- static_cast(__float2bfloat16(v0))) |
- (static_cast(
- static_cast(__float2bfloat16(v1)))
- << 16) |
- (static_cast(
- static_cast(__float2bfloat16(v2)))
- << 32) |
- (static_cast(
- static_cast(__float2bfloat16(v3)))
- << 48);
- }
- *reinterpret_cast(out + static_cast(row) * n + col0) =
- packed;
-}
-
-// ---------------------------------------------------------------------------
-// Direct 8-byte LDS fragment load (lineage-validated on gfx928/DTK 26.04 for
-// m16n16k32 int8 fragments). du_load_matrix_sync's int8 loaders assign
-// x[0..7] = 8 consecutive bytes at (lane & 15) * ldm + ((lane >> 4) << 3) for
-// matrix_b col_major, but the compiler lowers that to per-byte ds_read_u8 +
-// mask/OR reassembly VALU. Writing the same 8 bytes directly into the
-// fragment storage keeps the operand bit pattern identical (exact int32
-// accumulation unchanged) and lets the compiler emit one ds_read2_b64 per
-// fragment straight into the v_mmac operand.
-// ---------------------------------------------------------------------------
-__device__ __forceinline__ void load_fragment8(
- DUFragmentBase& frag,
- const signed char* __restrict__ base,
- int stride,
- int lane) {
- const int off = (lane & 15) * stride + ((lane >> 4) << 3);
- *reinterpret_cast(&frag.x[0]) =
- *reinterpret_cast(base + off);
-}
-
-// ---------------------------------------------------------------------------
-// Iteration 5 (packing round): 8-byte fragment load for the swizzled plane
-// LDS layout b_tile[kc][n][8] (kc = 8-byte k sub-chunk, n row stride 8 B,
-// plane stride 1024 B). Lane l reads the same 8 consecutive k bytes as
-// load_fragment8 (n row = lane&15, k chunk = lane>>4) but from
-// (lane >> 4) * 1024 + (lane & 15) * 8
-// relative to the fragment origin, so the 64 lanes hit 16 distinct bank
-// pairs per 128-B phase (all 32 banks once) -> zero bank conflicts, the
-// 4-cycle minimum for a 512-B fragment.
-// ---------------------------------------------------------------------------
-__device__ __forceinline__ void load_fragment8_plane(
- DUFragmentBase& frag,
- const signed char* __restrict__ base,
- int lane) {
- const int off = ((lane >> 4) << 10) + ((lane & 15) << 3);
- *reinterpret_cast(&frag.x[0]) =
- *reinterpret_cast(base + off);
-}
-
-// ---------------------------------------------------------------------------
-// Plane-layout 8-byte fragment load with a parameterized plane stride
-// (plane_stride = kBlockN * 8 bytes). Identical to load_fragment8_plane for
-// kBlockN == 128 (plane stride 1024 B); for narrower tiles the plane is
-// smaller but the bank pattern is unchanged: lane l reads the 8 bytes at
-// (lane >> 4) * plane_stride + (lane & 15) * 8, so each 16-lane group of a
-// 128-B phase covers all 32 banks exactly once (zero bank conflicts) for any
-// plane stride.
-// ---------------------------------------------------------------------------
-__device__ __forceinline__ void load_fragment8_plane_stride(
- DUFragmentBase& frag,
- const signed char* __restrict__ base,
- int plane_stride,
- int lane) {
- const int off = ((lane >> 4) * plane_stride) + ((lane & 15) << 3);
- *reinterpret_cast(&frag.x[0]) =
- *reinterpret_cast(base + off);
+ const uint64_t packed =
+ static_cast(bf16_rne_u16(v0)) |
+ (static_cast(bf16_rne_u16(v1)) << 16) |
+ (static_cast(bf16_rne_u16(v2)) << 32) |
+ (static_cast(bf16_rne_u16(v3)) << 48);
+ *reinterpret_cast(
+ out + static_cast(row) * n + col0) = packed;
}
-// ---------------------------------------------------------------------------
-// Compile-time B-fragment layout selector. row_major / col_major are tag
-// types in du::dumma, so the layout template argument must be selected as a
-// type (a conditional expression over type names is not a valid template
-// argument). kNMajorB == true (n-major packed B, exact (2048, 2048)) ->
-// col_major fragments (8 contiguous k bytes per lane, load_fragment8);
-// kNMajorB == false (raw [K, N] row-major B) -> row_major fragments (the
-// accepted-best loader).
-// ---------------------------------------------------------------------------
-template
-struct b_frag_layout {
- using type = row_major;
-};
-template <>
-struct b_frag_layout {
- using type = col_major;
-};
-
// ---------------------------------------------------------------------------
// Large-M prefill: 128x64 output tile per block, four wavefronts of 64 lanes.
// Each wave owns a 64x32 quadrant (eight m16n16k32 int32 accumulators); the
// block cooperatively stages A[128,64] from x_q (row-major, stride k) and
-// B[64,64] into a single-buffered LDS stage. Two barriers per stage: one
-// after the cooperative load, one before the next stage overwrites LDS.
-// A fragments use the library du_load_matrix_sync row-major loader. B is
-// templated on its staged layout:
-// * kNMajorB == true: B is packed n-major (packed[n*K + k]) and staged
-// n-major (n row, 80 B stride, k contiguous); each B fragment is 8
-// contiguous k bytes per lane, loaded by load_fragment8 (one 8-byte LDS
-// read per fragment), eliminating the ds_read_u8 byte-gather + VALU
-// reassembly of the accepted-best kernel.
-// * kNMajorB == false: B stays raw [K, N] row-major, staged k-major with
-// row-major library fragments (accepted-best path, unchanged).
-// Dispatch guarantees m % 128 == 0, n % 64 == 0, k % 64 == 0, so every
-// global load/store is in-bounds and 16-byte aligned.
+// B[64,64] from the raw [K, N] row-major weight into a single-buffered LDS
+// stage. Two barriers per 64-K stage: one after the cooperative load, one
+// before the next stage overwrites the buffer. Dispatch guarantees
+// m % 128 == 0, n % 64 == 0, k % 64 == 0, so every global load/store is
+// in-bounds and 16-byte aligned.
// ---------------------------------------------------------------------------
-template
-__global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x64_kernel(
- const int8_t* __restrict__ x_q,
- const int8_t* __restrict__ weight,
- const float* __restrict__ x_scale,
- const float* __restrict__ weight_scale,
- bf16_t* __restrict__ out,
- int m,
- int n,
- int k) {
+__global__ __launch_bounds__(kBlockThreads) void
+w8a8_dumma_prefill_128x64_kernel(const int8_t* __restrict__ x_q,
+ const int8_t* __restrict__ weight,
+ const float* __restrict__ x_scale,
+ const float* __restrict__ weight_scale,
+ bf16_t* __restrict__ out,
+ int m,
+ int n,
+ int k) {
const int tid = static_cast(threadIdx.x);
const int wave = tid / kWaveSize;
const int lane = tid % kWaveSize;
- const int wave_row = wave >> 1;
- const int wave_col = wave & 1;
+ const int wave_row = wave >> 1; // 0 or 1
+ const int wave_col = wave & 1; // 0 or 1
const int m0 = static_cast(blockIdx.y) * kBlockM;
const int n0 = static_cast(blockIdx.x) * kBlockN;
const int local_row = wave_row * 64;
const int local_col = wave_col * 32;
- // Single-buffered stage: A[128, 80] + B[64, 80] = 15,360 B/block
- // (4 blocks/CU fit the 64 KiB LDS budget; the padded 80-byte strides are
- // 16-byte-aligned and break the 64-byte LDS bank periodicity).
+ // Single-buffered stage: A[128, 80] + B[64, 80] = 15,360 B/block (four
+ // blocks fit the 64 KiB per-CU LDS budget; the padded 80-byte strides are
+ // 16-byte aligned and break the 64-byte LDS bank periodicity).
__shared__ __align__(16) int8_t a_tile[kBlockM * kAStride];
__shared__ __align__(16) int8_t b_tile[kStageK * kBStride];
DUFragment
a_frag0, a_frag1;
- DUFragment::type>
+ DUFragment
b_frag0, b_frag1;
DUFragment
acc00, acc01, acc10, acc11, acc20, acc21, acc30, acc31;
@@ -602,21 +622,18 @@ __global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x64_kerne
du_fill_fragment(acc30, 0);
du_fill_fragment(acc31, 0);
- // Cooperative staging: A[128,64] is 512 int4s (two per thread), B[64,64]
- // is 256 int4s (one per thread).
- // A: thread tid owns row (tid*16)/64 and 16-byte column group
+ // Cooperative staging: A[128,64] is 512 int4s (two per thread), B[64,64] is
+ // 256 int4s (one per thread).
+ // A: thread tid owns row (tid*16)/64 and the 16-byte K-column group
// (tid*16)%64 for rows [0,64) and [64,128).
- // B (kNMajorB == false): raw [K, N] row-major; thread tid owns K row
- // tid>>2 and the 16-byte N-column group (tid&3)*16 (16 consecutive N
- // values at a fixed K row -> one aligned int4 in global and in LDS).
- // B (kNMajorB == true): packed n-major; thread tid owns N row tid>>2 and
- // the 16-byte K-column group (tid&3)*16 (16 consecutive K values at a
- // fixed N row -> one aligned int4 in global and in the n-major LDS).
+ // B: raw [K, N] row-major; thread tid owns K row tid>>2 and the 16-byte
+ // N-column group (tid&3)*16 (16 consecutive N values at a fixed K row,
+ // one aligned int4 in global and in LDS).
const int vector_byte_offset = tid * static_cast(sizeof(int4));
const int stage_row = vector_byte_offset / kStageK; // 0..63
const int stage_col = vector_byte_offset - stage_row * kStageK;
- const int b_k = tid >> 2; // K row (false) / N row (true)
- const int b_nc = (tid & 3) * 16; // N group (false) / K group (true)
+ const int b_k = tid >> 2; // K row 0..63
+ const int b_nc = (tid & 3) * 16; // N-column group 0/16/32/48
for (int k0 = 0; k0 < k; k0 += kStageK) {
*reinterpret_cast(a_tile + stage_row * kAStride + stage_col) =
@@ -625,43 +642,23 @@ __global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x64_kerne
*reinterpret_cast(a_tile +
(stage_row + kBlockM / 2) * kAStride + stage_col) =
*reinterpret_cast(
- x_q + static_cast(m0 + stage_row + kBlockM / 2) * k +
- k0 + stage_col);
- if constexpr (kNMajorB) {
- // n-major packed B: 16 consecutive k bytes of the (n0 + b_k) row.
- *reinterpret_cast(b_tile + b_k * kBStride + b_nc) =
- *reinterpret_cast(
- weight + static_cast(n0 + b_k) * k + k0 + b_nc);
- } else {
- // raw [K, N] row-major B: 16 consecutive n bytes of the (k0 + b_k) row.
- *reinterpret_cast(b_tile + b_k * kBStride + b_nc) =
- *reinterpret_cast(
- weight + static_cast(k0 + b_k) * n + n0 + b_nc);
- }
+ x_q + static_cast(m0 + stage_row + kBlockM / 2) * k + k0 +
+ stage_col);
+ *reinterpret_cast(b_tile + b_k * kBStride + b_nc) =
+ *reinterpret_cast(
+ weight + static_cast(k0 + b_k) * n + n0 + b_nc);
__syncthreads();
// Each wave consumes its 64x32 quadrant: eight m16n16k32 MMACs per kk.
- // Accumulation order: k0-outer over 64-K stages, kk-inner (kk=0 then
- // kk=32), matching the reference int32 accumulation.
+ // Accumulation order: k0-outer over 64-K stages, kk-inner (kk = 0 then
+ // 32), matching the reference int32 accumulation exactly.
#pragma unroll
for (int kk = 0; kk < kStageK; kk += kTileK) {
- if constexpr (kNMajorB) {
- // n-major LDS tile: n rows stride kBStride, k contiguous within a
- // row; col_major fragment slots are 8 consecutive k bytes at
- // (lane & 15) * kBStride + ((lane >> 4) << 3) relative to the
- // fragment origin -> one ds_read2_b64 per fragment per lane.
- load_fragment8(b_frag0, b_tile + local_col * kBStride + kk, kBStride,
- lane);
- load_fragment8(b_frag1,
- b_tile + (local_col + kTileN) * kBStride + kk,
- kBStride, lane);
- } else {
- du_load_matrix_sync(b_frag0, b_tile + kk * kBStride + local_col,
- kBStride);
- du_load_matrix_sync(b_frag1,
- b_tile + kk * kBStride + local_col + kTileN,
- kBStride);
- }
+ du_load_matrix_sync(b_frag0, b_tile + kk * kBStride + local_col,
+ kBStride);
+ du_load_matrix_sync(b_frag1,
+ b_tile + kk * kBStride + local_col + kTileN,
+ kBStride);
du_load_matrix_sync(a_frag0, a_tile + local_row * kAStride + kk,
kAStride);
du_load_matrix_sync(a_frag1,
@@ -689,101 +686,84 @@ __global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x64_kerne
const int base_row = m0 + local_row;
const int base_col = n0 + local_col;
- store_prefill_fragment(acc00, x_scale, weight_scale, out, base_row,
- base_col, m, n, lane);
- store_prefill_fragment(acc01, x_scale, weight_scale, out, base_row,
- base_col + kTileN, m, n, lane);
- store_prefill_fragment(acc10, x_scale, weight_scale, out,
- base_row + kTileM, base_col, m, n, lane);
- store_prefill_fragment(acc11, x_scale, weight_scale, out,
- base_row + kTileM, base_col + kTileN, m, n, lane);
- store_prefill_fragment(acc20, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col, m, n, lane);
- store_prefill_fragment(acc21, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col + kTileN, m, n,
- lane);
- store_prefill_fragment(acc30, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col, m, n, lane);
- store_prefill_fragment(acc31, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col + kTileN, m, n,
- lane);
+ store_fragment_scaled(acc00, x_scale, weight_scale, out, base_row, base_col,
+ m, n, lane);
+ store_fragment_scaled(acc01, x_scale, weight_scale, out, base_row,
+ base_col + kTileN, m, n, lane);
+ store_fragment_scaled(acc10, x_scale, weight_scale, out,
+ base_row + kTileM, base_col, m, n, lane);
+ store_fragment_scaled(acc11, x_scale, weight_scale, out,
+ base_row + kTileM, base_col + kTileN, m, n, lane);
+ store_fragment_scaled(acc20, x_scale, weight_scale, out,
+ base_row + 2 * kTileM, base_col, m, n, lane);
+ store_fragment_scaled(acc21, x_scale, weight_scale, out,
+ base_row + 2 * kTileM, base_col + kTileN, m, n, lane);
+ store_fragment_scaled(acc30, x_scale, weight_scale, out,
+ base_row + 3 * kTileM, base_col, m, n, lane);
+ store_fragment_scaled(acc31, x_scale, weight_scale, out,
+ base_row + 3 * kTileM, base_col + kTileN, m, n, lane);
}
// ---------------------------------------------------------------------------
-// Iteration 4 (tile-aspect round) + iteration 5 (packing round): 128x128
-// tile for the exact q_b_proj (k, n) == (2048, 2048) swizzled-pack B arm.
-// Four wavefronts of 64 lanes; each wave owns a 64x64 quadrant (four A x
-// four B m16n16k32 fragments = 16 MMACs per kk, 16 int32 accumulator
-// fragments). The block cooperatively stages A[128,64] (2 int4/thread,
-// 80-byte-strided rows) and B[128,64] from the iteration-5 swizzled pack (2
-// int4/thread, 16-byte plane chunks) into a single-buffered LDS stage:
-// A[128,80] (10,240 B) + B plane tile [8][128][8] (8,192 B) = 18,432
-// B/block (2 blocks/CU = 36,864 B <= 64 KiB); two __syncthreads per stage.
-// A fragments use the library du_load_matrix_sync row-major loader
-// (4 x ds_read2_b32 per kk per wave), B fragments use load_fragment8_plane
-// (lane-linear 8-byte reads, zero LDS bank conflicts). A-side global
-// re-read halves (A reuse 32 -> 16, 268 -> 134 MiB) while B reuse stays 32
-// (134 MiB): per-byte A:B tile traffic balances 1:1, total 402 -> 268 MiB
-// (-33%). Grid dim3(N/128, M/128) = 512 blocks (~4.3/CU). Dispatch
-// guarantees m % 128 == 0, n == 2048, k % 64 == 0, so every global
-// load/store is in-bounds and 16-byte aligned. Accumulation order stays
-// k0-outer / kk-inner with the same element-to-slot fragment mapping, so the
-// int32 accumulation is bit-identical to the accepted kernel.
+// Exact-shape (k, n) == (512, 3584) large-M prefill: 128x128 macro-tile per
+// block, EIGHT gfx928 wavefronts of 64 lanes (512 threads). Each wave owns a
+// 32x64 quadrant = kAccM 2 A fragments x kAccN 4 B fragments = eight
+// m16n16k32 int32 accumulators (8 MMACs per kk). Cooperative A+B LDS staging
+// is retained (each A row-block feeds 4 B fragments x 2 kk, each B fragment
+// feeds 2 A row-blocks x 2 kk) and every fragment read is one direct 8-byte
+// LDS load. A comes from the identity [M, K] x_q (row-major, stride k) and
+// keeps the dense swizzled 8-byte-group LDS layout; B comes from the
+// iteration-5 one-time tile-contiguous 64-K-stage pack
+// packed[(((k0*(n/128)+nt)*8+kc)*128+row)*8+b] = raw[kk*n+col] and uses the
+// plane LDS layout [kc][128][8], so its staging is a plain contiguous 16-B
+// vector copy and its fragment loads are lane-linear and bank-conflict-free.
+// Two barriers per 64-K stage, single-buffered, int32 accumulation resident
+// through the whole K loop, no split-K (grid dim3(28, 32) = 896 blocks over
+// 120 CUs, ~3.7 waves of the 240 co-resident blocks). Dispatch guarantees
+// k == 512, n == 3584, m >= 128 and m % 128 == 0, so every global load is
+// in-bounds and 16-byte aligned. The two-argument __launch_bounds__ caps the
+// register file at the 64-VGPR 2-blocks/CU residency point for 512-thread
+// blocks (65536 VGPR per CU / (2 x 512 lanes) = 64): 2 blocks x 8 waves =
+// 16 waves/CU = 4 waves/SIMD instead of the measured 120-VGPR 128x128
+// four-wave form at 2 waves/SIMD. The iteration-6 epilogue is kept
+// register-only: 2 batched per-row scale loads, one aligned float4 per-column
+// scale load per fragment and one 8-byte coalesced store per lane per
+// fragment, so it adds no LDS, no barrier and no occupancy pressure.
// ---------------------------------------------------------------------------
-__global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x128_kernel(
- const int8_t* __restrict__ x_q,
- const int8_t* __restrict__ weight,
- const float* __restrict__ x_scale,
- const float* __restrict__ weight_scale,
- bf16_t* __restrict__ out,
- int m,
- int n,
- int k) {
- constexpr int kBlockM128 = 128;
- constexpr int kBlockN128 = 128;
- constexpr int kTileStride = kStageK + kBPad; // 64 -> 80-byte row stride
- constexpr int kTileBytes = kBlockM128 * kTileStride; // 10,240 B per tile
- // Iteration 5 (packing round): B tile is the swizzled plane layout
- // [kc][n][8] (8 planes of 128 n rows x 8 B, plane stride 1024 B) so each
- // lane's 8-byte fragment chunk is lane-linear and bank-conflict-free.
- constexpr int kBTileBytes = 8 * kBlockN128 * 8; // 8,192 B/block
- constexpr int kBTilePlane = 1024; // 128 n rows x 8 B
-
+__global__ __launch_bounds__(kNMBlockThreads, 2) void
+w8a8_dumma_prefill_128x128_nmajor_kernel(const int8_t* __restrict__ x_q,
+ const int8_t* __restrict__ weight,
+ const float* __restrict__ x_scale,
+ const float* __restrict__ weight_scale,
+ bf16_t* __restrict__ out,
+ int m,
+ int n,
+ int k) {
+ // Iteration 10 (consolidate round): the dispatch guard for this arm is
+ // is_tile_major_pack(k, n) && m >= kNMBlockM && m % kNMBlockM == 0, so the
+ // 128-row tile is always fully in bounds and no tail-M masking is needed.
+ (void)m;
const int tid = static_cast(threadIdx.x);
- const int wave = tid / kWaveSize;
+ const int wave = tid / kWaveSize; // 0..7
const int lane = tid % kWaveSize;
- const int wave_row = wave >> 1;
- const int wave_col = wave & 1;
- const int m0 = static_cast(blockIdx.y) * kBlockM128;
- const int n0 = static_cast(blockIdx.x) * kBlockN128;
- const int local_row = wave_row * (kBlockM128 / 2); // 0 or 64
- const int local_col = wave_col * (kBlockN128 / 2); // 0 or 64
- const int base_row = m0 + local_row;
- const int base_col = n0 + local_col;
+ const int wave_row = wave >> 1; // 0..3
+ const int wave_col = wave & 1; // 0..1
+ const int m0 = static_cast(blockIdx.y) * kNMBlockM;
+ const int n0 = static_cast(blockIdx.x) * kNMBlockN;
+ const int local_row = wave_row * kNMQuadM; // 0, 32, 64, 96
+ const int local_col = wave_col * kNMQuadN; // 0 or 64
- // Iteration 6 (epilogue round): per-lane register-batched scales. Each
- // lane owns 4 rows (base_row + 16*i + (lane & 15), i = 0..3) and 4
- // weight_scale float4s (base_col + 16*j + 4*(lane >> 4), j = 0..3) across
- // its 16 fragments; they are loaded exactly once on the last K stage
- // before the final protective barrier, so the coalesced epilogue below is
- // pure compute + 16 eight-byte stores with no interleaved vmem loads.
- float xs_m[4];
- float4 ws_m[4];
-
- // Single-buffered stage: A[128,80] (10,240 B) + B plane tile [8][128][8]
- // (8,192 B) = 18,432 B/block (2 blocks/CU = 36,864 B <= 64 KiB; the A rows
- // keep the padded 80-byte stride, the B tile is plane-swizzled so fragment
- // reads are lane-linear with zero bank conflicts).
- __shared__ __align__(16) int8_t a_tile[kTileBytes];
- __shared__ __align__(16) int8_t b_tile[kBTileBytes];
+ // A 8 KiB (dense swizzled 8-byte groups) + B 8 KiB (iteration-5
+ // tile-contiguous plane [kc][128][8]) = 16,384 B/block.
+ __shared__ __align__(16) int8_t a_tile[kNMBlockM * kStageK];
+ __shared__ __align__(16) int8_t b_tile[kNMBlockN * kStageK];
DUFragment
- a_frag0, a_frag1, a_frag2, a_frag3;
+ a_frag0, a_frag1;
DUFragment
b_frag0, b_frag1, b_frag2, b_frag3;
DUFragment
- acc00, acc01, acc02, acc03, acc10, acc11, acc12, acc13,
- acc20, acc21, acc22, acc23, acc30, acc31, acc32, acc33;
+ acc00, acc01, acc02, acc03, acc10, acc11, acc12, acc13;
du_fill_fragment(acc00, 0);
du_fill_fragment(acc01, 0);
du_fill_fragment(acc02, 0);
@@ -792,102 +772,59 @@ __global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_prefill_128x128_kern
du_fill_fragment(acc11, 0);
du_fill_fragment(acc12, 0);
du_fill_fragment(acc13, 0);
- du_fill_fragment(acc20, 0);
- du_fill_fragment(acc21, 0);
- du_fill_fragment(acc22, 0);
- du_fill_fragment(acc23, 0);
- du_fill_fragment(acc30, 0);
- du_fill_fragment(acc31, 0);
- du_fill_fragment(acc32, 0);
- du_fill_fragment(acc33, 0);
- // Cooperative staging: A[128,64] is 512 int4s (two per thread) and
- // B[128,64] from the iteration-5 swizzled pack is 512 int4s (two per
- // thread). Thread tid owns linear int4 slots tid and tid + 256; for A a
- // slot maps to row linear >> 2 and the 16-byte column group
- // (linear & 3) * 16 (an M row of x_q, one aligned int4 in global and in
- // LDS); for B a slot maps to plane linear >> 6 and n-pair linear & 63
- // (two n rows of one 8-byte k sub-chunk, one aligned int4 in the pack and
- // in the plane LDS tile).
+ // Cooperative staging (iteration-8 resource round): with 512 threads each
+ // thread owns exactly ONE 16-byte chunk of A and ONE 16-byte chunk of B, so
+ // the staging payload per thread halves (1 int4 + 1 int4 instead of
+ // 2 + 2) - this is what keeps the live-range budget compatible with the
+ // 64-VGPR / 2-blocks/CU occupancy point.
+ // A: thread tid owns row tid >> 2 and the 16-byte K column group
+ // (tid & 3) * 16; the swizzled 8-byte-group LDS store is unchanged.
+ // B (iteration-5 packing round): the tile-contiguous pack lays each
+ // (64-K stage, 128-N tile) B tile out as one contiguous 8192-B region
+ // already in the LDS plane order [kc][128 N rows][8 bytes], so staging
+ // is a plain 1 x 16-B contiguous vector copy per thread.
+ const int vector_byte_offset = tid * static_cast(sizeof(int4));
+ const int stage_row = vector_byte_offset / kStageK; // 0..127
+ const int stage_col = vector_byte_offset - stage_row * kStageK; // 0,16,32,48
+ const int b_tile_off = vector_byte_offset; // 0..8176
+
+ // Iteration 13 (consolidate round): one-stage register prefetch of the
+ // cooperative staging payload. The A row base and the B per-stage stride are
+ // loop-invariant, so the next 64-K A/B chunks are loaded while the current
+ // stage's fragment reads and MMACs are still draining; at the consumer side
+ // the load is already in flight across the trailing __syncthreads() instead
+ // of being issued and drained within a handful of instructions at the stage
+ // top. The loaded values, their order, the LDS bytes and the int32
+ // accumulation are unchanged, so every stored bit is unchanged.
+ const int8_t* const a_stage_base =
+ x_q + static_cast(m0 + stage_row) * k + stage_col;
+ const int64_t b_stage_step =
+ static_cast(n / kNMBlockN) * (kNMBlockN * kStageK);
+ const int64_t b_stage_base =
+ static_cast(n0 / kNMBlockN) * (kNMBlockN * kStageK);
+ int4 a_chunk = *reinterpret_cast(a_stage_base);
+ int4 b_chunk =
+ *reinterpret_cast(weight + b_stage_base + b_tile_off);
+
for (int k0 = 0; k0 < k; k0 += kStageK) {
- const int kstage = k0 >> 6; // 64-k group index inside the swizzled pack
-#pragma unroll
- for (int i = 0; i < 2; ++i) {
- const int linear = tid + i * kBlockThreads;
- const int row = linear >> 2; // 0..127
- const int col = (linear & 3) * 16; // 0/16/32/48
- *reinterpret_cast