+ ${readOnly ? null : html`
+ `}
${(() => {
const v = variantIndex ? variantIndex.get(shape) : null;
if (!v) {
@@ -940,10 +1018,11 @@ export default function DcuKernelAutoOptDetail({ taskId, data }) {
+ ${readOnly ? html`
以只读方式嵌入:不提供发布/同步操作。
` : null}
${showSkills ? html`
- <${SkillLibrary} taskId=${taskId} onClose=${() => setShowSkills(false)} />
+ <${SkillLibrary} ctx=${ctx} onClose=${() => setShowSkills(false)} />
` : null}
`;
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 57fffc1a..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) {
@@ -468,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) {
@@ -488,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`
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_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
index f3f0b80c..faa21706 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_preflight.py
+++ b/metainfer/tasks/dcu_kernel_auto_opt/tests/test_gpu_preflight.py
@@ -2,10 +2,28 @@
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 ==================================
================================================================================================
@@ -17,6 +35,20 @@
======================================== 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)
@@ -27,28 +59,72 @@ def test_parse_hy_smi_rows():
assert parsed[1]["temp_c"] == 49.0
-def test_usability_is_free_vram_not_percentage():
- """The rule is `free >= 4GB`, so a 91%-used card is still usable."""
- ok = gp.check_gpu(0, {"vram_percent": 91.0, "util_percent": 0.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)
- too_full = gp.check_gpu(1, {"vram_percent": 94.0, "util_percent": 0.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 ok["usable"] is True and ok["free_gb"] == 5.76
- assert ok["measurement_suspect"] is False
- assert too_full["usable"] is False and too_full["free_gb"] == 3.84
- assert any("VRAM free" in r for r in too_full["reasons"])
- # a full device is a capacity problem, not a pollution problem
- assert too_full["measurement_suspect"] is False
- assert gp.DEFAULT_MIN_FREE_GB == 4.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_marks_suspect_without_blocking():
+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 True # usage is allowed (VRAM ok)
+ 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"])
- assert any("power" 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"])
@@ -63,28 +139,41 @@ def test_idle_foreign_holders_are_not_suspicion():
assert any("foreign" in n for n in idle["notes"])
-def test_preflight_prefers_clean_then_suspect_then_over_limit(monkeypatch):
+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": 55.0,
+ 1: {"vram_percent": 5.0, "util_percent": 0.0,
"power_w": 250.0},
- 2: {"vram_percent": 30.0, "util_percent": 3.0},
- 3: {"vram_percent": 97.0, "util_percent": 10.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, 2]
- assert plan["suspect_ids"] == [1]
+ assert plan["clean_ids"] == [0, 1]
+ assert plan["suspect_ids"] == []
+ assert plan["blocked_ids"] == [2, 3]
assert plan["over_limit_ids"] == [3]
- # usable devices first (clean before suspect); the over-limit device drops
- # out while any device is within the limit
- assert plan["preferred"] == [0, 2, 1]
- assert 3 not in plan["preferred"]
+ # only gate-passing devices are ever offered
+ assert plan["preferred"] == [0, 1]
+ assert "HCU 4.0%" in plan["gate_reasons"]["2"][0]
-def test_over_limit_devices_return_when_nothing_else_is_usable(monkeypatch):
+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},
@@ -93,7 +182,20 @@ def test_over_limit_devices_return_when_nothing_else_is_usable(monkeypatch):
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"] == [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):
@@ -119,3 +221,256 @@ def test_disabled_preflight_returns_default_order():
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_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_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_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_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/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/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/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 f5387ae1..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,457 +1,78 @@
-// @@variant shape=glm_tp8_fused_qkv_a_proj_m4096 commit=99b9a733eddfc7196f3d186b80affe4324351ac4 added=2026-09-11
-// median_us=883.7 p90_us=885.5 speedup=77.7 baseline_us=6.866e+04
-// source=test-9-11-1-b01955b3/iteration_001/glm_tp8_fused_qkv_a_proj_m4096
+// @@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] )
//
-// Round 1 (mandatory establish_arch) -- packed_weight architecture family:
+// 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.
//
-// Bottleneck hypothesis, from the trusted PMC of the current best source
-// (the correctness-first bootstrap: w8a8_dumma_prefill_tiled_kernel
-// <128, 64, 128>, identity [K, N] weight, library du_load_matrix_sync
-// loaders, LDS row strides 128 (A) / 64 (B), grid 41 x 32 = 1312 blocks,
-// 256 threads, 2 blocks/CU):
-// median 2920.88 us / p90 2962.82 us, lds_instructions 21,999,616,
-// lds_bank_conflicts 274,071,552 (~12.5 per LDS instruction),
-// lds_wait_instructions 50,181,819, valu_instructions 86,560,512,
-// vmem_read_instructions 1,574,400, and a counter-derived operator HBM
-// bandwidth of only 50.6 GB/s (147.8 MB per replay at a 90.7% L2 hit
-// rate). The exact-shape stage loop is LDS-issue/conflict bound, not
-// MMA- or HBM-bound: both LDS row strides alias every row onto one bank
-// phase (128 mod 32 == 0, 64 mod 32 == 0), and the row_major matrix_b
-// loader of du_load_matrix_sync lowers each m16n16k32 int8 B fragment to
-// 8 ds_read_u8 at 16-byte spacing plus a byte-reassembly chain.
+// (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.
//
-// Focused change (ONE architecture family -- packed_weight -- entirely
-// inside csrc/w8a8_gemm_hip.hip):
-// (1) launch_pack_w8a8_weight 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, same buffer allocation, so every captured
-// address is unchanged). Every other (K, N) keeps the identity
-// device-to-device copy.
-// (2) the exact-shape guard dispatches the new
-// w8a8_dumma_prefill_packedb_kernel<128, 64, 128>. It keeps the
-// EXACT bootstrap geometry -- grid (N/64) x (M/128) = 41 x 32 = 1312
-// blocks, 256 threads = 4 gfx928 wavefronts, one 64 x 32 quadrant of
-// eight m16n16k32 int32 accumulator fragments per wave, the same
-// single-buffered 128-K LDS stage with two __syncthreads per stage,
-// the same k0-outer / kk-inner ascending int32 accumulation order and
-// the same coalesced bf16 epilogue -- and changes only the operand
-// path to the packed-weight architecture:
-// * B is staged n-major (b_tile[n_row][k]) from the packed [N, K]
-// weight, so a col_major matrix_b fragment is 8 contiguous K
-// bytes per lane and every staging vector is one coalesced
-// 16-byte global read;
-// * both LDS tiles use the bank-safe row stride StageK + 8 = 136
-// bytes (34 dwords): 8-byte aligned with an odd 8-byte count
-// (136/8 = 17), so the 16 rows of each 16-lane ds_read_b64 phase
-// land on 16 distinct bank pairs (the conflict-free 4-phase
-// floor), and 136 % 16 == 8 keeps the two-half int64 staging
-// form;
-// * both fragment loaders are the explicit load_frag8: ONE aligned
-// 8-byte load per lane per fragment (row = lane & 15, k-quarter =
-// lane >> 4), replacing the byte-reassembly library path with a
-// single ds_read_b64 while producing identical fragment values.
-// (3) LDS/block grows 24,576 -> 26,112 B (A[128, 136] + B[64, 136]);
-// 2 x 26,112 = 52,224 <= 65,536 B keeps TWO resident blocks/CU =
-// 8 waves/CU, so grid_blocks (1312), waves_per_block (4) and the
-// active CU count (120) are unchanged.
-// (4) the scalar fallback decodes the n-major layout when
-// (k, n) == (6144, 2624) with ascending kk, so the paired M=2 /
-// M=16 API shapes with the same (N, K) stay byte-exact; the generic
-// <64, 64, 128> identity arm is untouched.
+// 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.
//
-// Correctness: the same weight elements reach the same fragment slots in
-// the same ascending k-chunk sequence (k0 outer over 48 stages of 128 K,
-// kk inner ascending over four m16n16k32 steps), so every int32
-// accumulator -- and therefore every bf16 output byte -- is identical to
-// the bootstrap kernel and to the scalar reference.
-//
-// Round 2 (mandated operand-reuse round: cooperative A+B LDS staging vs direct
-// loads, quantified A/B reuse per macro-tile, vectorized coalesced global loads
-// with a bank-safe LDS layout):
-//
-// Operand-reuse accounting for the accepted <128, 64, 128> macro-tile
-// (M=4096, N=2624, K=6144):
-// * per block the unique operand bytes are A = 128 x 6144 = 786,432 and
-// B = 6144 x 64 = 393,216; the 41 N-blocks re-read every A row-slab
-// (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): 1.55 GB/replay of staged
-// global traffic;
-// * inside one block each staged 16-B vector is re-read from LDS by the two
-// wave columns (A) / the two wave rows (B), i.e. exactly 2x wave-level
-// reuse per staged byte, so the LDS read traffic is 2x the staged bytes
-// (4.72 GB/replay, ~23 B/cycle/CU at the measured 1.271 GHz -- far below
-// the LDS bandwidth), and the counter-derived HBM rate is only
-// 108.4 GB/s (147.8 MB/replay at a 90.7% L2 hit rate). The kernel is
-// therefore NOT HBM- or LDS-bandwidth bound, so a direct per-wave global
-// fragment path is rejected up front: it would issue 4x more vmem
-// instructions (4 waves x 48 stages of 8-B fragment loads) and every 8-B
-// fragment load would be a row-strided multi-line access instead of the
-// coalesced 16-B staging vectors of the cooperative path. Cooperative
-// A+B LDS staging is RETAINED -- the round improves it instead.
-//
-// Remaining exposed cost of the retained path: the VMEM round trip in front
-// of every stage consume. The accepted source (PMC digest daf403b3...,
-// profiles/.../iteration2/pmc.json: profiled 1368.793 us, grid 1312 blocks,
-// 2 blocks/CU, 26,112 B LDS, lds_instructions 7,892,992 with
-// lds_wait 4,535,391, lds_bank_conflicts 0, vmem_read 1,574,400,
-// grbm_gui_active == grbm_count) issues per 128-K stage 24 wave-level
-// global_load_dwordx4 + 24 ds_write2_b64 and then, per the exact-source ISA,
-// strictly serializes global_load_dwordx4 -> s_waitcnt vmcnt(0) ->
-// ds_write2_b64 -> s_barrier -> 24x ds_read_b64 -> 32x v_mmac -> s_barrier:
-// with 8 resident waves/CU the GPU is 100% GRBM GUI-active yet issues only
-// ~4.6% of its slots, i.e. the stage boundary -- not HBM, LDS bandwidth or
-// MMA throughput -- sets the pace.
-//
-// Focused change (ONE mechanism -- the K stage pipeline -- entirely inside
-// csrc/w8a8_gemm_hip.hip): the exact-shape guard now dispatches
-// w8a8_dumma_prefill_packedb_db_kernel<128, 64, 64>. It keeps the SAME
-// macro-tile (grid 41 x 32 = 1312 blocks, 256 threads = 4 wavefronts, 64 x 32
-// quadrant of eight m16n16k32 int32 accumulators per wave), the SAME packed
-// n-major B (16-B coalesced staging vectors), the SAME explicit load_frag8
-// fragment reads, the SAME bank-safe LDS stride form and the SAME coalesced
-// epilogue -- and changes only how the K stages are staged:
-// * the single-buffered 128-K stage (TWO __syncthreads per stage, full
-// VMEM latency exposed in front of every consume) becomes a
-// DOUBLE-buffered 64-K stage with ONE __syncthreads per stage: the
-// stage-(s+1) global loads (2 int4 A + 1 int4 B per thread) are issued
-// into staging VGPRs BEFORE the stage-s consume and flushed into the idle
-// LDS buffer after it, so the stage-(s+1) VMEM latency overlaps the
-// stage-s LDS/v_mmac work.
-//
-// 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. kLdsStride = 72 =
-// StageK + 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 the accepted kernel --
-// and 72 % 16 == 8 keeps the two-half ds_write2_b64 staging form. 12 extra
-// VGPR hold the prefetch payload (~72 -> ~84, still <= the 128/thread budget
-// of 2 blocks/CU; 0 spills expected).
-//
-// 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 the accepted 48 x 128-K stages, with the identical
-// element-to-slot fragment mapping, so every int32 accumulator -- and every
-// bf16 output byte -- is bit-identical to the accepted kernel and to the
-// scalar reference. The n-major pack, the scalar-fallback n-major decode, the
-// generic <64, 64, 128> identity arm and every shape guard are untouched.
-//
-// Round 3 (mandated pipeline round: compare single and double buffering across
-// K tiles; retain double buffering only on ISA/PMC evidence of reduced VMEM
-// stalls without harmful LDS or occupancy growth):
-//
-// * Mandatory comparison (exact trusted PMC/ISA of the two accepted arms):
-// arm median us LDS B/blk blocks/CU VGPR
-// single 128-K <128,64,128> 1363.078 26,112 2/8 waves 72
-// double 64-K <128,64,64> 1153.970 27,648 2/8 waves 96
-// A single 64-K arm is not built: with NO prefetch it can only add
-// barriers (2 per 64-K stage = 192 per block vs the 128-K arm's 96) while
-// exposing the same VMEM round trip, so it is strictly dominated and
-// rejected by construction.
-// RETAINED -- double buffering. Evidence: median 1363.078 -> 1153.970 us
-// (-15.3%) and p90 1364.825 -> 1164.522 us (-14.7%); the exact-source PMC
-// of the double-buffered arm (digest ee609c3f..., iteration3/pmc.json)
-// shows lds_wait_instructions 4,562,210 -> 3,590,796 (-21.3%) with
-// grbm_gui_active == grbm_count (100% GUI-active), LDS 27,648 B/block
-// (2 x 27,648 = 55,296 <= 65,536 so the trusted occupancy-probe split 2
-// is unchanged = 8 waves/CU), arch_vgpr 96 with 0 spills (<= the
-// 128/thread budget of 2 blocks/CU) and sgpr 32. The ISA confirms the
-// mechanism: the stage-(s+1) global_load_dwordx4s are issued BEFORE the
-// stage-s ds_read_b64 / v_mmac burst and the vmcnt(0) wait lands on the
-// ds_write2_b64 flush AFTER the MMACs, while the 128-K arm serializes
-// global_load_dwordx4 -> s_waitcnt vmcnt(0) -> ds_write2_b64 -> s_barrier
-// with zero compute overlap. So the prefetch converted exposed VMEM
-// latency into overlapped work without trading away occupancy.
-// LDS-growth verdict: the +1,536 B/block stays inside the 2-block budget
-// and lds_wait FELL (the LDS path is not harmed in the net); the new
-// lds_bank_conflicts 12,091,392 is exactly 6 x 4 x 4 x 96 x 1312 = one
-// conflict per 16-lane phase of every merged ds_read2_b64, i.e. the
-// bank-preserving alias of the +16-row second fragment offset
-// (16 x 72 = 1152 B = 9 x 128 B; for ANY 8-aligned uniform row stride S,
-// 16 rows = 4S dwords with 4S % 32 == 0). The unmerged ds_read_b64
-// control at the same 72/136-byte strides measured 0 conflicts, so this
-// is an alias of the merged read form, not a layout defect.
-// Every other K-tile pipeline is rejected inside the 65,536 B budget:
-// StageK=96 double buffer = 39,936 B/block, StageK=128 double buffer =
-// 52,224 B/block and 3 x 64-K = 41,472 B/block all exceed the 32,768
-// B/block that 2 resident blocks/CU permit (each would drop to 1
-// block/CU = 4 waves/CU, a 2x occupancy loss, explicitly excluded), and a
-// 32-K stage leaves only 128 B staging vectors per block (half the 256
-// threads idle) while carrying a single m16n16k32 step per stage (too
-// little consume work to amortize the barrier). StageK=64 x 2 buffers is
-// therefore the unique 2-resident-block K pipeline and is retained.
-//
-// * Focused change (ONE mechanism -- the B operand's global-memory layout --
-// entirely inside csrc/w8a8_gemm_hip.hip). The pipeline round's own PMC
-// identifies the next exposed cost: the accepted source reads
-// 170,221,184 B/replay at only an 86.938% L2 hit rate (3,250,848 misses)
-// and 166.147 GB/s counter-derived HBM, while the unique operands are only
-// 25.2 MB (A) + 16.1 MB (B). The B side causes it: the n-major pack
-// packed[n*K + kk] lays a [64, 64] B tile out as 64 separate 64-byte
-// column chunks at a 6,144-byte stride, so every one of the 8 warp-level
-// 16-byte staging loads touches 32 half-used 128-byte lines.
-// launch_pack_w8a8_weight now packs the exact (k, n) == (6144, 2624)
-// weight once, outside the timed region and out of Graph capture, into 41
-// B 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]
-// (byte-exact bijection 41 x 384 x 64 x 16 = k*n = 16,121,856; same byte
-// count, same buffer and same captured addresses; every other (K, N)
-// keeps the identity copy). Each [64, 64] B tile of a 64-K stage is then
-// FOUR contiguous 1-KiB slabs (one per 16-k-run) and every staging vector
-// is a 16-byte-aligned contiguous k-run of one column, so a warp's 32
-// x 16-byte B staging loads touch 4 FULLY-USED 128-byte lines instead of
-// 32 half-used ones (8x fewer line requests per warp-load; the tile's line
-// footprint is 32 fully-used lines instead of 64 half-used ones).
-// The ONLY kernel change is the two B global-load sites (prologue +
-// steady-state prefetch) of the dispatched
-// w8a8_dumma_prefill_packedb_db_kernel<128, 64, 64>: byte
-// (n0>>6)*(k*64) + ((k0>>4)+k16)*1024 + n_row*16 replaces
-// (n0+n_row)*k + k0 + k16*16. The LDS staging layout b_tile[2][64][72]
-// (two-half ds_write2_b64), the col_major load_frag8 consumes, the
-// k0-outer/kk-inner ascending int32 accumulation order, the coalesced
-// epilogue, the grid (41 x 32 = 1312 blocks), the block size (256 = 4
-// wavefronts), the retained double-buffered 64-K pipeline (one
-// __syncthreads per stage), LDS/block 27,648 B (2 blocks/CU = 8 waves/CU)
-// and the VGPR budget are ALL unchanged, so this is a pure global-layout
-// A/B.
-// * Correctness: the same weight elements reach the same fragment slots in
-// the same ascending k-chunk sequence, so every int32 accumulator -- and
-// every bf16 output byte -- is bit-identical to the accepted kernel and to
-// the scalar reference. The scalar fallback for (k, n) == (6144, 2624)
-// decodes the panel layout in ascending 16-byte k-runs (the paired M=2 /
-// M=16 API shapes with the same (N, K) stay byte-exact); the generic
-// identity arm, the round-2 n-major pack kernel (kept compiled as the
-// layout control) and every shape guard are untouched. Packing stays
-// outside the timed region and out of Graph capture.
-// * Predicted deltas: vmem_read_instructions unchanged (1,574,400), LDS
-// counters unchanged (lds_instructions 4,901,632, lds_wait 3,590,796,
-// conflicts at the merged-read alias floor 12,091,392), L2 misses strictly
-// below 3,250,848 with the hit rate up, operator read bytes strictly below
-// 170,221,184 (B line fetches halve), median < 1153.970 us.
-// * Superseded by Round 5: the B-panel pack and the b_tile[2][64][72] LDS
-// form below are replaced by the fragment-interleaved swizzle.
-//
-// Round 5 (mandated packing round: one weight packing/swizzle that makes each
-// DUMMA B tile vector-loadable and LDS-bank-safe):
-//
-// * Baseline: the accepted double-buffered <128, 64, 64> kernel (official
-// median 937.431 us / p90 939.452 us; exact-source PMC
-// profiles/.../iteration4/pmc.json: profiled 889.453 us, grid 41x32 =
-// 1312 blocks, 2 blocks/CU = 8 waves/CU, 27,648 B LDS, arch_vgpr 96 with
-// 0 spills, lds_instructions 4,901,632 = 9.7 per wave-stage (6 merged
-// ds_read2_b64 + 3 ds_write2_b64), lds_wait 3,828,676,
-// lds_bank_conflicts 12,091,392 = 6 x 4 x 4 x 96 x 1312, vmem_read
-// 1,574,400, L2 hit 90.39%). The surviving exposed LDS cost is the
-// b_tile row-major layout: each col_major B fragment is one ds_read_b64
-// at a 72-byte row stride, and the compiler merges the fragment pairs
-// into ds_read2_b64 whose +16-row second offset (16 x 72 = 1152 B =
-// 9 x 128 B) aliases the same bank pairs -- one conflict per 16-lane
-// phase, exactly the measured 12,091,392 conflicts.
-// * Mechanism (ONE packing/swizzle, entirely in this file): replace the
-// Round-3 B-panel pack with a fragment-interleaved swizzle. For each
-// 32-k x 16-n DUMMA B fragment block the 4 k8 groups x 16 n-rows x 8 B
-// are stored n-interleaved as [k8][n][8] (512 B per fragment block),
-// 32-k major inside the 64-K stage:
-// packed[((kk>>5)*n16g + (col>>4))*512 +
-// (((kk>>3)&3)*16 + (col&15))*8 + (kk&7)] = raw[kk*n + col],
-// with n16g = n>>4 (164 for the exact shape). Byte-exact bijection
-// 192 x 164 x 512 = k*n = 16,121,856: same byte count, same buffer and
-// same captured addresses (graph-stable), and packing still runs once,
-// outside the timed region and out of Graph capture.
-// * Kernel-side effect: every 16-byte B staging vector is now 2 n-rows x
-// 8 B of one k8 of one fragment block, so the staging store is ONE
-// 16-byte-aligned int4 (ds_write_b128) flat copy per thread (256 chunks
-// = one per thread per 64-K stage; a warp reads 8 fully-used 128-byte
-// lines) and every col_major B fragment is load_frag8_swz: lane
-// (n = lane&15, k8 = lane>>4) reads p + k8*128 + n*8, i.e. one
-// ds_read_b64 whose 16-lane phase spans exactly one 128-byte k8 block --
-// all 32 banks exactly once per phase, conflict-free by construction.
-// The A tile keeps its bank-safe 72-byte stride, and the double-buffered
-// 64-K pipeline, the k0-outer/kk-inner ascending int32 order, the
-// coalesced epilogue, the grid (41x32 = 1312 blocks), the block
-// (256 threads = 4 wavefronts) and the occupancy (2 blocks/CU = 8
-// waves/CU) are all unchanged, so this is a pure packing/swizzle A/B.
-// * LDS: b_tile shrinks from 2 x [64, 72] = 9,216 B to 2 x 4,096 B; the
-// block total is 18,432 (A) + 8,192 (B) = 26,624 B, so 2 x 26,624 =
-// 53,248 <= 65,536 B keeps 2 resident blocks/CU = 8 waves/CU.
-// * Correctness: the same weight elements reach the same fragment slots in
-// the same k0-outer/kk-inner ascending k-chunk sequence, so every int32
-// accumulator -- and every bf16 output byte -- is bit-identical to the
-// accepted kernel and to the scalar reference. The scalar fallback for
-// (k, n) == (6144, 2624) decodes the swizzle element-wise in ascending
-// kk (the paired M=2 / M=16 API shapes stay byte-exact); the generic
-// identity arm, the Round-2 n-major pack control and every shape guard
-// are untouched.
-// * Predicted deltas: vmem_read_instructions unchanged (~1,574,400),
-// LDS/block 27,648 -> 26,624 B, lds_bank_conflicts strictly below
-// 12,091,392, lds_wait below 3,828,676, median < 937.431 us.
-// Falsifiable: retained only if median_us < 937.4308013916016,
-// mismatch_count == 0, max_abs_error == 0.0,
-// graph_capture_passed == true, the exact code object shows 0 spills /
-// LDS == 26,624 B / VGPR <= 128 (2 blocks/CU), and the PMC shows
-// lds_bank_conflicts and lds_wait strictly below the baseline with
-// vmem_read not growing; any regression restores the iteration-3 source.
-//
-// Round 6 (mandated epilogue round: fuse per-row/per-column scales, bf16
-// conversion and the final coalesced store into the compute kernel; remove any
-// unnecessary workspace/combine pass):
-//
-// * Fusion audit (nothing to remove, so the fused epilogue is made cheaper).
-// The accepted source ALREADY fuses everything: store_prefill_fragment_
-// coalesced applies x_scale[row](per-row) * weight_scale[col](per-column),
-// converts to bf16 and writes one coalesced 8-byte store per lane, and the
-// exact-shape operator replay is a SINGLE dispatch -- the trusted PMC
-// (profiles/glm_tp8_fused_qkv_a_proj_m4096/iteration6/pmc.json, digest
-// 7e5731dc..., compile key 120c807e...) reports operator_aggregate
-// kernel_count == 1, scratch_bytes == 0, vmem_write_instructions 41,984 =
-// 8 x 4 x 1312 and write_bytes_per_operator_replay 21,500,384 = m*n*2, i.e.
-// there is no workspace, no split-K combine pass and no second Graph node
-// to remove; workspace_bytes is unused and workspace_split_k_capacity is
-// zero for this shape. The remaining epilogue cost is therefore the
-// conversion itself.
-// * Measured cost of the fused epilogue (exact gfx928 code object of the
-// current best, iteration6/current-best-isa/isa.txt). The dispatched
-// kernel is ~1,920 instructions; the steady-state K-stage loop body is 107
-// instructions (one `s_branch 65428` back-edge at 0x6BB8 -> 0x6A0C) and
-// the prologue is ~259, so the fused epilogue is ~1,490 static
-// instructions -- roughly 13% of the dynamic instruction stream over
-// 96 stages. The dominant term is that this DTK lowers EACH of the 32
-// per-lane __float2bfloat16 calls to a ~13-instruction EXEC-masked
-// Inf/NaN sequence:
-// v_and_b32 (b & 0x7f800000) / v_cmp_ne_u32 / s_and_saveexec_b64 /
-// s_xor_b64 / v_bfe_u32 (bit 16) / v_add3_u32 (b + lsb + 0x7fff) /
-// s_andn2_saveexec_b64 / v_mov_b32 0 / v_or_b32 (b | 0x10000) /
-// v_cmp_eq_u32_sdwa / v_cndmask_b64 / s_or_b64 (restore EXEC)
-// so ~420 instructions per lane plus 4 x (3 shifts + 3 ors) of packing,
-// with 32 dependent SALU EXEC save/restore pairs that cannot be software-
-// pipelined and serialise the wavefront (the epilogue is the only place
-// VALU issues drop while the exact-shape kernel is 100% GRBM GUI-active).
-// * Focused change (ONE mechanism -- the bf16 conversion of the fused
-// epilogue -- entirely inside csrc/w8a8_gemm_hip.hip). Replace the library
-// conversion + shift/or pack in store_prefill_fragment_coalesced with the
-// branch-free round-to-nearest-even integer identity
-// bf16_rne_high16(v) = (b + 0x7fff + ((b >> 16) & 1)) >> 16,
-// b = __float_as_uint(v),
-// which the disassembly proves is exactly what the normal (finite) path of
-// the library conversion computes (v_bfe bit 16, v_add3 b + lsb + 0x7fff).
-// 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 (RNE), max-float still
-// rounds to +/-Inf and 0.0/-0.0 and subnormal encodings are preserved, so
-// every 16-bit payload is identical. Every value converted here is finite
-// by construction (|int32 dot| < 2^31 times finite dequant scales), so the
-// removed Inf/NaN special case is unreachable. Four values now pack into
-// one 8-byte word with three VALU ops (v_bfe/v_and + v_add3 + v_lshrrev)
-// instead of ~13 EXEC-bracketed instructions each, and the 32 EXEC
-// save/restore chains disappear. The epilogue structure, the per-row
-// x_scale gather, the per-column float4 weight_scale gather, the register
-// 4x4 shuffle transpose, the multiply order float(dot) * x_scale *
-// weight_scale, the 8-byte-aligned coalesced store address, the grid
-// (41 x 32 = 1312), the block (256 = 4 wavefronts), the double-buffered
-// 64-K pipeline, the swizzled packed B, LDS/block 26,624 B (2 blocks/CU =
-// 8 waves/CU) and the 65,536 B workspace budget are ALL unchanged, so this
-// is a pure conversion A/B. The scalar fallback keeps __float2bfloat16 (it
-// is not on the timed path); the generic identity arm and every shape
-// guard are untouched.
-// * Predicted deltas: grid_blocks 1312, LDS/block 26,624 B, 2 blocks/CU =
-// 8 waves/CU, vmem_read 1,574,400 and vmem_write 41,984 unchanged;
-// valu_instructions strictly below 23,920,384 and the epilogue's static
-// EXEC/SALU count down by >1,000; median < 918.191 us with no p90
-// regression. Falsifiable: retained only if median_us <= 909.0 us
-// (>= 1% round acceptance) and p90_us <= 919.530 us, mismatch_count == 0,
-// max_abs_error == 0.0, graph_capture_passed == true, the exact code
-// object shows 0 spills / LDS == 26,624 B / VGPR <= 128 (2 blocks/CU) and
-// no __float2bfloat16 EXEC sequence remains in the epilogue; otherwise the
-// iteration-5 source is restored.
-//
-// Round 9 (late ISA-diagnosis round; the control-plane plateau gate is CLOSED
-// for this round -- phase hip_only, plateau false, raw inline asm and ISA
-// Skills forbidden -- so this stays a HIP-only round):
-//
-// * Bottleneck evidence (exact trusted gfx928 ISA of the current best,
-// profiles/glm_tp8_fused_qkv_a_proj_m4096/iteration6/current-best-isa/
-// isa.txt, digest 7e5731dc...; the round-6 source differs from that
-// profiled official best only in the bf16 conversion, so the steady-state
-// loop analysed here is the dispatched loop). The steady-state stage body
-// is 107 instructions (0x6A0C..0x6BB8, one back-edge `s_branch 65428` at
-// 0x6BB8 -> 0x6A0C) and roughly 14 of them are pure A-row tail
-// predication:
-// 0x6A0C..0x6A20 six unconditional `v_mov_b32` zero-fills of the two
-// int4 staging vectors (v33/v34 and v37..v40),
-// 0x6A24..0x6A34 `s_and_saveexec_b64` / `s_cbranch_execz` /
-// `global_load_dwordx4` / `s_or_b64` for A vector 0,
-// 0x6A38..0x6A50 two more `v_mov_b32` zero-fills (v35/v36) plus the
-// same EXEC save/restore + uniform branch for A
-// vector 1.
-// The guard source `a_pre[v] = (g_row < m) ? load : int4{0,0,0,0}` is
-// never false for the dispatched M=4096 (M / BM = 4096 / 128 = 32 exact),
-// yet its zero-fills execute unconditionally in every one of the 96
-// stages and its two EXEC read-modify-writes serialize the wavefront in
-// front of the very global loads that prefetch the next tile. That is
-// ~14/107 = 13% of the hot loop's instruction stream, on top of the
-// epilogue's per-fragment `row >= m` EXEC bracket (s_and_saveexec /
-// s_xor_b64 / s_andn2_saveexec_b64 / s_cbranch_execz / s_or_b64 around
-// each of the eight unrolled fragment stores), for a kernel the PMC shows
-// 100% GRBM GUI-active yet issuing only ~0.42 instructions per cycle per
-// CU (lds_wait_instructions 5,302,705 against lds_instructions
-// 4,901,632).
-// * Focused change (ONE mechanism -- exact-shape full-tile specialization,
-// entirely inside csrc/w8a8_gemm_hip.hip): the dispatched kernel gains a
-// compile-time `bool kFullTile` template parameter. The exact-shape guard
-// keeps `(m >= 128, n == 2624, k == 6144)` and the same grid (N/64 x
-// ceil(M/128)) and block (256 = 4 wavefronts), and selects
-// `kFullTile == true` exactly when the caller's m is a multiple of
-// BM = 128 -- i.e. when every block tile lies fully inside the output (the
-// assigned M=4096 case). In that instantiation the A staging loads
-// (prologue AND steady state) drop the `(g_row < m)` select / zero-fill
-// and the epilogue drops the per-fragment `row >= m` early return, so the
-// loop body keeps only the address math and the three coalesced
-// `global_load_dwordx4`. Every other m keeps the guarded instantiation
-// byte-for-byte (zero-fill, epilogue row mask, tail-block coverage), and
-// the paired M=2/M=16 API shapes with the same (N, K) still take the
-// scalar fallback because the `m >= 128` guard is untouched.
-// * Why this is the remaining lever and not a new pipeline: rounds 7
-// (deeper register prefetch / stage pairs) and 8 (smaller macro-tile, more
-// resident waves) both regressed, so the double-buffered 64-K stage (ONE
-// __syncthreads per stage, LDS 26,624 B/block, 2 blocks/CU = 8 waves/CU),
-// the 128 x 64 macro-tile, the fragment-interleaved packed B swizzle, the
-// load_frag8 / load_frag8_swz consumes, the k0-outer / kk-inner ascending
-// m16n16k32 int32 accumulation order and the fused branch-free bf16
-// epilogue are ALL retained unchanged. The round removes instructions that
-// are provably dead for this shape without touching any resource, any
-// memory access order or any MMA input.
-// * Correctness: for m % 128 == 0 no block tile crosses m, so the removed
-// guard predicate is never false and the removed zero-fill never reaches an
-// accumulator; the surviving assignments, the ascending k-chunk sequence
-// and the exact `float(dot) * x_scale * weight_scale` multiply order are
-// bit-identical to the guarded path (mismatch 0 / max_abs_error 0.0
-// expected). The guarded instantiation preserves the exact current
-// semantics for every m % 128 != 0.
-// * Predicted deltas: steady-state loop 107 -> ~93 instructions (the eight
-// zero-fills and the two EXEC save/restore + uniform-branch pairs removed),
-// valu_instructions strictly below 23,920,384, grid_blocks 1312,
-// block_threads 256, LDS/block 26,624 B, arch_vgpr <= 128 with 0 spills,
-// vmem_read_instructions 1,574,400, vmem_write_instructions 41,984 and the
-// epilogue store count all unchanged. Falsifiable: retained only if
-// median_us < 912.7147674560547 (the current shadow base) with no p90
-// regression, mismatch_count == 0, max_abs_error == 0.0,
-// graph_capture_passed == true, the exact code object shows 0 spills /
-// LDS == 26,624 B / VGPR <= 128 (2 blocks/CU) and the kFullTile symbol's
-// loop body no longer contains the A-guard
-// `s_and_saveexec`/`s_cbranch_execz`/zero-fill sequence; any median/p90
-// regression, mismatch, spill or occupancy loss restores the round-6
-// 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).
@@ -470,6 +91,12 @@ 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;
@@ -477,53 +104,38 @@ constexpr int kExactK = 6144;
using namespace du::dumma;
-// Coalesced accumulator epilogue for gfx928 INT8 m16n16k32: 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 -> offset a multiple of 8) are
-// aligned.
-//
-// Round 6: the fused bf16 conversion itself is made branch-free. The exact
-// gfx928 code object of the previous best lowers each of the 32 per-lane
-// __float2bfloat16 calls to a ~13-instruction sequence that brackets an
-// Inf/NaN special case with EXEC save/restore:
-// v_and_b32 (b & 0x7f800000) / v_cmp_ne_u32 / s_and_saveexec_b64 / s_xor_b64
-// / v_bfe_u32 (bit 16) / v_add3_u32 (b + lsb + 0x7fff) / s_andn2_saveexec_b64
-// / v_mov_b32 0 / v_or_b32 (b | 0x10000) / v_cmp_eq_u32_sdwa / v_cndmask_b64
-// / s_or_b64 (restore EXEC)
-// The EXEC manipulation cannot be software-pipelined and serialises the
-// wavefront. For every FINITE fp32 the library conversion is exactly
-// round-to-nearest-even at bit 16, which is the branch-free integer identity
-// bf16 = (b + 0x7fff + ((b >> 16) & 1)) >> 16, b = __float_as_uint(v)
-// (adding 0x7fff + the retained bit-16 LSB never carries out of the sign, and
-// the +((b>>16)&1) term breaks exact ties toward an even LSB, so max-float
-// still rounds to +/-Inf and -0.0/0.0 are preserved). Every value converted
-// here is finite by construction: |int32 dot| < 2^31 and the dequant scales
-// are finite, so no Inf/NaN can reach the conversion. This yields the exact
-// same 16-bit payloads as __float2bfloat16 while dropping the EXEC sequence
-// and the separate shift/or pack (four bf16 pairs become three VALU ops per
-// 8-byte word instead of 4 x 13 + 3).
+// 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;
}
-// Round 9 adds the compile-time `kGuarded` flag: kGuarded == false is the
-// exact-shape full-tile specialization (m % BM == 0), for which every row of
-// every block tile is strictly inside the output, so the `row >= m` early
-// return -- which the exact code object lowers to an EXEC save/restore pair
-// plus a uniform branch around each of the eight unrolled fragment stores --
-// is removed at compile time. The default kGuarded == true keeps the exact
-// current semantics for every other m (including the generic arm and the
-// single-buffered siblings, which never pass the flag).
+// 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,
@@ -567,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);
@@ -588,15 +200,14 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced(
packed;
}
-// 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
-// -> 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.
+// 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,
@@ -609,17 +220,17 @@ __device__ __forceinline__ void load_frag8(
reinterpret_cast(f.x)[0] = v;
}
-// Round-5 fragment-interleaved B swizzle helpers.
+// Fragment-interleaved B swizzle helpers.
//
// 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
+// (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 Round-3 layout delivered.
+// 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,
@@ -630,12 +241,11 @@ __device__ __forceinline__ int64_t packed_b_swz_offset(
static_cast(k8) * 128 + (n_pair << 4);
}
-// Round-5 col_major B fragment loader on the swizzled tile: lane (n = lane&15,
+// 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.
+// load_frag8, so the v_mmac operands and the int32 accumulation stay bit-exact.
template
__device__ __forceinline__ void load_frag8_swz(
Frag& f,
@@ -648,27 +258,50 @@ __device__ __forceinline__ void load_frag8_swz(
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 =
+// 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)).
//
// 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. 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 above.
+// 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(
@@ -787,282 +420,146 @@ w8a8_dumma_prefill_tiled_kernel(
}
}
-// Round-1 packed-weight architecture kernel, kept as the single-buffered
-// 128-K sibling of the dispatched <128, 64, 64> double-buffered kernel below
-// (Round 2 moved the exact-shape dispatch to that sibling; Round 3 keeps this
-// kernel compiled but undispatched as the n-major layout / single-buffer
-// control, and only the double-buffered kernel reads the packed B; Round 5
-// replaces that packed B layout with the fragment-interleaved swizzle, so this
-// n-major control no longer matches the pack either). It is defined for
-// the n-major packed_weight[n * K + kk] = raw[kk * N + n] layout produced by
-// launch_pack_w8a8_weight. SAME geometry as the bootstrap
-// w8a8_dumma_prefill_tiled_kernel<128, 64, 128>:
-// grid (N/64) x (M/128) = 41 x 32 = 1312 blocks, 256 threads = 4 wavefronts,
-// one 64 x 32 quadrant of eight m16n16k32 int32 accumulators per wave,
-// single-buffered 128-K LDS stage with TWO __syncthreads per stage, same
-// k0-outer/kk-inner ascending int32 accumulation order, same coalesced
-// fragment epilogue. The two differences, both inside the operand path:
-// * B is staged n-major (b_tile[n_row][kk] from the packed [N, K] weight)
-// so col_major B fragments are 8 contiguous K bytes per lane and staging
-// is one coalesced 16-byte vector per (n_row, k16) pair;
-// * A and B LDS row strides are 136 B (StageK + 8): 8-byte aligned with an
-// odd 8-byte count (136/8 = 17) so the 16 rows of each 16-lane ds_read_b64
-// phase land on 16 distinct bank pairs (the conflict-free 4-phase floor),
-// and 136 % 16 == 8 keeps the two-half int64 staging form.
-// 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 bootstrap 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]
- 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 for StageK = 128.
- // 8-aligned with an odd 8-byte count -> every lane's 8-byte fragment is a
- // ds_read_b64 at the 4-phase conflict-free floor (see the kernel comment
- // above); 8 mod 16 keeps the two-half staging store form.
- 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 (the dispatch guarantees
- // k % StageK == 0 and n % BN == 0). 8 vectors per StageK row for 128.
- 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));
-
- for (int k0 = 0; k0 < k; k0 += StageK) {
- // Stage A[BM, StageK] row-major with the bank-safe row stride: 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: the row stride is 8 mod 16, so odd rows are not
- // 16-byte aligned and a single int4 LDS 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,
- static_cast(kLdsStride),
- lane);
- }
-#pragma unroll
- for (int i = 0; i < kWaveN16; ++i) {
- load_frag8(
- b_frag[i],
- b_tile + (local_col + i * kTileN) * kLdsStride + kk,
- static_cast(kLdsStride),
- 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]);
- }
- }
- }
- // Single buffer: every wavefront must finish reading LDS before the next
- // stage's stores overwrite it.
- __syncthreads();
- }
-
- 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);
- }
- }
-}
-
-// Round 9 exact-shape A-staging load. For the dispatched M=4096 every 128-row
-// block tile lies fully inside the output, so the `(g_row < m)` row guard is
-// dead: the exact code object lowers it to eight unconditional zero-fill
-// `v_mov_b32` plus two EXEC save/restore (`s_and_saveexec_b64` /
-// `s_cbranch_execz` / `s_or_b64`) uniform-branch pairs per K stage, i.e. ~14
-// of the 107 steady-state loop instructions and an EXEC serialization in
-// front of the very `global_load_dwordx4`s 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 to the current kernel for every m % 128 != 0.
+// 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__ x_q,
- int64_t offset,
+ const int8_t* __restrict__ p,
int m,
int g_row) {
if (kFullTile) {
- return *reinterpret_cast(x_q + offset);
+ return *reinterpret_cast(p);
}
- return (g_row < m) ? *reinterpret_cast(x_q + offset)
- : int4{0, 0, 0, 0};
+ return (g_row < m) ? *reinterpret_cast(p) : int4{0, 0, 0, 0};
+}
+
+// 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);
}
-// Round-2 double-buffered packed-B prefill kernel, dispatched ONLY for
-// (m >= 128, n == 2624, k == 6144). SAME macro-tile and SAME operand path as
-// the accepted w8a8_dumma_prefill_packedb_kernel<128, 64, 128>: grid
-// (N/64) x (M/128) = 41 x 32 = 1312 blocks, 256 threads = 4 wavefronts, one
-// 64 x 32 quadrant of eight m16n16k32 int32 accumulators per wave, packed
-// n-major B, explicit load_frag8 reads, coalesced epilogue. The ONLY change is
-// the K pipeline (see the Round-2 section of the file header):
-// * 64-K DOUBLE-buffered stages (2 x 13,824 = 27,648 B LDS/block, still
-// 2 blocks/CU = 8 waves/CU): the stage-(s+1) global loads are issued into
-// 3 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 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, 96, ... as the accepted 48 x 128-K stages, 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
+// 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.
//
-// Round 5 additionally reads the packed B operand from the fragment-
-// interleaved swizzle (see w8a8_pack_bswz_i8_kernel and the Round-5 header
-// section): the LDS b_tile is now the SAME [k32][n16][k8][n][8] layout as the
-// pack (2 x 256 16-byte chunks 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.
+// 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.
//
-// Round 9 adds the compile-time `kFullTile` flag: it 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 (see
-// load_a_stage_vec) and the epilogue drops the per-fragment row mask (see
-// store_prefill_fragment_coalesced); false keeps the guarded form
-// byte-for-byte. Geometry, resources, operand path and accumulation order are
-// identical in both instantiations.
-template
-__global__ __launch_bounds__(kThreadsPerBlock) void
+// 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>>4][k>>5][512 B] swizzled B
- // fragment blocks (round-5 layout)
+ // fragment blocks
const float* __restrict__ x_scale, // [M]
const float* __restrict__ weight_scale, // [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");
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");
- // Round 5: the packed B weight is the fragment-interleaved swizzle (see
- // w8a8_pack_bswz_i8_kernel), defined for BN == 64 and StageK == 64 (one
- // 32-k x 16-n fragment block per 512 B; 2 k32 blocks x 4 n16 groups x 4 k8
- // groups x 8 n-pairs = 256 16-byte chunks per stage, one per thread). Only
- // <128, 64, 64> is ever dispatched by the exact-shape guard, so the chunk
- // index math below is exact.
- static_assert(BN == 64 && StageK == 64,
- "swizzled B pack requires BN == 64 and StageK == 64");
+ // 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];
- // Round-5 swizzled B tile: flat [k32][n16][k8][n_pair][16 B] = 4,096 B per
- // 64-K stage, byte-identical to the packed global layout of that stage.
+ // 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
@@ -1079,56 +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");
- // Round 5: 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 == kThreadsPerBlock,
+ // 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];
+ // 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] = load_a_stage_vec(
- x_q, static_cast(g_row) * k + k0 + k16 * 16, m, g_row);
+ 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;
- // Round 5: the 16-byte chunk (k32 = vec>>7, n16 = (vec>>5)&3,
+ // 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;
- // consecutive lanes read consecutive 16-byte chunks, so a warp's
+ // (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(
- packed_w +
- packed_b_swz_offset((k0 >> 5) + (vec >> 7),
- (n0 >> 4) + ((vec >> 5) & 3), n16g,
- (vec >> 3) & 3, vec & 7));
+ 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]);
@@ -1139,112 +667,157 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
#pragma unroll
for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- // Round 5: 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.
+ 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 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.
- for (int s = 0; s < kStages - 1; ++s) {
- // Prefetch stage s+1: issue the global loads (held in VGPRs 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] = load_a_stage_vec(
- x_q, static_cast(g_row) * k + k0 + k16 * 16, m, g_row);
- }
+ 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;
- // Round 5: the 16-byte chunk (k32 = vec>>7, n16 = (vec>>5)&3,
- // 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;
- // 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(
- packed_w +
- packed_b_swz_offset((k0 >> 5) + (vec >> 7),
- (n0 >> 4) + ((vec >> 5) & 3), n16g,
- (vec >> 3) & 3, vec & 7));
- }
+ 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);
+ // ---- 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,
- static_cast(kLdsStride),
- lane);
- }
+ 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) {
- // Round 5: 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[s & 1] +
- ((kk >> 5) * (BN / 16) + wave_col * kWaveN16 + i) * 512,
- lane);
- }
+ 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 < kAVecPerThread; ++v) {
+ a_src[v] += StageK;
+ }
#pragma unroll
- for (int v = 0; v < kBVecPerThread; ++v) {
- const int vec = tid + v * kThreadsPerBlock;
- // Round 5: flat 16-byte swizzled-tile commit (ds_write_b128).
- *reinterpret_cast(b_tile[buf] + vec * 16) = b_pre[v];
+ 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_row = wave_row * kRowsPerWave;
#pragma unroll
for (int kk = 0; kk < StageK; kk += kTileK) {
#pragma unroll
@@ -1257,9 +830,6 @@ w8a8_dumma_prefill_packedb_db_kernel(
}
#pragma unroll
for (int i = 0; i < kWaveN16; ++i) {
- // Round 5: 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[s & 1] +
@@ -1276,14 +846,14 @@ 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) {
- // Round 9: kGuarded == !kFullTile -- the full-tile instantiation skips
- // the dead per-fragment `row >= m` EXEC bracket (see the helper above).
+ // 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);
@@ -1292,26 +862,21 @@ 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/packed paths). One output element per
+// 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.
//
-// Round 1: for the exact (k, n) == (6144, 2624) the packed_weight buffer
-// holds a packed layout (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). The int32 accumulation stays
-// kk-ascending 0..k-1 in both arms; integer addition is exact and commutative,
-// so the int32 dot -- and the paired M=2 / M=16 API shapes -- remain
-// byte-exact against the identity-layout reference. Round 5 replaces the
-// exact-shape pack with the fragment-interleaved swizzle (see
-// w8a8_pack_bswz_i8_kernel), so the exact-shape decode below walks
+// 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.
+// 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, // packed weight (swizzled/identity)
@@ -1333,8 +898,8 @@ __global__ __launch_bounds__(256) void w8a8_scalar_gemm_kernel(
const int8_t* a_row = a + static_cast(row) * k;
int32_t acc = 0;
if (packed_swz) {
- // Round-5 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.
+ // 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 =
@@ -1372,46 +937,19 @@ __global__ __launch_bounds__(256) void w8a8_pack_identity_i8_kernel(
}
}
-// Round-1 exact-shape n-major pack: 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] so every col_major B fragment is 8
-// contiguous K bytes in the packed buffer and every staging vector is one
-// coalesced 16-byte read. Same byte count and buffer as the identity pack, so
-// the captured addresses are unchanged. One thread per output byte; runs once
-// during weight preparation.
-__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];
- }
-}
-
-// Round-5 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 512-byte 32-k x 16-n DUMMA B fragment blocks laid out
-// [k32][n16][k8][n][8]:
+// 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 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/n-major/panel packs, 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). The round-2 n-major pack kernel above stays compiled as the
-// layout control.
+// 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,
@@ -1437,8 +975,8 @@ __global__ __launch_bounds__(256) void w8a8_pack_bswz_i8_kernel(
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).
+ // 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];
}
@@ -1462,10 +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. The workspace is unused for this shape
-// (single-buffered, no split-K): m*n*4 = 42,990,080 B exceeds the 16 MiB
-// workspace budget, so workspace_split_k_capacity is zero.
+// 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,
@@ -1486,41 +1024,23 @@ 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). The m >= 128 guard keeps the paired M=2/M=16
- // API shapes with the same (N, K) on the scalar fallback below. Round 2
- // keeps the accepted packed-weight architecture (the exact-fit 2-D macro-tile
- // <128, 64, ...> with the packed weight produced once, out of the
- // timed region, by launch_pack_w8a8_weight, the bank-safe LDS row strides and
- // the explicit 8-byte load_frag8 fragment loads) and changes ONLY the K-stage
- // pipeline: the dispatch goes to the 64-K double-buffered prefetch kernel
- // <128, 64, 64> instead of the single-buffered <128, 64, 128> one, so the
- // stage-(s+1) VMEM latency overlaps the stage-s LDS/v_mmac work. Round 3
- // retains that double-buffered 64-K pipeline unchanged (it is the unique
- // 2-resident-block K pipeline inside the 65,536 B LDS budget: 2 x 27,648 =
- // 55,296 <= 65,536) and changes only the packed B layout consumed by the two
- // B global-load sites -- the weight is now the B-panel layout, so a warp's B
- // staging loads touch 4 fully-used 128-byte lines instead of 32 half-used
- // ones. Round 5 keeps the kernel, geometry and occupancy identical and
- // changes only the packed B layout again: the B-panel pack is replaced by the
- // fragment-interleaved swizzle, so the B stage is a flat 16-byte copy and
- // every col_major B fragment is one conflict-free ds_read_b64
- // (load_frag8_swz). M = 32*128 and N = 41*64 are both exact, so no tail
- // predication is introduced. Round 9 keeps that kernel, geometry and
- // occupancy and only selects a compile-time tail-predication specialization:
- // when m is a multiple of BM = 128 every block tile is fully inside the
- // output, so the instantiation with kFullTile == true drops the dead A-row
- // zero-fill/select in both staging sites and the dead epilogue row mask
- // (see the Round-9 header section); every other m keeps the guarded
- // instantiation byte-for-byte.
+ // (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));
- if ((m & 127) == 0) {
+ 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<128, 64, 64, true>),
+ w8a8_dumma_prefill_packedb_db_kernel<
+ kPackedBM, kPackedBN, 64, kPackedThreads, true>),
grid,
block,
0,
@@ -1536,7 +1056,8 @@ extern "C" void launch_w8a8_gemm(
} else {
hipLaunchKernelGGL(
HIP_KERNEL_NAME(
- w8a8_dumma_prefill_packedb_db_kernel<128, 64, 64, false>),
+ w8a8_dumma_prefill_packedb_db_kernel<
+ kPackedBM, kPackedBN, 64, kPackedThreads, false>),
grid,
block,
0,
@@ -1555,8 +1076,8 @@ extern "C" void launch_w8a8_gemm(
// 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.
+ // 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),
@@ -1608,7 +1129,7 @@ extern "C" void launch_w8a8_gemm(
k);
}
-// Stable host launch symbol consumed by csrc/bindings.cpp. Round 5: the exact
+// 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)]
@@ -1628,7 +1149,7 @@ extern "C" void launch_pack_w8a8_weight(
constexpr int kPackThreads = 256;
const int64_t weight_count = static_cast(k) * n;
if (k == kExactK && n == kExactN) {
- // Round 5: one thread packs one 16-byte chunk (16 bytes = 2 n-rows x 8 B).
+ // 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) {
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/o_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/o_proj.hip
index 49710bb2..2a553a70 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/o_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/o_proj.hip
@@ -1,76 +1,147 @@
-// @@variant shape=glm_tp8_o_proj_m4096 commit=486f3ec14b6d88c045aa35a5c85d1d44816c8c2d added=2026-08-31
-// median_us=759.9 p90_us=760.8 speedup=72.08 baseline_us=5.478e+04
-// source=glm5-2-dsh-tp8-m4096-1-e6a280a2
-// @@variant shape=glm_tp8_o_proj_m4096
+// @@variant shape=glm_tp8_o_proj_m4096 commit=41d2ef9da654ecb15154a974fa11d84df5e4fc5d added=2026-09-14
+// median_us=651.5 p90_us=654.2 speedup=84.07 baseline_us=5.478e+04
+// source=test-9-11-1/iteration_002/glm_tp8_o_proj_m4096
// INT8 W8A8 GEMM HIP implementation for Hygon K500SM_AI / gfx928.
//
-// Worker: worker_2 (physical GPU 2), assigned shape:
+// Worker: worker_0 (physical GPU 0), assigned shape:
// glm_tp8_o_proj_m4096 : M=4096, N=6144, K=2048
//
-// Bootstrap strategy (iteration 1, correctness-first but a usable profiling
-// baseline - a large-Prefill scalar K loop would not be):
-// * Large-prefill path (exact (m, n, k) == (4096, 6144, 2048)):
-// 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 accumulator fragments;
-// the block cooperatively vector-loads A[128,64] and B[64,64] into one
-// single-buffered 64-K 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. Grid dim3(96, 32) = 3072 blocks
-// dwarfs the 120 CUs, so no split-K is needed.
-// The weight is read in the identity [K, N] int8 layout (bootstrap:
-// launch_pack_w8a8_weight is an identity device-to-device copy for
-// every (k, n)). A is staged m-major into a_tile[m][k] and its eight
-// fragment k-values per lane are contiguous, so A fragments use the
-// direct 8-byte LDS fill (load_fragment8, the exact pattern validated
-// on the worker-29 TP4 gate_up lineage iteration 13 and the TP8 hy3
-// o_proj lineage iteration 12). B is staged k-major into
-// b_tile[k][n] and its fragments use the library du_load_matrix_sync
-// row_major loader (8 per-byte LDS reads strided by ldm per lane -
-// correct, direct, and simple; the packed n-major layout + 8-byte B
-// fragment fill is a later-round optimization once the packed layout
-// is introduced).
-// * Generic scalar int8/int32 fallback for every unmatched (m, n, k),
-// including all small-M API cases (M=2, M=16) and M=3072 with the same
-// (K, N) = (2048, 6144); the fallback reads the identity [K, N] layout
-// for every (k, n).
-// * launch_pack_w8a8_weight: identity device-to-device copy of the raw
-// [K, N] weight and the [N] weight_scale for every (k, n) (bootstrap;
-// later rounds may change only the pack HIP implementation and the
-// matching GEMM interpretation). Packing never happens inside the
-// timed GEMM.
-// Iteration 5 (packing round): for the exact (k, n) == (2048, 6144)
-// launch_pack_w8a8_weight now produces the n-major [N, K] layout
-// packed[n*K + kk] = raw[kk*N + n] (same byte count and buffer ->
-// graph-stable), and the exact-shape dispatch uses the new
-// w8a8_dumma_256x64x64_packedb_kernel, which stages B n-major into
-// bank-safe 72-byte LDS rows and fills each col_major B fragment with
-// one contiguous 8-byte ds_read_b64 (zero per-byte reads, zero
-// reassembly VALU); the scalar fallback decodes the n-major layout
-// when (k, n) matches so the paired M=2/M=16 and M=3072 API shapes
-// stay byte-exact. Every other (k, n) keeps the identity copy.
-// Iteration 6 (epilogue round): the fused scale->bf16->coalesced
-// store epilogue of w8a8_dumma_256x64x64_packedb_kernel is
-// restructured to register-batch the per-row and per-column scales:
-// each lane loads its 8 x_scale rows and its 2 weight_scale float4
-// columns ONCE (issued on the last K stage, before the final
-// protective barrier) and the 16 per-fragment coalesced stores reuse
-// them; the previous per-call epilogue re-loaded the scales inside
-// each store (PMC vmem_read surplus ~153,600 wavefront loads above
-// the pure staging floor). Output bits are identical (same transpose,
-// same float(dot) * x_scale[row] * weight_scale[col] order, same bf16
-// rounding). No workspace/combine pass exists: the workspace
-// argument stays unused and the epilogue is fully fused in-kernel.
+// Round 1 (establish_arch) established ONE complete, correct DUMMA
+// architecture family and measured median 753.83 us:
+//
+// * family: packed_weight + staged A/B, 256x64 output macro-tile, native
+// INT8 DUMMA m16n16k32 with int32 accumulation, four 64-lane wavefronts
+// (256 threads);
+// * weight packing (outside the timed region, before Graph capture, via the
+// ABI-preserving launch_pack_w8a8_weight): for the exact (k, n) ==
+// (2048, 6144) the raw logical [K, N] int8 weight is transposed once into
+// the n-major [N, K] layout packed[col * K + kk] = raw[kk * N + col].
+// Same byte count and same caller-owned buffer as the identity layout, so
+// every captured address is graph-stable. Every other (k, n) keeps the
+// identity device-to-device copy.
+// * fragment fills: the m16n16k32 matrix_a slot for lane (r = lane & 15,
+// kq = lane >> 4) is eight consecutive k bytes at fixed row, and the
+// matrix_b (col_major) slot is eight consecutive k bytes at fixed n on the
+// n-major tile, so BOTH operands use one contiguous 8-byte ds_read_b64 per
+// fragment (load_fragment8) with no per-byte reads and no mask/OR
+// reassembly VALU.
+// * LDS bank safety: both padded strides are odd multiples of 8 bytes
+// (A = 88 = 22 dwords, B = 72 = 18 dwords). 22*r mod 32 and 18*r mod 32
+// over r = 0..15 are permutations of the 16 even dword starts, so each
+// 16-lane ds_read_b64 fragment phase lands on 16 distinct bank pairs.
+// Both strides are 8 mod 16, so the 16-byte cooperative staging stores are
+// emitted as two int64 halves (the validated ds_write2_b64 pattern).
+// * fused coalesced epilogue: the 4x4 shfl_xor register transpose keeps the
+// exact per-element multiply order float(dot) * x_scale[row] *
+// weight_scale[col] and bf16 rounding, then one 8-byte store per lane per
+// fragment (100% store-sector efficiency). Each lane's x_scale rows and
+// weight_scale float4 columns are loaded ONCE into registers on the last
+// K stage, issued before the final protective barrier, so the epilogue
+// reads no scale inside the store chains. The epilogue is fully fused:
+// no workspace pass, no separate combine kernel.
+//
+// Round 2 (mandatory decision: operand-reuse round - compare direct loads
+// with cooperative A+B LDS staging, quantify A/B reuse per macro-tile, use
+// vectorized coalesced global loads with a bank-safe LDS layout):
+//
+// FALSIFIABLE BOTTLENECK HYPOTHESIS. The round-1 kernel is bound by the
+// L2/global operand traffic its macro-tile reuse ratio forces, not by the
+// LDS fragment pipe, not by MMAC issue, and not by HBM. With a
+// (tileM, tileN) = (256, 64) macro-tile every A byte is re-read N/tileN =
+// 96x and every B byte M/tileM = 16x by the 1,536-block grid, i.e.
+// A 805.3 MB + B 201.3 MB = 1,006.6 MB of global reads per operator replay.
+// The exact-source PMC record (digest acb11003..., grid 1,536 x 256
+// threads, 151 VGPR, 27,136 B LDS, 957,746 grbm_gui_active cycles at
+// ~1.29 GHz) shows only 99.3 MB of HBM reads and 50.3 MB of HBM writes
+// (198.5 GB/s counter-derived HBM) against 1,044,480 vmem_read_instructions
+// of 16-B-per-lane coalesced staging, i.e. ~1.07 GB of L2 read traffic in
+// 740.5 us (~1.4 TB/s). HBM is far from its ceiling, so the traffic that
+// matters is the L2-resident operand re-read, and the only structural
+// lever on it is the macro-tile reuse ratio 1/tileM + 1/tileN.
+//
+// Direct fragment loads vs cooperative A+B LDS staging. Direct global
+// fragment fills were rejected without measuring them: the m16n16k32
+// fragment owns the k dimension and lane index, so lane (r, kq) needs
+// 8 consecutive k bytes at one row; in a 64-lane wave that is 16 different
+// rows 2,048 B apart per 16-lane group (8 B useful out of a 32-B sector),
+// and each wave reads the same A fragment once per n-group. Per block that
+// is >= 4x the L2 sectors of the 16-B-per-thread coalesced cooperative
+// staging (which touches every operand byte exactly once per block), before
+// counting the duplicated A reads. Cooperative A+B staging therefore
+// stays, and this round moves the reuse lever: DOUBLE THE N MACRO-TILE.
+//
+// Change (ONE bounded factor: the macro-tile reuse geometry). The output
+// macro-tile becomes 256x128 at the same 256 threads / 4 wavefronts; the
+// 2(m) x 2(n) wave split gives every wave a 128x64 quadrant = 8 m-groups x
+// 4 n-groups = 32 m16n16k32 int32 accumulators (the A quadrant height and
+// the 88-byte A stride are unchanged, so the eight A fragment fills per kk
+// keep the exact round-1 addresses; the four B fragment fills per kk are the
+// only new ones). Global staging stays per-thread 16-byte coalesced int4
+// loads, 6 per thread per stage (A rows [0,64),[64,128),[128,192),[192,256)
+// at k-chunk (tid&3)*16 -- unchanged from round 1 -- plus TWO B bands now:
+// n = b_row0 and b_row0 + 64 at the same k-chunk). The single-buffered
+// 64-K LDS stage becomes A[256,88] = 22,528 B + B[128,72] = 9,216 B =
+// 31,744 B/block, both strides unchanged and still bank-safe, and the two
+// int64-half ds_write2_b64 staging stores are unchanged.
+//
+// Quantified reuse per macro-tile (bytes of global operand read per
+// operator replay, M*K*(N/tileN) + N*K*(M/tileM)):
+// (256, 64) -> A 805.3 MB + B 201.3 MB = 1,006.6 MB (round 1)
+// (256, 128) -> A 402.7 MB + B 201.3 MB = 604.0 MB (this round, -40%)
+// (128, 128) -> A 402.7 MB + B 402.7 MB = 805.3 MB (rejected: the
+// halved tileM double-charges B, only -20% total)
+// (512, 64) -> A 805.3 MB + B 100.7 MB = 906.0 MB (rejected: the
+// unchanged tileN leaves the dominant A term untouched)
+// LDS fragment reads per MMAC also fall from (8+2)/16 = 0.625 to
+// (8+4)/32 = 0.375 per kk, and the block count falls 1,536 -> 768 (still
+// 6.4 blocks/CU on the 120-CU part, so split-K remains unnecessary and the
+// caller workspace stays unused; the trusted occupancy-probe split
+// candidates [2,3,4,5] are irrelevant to a grid this wide).
+//
+// Expected effect: if the L2 operand-traffic term is the limiter, median
+// moves from ~754 us toward ~500-620 us; if the kernel is instead at the
+// MMAC/LDS-issue floor, the larger tile stays within a few percent (the
+// extra 64 accumulator VGPRs and the 768-block tail are the only costs).
+//
+// Round 8 (mandatory decision: fix_build - repair the previous round's
+// compile/API failure with the smallest source-owned change; no architecture
+// change, no redesign). Rounds 6 and 7 produced no built candidate because the
+// worker process was terminated by an infrastructure rate limit, and the
+// workspace source was restored to the accepted round-2 kernel. Round 6's
+// intended epilogue change is re-applied with the compile-safety correction
+// demanded by this round. Preserved byte-for-byte: the 256x128 macro-tile,
+// 256 threads / 4 wavefronts, 88-byte A and 72-byte B padded strides, the
+// single-buffered 64-K LDS stage (A 22,528 B + B 9,216 B = 31,744 B), the
+// one-ds_read_b64-per-fragment fills, the 32 int32 m16n16k32 accumulators and
+// their exact accumulation order, the two __syncthreads per stage, the exact
+// (m,n,k) dispatch guard, the packed n-major weight path and the scalar
+// generic fallback. The only edit is the epilogue conversion. hip's
+// __float2bfloat16 lowers on this DTK to a NaN-guarded exec-masked sequence
+// per element (~12 extra SALU/VALU plus a divergent branch region for each of
+// the 128 conversions per lane); it is replaced by the branchless
+// finite-domain RNE helper bf16_rne_finite_bits, which stores the identical
+// 16-bit payload for every finite value this epilogue can produce and can
+// differ only for NaN/Inf. The 4x4 shfl_xor transpose, the per-element
+// multiply order float(dot) * x_scale[row] * weight_scale[col], the tail-row
+// guard, the 8-byte coalesced store and the dispatch are unchanged.
+//
+// Generic scalar int8/int32 fallback for every unmatched (m, n, k), including
+// all small-M API cases (M=2, M=16) and M=3072 with the same (K, N): it always
+// reproduces the exact int32 dot over K in ascending k order, then the fused
+// float scale and bf16 store. For (k, n) == (2048, 6144) it decodes the
+// n-major packed layout produced above; every other (k, n) reads the identity
+// [K, N] layout.
//
// Rules honored: headers in the known-good order (hip/hip_runtime.h,
-// hip/hip_bfloat16.h, du_mma.h); wavefront = 64 and blockDim a multiple of
-// 64; int32 accumulation with float conversion only for
+// hip/hip_bfloat16.h, du_mma.h); wavefront = 64 and blockDim a multiple of 64;
+// int32 accumulation with float conversion only for
// dot * x_scale[m] * weight_scale[n]; hipLaunchKernelGGL on the
-// caller-provided PyTorch stream; no stream-0 / synchronization / temporary
-// allocation in gemm_out; no NVIDIA wmma/mma.sync/PTX/warp-32 masks; all
-// barriers on paths reached by every thread in the block; explicit dispatch
-// for the assigned shape with a scalar generic fallback for everything else.
+// caller-provided PyTorch stream; no allocation / packing / synchronization /
+// default-stream launch inside the timed API; no NVIDIA wmma/mma.sync/PTX/
+// FP8/INT4 and no raw inline asm; all __syncthreads() on paths reached by
+// every block thread; coalesced lanes follow contiguous N addresses; single
+// 31,744 B LDS buffer (no double buffering); explicit exact-shape dispatch for
+// the assigned shape with a preserved scalar generic fallback.
#include
#include
@@ -89,62 +160,37 @@ constexpr int kWaveSize = 64; // gfx928 native wavefront
constexpr int kTargetM = 4096;
constexpr int kTargetN = 6144;
constexpr int kTargetK = 2048;
-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;
-
-// Iteration-4 tile-shape constants: 256x64 output tile (M-major doubling of
-// the accepted 128x64 tile) at 256 threads = 4 wavefronts, each wave owning a
-// 128x32 quadrant (16 m16n16k32 int32 accumulators).
-constexpr int kBlockM256 = 2 * kBlockM;
-
-// Iteration-5 packed-B constants: for the exact (k, n) == (2048, 6144) the
-// weight is packed once (outside the timed region) into the n-major [N, K]
-// layout packed[n*K + kk] = raw[kk*N + n]; B is then staged n-major into LDS
-// rows of 72 B (64 + 8). 72 is an odd 8-byte count, so the 16 lanes of each
-// ds_read_b64 fragment-load phase land on 16 distinct bank pairs (18r mod 32,
-// r = 0..15 -> all 16 even dword starts: the conflict-free 4-phase floor),
-// and 72 % 16 == 8, so staging stores use two int64 halves (ds_write2_b64)
-// instead of one misaligned int4.
-constexpr int kBPadPacked = 8; // 64 -> 72-byte n-major B LDS row stride
-constexpr int kBStridePacked = kStageK + kBPadPacked;
-// Iteration-9 (A-stride round) constant: packedb A LDS row stride 64 -> 88
-// bytes, replacing the shared 80-byte kAStride for w8a8_dumma_256x64x64_
-// packedb_kernel only (the identity-layout kernels keep kAStride = 80).
-// 88 = 22 dwords: 22r mod 32 is a permutation of the 16 even dword starts
-// over r = 0..15, so the 16 lanes of each ds_read_b64 fragment-load phase
-// land on 16 distinct bank pairs (the same conflict-free 4-phase floor the
-// packed-B stride 72 already uses); the current 80-byte (20-dword) stride
-// has 20r mod 32 with period 8, so rows r and r+8 alias onto one bank phase
-// and EVERY A fragment read pays a 2-way bank conflict (PMC
-// lds_bank_conflicts 20,447,232, deterministic across iterations 5..8).
-// 88 % 16 == 8, so the 16-byte staging stores go through two int64 halves
-// (ds_write2_b64, the exact accepted packed-B staging pattern) - the same
-// instruction count as the current ds_write_b128. A[256,88] = 22,528 B +
-// B[64,72] = 4,608 B = 27,136 B/block; 2 blocks/CU = 54,272 <= 65,536, so
-// co-residency, VGPR, barrier count and the global-load pattern are all
-// unchanged; only the physical LDS addresses move (same bytes stored, same
-// bytes read, fragment values bit-identical).
-constexpr int kAStridePacked = kStageK + 24; // 64 -> 88-byte A LDS row stride
+// Round 2: 256x128 macro-tile, four wavefronts, single-buffered 64-K stage.
+// Doubling the N macro-tile halves the A re-read term M*K*(N/tileN) while the
+// unchanged tileM keeps the B term N*K*(M/tileM); total global operand reads
+// fall 1,006.6 -> 604.0 MB per replay (-40%).
+constexpr int kBlockM = 256;
+constexpr int kBlockN = 128;
+constexpr int kStageK = 64;
+constexpr int kBlockThreads = 4 * kWaveSize; // 256
+
+// A[256, 88]: m-major, 88 = 22 dwords (odd 8-byte count). 22*r mod 32 over
+// r = 0..15 is a permutation of the 16 even dword starts, so every 16-lane
+// ds_read_b64 A fragment phase is conflict-free; 88 % 16 == 8 forces the
+// two-int64-half (ds_write2_b64) staging store.
+constexpr int kAStride = kStageK + 24; // 64 -> 88-byte A LDS row stride
+// B[128, 72]: n-major (row index = n), 72 = 18 dwords. 18*r mod 32 over
+// r = 0..15 is the same conflict-free permutation; 72 % 16 == 8 gives the same
+// two-int64-half staging store.
+constexpr int kBStride = kStageK + 8; // 64 -> 72-byte B LDS row stride
using bf16_t = hip_bfloat16;
-// Direct 8-byte LDS fragment fill for matrix_a on the m-major A tile.
-// The m16n16k32 matrix_a row_major loader assigns x[i] =
-// p[row*ldm + col_group*8 + i] (row = lane & 15, col_group = lane >> 4),
-// i.e. eight consecutive k-values per lane; the library path makes the
-// compiler emit per-byte ds_read_u8 loads plus a mask/OR reassembly chain
-// per dword that is an arithmetic identity on the v_mmac A operand. Writing
-// the same 8 bytes straight into the fragment storage keeps the operand bit
-// pattern and the int32 accumulation identical and lets the compiler feed a
-// single 8-byte LDS read to the v_mmac. This is the exact pattern validated
-// on the worker-29 TP4 gate_up lineage (accepted iteration 13) and the TP8
-// hy3 o_proj lineage (accepted iteration 12).
+// Direct 8-byte LDS fragment fill. The m16n16k32 fragment layout assigns
+// x[i] = p[(row++) * ldm + col_group*8 + i] (row = lane & 15,
+// col_group = lane >> 4), i.e. eight consecutive 8-bit elements per lane on a
+// row-major tile. On the m-major A tile these are eight consecutive k bytes of
+// one row; on the n-major B tile (row = n) they are eight consecutive k bytes
+// of one column. Writing those bytes straight into the fragment storage keeps
+// the operand bit pattern, the v_mmac inputs and the int32 accumulation order
+// identical to the library loader while emitting one ds_read_b64 instead of
+// eight per-byte ds_read_u8 plus a mask/OR reassembly chain per dword.
template
__device__ __forceinline__ void load_fragment8(
Frag& frag,
@@ -156,103 +202,46 @@ __device__ __forceinline__ void load_fragment8(
*reinterpret_cast(base + off);
}
-// Coalesced fragment store. The m16n16k32 accumulator lane mapping
-// (row = lane & 15, column group c4 = lane >> 4, frag.x[i] -> column
-// c4 + 4*i) gives each lane four elements strided by 4 columns, so a direct
-// store would be four 2-byte scalar stores per lane whose wavefront
-// addresses touch each 32-B sector at 25% utilization. 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), 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). This DTK lowers
-// __shfl_xor to ds_bpermute at the block tail where the LDS pipe is idle.
-// 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
-// bf16 rounding are unchanged, so the stored bits are identical to the
-// direct per-element store. 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 16 and n*2 a
-// multiple of 8, so the float4 weight_scale load and the 8-byte store are
-// aligned.
-template
-__device__ __forceinline__ void store_prefill_fragment_coalesced(
- const AccFragment& frag,
- const float* __restrict__ x_scale,
- 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];
- const int x2 = frag.x[2];
- const int x3 = frag.x[3];
- // Step 1: transpose the 2x2 in-block pairs (c4 bit0 <-> element bit0).
- const int s0 = __shfl_xor(x0, 16, kWaveSize);
- const int s1 = __shfl_xor(x1, 16, kWaveSize);
- const int s2 = __shfl_xor(x2, 16, kWaveSize);
- const int s3 = __shfl_xor(x3, 16, kWaveSize);
- const bool lo = (c4 & 1) == 0;
- const int a0 = lo ? x0 : s1;
- const int a1 = lo ? s0 : x1;
- const int a2 = lo ? x2 : s3;
- const int a3 = lo ? s2 : x3;
- // Step 2: transpose the 2x2 blocks (c4 bit1 <-> element bit1).
- const int t0 = __shfl_xor(a0, 32, kWaveSize);
- const int t1 = __shfl_xor(a1, 32, kWaveSize);
- const int t2 = __shfl_xor(a2, 32, kWaveSize);
- const int t3 = __shfl_xor(a3, 32, kWaveSize);
- const bool hi = (c4 & 2) == 0;
- const int f0 = hi ? a0 : t2;
- const int f1 = hi ? a1 : t3;
- 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).
- const int col0 = base_col + 4 * c4;
- const float xs = x_scale[row];
- 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;
- const uint64_t 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;
+// Round 8 (fix_build / epilogue): plain round-to-nearest-even fp32 -> bf16
+// conversion without hip's __float2bfloat16 NaN/Inf signaling-bit fixup. The
+// arithmetic (u + 0x7fff + ((u >> 16) & 1)) >> 16 is exactly the finite-domain
+// path of DTK's float_2_bfloatraw (exponent bits != all-ones), so every finite
+// fp32 input - including +/-0, denormals, and values whose RNE carries into or
+// out of the bf16 exponent field - produces the identical 16-bit pattern as
+// __float2bfloat16; the outputs can differ only for NaN/Inf inputs, which this
+// kernel's epilogue can never produce (|int32 dot| <= K*127^2 and both scale
+// inputs are finite, so the scaled accumulator is always finite). The helper
+// is spelled with a union rather than __float_as_uint so it depends on no
+// optional device intrinsic; this is the validated formulation from the
+// accepted hy3/TP8/M4096 o_proj lineage, not a new architecture.
+__device__ __forceinline__ unsigned short bf16_rne_finite_bits(float f) {
+ union {
+ float fp32;
+ unsigned int u32;
+ } u = {f};
+ u.u32 += 0x7fffu + ((u.u32 >> 16) & 1u); // round half to even at bit 16
+ return static_cast(u.u32 >> 16);
}
-// Register-batched variant of the coalesced fragment store (iteration 6,
-// epilogue round). Identical to store_prefill_fragment_coalesced in every
-// output bit: same 4x4 shuffle transpose, same per-element multiply order
-// float(dot) * x_scale[row] * weight_scale[col], same __float2bfloat16
-// rounding, same single 8-byte coalesced store. The only difference is that
-// the per-row x_scale and per-column weight_scale values are passed in from
-// registers: the caller loads each lane's 8 x_scale rows and 2 weight_scale
-// float4 columns exactly once per wave (the per-call epilogue re-loaded the
-// scales inside each of the 16 fragment stores), so the stored bits are
-// unchanged while the epilogue vmem read count per lane per wave drops from
-// up to 8 x_scale + 16 float4 loads to exactly 8 + 2 and the loads are
-// independent of (and issue before) the shuffle/convert/store chains.
+// Fused coalesced fragment store with register-batched scales. The m16n16k32
+// accumulator lane mapping (row = lane & 15, column group c4 = lane >> 4,
+// frag.x[i] -> column c4 + 4*i) gives each lane four elements strided by 4
+// columns, so a direct store would be four 2-byte scalar stores per lane whose
+// wavefront addresses touch each 32-B sector at 25% utilization. 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), 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 multiply order
+// (float(dot) * x_scale[row] * weight_scale[col]) and the bf16 rounding are
+// unchanged, so the stored bits are identical to a direct per-element store.
+// xs is this lane's x_scale[row] and ws the float4 at
+// weight_scale + base_col + 4*c4, both preloaded once by the caller. The
+// row >= m guard is wavefront-uniform (all 64 lanes share the same 16-row
+// window), so the shuffles never mix active and inactive lanes; base_col is a
+// multiple of 16 and n*2 a multiple of 8, so the float4 scale load and the
+// 8-byte store are aligned.
template
__device__ __forceinline__ void store_prefill_fragment_coalesced_scaled(
const AccFragment& frag,
@@ -294,612 +283,42 @@ __device__ __forceinline__ void store_prefill_fragment_coalesced_scaled(
const int f2 = hi ? t0 : a2;
const int f3 = hi ? t1 : a3;
- // Lane (r, c4) 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*2 is a multiple of 8).
const int col0 = base_col + 4 * c4;
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;
const uint64_t 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);
+ static_cast(bf16_rne_finite_bits(v0)) |
+ (static_cast(bf16_rne_finite_bits(v1)) << 16) |
+ (static_cast(bf16_rne_finite_bits(v2)) << 32) |
+ (static_cast(bf16_rne_finite_bits(v3)) << 48);
*reinterpret_cast(out + static_cast(row) * n + col0) =
packed;
}
// ---------------------------------------------------------------------------
-// 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] from the identity [K, N] weight into a single-buffered LDS
-// buffer. Two barriers per stage: one after the cooperative load, one before
-// the next stage overwrites LDS. Direct fragment epilogue (coalesced store).
+// Exact-shape GEMM: 256x128 output tile per block, four wavefronts.
//
-// Stage layout (bootstrap, identity [K, N] weight):
-// A[128, 80]: m-major, 512 int4s; thread t loads int4 #0 at
-// a_tile[(t/4)*80 + (t%4)*16] from x_q[(m0 + t/4)*k + k0 +
-// (t%4)*16] and int4 #1 at row +64 (rows 64..127).
-// B[64, 80]: k-major, 256 int4s; thread t loads int4 at
-// b_tile[(t/4)*80 + (t%4)*16] from
-// weight[(k0 + t/4)*N + n0 + (t%4)*16] (16 consecutive n at
-// fixed k; one int4 per thread). B fragments are loaded with
-// the library row_major loader (ldm = 80): lane (r, c4) holds
-// x[i] = b_tile[(kk + r + i)*80 + local_col + c4], i.e. eight
-// k-rows strided by ldm - correct for the identity [K, N]
-// layout; later rounds can switch to an n-major packed B with
-// the 8-byte fragment fill.
-// ---------------------------------------------------------------------------
-__global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_128x64x64_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 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 64-K 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).
- __shared__ __align__(16) int8_t a_tile[kBlockM * kAStride];
- __shared__ __align__(16) int8_t b_tile[kBlockN * kBStride];
-
- DUFragment
- a_frag0, a_frag1;
- DUFragment
- b_frag0, b_frag1;
- DUFragment
- acc00, acc01, acc10, acc11, acc20, acc21, acc30, acc31;
- du_fill_fragment(acc00, 0);
- du_fill_fragment(acc01, 0);
- du_fill_fragment(acc10, 0);
- du_fill_fragment(acc11, 0);
- du_fill_fragment(acc20, 0);
- du_fill_fragment(acc21, 0);
- du_fill_fragment(acc30, 0);
- du_fill_fragment(acc31, 0);
-
- // Cooperative staging: each thread owns one 16-byte vector in A rows
- // [0,64), one in A rows [64,128), and one 16-byte vector of B.
- // A[128,64] is 512 int4s; B[64,64] is 256 int4s.
- const int vector_byte_offset = tid * static_cast(sizeof(int4));
- const int stage_row = vector_byte_offset / kStageK; // tid / 4, 0..63
- const int stage_col = vector_byte_offset - stage_row * kStageK; // (tid%4)*16
- const int b_row = tid >> 2; // k row of this thread's B int4
- const int b_n16 = (tid & 3) * 16; // n offset (stage-local) of B int4
-
- for (int k0 = 0; k0 < k; k0 += kStageK) {
- *reinterpret_cast(a_tile + stage_row * kAStride + stage_col) =
- *reinterpret_cast(
- x_q + static_cast(m0 + stage_row) * k + k0 + stage_col);
- *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);
- *reinterpret_cast(b_tile + b_row * kBStride + b_n16) =
- *reinterpret_cast(
- weight + (static_cast(k0 + b_row) * n + n0 + b_n16));
- __syncthreads();
-
- // Each wave consumes its 64x32 quadrant: eight m16n16k32 MMACs per kk.
-#pragma unroll
- for (int kk = 0; kk < kStageK; kk += kTileK) {
- // A fragments: direct 8-byte LDS fill (m-major A tile; each lane's
- // eight fragment k-values are contiguous).
- load_fragment8(a_frag0, a_tile + local_row * kAStride + kk, kAStride,
- lane);
- load_fragment8(a_frag1, a_tile + (local_row + kTileM) * kAStride + kk,
- kAStride, lane);
- // B fragments: library row_major loader on the k-major [K, N] LDS
- // tile (bootstrap, identity layout; ldm = 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_mma_sync(acc00, a_frag0, b_frag0, acc00);
- du_mma_sync(acc01, a_frag0, b_frag1, acc01);
- du_mma_sync(acc10, a_frag1, b_frag0, acc10);
- du_mma_sync(acc11, a_frag1, b_frag1, acc11);
- load_fragment8(
- a_frag0, a_tile + (local_row + 2 * kTileM) * kAStride + kk,
- kAStride, lane);
- load_fragment8(
- a_frag1, a_tile + (local_row + 3 * kTileM) * kAStride + kk,
- kAStride, lane);
- du_mma_sync(acc20, a_frag0, b_frag0, acc20);
- du_mma_sync(acc21, a_frag0, b_frag1, acc21);
- du_mma_sync(acc30, a_frag1, b_frag0, acc30);
- du_mma_sync(acc31, a_frag1, b_frag1, acc31);
- }
-
- // Protect the LDS buffer from the next stage's cooperative overwrite.
- __syncthreads();
- }
-
- const int base_row = m0 + local_row;
- const int base_col = n0 + local_col;
- // Coalesced epilogue store (one 8-byte store per lane per fragment; scales
- // + bf16 conversion fused in-kernel, no workspace pass).
- store_prefill_fragment_coalesced(acc00, x_scale, weight_scale, out,
- base_row, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc01, x_scale, weight_scale, out,
- base_row, base_col + kTileN, m, n, lane);
- store_prefill_fragment_coalesced(acc10, x_scale, weight_scale, out,
- base_row + kTileM, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc11, x_scale, weight_scale, out,
- base_row + kTileM, base_col + kTileN, m, n,
- lane);
- store_prefill_fragment_coalesced(acc20, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc21, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc30, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc31, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col + kTileN,
- m, n, lane);
-}
-
-// ---------------------------------------------------------------------------
-// Iteration-3 pipeline-control variant: double-buffered K-stage of the same
-// 128x64 tile / 64-K stage / fragment loads / epilogue as the accepted
-// single-buffered kernel. Stage s+1's three global loads (2 A int4s + 1 B
-// int4 per thread) are issued BEFORE the stage s MMAC burst and the loaded
-// vectors stay in VGPR; the publish (ds_write into the idle buffer) happens
-// AFTER the burst, so the compiler's vmcnt(0) wait sits behind 16 v_mmac +
-// ~70 LDS reads instead of directly behind the load issue (the accepted
-// single-buffered kernel waits immediately after 3 loads per wave per
-// stage). One __syncthreads per stage (32/block vs 64/block: the barrier
-// doubles as read-done for stage s+1 and write-done for stage s; the last
-// stage needs no fence before the register epilogue). Costs: LDS
-// 15,360 -> 30,720 B/block (A[2][128,80] + B[2][64,80]), residency
-// 4 -> 2 blocks/CU = 8 waves/CU; VGPR cap at 2 blocks/CU is 128, so the
-// ~12-VGPR loop-carried prefetch payload cannot change residency. Retained
-// only if ISA/PMC evidence shows the vmcnt wait moved off the load->ds_write
-// path (covered by the MMAC burst) and the measured median/p90 beat the
-// 1162.15/1164.03 guard.
-// ---------------------------------------------------------------------------
-__global__ __launch_bounds__(kBlockThreads) void w8a8_dumma_128x64x64_db_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 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;
-
- // Double-buffered 64-K stage: A[2][128,80] + B[2][64,80] = 30,720 B/block
- // (2 blocks/CU fit the 64 KiB LDS budget; padded 80-byte strides and bank
- // phases unchanged from the accepted kernel).
- __shared__ __align__(16) int8_t a_tile[2][kBlockM * kAStride];
- __shared__ __align__(16) int8_t b_tile[2][kBlockN * kBStride];
-
- DUFragment
- a_frag0, a_frag1;
- DUFragment
- b_frag0, b_frag1;
- DUFragment
- acc00, acc01, acc10, acc11, acc20, acc21, acc30, acc31;
- du_fill_fragment(acc00, 0);
- du_fill_fragment(acc01, 0);
- du_fill_fragment(acc10, 0);
- du_fill_fragment(acc11, 0);
- du_fill_fragment(acc20, 0);
- du_fill_fragment(acc21, 0);
- du_fill_fragment(acc30, 0);
- du_fill_fragment(acc31, 0);
-
- // Cooperative staging indexing is identical to the accepted kernel: each
- // thread owns two 16-byte A vectors (rows [0,64) and [64,128)) and one
- // 16-byte B vector.
- const int vector_byte_offset = tid * static_cast(sizeof(int4));
- const int stage_row = vector_byte_offset / kStageK; // tid / 4, 0..63
- const int stage_col = vector_byte_offset - stage_row * kStageK; // (tid%4)*16
- const int b_row = tid >> 2; // k row of this thread's B int4
- const int b_n16 = (tid & 3) * 16; // n offset (stage-local) of B int4
-
- const int k_stages = k / kStageK;
-
- // Loop-carried prefetch payload (stage s+1 global vectors in VGPR).
- int4 pa0, pa1, pb0;
-
- // Prologue: fetch stage 0 into registers, publish into buffer 0, fence.
- {
- const int k0 = 0;
- pa0 = *reinterpret_cast(
- x_q + static_cast(m0 + stage_row) * k + k0 + stage_col);
- pa1 = *reinterpret_cast(
- x_q + static_cast(m0 + stage_row + kBlockM / 2) * k + k0 +
- stage_col);
- pb0 = *reinterpret_cast(
- weight + (static_cast(k0 + b_row) * n + n0 + b_n16));
- *reinterpret_cast(a_tile[0] + stage_row * kAStride + stage_col) =
- pa0;
- *reinterpret_cast(a_tile[0] +
- (stage_row + kBlockM / 2) * kAStride +
- stage_col) = pa1;
- *reinterpret_cast(b_tile[0] + b_row * kBStride + b_n16) = pb0;
- __syncthreads();
- }
-
- for (int s = 0; s < k_stages; ++s) {
- const int cur = s & 1;
- const int nxt = cur ^ 1;
-
- // Prefetch stage s+1: issue the global loads BEFORE the stage s MMAC
- // burst; the loaded vectors stay in VGPR and the compiler's vmcnt(0) wait
- // is inserted only at the publish below (covered by the burst).
- if (s + 1 < k_stages) {
- const int k0 = (s + 1) * kStageK;
- pa0 = *reinterpret_cast(
- x_q + static_cast(m0 + stage_row) * k + k0 + stage_col);
- pa1 = *reinterpret_cast(
- x_q + static_cast(m0 + stage_row + kBlockM / 2) * k + k0 +
- stage_col);
- pb0 = *reinterpret_cast(
- weight + (static_cast(k0 + b_row) * n + n0 + b_n16));
- }
-
- // Consume stage s from the current buffer: bit-identical fragment values,
- // MMAC sequence, int32 accumulation order and lane mapping to the
- // accepted kernel (only the buffer base changes).
-#pragma unroll
- for (int kk = 0; kk < kStageK; kk += kTileK) {
- load_fragment8(a_frag0, a_tile[cur] + local_row * kAStride + kk,
- kAStride, lane);
- load_fragment8(a_frag1, a_tile[cur] + (local_row + kTileM) * kAStride +
- kk,
- kAStride, lane);
- du_load_matrix_sync(b_frag0,
- b_tile[cur] + kk * kBStride + local_col, kBStride);
- du_load_matrix_sync(b_frag1,
- b_tile[cur] + kk * kBStride + local_col + kTileN,
- kBStride);
- du_mma_sync(acc00, a_frag0, b_frag0, acc00);
- du_mma_sync(acc01, a_frag0, b_frag1, acc01);
- du_mma_sync(acc10, a_frag1, b_frag0, acc10);
- du_mma_sync(acc11, a_frag1, b_frag1, acc11);
- load_fragment8(
- a_frag0, a_tile[cur] + (local_row + 2 * kTileM) * kAStride + kk,
- kAStride, lane);
- load_fragment8(
- a_frag1, a_tile[cur] + (local_row + 3 * kTileM) * kAStride + kk,
- kAStride, lane);
- du_mma_sync(acc20, a_frag0, b_frag0, acc20);
- du_mma_sync(acc21, a_frag0, b_frag1, acc21);
- du_mma_sync(acc30, a_frag1, b_frag0, acc30);
- du_mma_sync(acc31, a_frag1, b_frag1, acc31);
- }
-
- // Publish stage s+1 into the idle buffer after the burst (vmcnt wait
- // covered by the burst), then one fence: read-done for s+1 and
- // write-done for s (buffer cur is free for the s+2 publish).
- if (s + 1 < k_stages) {
- *reinterpret_cast(a_tile[nxt] + stage_row * kAStride +
- stage_col) = pa0;
- *reinterpret_cast(a_tile[nxt] +
- (stage_row + kBlockM / 2) * kAStride +
- stage_col) = pa1;
- *reinterpret_cast(b_tile[nxt] + b_row * kBStride + b_n16) = pb0;
- __syncthreads();
- }
- }
-
- const int base_row = m0 + local_row;
- const int base_col = n0 + local_col;
- store_prefill_fragment_coalesced(acc00, x_scale, weight_scale, out,
- base_row, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc01, x_scale, weight_scale, out,
- base_row, base_col + kTileN, m, n, lane);
- store_prefill_fragment_coalesced(acc10, x_scale, weight_scale, out,
- base_row + kTileM, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc11, x_scale, weight_scale, out,
- base_row + kTileM, base_col + kTileN, m, n,
- lane);
- store_prefill_fragment_coalesced(acc20, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc21, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc30, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc31, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col + kTileN,
- m, n, lane);
-}
-
-// ---------------------------------------------------------------------------
-// Iteration-4 tile-shape variant: 256x64 output tile (M-major doubling of the
-// accepted 128x64 tile) at 256 threads = 4 wavefronts, each wave owning a
-// 128x32 quadrant (sixteen m16n16k32 int32 accumulators). This is the
-// measured tile axis pushed one step further: per-FLOP B-fragment LDS reads
-// scale as 1/(quadrant m-height) (the 64x128 flip's quadrants were 32x64 and
-// lost +8% with a doubled per-FLOP B cost), so doubling the quadrant height
-// to 128 HALVES the per-FLOP B ds_read_u8 + reassembly VALU - the measured
-// LDS-read/wait bottleneck (PMC lds_wait ~= lds_instructions) - while keeping
-// the same 64-column quadrant width (B library row_major loader ldm = 80).
-// B reuse doubles (each staged B element is consumed by all 256 block rows),
-// so the grid-level B re-read drops from M/128 = 32x to M/256 = 16x (staged B
-// traffic 402.7 -> 201.3 MB/replay; A re-read stays N/64 = 96x = 805.3 MB;
-// total 1.208 -> 1.007 GB, -17%). Single-buffered 64-K stage: A[256,80] +
-// B[64,80] = 25,600 B/block -> 2 blocks/CU = 8 waves/CU, the SAME
-// co-residency structure as the accepted double-buffered 128x64 kernel
-// (2 blocks x 4 waves), so the comparison is not residency-confounded.
-// Two __syncthreads per stage, direct 8-byte A fragment fills, fused
-// coalesced scale->bf16 epilogue, row>=m guard. Grid dim3(96,16) = 1536
-// blocks = 12.8 blocks/CU, so the MxN output-tile grid still dwarfs the
-// 120 CUs. Fragment values, MMAC sequence, int32 accumulation order
-// (k0-outer 64, kk-inner 32, hardware k-order in m16n16k32),
-// element-to-lane mapping, barrier count and the fused epilogue are
-// bit-identical to the accepted kernels (only the tile height, the quadrant
-// split and the staging vectorization change), so exact bf16 equality
-// (mismatch 0 / max_abs_error 0.0) is expected without tolerance debate.
-// ---------------------------------------------------------------------------
-__global__ __launch_bounds__(kBlockThreads)
- void w8a8_dumma_256x64x64_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;
- // 4 waves: wave_row = (wave>>1)*128 covers rows 0..255 in two 128-row bands
- // (waves 2w, 2w+1 share the band), wave_col picks the 32-col half.
- const int wave_row = (wave >> 1) * (kBlockM256 / 2);
- const int wave_col = wave & 1;
- const int m0 = static_cast(blockIdx.y) * kBlockM256;
- const int n0 = static_cast(blockIdx.x) * kBlockN;
- const int local_row = wave_row;
- const int local_col = wave_col * 32;
-
- // Single-buffered 64-K stage: A[256, 80] + B[64, 80] = 25,600 B/block
- // (2 blocks/CU fit the 64 KiB LDS budget; padded 80-byte strides and bank
- // phases unchanged from the accepted kernels).
- __shared__ __align__(16) int8_t a_tile[kBlockM256 * kAStride];
- __shared__ __align__(16) int8_t b_tile[kBlockN * kBStride];
-
- DUFragment
- a_frag0, a_frag1, a_frag2, a_frag3, a_frag4, a_frag5, a_frag6, a_frag7;
- DUFragment
- b_frag0, b_frag1;
- // Sixteen m16n16k32 int32 accumulators (8 m-groups x 2 n-groups per wave).
- DUFragment
- acc00, acc01, acc10, acc11, acc20, acc21, acc30, acc31,
- acc40, acc41, acc50, acc51, acc60, acc61, acc70, acc71;
- du_fill_fragment(acc00, 0);
- du_fill_fragment(acc01, 0);
- du_fill_fragment(acc10, 0);
- du_fill_fragment(acc11, 0);
- du_fill_fragment(acc20, 0);
- du_fill_fragment(acc21, 0);
- du_fill_fragment(acc30, 0);
- du_fill_fragment(acc31, 0);
- du_fill_fragment(acc40, 0);
- du_fill_fragment(acc41, 0);
- du_fill_fragment(acc50, 0);
- du_fill_fragment(acc51, 0);
- du_fill_fragment(acc60, 0);
- du_fill_fragment(acc61, 0);
- du_fill_fragment(acc70, 0);
- du_fill_fragment(acc71, 0);
-
- // Cooperative staging: A[256,64] is 1,024 int4s (4 per thread: rows
- // [0,64), [64,128), [128,192), [192,256)) and B[64,64] is 256 int4s (one
- // per thread), 5 loads per thread per stage.
- const int a_row0 = tid >> 2; // 0..63
- const int a_col = (tid & 3) * 16; // 0,16,32,48
-
- for (int k0 = 0; k0 < k; k0 += kStageK) {
- *reinterpret_cast(a_tile + a_row0 * kAStride + a_col) =
- *reinterpret_cast(
- x_q + static_cast(m0 + a_row0) * k + k0 + a_col);
- *reinterpret_cast(a_tile +
- (a_row0 + kBlockM256 / 4) * kAStride + a_col) =
- *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + kBlockM256 / 4) * k +
- k0 + a_col);
- *reinterpret_cast(a_tile +
- (a_row0 + kBlockM256 / 2) * kAStride + a_col) =
- *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + kBlockM256 / 2) * k +
- k0 + a_col);
- *reinterpret_cast(a_tile +
- (a_row0 + 3 * kBlockM256 / 4) * kAStride +
- a_col) =
- *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + 3 * kBlockM256 / 4) * k +
- k0 + a_col);
- const int b_row = tid >> 2; // k row of this thread's B int4, 0..63
- const int b_n16 = (tid & 3) * 16;
- *reinterpret_cast(b_tile + b_row * kBStride + b_n16) =
- *reinterpret_cast(
- weight + (static_cast(k0 + b_row) * n + n0 + b_n16));
- __syncthreads();
-
- // Each wave consumes its 128x32 quadrant: eight m-groups x two n-groups,
- // 16 m16n16k32 MMACs per kk. The A fragment loads and the MMAC sequence
- // keep the exact int32 accumulation order (k0-outer 64, kk-inner 32,
- // hardware k-order in m16n16k32) and the element-to-lane mapping of the
- // accepted kernels; only the number of m-groups per wave changes.
-#pragma unroll
- for (int kk = 0; kk < kStageK; kk += kTileK) {
- load_fragment8(a_frag0, a_tile + local_row * kAStride + kk, kAStride,
- lane);
- load_fragment8(a_frag1, a_tile + (local_row + kTileM) * kAStride + kk,
- kAStride, lane);
- load_fragment8(a_frag2, a_tile + (local_row + 2 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- load_fragment8(a_frag3, a_tile + (local_row + 3 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- 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_mma_sync(acc00, a_frag0, b_frag0, acc00);
- du_mma_sync(acc01, a_frag0, b_frag1, acc01);
- du_mma_sync(acc10, a_frag1, b_frag0, acc10);
- du_mma_sync(acc11, a_frag1, b_frag1, acc11);
- du_mma_sync(acc20, a_frag2, b_frag0, acc20);
- du_mma_sync(acc21, a_frag2, b_frag1, acc21);
- du_mma_sync(acc30, a_frag3, b_frag0, acc30);
- du_mma_sync(acc31, a_frag3, b_frag1, acc31);
- load_fragment8(a_frag4, a_tile + (local_row + 4 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- load_fragment8(a_frag5, a_tile + (local_row + 5 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- load_fragment8(a_frag6, a_tile + (local_row + 6 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- load_fragment8(a_frag7, a_tile + (local_row + 7 * kTileM) * kAStride +
- kk,
- kAStride, lane);
- du_mma_sync(acc40, a_frag4, b_frag0, acc40);
- du_mma_sync(acc41, a_frag4, b_frag1, acc41);
- du_mma_sync(acc50, a_frag5, b_frag0, acc50);
- du_mma_sync(acc51, a_frag5, b_frag1, acc51);
- du_mma_sync(acc60, a_frag6, b_frag0, acc60);
- du_mma_sync(acc61, a_frag6, b_frag1, acc61);
- du_mma_sync(acc70, a_frag7, b_frag0, acc70);
- du_mma_sync(acc71, a_frag7, b_frag1, acc71);
- }
-
- // Protect the LDS buffer from the next stage's cooperative overwrite.
- __syncthreads();
- }
-
- const int base_row = m0 + local_row;
- const int base_col = n0 + local_col;
- store_prefill_fragment_coalesced(acc00, x_scale, weight_scale, out,
- base_row, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc01, x_scale, weight_scale, out,
- base_row, base_col + kTileN, m, n, lane);
- store_prefill_fragment_coalesced(acc10, x_scale, weight_scale, out,
- base_row + kTileM, base_col, m, n, lane);
- store_prefill_fragment_coalesced(acc11, x_scale, weight_scale, out,
- base_row + kTileM, base_col + kTileN, m, n,
- lane);
- store_prefill_fragment_coalesced(acc20, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc21, x_scale, weight_scale, out,
- base_row + 2 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc30, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc31, x_scale, weight_scale, out,
- base_row + 3 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc40, x_scale, weight_scale, out,
- base_row + 4 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc41, x_scale, weight_scale, out,
- base_row + 4 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc50, x_scale, weight_scale, out,
- base_row + 5 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc51, x_scale, weight_scale, out,
- base_row + 5 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc60, x_scale, weight_scale, out,
- base_row + 6 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc61, x_scale, weight_scale, out,
- base_row + 6 * kTileM, base_col + kTileN,
- m, n, lane);
- store_prefill_fragment_coalesced(acc70, x_scale, weight_scale, out,
- base_row + 7 * kTileM, base_col, m, n,
- lane);
- store_prefill_fragment_coalesced(acc71, x_scale, weight_scale, out,
- base_row + 7 * kTileM, base_col + kTileN,
- m, n, lane);
-}
-
-// ---------------------------------------------------------------------------
-// Iteration-5 packed-B variant of the accepted 256x64 kernel: the B operand
-// switches to the one-time n-major [N, K] pack (packed[n*K + kk] =
-// raw[kk*N + n], produced by launch_pack_w8a8_weight outside the timed
-// region and Graph capture; same byte count and same buffer, so captured
-// addresses are unchanged) and the B fragment path switches from the
-// library row_major loader (8 per-byte ds_read_u8 strided by ldm per lane
-// plus mask/OR reassembly VALU per dword - the measured LDS-read axis) to
-// the lineage-validated 8-byte fragment fill (load_fragment8, one
-// ds_read_b64 per col_major fragment, zero reassembly VALU) on the n-major
-// staged tile. The m16n16k32 matrix_b hardware slot for lane
-// (r = lane & 15, kq = lane >> 4) is B[k = kk + kq*8 + i][n = local_col +
-// i*16 + r] - eight consecutive k bytes at fixed n - so staging B n-major
-// (b_tile[n][k]) makes each lane's eight fragment bytes CONTIGUOUS in LDS:
-// the fill reads the same elements into the same fragment slots with the
-// same v_mmac inputs and the exact int32 accumulation order (k0-outer 64,
-// kk-inner 32, hardware k-order), so exact bf16 equality (mismatch 0 /
-// max_abs_error 0.0) is expected without tolerance debate. LDS bank
-// safety: the n-major B tile uses the 72-byte row stride (odd 8-byte
-// count) so the 16 lanes of each ds_read_b64 phase read 18r mod 32
-// (r = 0..15) distinct bank pairs - conflict-free; 72 % 16 == 8 forces the
-// two-int64-half staging stores (ds_write2_b64). Iteration 9 gives the A
-// tile the same conflict-free floor: the A row stride is 88 bytes (22
-// dwords; 22r mod 32 over r = 0..15 is a permutation of the 16 even dword
-// starts), replacing the 80-byte stride's 2-way class (20r mod 32, period
-// 8, rows r and r+8 alias), and A staging uses the same two-int64-half
-// ds_write2_b64 pattern (88 % 16 == 8). Everything else is bit-identical
-// to the accepted kernel: same bytes staged/read (only the physical LDS
-// addresses move), same 128x32 quadrant per wave, same MMAC sequence, two
-// __syncthreads per stage, fused coalesced scale->bf16 epilogue, row>=m
-// guard, grid dim3(96,16) = 1536 blocks. LDS 22,528 (A) + 4,608 (B) =
-// 27,136 B/block -> 2 blocks/CU = 8 waves/CU, the same co-residency
-// structure as the accepted kernel.
+// 4 waves: wave_row = (wave >> 1) * 128 covers rows 0..255 in two 128-row
+// bands (waves 2w and 2w+1 share the band), wave_col = (wave & 1) * 64 selects
+// the 64-column half. Each wave consumes a 128x64 quadrant: eight m-groups x
+// four n-groups = thirty-two m16n16k32 int32 accumulators. The eight A
+// fragment addresses per kk and the 88-byte A stride are byte-identical to
+// round 1; only the n-major B tile and its four fragment fills are new.
+//
+// Staging (6 int4 global loads per thread per stage):
+// A[256, 88]: thread t owns 16-byte A chunks at (row = t >> 2,
+// col = (t & 3) * 16) in each of the four 64-row bands; each
+// store is two int64 halves because 88 % 16 == 8.
+// B[128, 72]: thread t owns the 16-byte k-chunk at (n_row = t >> 2,
+// k16 = (t & 3) * 16) and at n_row + 64 of the packed n-major
+// weight; each store is two int64 halves because 72 % 16 == 8.
// ---------------------------------------------------------------------------
__global__ __launch_bounds__(kBlockThreads)
- void w8a8_dumma_256x64x64_packedb_kernel(
+ void w8a8_dumma_256x128x64_packedb_kernel(
const int8_t* __restrict__ x_q,
const int8_t* __restrict__ packed_w, // [N, K] n-major (exact shape)
const float* __restrict__ x_scale,
@@ -911,100 +330,105 @@ __global__ __launch_bounds__(kBlockThreads)
const int tid = static_cast(threadIdx.x);
const int wave = tid / kWaveSize;
const int lane = tid % kWaveSize;
- // 4 waves: wave_row = (wave>>1)*128 covers rows 0..255 in two 128-row bands
- // (waves 2w, 2w+1 share the band), wave_col picks the 32-col half.
- const int wave_row = (wave >> 1) * (kBlockM256 / 2);
- const int wave_col = wave & 1;
- const int m0 = static_cast(blockIdx.y) * kBlockM256;
+ const int wave_row = (wave >> 1) * (kBlockM / 2);
+ const int wave_col = (wave & 1) * (kBlockN / 2);
+ const int m0 = static_cast(blockIdx.y) * kBlockM;
const int n0 = static_cast(blockIdx.x) * kBlockN;
const int local_row = wave_row;
- const int local_col = wave_col * 32;
+ const int local_col = wave_col;
- // Single-buffered 64-K stage: A[256, 88] (iteration-9 stride) + B[64, 72]
- // n-major = 27,136 B/block (2 blocks/CU fit the 64 KiB LDS budget).
- __shared__ __align__(16) int8_t a_tile[kBlockM256 * kAStridePacked];
- __shared__ __align__(16) int8_t b_tile[kBlockN * kBStridePacked];
+ // Single-buffered 64-K stage: A[256, 88] + B[128, 72] = 31,744 B/block.
+ __shared__ __align__(16) int8_t a_tile[kBlockM * kAStride];
+ __shared__ __align__(16) int8_t b_tile[kBlockN * kBStride];
DUFragment
a_frag0, a_frag1, a_frag2, a_frag3, a_frag4, a_frag5, a_frag6, a_frag7;
DUFragment
- b_frag0, b_frag1;
- // Sixteen m16n16k32 int32 accumulators (8 m-groups x 2 n-groups per wave).
+ b_frag0, b_frag1, b_frag2, b_frag3;
+ // Thirty-two m16n16k32 int32 accumulators (8 m-groups x 4 n-groups per wave).
DUFragment
- acc00, acc01, acc10, acc11, acc20, acc21, acc30, acc31,
- acc40, acc41, acc50, acc51, acc60, acc61, acc70, acc71;
+ acc00, acc01, acc02, acc03, acc10, acc11, acc12, acc13,
+ acc20, acc21, acc22, acc23, acc30, acc31, acc32, acc33,
+ acc40, acc41, acc42, acc43, acc50, acc51, acc52, acc53,
+ acc60, acc61, acc62, acc63, acc70, acc71, acc72, acc73;
du_fill_fragment(acc00, 0);
du_fill_fragment(acc01, 0);
+ du_fill_fragment(acc02, 0);
+ du_fill_fragment(acc03, 0);
du_fill_fragment(acc10, 0);
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);
du_fill_fragment(acc40, 0);
du_fill_fragment(acc41, 0);
+ du_fill_fragment(acc42, 0);
+ du_fill_fragment(acc43, 0);
du_fill_fragment(acc50, 0);
du_fill_fragment(acc51, 0);
+ du_fill_fragment(acc52, 0);
+ du_fill_fragment(acc53, 0);
du_fill_fragment(acc60, 0);
du_fill_fragment(acc61, 0);
+ du_fill_fragment(acc62, 0);
+ du_fill_fragment(acc63, 0);
du_fill_fragment(acc70, 0);
du_fill_fragment(acc71, 0);
+ du_fill_fragment(acc72, 0);
+ du_fill_fragment(acc73, 0);
- // Cooperative staging: A[256,64] is 1,024 int4s (4 per thread: rows
- // [0,64), [64,128), [128,192), [192,256)) and B[64,64] is 256 int4s (one
- // per thread), 5 loads per thread per stage.
const int a_row0 = tid >> 2; // 0..63
- const int a_col = (tid & 3) * 16; // 0,16,32,48
- // B staging is n-major: thread owns the 16-byte k-chunk at
- // (n_row = tid >> 2, k16 = (tid & 3) * 16).
+ const int a_col = (tid & 3) * 16; // 0, 16, 32, 48
const int b_n_row = tid >> 2; // 0..63
- const int b_k16 = (tid & 3) * 16; // 0,16,32,48
+ const int b_k16 = (tid & 3) * 16; // 0, 16, 32, 48
const int base_row = m0 + local_row;
const int base_col = n0 + local_col;
- // Iteration-6 (epilogue round) register-batched scales: each lane's 16
- // fragment stores cover the 8 m-groups (rows base_row + r + 16*i,
- // r = lane & 15) x 2 n-groups (cols base_col + 4*c4 and
- // base_col + 16 + 4*c4, c4 = lane >> 4), so each lane needs exactly 8
- // x_scale values and 2 weight_scale float4s. They are loaded ONCE on the
- // last K stage, issued BEFORE the final protective barrier (uniform
- // branch: k0 is block-uniform, so the barrier below is reached by every
- // thread), so the vmem latency overlaps the barrier and the 16
- // shuffle/convert/store chains; the previous per-call epilogue re-loaded
- // the scales inside each of the 16 fragment stores (PMC vmem_read
- // surplus ~24-25 wavefront loads/wave above the 5/thread/stage staging
- // floor matches ~8 x_scale + ~16 float4 re-loads per lane per wave).
+
+ // Epilogue scale registers: each lane's 32 fragment stores cover the 8
+ // m-groups (rows base_row + 16*i + r, r = lane & 15) x 4 n-groups
+ // (cols base_col + 16*j + 4*c4, c4 = lane >> 4), so each lane needs exactly
+ // 8 x_scale values and 4 weight_scale float4s. They are loaded once on the
+ // last K stage (uniform branch: k0 is block-uniform, so the following
+ // barrier is reached by every thread) and reused by all thirty-two coalesced
+ // stores; xs_m is zeroed for out-of-range tail rows.
float xs_m[8];
- float4 ws0, ws1;
+ float4 ws0, ws1, ws2, ws3;
for (int k0 = 0; k0 < k; k0 += kStageK) {
const int4 av0 = *reinterpret_cast(
x_q + static_cast(m0 + a_row0) * k + k0 + a_col);
const int4 av1 = *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + kBlockM256 / 4) * k + k0 +
+ x_q + static_cast(m0 + a_row0 + kBlockM / 4) * k + k0 +
a_col);
const int4 av2 = *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + kBlockM256 / 2) * k + k0 +
+ x_q + static_cast(m0 + a_row0 + kBlockM / 2) * k + k0 +
a_col);
const int4 av3 = *reinterpret_cast(
- x_q + static_cast(m0 + a_row0 + 3 * kBlockM256 / 4) * k +
- k0 + a_col);
+ x_q + static_cast(m0 + a_row0 + 3 * kBlockM / 4) * k + k0 +
+ a_col);
// A rows are 88 B apart (8 mod 16), so each 16-byte staging store is two
- // int64 halves (ds_write2_b64, one instruction) - the exact accepted
- // packed-B staging pattern below; a single int4 store would be
- // misaligned on odd rows.
+ // int64 halves (ds_write2_b64); a single int4 store would be misaligned on
+ // odd rows.
const int64_t* asrc0 = reinterpret_cast(&av0);
const int64_t* asrc1 = reinterpret_cast(&av1);
const int64_t* asrc2 = reinterpret_cast(&av2);
const int64_t* asrc3 = reinterpret_cast(&av3);
int64_t* adst0 = reinterpret_cast(
- a_tile + a_row0 * kAStridePacked + a_col);
+ a_tile + a_row0 * kAStride + a_col);
int64_t* adst1 = reinterpret_cast(
- a_tile + (a_row0 + kBlockM256 / 4) * kAStridePacked + a_col);
+ a_tile + (a_row0 + kBlockM / 4) * kAStride + a_col);
int64_t* adst2 = reinterpret_cast(
- a_tile + (a_row0 + kBlockM256 / 2) * kAStridePacked + a_col);
+ a_tile + (a_row0 + kBlockM / 2) * kAStride + a_col);
int64_t* adst3 = reinterpret_cast(
- a_tile + (a_row0 + 3 * kBlockM256 / 4) * kAStridePacked + a_col);
+ a_tile + (a_row0 + 3 * kBlockM / 4) * kAStride + a_col);
adst0[0] = asrc0[0];
adst0[1] = asrc0[1];
adst1[0] = asrc1[0];
@@ -1013,78 +437,95 @@ __global__ __launch_bounds__(kBlockThreads)
adst2[1] = asrc2[1];
adst3[0] = asrc3[0];
adst3[1] = asrc3[1];
- // Stage B[64, 64] n-major from the packed [N, K] weight: 16 consecutive
- // k bytes per thread at fixed n (one coalesced int4 global load). Row
- // stride 72 is 8 mod 16, so the LDS store is two int64 halves
- // (ds_write2_b64) - a single int4 store would be misaligned on odd rows.
- const int4 bv = *reinterpret_cast(
- packed_w +
- (static_cast(n0 + b_n_row) * k + k0 + b_k16));
- const int64_t* bsrc64 = reinterpret_cast(&bv);
- int64_t* bdst64 = reinterpret_cast(
- b_tile + b_n_row * kBStridePacked + b_k16);
- bdst64[0] = bsrc64[0];
- bdst64[1] = bsrc64[1];
+ // B staging is n-major: 16 consecutive k bytes at fixed n (one coalesced
+ // int4 global load from the packed [N, K] weight) in two 64-row bands;
+ // row stride 72 is 8 mod 16, so each LDS store is two int64 halves.
+ const int4 bv0 = *reinterpret_cast(
+ packed_w + static_cast(n0 + b_n_row) * k + k0 + b_k16);
+ const int4 bv1 = *reinterpret_cast(
+ packed_w + static_cast(n0 + b_n_row + kBlockN / 2) * k + k0 +
+ b_k16);
+ const int64_t* bsrc0 = reinterpret_cast(&bv0);
+ const int64_t* bsrc1 = reinterpret_cast(&bv1);
+ int64_t* bdst0 = reinterpret_cast(
+ b_tile + b_n_row * kBStride + b_k16);
+ int64_t* bdst1 = reinterpret_cast(
+ b_tile + (b_n_row + kBlockN / 2) * kBStride + b_k16);
+ bdst0[0] = bsrc0[0];
+ bdst0[1] = bsrc0[1];
+ bdst1[0] = bsrc1[0];
+ bdst1[1] = bsrc1[1];
__syncthreads();
- // Each wave consumes its 128x32 quadrant: eight m-groups x two n-groups,
- // 16 m16n16k32 MMACs per kk. The A fragment loads and the MMAC sequence
- // keep the exact int32 accumulation order (k0-outer 64, kk-inner 32,
- // hardware k-order in m16n16k32) and the element-to-lane mapping of the
- // accepted kernels; only the B fragment fill (one 8-byte contiguous LDS
- // read per col_major fragment on the n-major tile) changes.
+ // Each wave consumes its 128x64 quadrant: eight m-groups x four n-groups,
+ // 32 m16n16k32 MMACs per kk, in the exact int32 accumulation order
+ // (k0-outer 64, kk-inner 32, hardware k-order inside m16n16k32). The B
+ // fragments are loaded once per kk and reused by all eight m-groups; the A
+ // fragments are consumed in two batches of four to bound their live range
+ // to eight VGPRs at a time.
#pragma unroll
for (int kk = 0; kk < kStageK; kk += kTileK) {
- load_fragment8(a_frag0, a_tile + local_row * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag1,
- a_tile + (local_row + kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag2,
- a_tile + (local_row + 2 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag3,
- a_tile + (local_row + 3 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(b_frag0, b_tile + local_col * kBStridePacked + kk,
- kBStridePacked, lane);
- load_fragment8(b_frag1,
- b_tile + (local_col + kTileN) * kBStridePacked + kk,
- kBStridePacked, 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);
+ load_fragment8(b_frag2, b_tile + (local_col + 2 * kTileN) * kBStride + kk,
+ kBStride, lane);
+ load_fragment8(b_frag3, b_tile + (local_col + 3 * kTileN) * kBStride + kk,
+ kBStride, lane);
+ load_fragment8(a_frag0, a_tile + local_row * kAStride + kk, kAStride,
+ lane);
+ load_fragment8(a_frag1, a_tile + (local_row + kTileM) * kAStride + kk,
+ kAStride, lane);
+ load_fragment8(a_frag2, a_tile + (local_row + 2 * kTileM) * kAStride + kk,
+ kAStride, lane);
+ load_fragment8(a_frag3, a_tile + (local_row + 3 * kTileM) * kAStride + kk,
+ kAStride, lane);
du_mma_sync(acc00, a_frag0, b_frag0, acc00);
du_mma_sync(acc01, a_frag0, b_frag1, acc01);
+ du_mma_sync(acc02, a_frag0, b_frag2, acc02);
+ du_mma_sync(acc03, a_frag0, b_frag3, acc03);
du_mma_sync(acc10, a_frag1, b_frag0, acc10);
du_mma_sync(acc11, a_frag1, b_frag1, acc11);
+ du_mma_sync(acc12, a_frag1, b_frag2, acc12);
+ du_mma_sync(acc13, a_frag1, b_frag3, acc13);
du_mma_sync(acc20, a_frag2, b_frag0, acc20);
du_mma_sync(acc21, a_frag2, b_frag1, acc21);
+ du_mma_sync(acc22, a_frag2, b_frag2, acc22);
+ du_mma_sync(acc23, a_frag2, b_frag3, acc23);
du_mma_sync(acc30, a_frag3, b_frag0, acc30);
du_mma_sync(acc31, a_frag3, b_frag1, acc31);
- load_fragment8(a_frag4,
- a_tile + (local_row + 4 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag5,
- a_tile + (local_row + 5 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag6,
- a_tile + (local_row + 6 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
- load_fragment8(a_frag7,
- a_tile + (local_row + 7 * kTileM) * kAStridePacked + kk,
- kAStridePacked, lane);
+ du_mma_sync(acc32, a_frag3, b_frag2, acc32);
+ du_mma_sync(acc33, a_frag3, b_frag3, acc33);
+ load_fragment8(a_frag4, a_tile + (local_row + 4 * kTileM) * kAStride + kk,
+ kAStride, lane);
+ load_fragment8(a_frag5, a_tile + (local_row + 5 * kTileM) * kAStride + kk,
+ kAStride, lane);
+ load_fragment8(a_frag6, a_tile + (local_row + 6 * kTileM) * kAStride + kk,
+ kAStride, lane);
+ load_fragment8(a_frag7, a_tile + (local_row + 7 * kTileM) * kAStride + kk,
+ kAStride, lane);
du_mma_sync(acc40, a_frag4, b_frag0, acc40);
du_mma_sync(acc41, a_frag4, b_frag1, acc41);
+ du_mma_sync(acc42, a_frag4, b_frag2, acc42);
+ du_mma_sync(acc43, a_frag4, b_frag3, acc43);
du_mma_sync(acc50, a_frag5, b_frag0, acc50);
du_mma_sync(acc51, a_frag5, b_frag1, acc51);
+ du_mma_sync(acc52, a_frag5, b_frag2, acc52);
+ du_mma_sync(acc53, a_frag5, b_frag3, acc53);
du_mma_sync(acc60, a_frag6, b_frag0, acc60);
du_mma_sync(acc61, a_frag6, b_frag1, acc61);
+ du_mma_sync(acc62, a_frag6, b_frag2, acc62);
+ du_mma_sync(acc63, a_frag6, b_frag3, acc63);
du_mma_sync(acc70, a_frag7, b_frag0, acc70);
du_mma_sync(acc71, a_frag7, b_frag1, acc71);
+ du_mma_sync(acc72, a_frag7, b_frag2, acc72);
+ du_mma_sync(acc73, a_frag7, b_frag3, acc73);
}
- // Last stage only: prefetch the epilogue's per-row x_scale and
- // per-column weight_scale values into registers (uniform branch: k0 is
- // block-uniform, so the barrier below is reached by every thread; the
- // row < m guard keeps the x_scale reads in bounds for any m tail).
+ // Last stage only: prefetch the epilogue's per-row x_scale and per-column
+ // weight_scale values into registers (uniform branch; the row < m guard
+ // keeps the x_scale reads in bounds for any m tail).
if (k0 + kStageK >= k) {
const int r = lane & 15;
const int c4 = lane >> 4;
@@ -1097,7 +538,12 @@ __global__ __launch_bounds__(kBlockThreads)
4 * c4);
ws1 = *reinterpret_cast(weight_scale + base_col +
kTileN + 4 * c4);
+ ws2 = *reinterpret_cast(weight_scale + base_col +
+ 2 * kTileN + 4 * c4);
+ ws3 = *reinterpret_cast(weight_scale + base_col +
+ 3 * kTileN + 4 * c4);
}
+
// Protect the LDS buffer from the next stage's cooperative overwrite.
__syncthreads();
}
@@ -1106,58 +552,105 @@ __global__ __launch_bounds__(kBlockThreads)
base_col, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc01, xs_m[0], ws1, out, base_row,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc02, xs_m[0], ws2, out, base_row,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc03, xs_m[0], ws3, out, base_row,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc10, xs_m[1], ws0, out,
base_row + kTileM, base_col, m, n,
lane);
store_prefill_fragment_coalesced_scaled(acc11, xs_m[1], ws1, out,
base_row + kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc12, xs_m[1], ws2, out,
+ base_row + kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc13, xs_m[1], ws3, out,
+ base_row + kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc20, xs_m[2], ws0, out,
base_row + 2 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc21, xs_m[2], ws1, out,
base_row + 2 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc22, xs_m[2], ws2, out,
+ base_row + 2 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc23, xs_m[2], ws3, out,
+ base_row + 2 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc30, xs_m[3], ws0, out,
base_row + 3 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc31, xs_m[3], ws1, out,
base_row + 3 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc32, xs_m[3], ws2, out,
+ base_row + 3 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc33, xs_m[3], ws3, out,
+ base_row + 3 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc40, xs_m[4], ws0, out,
base_row + 4 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc41, xs_m[4], ws1, out,
base_row + 4 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc42, xs_m[4], ws2, out,
+ base_row + 4 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc43, xs_m[4], ws3, out,
+ base_row + 4 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc50, xs_m[5], ws0, out,
base_row + 5 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc51, xs_m[5], ws1, out,
base_row + 5 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc52, xs_m[5], ws2, out,
+ base_row + 5 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc53, xs_m[5], ws3, out,
+ base_row + 5 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc60, xs_m[6], ws0, out,
base_row + 6 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc61, xs_m[6], ws1, out,
base_row + 6 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc62, xs_m[6], ws2, out,
+ base_row + 6 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc63, xs_m[6], ws3, out,
+ base_row + 6 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
store_prefill_fragment_coalesced_scaled(acc70, xs_m[7], ws0, out,
base_row + 7 * kTileM, base_col, m,
n, lane);
store_prefill_fragment_coalesced_scaled(acc71, xs_m[7], ws1, out,
base_row + 7 * kTileM,
base_col + kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc72, xs_m[7], ws2, out,
+ base_row + 7 * kTileM,
+ base_col + 2 * kTileN, m, n, lane);
+ store_prefill_fragment_coalesced_scaled(acc73, xs_m[7], ws3, out,
+ base_row + 7 * kTileM,
+ base_col + 3 * kTileN, m, n, lane);
}
// ---------------------------------------------------------------------------
-// Generic scalar fallback: one thread per output element. Exact int32 dot
-// over K, then fused x_scale * weight_scale, then bf16 store. Handles every
-// unmatched (m, n, k), including all small-M API cases (M=2, M=16) and any
-// M with the same (K, N). For the exact (k, n) == (2048, 6144) the weight is
-// the n-major packed [N, K] layout (produced once out of the timed region),
-// which the fallback decodes; every other (k, n) keeps the identity [K, N]
-// layout.
+// Generic scalar fallback: one thread per output element. Exact int32 dot
+// over K in ascending k order, then fused x_scale * weight_scale, then bf16
+// store. Handles every unmatched (m, n, k), including all small-M API cases
+// (M=2, M=16) and any M with the same (K, N). For the exact
+// (k, n) == (2048, 6144) the packed_weight buffer holds the n-major transpose
+// packed[col * K + kk] = raw[kk * N + col] (produced once outside the timed
+// region), which this fallback decodes; every other (k, n) keeps the identity
+// [K, N] row-major layout.
// ---------------------------------------------------------------------------
__global__ __launch_bounds__(256) void w8a8_gemm_scalar_fallback_kernel(
const int8_t* __restrict__ x_q,
@@ -1178,12 +671,6 @@ __global__ __launch_bounds__(256) void w8a8_gemm_scalar_fallback_kernel(
const int col = static_cast(linear - static_cast(row) * n);
int32_t acc = 0;
const int8_t* a_row = x_q + static_cast(row) * k;
- // Iteration 5: for the exact (k, n) == (2048, 6144) the packed_weight
- // buffer holds the n-major transpose packed[col*k + kk] = raw[kk*n + col]
- // (produced once out of the timed region), so the fallback decodes that
- // layout there (keeps the paired M=2/M=16 and M=3072 API shapes with the
- // same (K, N) byte-exact); every other (k, n) keeps the identity [K, N]
- // row-major copy.
const bool packed_nmajor = (k == kTargetK && n == kTargetN);
const int8_t* b_ptr =
packed_nmajor ? weight + static_cast(col) * k : weight + col;
@@ -1198,16 +685,13 @@ __global__ __launch_bounds__(256) void w8a8_gemm_scalar_fallback_kernel(
}
// ---------------------------------------------------------------------------
-// Identity packing (bootstrap; outside the timed region and outside Graph
-// capture; the packed buffer keeps the same byte count K*N and the same
-// allocated address, so the layout is graph-stable). launch_pack_w8a8_weight
-// is an identity device-to-device copy for every (k, n): the packed weight
-// is byte-identical to the raw [K, N] weight and the packed scale is
-// byte-identical to the raw [N] scale. Later Parallel explore rounds may
-// replace these kernels with a real layout transform and must then update
-// the matching GEMM interpretation; the identity D2D copy is the generic
-// fallback for unmatched (k, n).
+// Weight packing, executed once by the trusted binding before Graph capture
+// and never inside launch_w8a8_gemm. The packed buffer keeps the same byte
+// count as the raw [K, N] weight, so captured addresses are graph-stable.
// ---------------------------------------------------------------------------
+
+// Generic identity device-to-device copy for every (k, n) that has no
+// specialized packing.
__global__ __launch_bounds__(256) void w8a8_pack_identity_kernel(
const int8_t* __restrict__ raw,
int8_t* __restrict__ packed,
@@ -1219,14 +703,8 @@ __global__ __launch_bounds__(256) void w8a8_pack_identity_kernel(
}
}
-// Exact-shape n-major pack (iteration 5): for (k, n) == (2048, 6144) the
-// weight is transposed once, outside the timed region and out of Graph
-// capture, into packed[col * k + kk] = raw[kk * n + col] so every col_major
-// B fragment is 8 contiguous K bytes in the packed buffer (and in the
-// staged n-major LDS tile). Same byte count and same buffer as the
-// identity pack, so captured addresses are unchanged; correctness is
-// checked against the raw logical [K, N] weight. One thread per output
-// byte; runs once during weight prep.
+// Exact-shape n-major transpose for (k, n) == (2048, 6144):
+// packed[col * k + kk] = raw[kk * n + col]. One thread per output byte.
__global__ __launch_bounds__(256) void w8a8_pack_nmajor_i8_kernel(
const int8_t* __restrict__ raw,
int8_t* __restrict__ packed,
@@ -1244,6 +722,7 @@ __global__ __launch_bounds__(256) void w8a8_pack_nmajor_i8_kernel(
}
}
+// The [N] weight scale needs no layout change.
__global__ __launch_bounds__(256) void w8a8_pack_scale_identity_kernel(
const float* __restrict__ raw,
float* __restrict__ packed,
@@ -1278,19 +757,17 @@ extern "C" void launch_w8a8_gemm(
(void)workspace_bytes;
auto* out_bf16 = reinterpret_cast(out);
- // Explicit dispatch. The assigned shape (M=4096, N=6144, K=2048) takes the
- // iteration-5 packed-B DUMMA 256x64 path (w8a8_dumma_256x64x64_packedb_kernel
- // consumes the n-major packed [N, K] weight produced by
- // launch_pack_w8a8_weight; the identity-layout kernels stay compiled for
- // re-dispatch only together with an identity pack revert); every other
+ // Explicit exact-shape dispatch. The assigned shape (M=4096, N=6144,
+ // K=2048) takes the 256x128 packed-B DUMMA path, which consumes the n-major
+ // packed [N, K] weight produced by launch_pack_w8a8_weight. Every other
// (m, n, k) - including small-M API cases (M=2, M=16) and M=3072 with the
- // same (K, N) - takes the scalar fallback, which decodes the n-major
- // packed layout when (k, n) == (2048, 6144) and the identity [K, N]
- // layout otherwise.
+ // same (K, N) - takes the scalar fallback, which decodes the n-major packed
+ // layout when (k, n) == (2048, 6144) and the identity [K, N] layout
+ // otherwise.
if (m == kTargetM && n == kTargetN && k == kTargetK) {
- const dim3 grid(kTargetN / kBlockN, kTargetM / kBlockM256);
+ const dim3 grid(kTargetN / kBlockN, kTargetM / kBlockM);
const dim3 block(kBlockThreads);
- hipLaunchKernelGGL(w8a8_dumma_256x64x64_packedb_kernel,
+ hipLaunchKernelGGL(w8a8_dumma_256x128x64_packedb_kernel,
grid, block, 0, stream,
a, b, x_scale, weight_scale, out_bf16, m, n, k);
} else {
@@ -1313,13 +790,12 @@ extern "C" void launch_pack_w8a8_weight(
int k,
int n,
hipStream_t stream) {
- // Iteration 5: the exact (k, n) == (2048, 6144) weight is packed once,
- // outside the timed region and out of Graph capture, into the n-major
- // [N, K] layout packed[col*k + kk] = raw[kk*n + col] (same byte count and
- // same buffer -> graph-stable addresses) so every col_major B fragment is
- // 8 contiguous K bytes; every other (k, n) keeps the identity
- // device-to-device copy (the scalar fallback decodes the n-major layout
- // when (k, n) matches and the identity layout otherwise).
+ // For the exact (k, n) == (2048, 6144) the logical [K, N] weight is
+ // transposed once, outside the timed region and outside Graph capture, into
+ // the n-major [N, K] layout packed[col * k + kk] = raw[kk * n + col] (same
+ // byte count and same caller-owned buffer -> graph-stable addresses), so
+ // every col_major B fragment is 8 contiguous K bytes; every other (k, n)
+ // keeps the identity device-to-device copy.
constexpr int kBlock = 256;
const dim3 block(kBlock);
diff --git a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/shared_down_proj.hip b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/shared_down_proj.hip
index a1b49be9..e07786c9 100644
--- a/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/shared_down_proj.hip
+++ b/metainfer/tasks/dcu_kernel_auto_opt/variant/int8w8a8-gemm/glm52/TP8/M4096/shared_down_proj.hip
@@ -1,313 +1,266 @@
-// @@variant shape=glm_tp8_shared_down_proj_m4096 commit=9b56135e1fec933006c477066840f9c84f497cef added=2026-08-31
-// median_us=278.4 p90_us=279.2 speedup=63.57 baseline_us=1.77e+04
-// source=glm5-2-dsh-tp8-m4096-1-e6a280a2
+// @@variant shape=glm_tp8_shared_down_proj_m4096 commit=44018b50c8c0557effd9aa58807ccc6a69dead1f added=2026-09-14
+// median_us=146.7 p90_us=147.2 speedup=120.7 baseline_us=1.77e+04
+// source=test-9-11-1/iteration_002/glm_tp8_shared_down_proj_m4096 (retested A/B)
// MetaInfer W8A8 INT8 GEMM backend for Hygon gfx928 (K500SM_AI).
-// worker_3 (iteration 3), assigned shapes:
-// glm_tp8_shared_gate_up_proj_m4096 : (M, N, K) = (4096, 512, 6144)
-// glm_tp8_shared_down_proj_m4096 : (M, N, K) = (4096, 6144, 256)
//
-// This file provides the two stable host launch symbols consumed by
-// csrc/bindings.cpp:
-// * launch_w8a8_gemm(...) - timed, graph-safe GEMM on the caller's
-// HIP stream; no allocation, packing,
-// autotuning, or synchronization. Only the
-// caller-provided out and workspace are
-// touched (workspace is unused this round).
-// * launch_pack_w8a8_weight(...) - out-of-timed-region packing. Identity
-// device-to-device copy for every (k, n)
-// pair EXCEPT the assigned gate_up weight
-// (k, n) == (6144, 512), which is packed
-// once as a [K, N] -> [N, K] transpose
-// (n-major rows, k contiguous) so the timed
-// GEMM consumes B fragments as
-// k-contiguous bytes. The pack stays
-// out-of-timed, out-of-Graph, same byte
-// count, same graph-stable buffer.
-//
-// Bootstrap strategy (correctness-first, usable profiling baseline):
-// * Native INT8 DUMMA m16n16k32 tiled kernel with int32 accumulation for
-// every assigned shape with M >= 128 (both assigned shapes qualify).
-// * 2-D macro-tile family (64x64, 64x128, 128x64) with one 32x32 quadrant
-// per wavefront (four m16n16k32 int32 accumulators kept resident over the
-// whole K loop). Single-buffered K loop (K stage 128, two barriers per
-// stage): no split-K, no raw asm, no speculative double-buffer pipeline.
+// worker_0 (iteration 2, operand-reuse round). Assigned shape:
+// glm_tp8_shared_down_proj_m4096 : (M, N, K) = (4096, 6144, 256)
//
-// Round 1 (baseline): the 64x128 tile measured 517.4 us median / 49.81 TOPS;
-// routing the exact gate_up shape to the 64x64 tile measured 432.5 us
-// (accepted; 19.6% improvement).
+// This file provides the two stable host launch symbols consumed by the
+// trusted csrc/bindings.cpp:
+// * launch_w8a8_gemm(...) - timed, graph-safe GEMM on the caller's
+// HIP stream. No allocation, no packing, no
+// autotuning, no host/device synchronization
+// and no default-stream launch. Only the
+// caller-provided out is written; the
+// caller-provided workspace is unused
+// because there is no split-K.
+// * launch_pack_w8a8_weight(...) - out-of-timed-region packing. The exact
+// down_proj (k, n) == (256, 6144) pair is
+// packed once into the round-5 panel layout
+// [N/128][K/64][128][64] (n-major rows of
+// 64 k-bytes per panel) so the timed GEMM
+// consumes each 128x64 B stage tile as ONE
+// contiguous 8 KiB region and each lane's
+// fragment stays k-contiguous; every other
+// (k, n) stays the identity device-to-device
+// copy of the logical [K, N] weight.
//
-// Round 2 (operand-reuse, mandated decision): cooperative A+B LDS staging is
-// kept and direct per-wave global loads are rejected (A is reused 64x per
-// macro-tile via kBlockN and B 64x via kBlockM; direct loads would refetch
-// both operands once per wavefront, 4x the 28.3 MB staged global traffic ->
-// ~1.73M vmem_read instructions vs 432K, and the TP4 gate_up direct-A arm
-// measured 7.9x slower). The focused change makes the staged B layout
-// bank-safe for the fragment consumers: the gate_up weight is packed once
-// [K,N] -> [N,K] (n-major, k contiguous) and B fragments are loaded
-// col_major from a 144-B (128+16) row-strided LDS tile, replacing the
-// byte-scattered row_major B reads on the identity [K,N] layout that are the
-// 17.3M PMC bank-conflict source. Global loads stay 16-byte vectorized and
-// coalesced on both operands. LDS stays 19,456 B/block -> 3 blocks/CU
-// (58,368 B, unchanged; round-1's "4 blocks/CU" estimate was impossible on
-// both LDS (4x19,456 = 77,824 B > 64 KiB) and VGPR (4x256x72 = 73,728 >
-// 65,536)). down_proj keeps the identity 64x128 arm; every unmatched shape
-// keeps the generic arms and the scalar fallback (which decodes the packed
-// [N,K] layout when (k, n) == (6144, 512) so the paired M=2/M=16 gate_up
-// tails stay correct).
-// * A[kBlockM, kStageK] and B staged in bank-skewed LDS (row strides
-// 144/80/144 = 16-byte-aligned non-power-of-two), consumed with the
-// library du_load_matrix_sync loaders: A row_major (unchanged, load-
-// bearing on every shape), B row_major for the identity [K, N] layout
-// (down_proj + unmatched shapes) and B col_major for the gate_up
-// [N, K] packed layout (round 2).
+// Round 1 (bootstrap, accepted at 313.66 us median): native INT8 DUMMA
+// m16n16k32 tiled kernel with exact int32 accumulation for every M >= 128.
+// * 64x128 macro-tile: 8 wavefronts, one 32x32 quadrant per wavefront, four
+// m16n16k32 int32 accumulators kept resident across the whole K loop. A and
+// B are staged cooperatively into bank-skewed LDS and consumed with the
+// library du_load_matrix_sync fragment loaders.
+// * SINGLE-buffered K loop, K stage 128 (A[64,128] + B[128,128] = 27,648
+// B/block -> 2 blocks/CU): two barriers per stage, no split-K, no raw asm,
+// no speculative double-buffer pipeline. K=256 is two stages.
// * Fused direct fragment -> x_scale -> weight_scale -> bf16 epilogue using
// the verified gfx928 accumulator ownership (row = lane & 15,
-// col_mod4 = lane >> 4, x[i] -> columns col_mod4 + 4*i). The int32 dot is
-// exact for every assigned shape (max |dot| = 6144*127*127 << 2^31), the
-// int32 -> float conversion is exact, and the single fp32 multiply
-// followed by one bfloat16 rounding matches the reference bit-for-bit.
-// * Tail M rows are zero-filled on load and masked on store, so any
-// M >= 128 is supported by the tiled path.
-// * All other (m, n, k) - including every M < 128 API case and the paired
-// M=2/M=16 shapes with the same (N, K) - go to a scalar int8/int32
-// fallback that decodes the [N, K] packed layout when (k, n) ==
-// (6144, 512) and the identity [K, N] layout otherwise.
-//
-// Round 3 (pipeline, mandated decision): compare single and double buffering
-// across K tiles. The exact (4096, 512, 6144) guard now routes to
-// w8a8_dumma_prefill_tiled_kernel<64, 64, 64, true, true>: a software-
-// pipelined 64-K double-buffered K loop (prologue + one barrier per stage,
-// 97 barriers for 96 stages) that issues stage s+1's global loads into the
-// idle LDS buffer before the stage-s MMAC burst, overlapping the
-// global-load latency that the round-2 single buffer serializes in front of
-// every stage (ISA-verified order: global_load_dwordx4 -> ds_write_b128 ->
-// s_waitcnt -> s_barrier -> ds_read2 x16 -> v_mmac x16 -> s_barrier, with
-// the next stage's loads gated behind the previous stage's barrier). The
-// 128-K double-buffered variant is arithmetically impossible at 3 blocks/CU
-// (2 x 19,456 = 38,912 B/block x 3 = 116,736 > 64 KiB -> 1 block/CU, a
-// harmful occupancy loss), so the pipelined comparison uses the 64-K
-// double buffer: LDS 20,480 B/block x 3 = 61,440 <= 64 KiB and
-// VGPR 64-72 x 256 x 3 <= 65,536, i.e. 3 blocks/CU preserved (no occupancy
-// confound). Everything else is byte-identical: same int4 staging pattern
-// (same 28.3 MB global traffic and ~432K vmem_read instructions), same
-// k0-outer/kk-inner int32 order (stage size changes only regroup exact
-// integer sums, so results stay bit-identical), same col_major B fragments
-// (stride 80 = 64+16 bank skew), same fused epilogue, same pack, same
-// generic single-buffered arms and scalar fallback. Decision rule (mandate):
-// retain double buffering only if PMC shows reduced VMEM/LDS stall evidence
-// (lds_wait_instructions drops below ~1.2M, lds_instructions stays ~1.97M,
-// occupancy stays 3 blocks/CU) AND the median improves >2%; otherwise the
-// round's evidence retains the single-buffered round-2 structure.
-//
-// Round 10 (fragment-load-form round, restarted from the accepted round-3
-// source after the round-9 infra kill): replace the library
-// du_load_matrix_sync fragment loads of the exact-shape gate_up kernel with
-// the lineage-validated explicit load_frag8 8-byte ds_read_b64 loader. The
-// exact (4096, 512, 6144) guard now routes to
-// w8a8_dumma_prefill_tiled_kernel<64, 64, 64, true, true, true> (new
-// kDirectFrag template arm). The round-3 ISA shows every fragment load as
-// ds_read2_b32 + per-byte v_and/v_or_b32_sdwa/v_or3 mask-OR "identity"
-// reassembly (~9.4M of the 15.43M VALU instructions per replay); loading the
-// same 8 bytes per lane as one int64 (ds_read_b64) removes that VALU and
-// halves the fragment bank-conflict factor (32-lane dword phases at stride
-// 80 -> 16-lane 8-byte phases). Element-to-slot fragment mapping, v_mmac
-// operands, exact k0-outer/kk-inner int32 accumulation, staging/global
-// traffic, barriers, tile, grid, LDS bytes, occupancy (3 blocks/CU) and the
-// fused epilogue are unchanged; results stay bit-identical. The generic arms
-// (<64,128,128>, <128,64,128>, <64,64,128>, all kBNMajor=false) keep the
-// library loaders and byte-identical codegen/resources (if constexpr).
-//
-// Round 11 (LDS bank-skew round, HIP-only; ISA policy keeps raw asm
-// disallowed because the HIP plateau rule - 8 valid HIP rounds within +/-2%
-// of the then-current best - is unmet). The accepted round-10 kernel still
-// shows PMC lds_bank_conflicts = 9,437,184, EXACTLY unchanged from the
-// round-3 code object (the -9.07% round-10 win came from VALU removal, not
-// from conflicts). 9,437,184 = 192 x 512 blocks x 96 stages: every
-// block-stage pays 192 conflict cycles = 128 from the 16 ds_read2_b64
-// fragment reads + 64 from the 8 ds_write_b128 staging stores (the 16-B/lane
-// store floor). The read half is a stride defect: at stride 80 = 20 dwords,
-// 20*r mod 32 has period 8, so every 16-lane read phase groups rows r and
-// r+8 onto the same bank pair - a 2-way conflict on ALL 4 sub-phases of
-// every fragment read. The exact (4096, 512, 6144) guard now routes to
-// w8a8_dumma_prefill_tiled_kernel<64, 64, 64, true, true, true, true> (new
-// kSkew8 template arm) which stages both operands at the down_proj-validated
-// 8-mod-16 row strides A = 88 (64+24) and B = 72 (64+8) instead of 80:
-// 22*r and 18*r mod 32 are distinct over r = 0..15 and gcd(22,32) =
-// gcd(18,32) = 2 does not divide 1, so every 16-lane read phase lands on 16
-// DISTINCT bank pairs - the fragment reads become conflict-free (128 of the
-// 192 cycles removed; only the 2-way 16-B/lane staging-store floor of 64
-// cycles remains -> lds_bank_conflicts expected ~3.15M). The 8-mod-16 row
-// starts forbid 16-byte staging stores, so each 16-byte int4 is staged as
-// two 8-byte halves (ds_write2_b64), the validated down_proj staging form
-// (its B tile is exactly stride 72; its A tile exactly stride 88). Global
-// loads (int4, same 432,128 vmem_read), pack layout, barrier structure
-// (97/block), tile/grid (512 blocks), LDS bytes (A[64,88] + B[64,72] x 2 =
-// 20,480 B/block -> 3 blocks/CU unchanged), VGPR, the k0-outer/kk-inner
-// int32 order, the element-to-slot fragment mapping and the fused epilogue
-// are untouched; only the LDS row strides and the staging store granularity
-// change, so results stay bit-identical. Generic arms (<64,128,128>,
-// <128,64,128>, <64,64,128>) keep kSkew8 = false and byte-identical
-// codegen/resources (if constexpr).
+// col_mod4 = lane >> 4, frag.x[i] -> columns col_mod4 + 4*i). The int32 dot
+// is exact for every allowed K and the fp32 left-associated multiply chain
+// matches the mathematical reference.
+// * A second 64x64 arm covers large-M shapes whose N is a multiple of 64 but
+// not of 128. Every other (m, n, k) - including every M < 128 API case and
+// the paired M=2/M=16 shapes with the same (N, K) - goes to the scalar
+// int8/int32 fallback.
//
-// Round 16 (lead-2 loop-carried payload round, HIP-only; raw asm stays
-// DISALLOWED - plateau=false, the recent-valid-improvement window
-// [-34.44%, +2.82%, +9.97%] does not meet three valid HIP rounds within
-// [-2%, +2%) of the then-current best, and no prior ISA-guided round
-// recorded a compiler limitation with target instructions). The accepted
-// round-11 exact code object (digest 9889820e, median 293.30 us / p90
-// 293.545) shows the software-pipeline intent is DEFEATED by the compiler's
-// wait placement: per stage the steady-state loop is
-// global_load_dwordx4 (A prefetch s+1) -> global_load_dwordx4 (B prefetch
-// s+1) -> s_waitcnt vmcnt(1) -> s_waitcnt vmcnt(0) <- waits IMMEDIATELY
-// -> 2 x ds_write2_b64 (staging stores) -> 4 x ds_read2_b64 (fragment
-// reads, stage s) -> 8 x v_mmac CONTIGUOUS (stage s) -> s_barrier -> loop.
-// i.e. the stage-(s+1) global loads are waited ~12-16 instructions after
-// issue, BEFORE the stage-s MMAC burst, so the full global-load latency
-// (~600-1000+ cycles under 120-CU L2 contention) sits on the per-stage
-// critical path with zero overlap (the round-12 analysis'
-// "prefetch global load -> ds_write2_b64 -> s_barrier -> ds_read2_b64 ->
-// v_mmac" chain; PMC: 6.0M VALU + 1.18M LDS + 465K VMEM wave-instructions
-// over 293 us x 120 CUs, i.e. ~95-99% issue-stalled). Iteration history
-// rules out the other levers on this exact shape: tile aspect (iter 4
-// 32x64: 749.9 us; iter 8 64x128 w8: 395.4 us), 2-wave quadrants (iter 6:
-// 629.0 us - wave count per CU is load-bearing), LDS-side fragment
-// prefetch (iter 7: 337.97 us, flat/worse), stage 32 / 6 blocks per CU
-// (iter 12: 447.4 us - per-stage fixed cost, not occupancy, is the lever).
-// The one untested mechanism that directly removes the exposed vmcnt wait:
-// a LEAD-2 LOOP-CARRIED REGISTER PAYLOAD (the TP4 o_proj lineage's
-// validated "stage-top publish" form). The exact (4096, 512, 6144) guard
-// now routes to w8a8_dumma_prefill_tiled_kernel<64,64,64,true,true,true,
-// true,true> (new kPayload template arm): each thread keeps the one A int4
-// + one B int4 of stage s+2 in VGPR across the stage-(s+1) barrier, issues
-// those global loads at the TOP of iteration s (no wait), and PUBLISHES
-// the payload carried from iteration s-1 into the idle LDS buffer at the
-// TOP of iteration s (before the stage-s MMAC burst). The first use of the
-// payload registers is the next iteration's publish, and the s_barrier
-// fences outstanding vmem, so the compiler can no longer place
-// s_waitcnt vmcnt(0) immediately after the loads: the wait lands after the
-// stage-s fragment reads + MMAC burst (immediately before the barrier), and
-// the global-load latency overlaps the whole compute segment instead of
-// serializing in front of it. One barrier per 64-K stage is preserved
-// (prologue + 96 = 97/block); the publish targets the idle buffer (WAR
-// retired by the previous barrier, RAW retired by the stage barrier), so no
-// second barrier is introduced. Everything else is byte-identical: 64x64
-// tile / 256 threads / 4 waves / 32x32 quadrant per wave, grid (8,64) =
-// 512 blocks, kSkew8 strides 88/72 with ds_write2_b64 staging (8 per
-// block-stage), same 16 ds_read2_b64 fragment reads per block-stage (same
-// lds_instructions ~1,179,648), same 20,480 B LDS/block -> 3 blocks/CU
-// (payload costs 8 VGPRs: ~54 total, 54 x 256 x 3 = 41,472 <= 65,536; no
-// spill risk at 3 blocks/CU), same int4 global loads (same ~432,128
-// vmem_read / 32,768 vmem_write), same k0-outer/kk-inner int32 order and
-// the same element-to-slot fragment mapping - the payload changes only WHEN
-// global bytes land in LDS, never which bytes or the accumulation order, so
-// results stay bit-identical (mismatch 0 / max_abs_error 0.0 expected).
-// Generic arms (<64,128,128>, <128,64,128>, <64,64,128>, all kBNMajor=
-// false) keep kPayload = false and byte-identical codegen/resources
-// (if constexpr).
+// Round 2 (this session, mandated operand-reuse decision): compare direct
+// per-wave global loads with cooperative A+B LDS staging on the exact
+// (4096, 6144, 256) shape.
+// * Reuse accounting for the 64x128 macro-tile with one 32x32 quadrant per
+// wave: every A element participates in kBlockN = 128 output MACs and every
+// B element in kBlockM = 64 MACs, so cooperative staging keeps ONE global
+// read and ONE LDS write per element while direct per-wave fragment loads
+// would re-read the A tile kBlockN/32 = 4 times and the B tile
+// kBlockM/32 = 2 times from global (per 64-K stage: 12,288 B staged vs
+// 32,768 B of fragment reads, a 2.67x global-traffic amplification on an
+// already L2-resident operand stream). Direct loads are rejected on that
+// amplification and on the sibling lineage's measured 7.9x direct-arm
+// regressions.
+// * The bootstrap exact-shape arm reads the identity [K, N] B layout with
+// byte-scattered row_major library fragment loads: its exact code object
+// and PMC show 3,686,400 LDS instructions, 7,864,320 LDS bank conflicts,
+// 5,131,888 LDS waits and 17.4M VALU per operator replay at 313.66 us
+// median (11% issue utilization).
+// * Focused change: the exact (m, n, k) == (4096, 6144, 256) guard now routes
+// to w8a8_dumma_proj_operand_reuse_kernel<64, 128, 64>, which keeps
+// cooperative A+B LDS staging, keeps 16-byte vectorized coalesced global
+// staging loads on both operands, and fixes the LDS consumer side:
+// - the one-time pack transposes B to the n-major [N, K] layout, so the
+// staged B tile is [kBlockN, kStageK] with k contiguous and each
+// matrix_b col_major fragment lane owns 8 contiguous bytes (the
+// identity-layout row_major B reads are the bank-conflict source);
+// round 5 keeps the identical staged LDS bytes but re-orders the
+// packed buffer into the [N/128][K/64][128][64] panel layout so each
+// block-stage B tile is one contiguous global region (see below);
+// - A and B fragments are read with the explicit 8-byte load_frag8
+// loader (ONE 8-byte LDS read per fragment) instead of the library
+// per-byte loaders plus reassembly VALU; the element-to-slot mapping
+// is identical, so the MMAC operands and accumulation are unchanged;
+// - LDS row strides are the bank-safe 8-mod-16 values A = 88 (64+24) and
+// B = 72 (64+8): 22*r and 18*r mod 32 are 16 distinct values over
+// r = 0..15, so every 16-lane ds_read_b64 fragment phase lands on 16
+// distinct bank pairs. The 8-mod-16 row starts forbid 16-byte stores,
+// so each staged int4 is written as two 8-byte halves;
+// - 64-K STAGE single buffer: 2 barriers/stage (8 per block for K=256),
+// LDS A[64,88] + B[128,72] = 14,848 B/block -> 2 blocks/CU (occupancy
+// unchanged from the bootstrap's 27,648 B / 2 blocks/CU), 64x128 tile,
+// 512 threads / 8 waves, grid (48, 64) = 3072 blocks.
+// The k0-outer / kk-inner int32 accumulation order, the element-to-slot
+// fragment mapping and the fused epilogue are unchanged, so results stay
+// bit-identical (mismatch 0 / max_abs_error 0.0 expected).
+// * Expected measured effect: lds_instructions 3,686,400 -> ~540,672,
+// lds_bank_conflicts 7,864,320 -> ~1,179,648, lds_wait 5,131,888 ->
+// ~162,664, VALU 17.4M -> ~6M, median 313.66 us -> ~275-290 us.
+// * Every other shape keeps the bootstrap arms. The scalar fallback decodes
+// the packed [N, K] layout for (k, n) == (256, 6144) so the paired M=2/M=16
+// API shapes stay correct, and the packed pair never reaches an
+// identity-layout tiled arm.
//
-// Round 2 (this session, glm_tp8_shared_down_proj_m4096 exact-shape arm):
-// the generic <64,128,128> identity-layout arm that served down_proj
-// (4096, 6144, 256) measured 314.72 us median with PMC
-// lds_bank_conflicts = 7,864,320, ~600 LDS instructions and heavy per-byte
-// fragment-reassembly VALU per block-stage (library row_major loaders on
-// the identity [K, N] B layout). The exact (4096, 6144, 256) guard now
-// routes to w8a8_dumma_prefill_tiled_kernel<64, 128, 64, true, true, true,
-// true> - the file's own lineage-validated gate_up machinery transplanted
-// to down_proj: B packed once [K,N] -> [N,K] n-major (kBNMajor, same
-// transpose kernel as gate_up, same byte count / graph-stable buffer,
-// out-of-timed/out-of-Graph), 64-K double-buffered K loop (kDoubleBuffer,
-// prologue + 4 stages for K=256, one barrier per stage), explicit 8-byte
-// load_frag8 fragment reads (kDirectFrag, one ds_read_b64 per fragment,
-// zero reassembly VALU) and the 8-mod-16 bank-skewed LDS strides A=88 /
-// B=72 (kSkew8, down_proj-validated, conflict-free 16-lane fragment
-// phases). LDS A[64,88] + B[128,72] x 2 = 29,696 B/block -> 2 blocks/CU
-// (occupancy unchanged vs the 27,648-B single buffer); grid (48, 64) =
-// 3072 blocks; the stage size 128->64 only regroups exact integer sums
-// (k0-outer / kk-inner int32 order over k = 0,32,64,96,128,160,192,224 is
-// unchanged) and load_frag8 fills the same x[0..7] slots with the same
-// little-endian bytes, so results stay bit-identical (mismatch 0 /
-// max_abs_error 0.0 expected).
+// Round 5 (this session, mandated PACKING round): test one weight
+// packing/swizzle that makes each DUMMA B tile vector-loadable and
+// LDS-bank-safe, with the pack outside timing and the graph-stable packed
+// layout validated.
+// * The accepted round-2 arm stages its 128x64 B tile from the plain
+// n-major [N, K] packed weight: per block-stage the 512 int4 loads are 16
+// n-rows x 4 int4 per 64-lane warp at a 256-B row stride, i.e. every warp
+// instruction touches 16 DISTINCT 128-B lines with only 64 B used in each
+// (the same line is completed by another stage of the same block). Still
+// 16-byte vector loads, but scattered at line granularity.
+// * Focused change: the one-time pack for (k, n) == (256, 6144) ONLY becomes
+// the panel layout [N/128][K/64][128][64] (n-major 64-k-byte rows per
+// panel, panels ordered n-tile outer / k-tile inner). Byte-exact index
+// relation: packed[((nt*(k/64)+kt)*128+nn)*64+kk64] ==
+// raw[(nt*128+nn)*k + kt*64+kk64]. Each block's per-stage B tile
+// (n0 == 128*blockIdx.x, k0 == 64*stage) is now ONE contiguous 8 KiB
+// region, so a 64-lane warp instruction reads 1 KiB contiguous (full
+// 128-B line utilization) instead of 16 half-used lines; the per-thread
+// address stays one base + one local_col*kStageK + v*16 offset.
+// * Everything else is byte-identical: same 64-K single-buffered stage with
+// two barriers, same 64x128 tile / 512 threads / 8 waves / 32x32 quadrant
+// per wave, same grid (48, 64) = 3072 blocks, same LDS A[64,88] +
+// B[128,72] = 14,848 B/block -> 2 blocks/CU, same ds_write2_b64 8-byte-half
+// staging stores and same conflict-free load_frag8 ds_read_b64 fragment
+// reads (the LDS-bank-safe strides 88/72 and the staged LDS bytes are
+// untouched - only the GLOBAL byte order of the packed buffer changes),
+// same k0-outer/kk-inner int32 order, same fused epilogue -> mismatch 0 /
+// max_abs_error 0.0 expected. The scalar fallback decodes the panel layout
+// for (k, n) == (256, 6144) so the paired M=2/M=16 tails stay correct;
+// every other (k, n) keeps the identity copy. Packing runs
+// out-of-timed/out-of-Graph in launch_pack_w8a8_weight, same byte count
+// (256*6144), same grid (n*k/16 threads), same graph-stable buffer.
+// * Falsifiable gate: median_us < 273.2886219024658 AND p90_us <
+// 275.1925849914551 (and mismatch 0 / graph capture true) -> the
+// line-scattered B staging was binding and the panel pack wins; otherwise
+// the plain n-major pack is retained and this round records that the B
+// stream is L2-resident (90.4% hit, same 32 sectors/warp either way) and
+// NOT binding on the latency-stalled critical path.
//
-// Round 3 (this session, glm_tp8_shared_down_proj_m4096 pipeline round):
-// mandated single vs double buffering comparison across K tiles. The
-// accepted round-2 arm (<64,128,64,true,true,true,true>, 286.53 us median /
-// 286.64 us p90, digest cd9791d0) is the 64-K DOUBLE-buffered pipeline; its
-// exact code object shows the software pipeline already works on this shape
-// (steady-state loop: ds_write2_b64 publish s+1 -> global_load_dwordx4 x2
-// for s+2 -> ds_read2_b64 x4 (fragment reads s+1) -> v_mmac x8 -> s_waitcnt
-// vmcnt(0) lgkmcnt(0) -> s_barrier: the stage-(s+2) loads stay in flight
-// across the fragment reads + MMAC burst and are waited only after it, i.e.
-// reduced VMEM stalls), and PMC shows no harmful LDS or occupancy growth vs
-// the pre-round-2 single buffer (29,696 vs 27,648 B/block, 2 blocks/CU
-// pinned; lds_wait 162,664 / lds_instructions 540,672 / conflicts
-// 1,179,648). The (4096, 6144, 256) guard now routes to the SINGLE-buffered
-// form <64,128,64,true,false,true,true>: identical stage size 64, identical
-// kBNMajor/kDirectFrag/kSkew8/tile/grid/epilogue, kDoubleBuffer = false, so
-// each stage serializes its global-load round trip (load -> wait ->
-// ds_write2_b64) in front of the MMAC burst with two barriers per stage
-// (8/block vs 5/block) and LDS halves to 14,848 B/block. The single-buffer
-// branch now threads kSkew8 and kDirectFrag through its staging/MMAC calls
-// (previously hardcoded to the library-loader defaults; generic arms keep
-// byte-identical codegen via if constexpr). Falsifiable gate: if the
-// single-buffered median is NOT strictly below 286.53368949890137 (p90
-// guard < 286.6385078430176), double buffering is retained with the ISA/PMC
-// evidence above; if it IS below, the single buffer wins on this shape
-// (LDS halved; VGPR may drop <= 42 and open 3 blocks/CU, the trusted
-// occupancy-probe set [2, 3, 4]). Results stay bit-identical either way
-// (same stage size, same k0-outer/kk-inner int32 order, same bytes).
+// Round 6 (this session, mandated EPILOGUE round): fuse per-row x per-column
+// scaling, bf16 conversion and a fully coalesced final store into the compute
+// kernel and remove any unnecessary workspace/combine pass.
+// * The accepted arm already fuses the scaling, conversion and store into
+// w8a8_dumma_proj_operand_reuse_kernel, and the operator never allocates or
+// reads the caller workspace (there is no split-K partial and no combine
+// kernel), so there is no pass to remove. What the round-2..5 epilogue did
+// NOT do is store coalesced: the gfx928 fragment ownership (row = lane & 15,
+// col = (lane >> 4) + 4*i) makes each 64-lane *bf16* store instruction
+// carry 16 rows x 4 columns, i.e. 16 distinct 128-B global lines with only
+// 8 useful bytes in each. The exact code object confirms it: 16
+// global_store_short_d16_hi per thread, interleaved with ~20 scalar
+// x_scale/weight_scale loads (the 21 static global loads) and s_waitcnt
+// vmcnt chains, and the counter-derived write traffic is 57,966,336 B for a
+// 50,331,648 B bf16 output (1.152x sector overhead) out of 61,997,440 B
+// total per operator replay.
+// * Focused change: the exact-arm epilogue now (a) cooperatively stages the
+// 64 x_scale and 128 weight_scale values into the retired A LDS tile with
+// one coalesced load per thread, (b) writes each thread's 16 scaled bf16
+// results into a 64-row x 64-column half of the retired B LDS tile with the
+// bank-safe 8-mod-16 element stride 68 (136 B), and (c) after one barrier
+// writes the half back with 16-byte stores - 8 threads cover one full
+// 128-B output line, so a 64-lane store instruction moves 8 complete lines
+// instead of 16 partial ones. Two column halves plus one scale-publish
+// barrier (5 extra barriers per block); the staged bytes are the same
+// values in the same logical [M, N] order, and the scale multiply order
+// ((float)acc * x_scale * weight_scale) and the bf16 conversion are
+// byte-identical, so results stay bit-exact (mismatch 0 / max_abs_error 0.0
+// expected).
+// * The A/B staging buffers are dead after the K loop, so the epilogue reuses
+// them: LDS stays 14,848 B/block -> 2 blocks/CU, the grid stays (48, 64) =
+// 3072 blocks and the tile stays 64x128 / 512 threads / 8 waves. Every
+// other shape keeps the bootstrap arms and the scalar fallback (including
+// the shared pair M=2/M=16 down_proj tails) untouched.
+// * Falsifiable gate: median_us < 270.556 us (>= 1% under the 273.2886219 us
+// official best) AND p90_us <= 275.1925849914551, with the coalesced store
+// cutting write_bytes_per_operator_replay from 57,966,336 toward the
+// 50,331,648 B theoretical minimum and replacing 393,216 vmem_write
+// instructions with ~98,304 (16 bytes each). If the median does not clear
+// the 1% bar, the direct 2-byte-scatter epilogue is restored.
//
-// Round 4 (this session, glm_tp8_shared_down_proj_m4096 tile-shape round):
-// the balanced 1:1 aspect probe <64,64,64,true,false,true,true> (256 thr =
-// 4 waves, 6144 blocks, LDS 10,240 B/block -> 4 blocks/CU, doubled A
-// staging) measured 297.15 us median / 297.97 us p90 and was REJECTED
-// (guard restored to the round-3 single-buffered <64,128,64,true,false,
-// true,true> winner, 281.65 us / 282.52 us): co-residency was already
-// saturated at 2 blocks/CU and the +33% A-side staging cost was not repaid.
+// Round 15 (final consolidate / HIP-only polish on the frozen winning arm):
+// move the exact arm's scale delivery from the retired-A LDS slab to
+// per-thread registers.
+// * Frozen winning architecture (accepted best 158.0229902267456 us median /
+// 158.49021911621094 us p90, source digest c548257a): 64x128 macro-tile,
+// 512 threads / 8 waves with a 32x32 quadrant per wave, 64-K
+// single-buffered two-barrier stages, round-5 panel pack
+// [N/128][K/64][128][64], grid (48,64) = 3072 blocks, A[64,88] + B[128,72]
+// = 14,848 B LDS -> 2 blocks/CU (arch_vgpr 48), round-6 LDS-staged
+// coalesced fused epilogue.
+// * Evidence this round acts on: the exact-source PMC (120 CU, 3072 x 512,
+// 147.532 us profiled) reports lds_instructions 1,631,232 with
+// lds_wait_instructions 1,486,046 at ~11% SIMD issue utilization, i.e. ~66
+// LDS instructions per wave of which ~60 stall; the epilogue's scale
+// read-back alone is ~10 of those LDS reads per wave (x_scale[row + r] and
+// the four weight_scale columns of each of the two column groups, once per
+// accumulator fragment), all of them on the post-K-loop critical path. The
+// per-block tail is otherwise serialized behind 5 barriers (scale publish +
+// publish/WAR per 64-column half).
+// * Focused change (one bounded mechanism, exact `(m,n,k) == (4096,6144,256)`
+// arm only): each thread loads its OWN two row scales (rows
+// local_row + (lane & 15) and local_row + kTileM + (lane & 15)) and eight
+// column scales (columns local_col + (lane >> 4) + 4*i for the two kTileN
+// column groups) into registers ONCE before the K loop, and
+// stage_epilogue_fragment consumes those registers instead of reading
+// xs_tile/ws_tile from the retired A tile. Consequences: (a) no
+// x_scale/weight_scale is staged into or read back from LDS anywhere in the
+// exact arm, (b) the per-block scale-publish barrier disappears (block
+// barriers 13 -> 12: 8 K-loop + 4 epilogue), (c) the 10 global scale loads
+// per thread are issued before the K loop, so their latency hides behind the
+// four 64-K stages instead of sitting in front of the staging multiplies.
+// * Byte-identical data path: the values are the same x_scale[m0 + row] /
+// weight_scale[n0 + col], the multiply order
+// ((float)acc * x_scale * weight_scale) and the bf16 conversion are
+// unchanged, and every staged LDS address (stride 68) and every coalesced
+// 16-byte global store address is unchanged, so mismatch 0 / max_abs_error
+// 0.0 are expected. The guard launches exactly grid(n/128, m/64) with
+// m == 4096 and n == 6144, so the registered rows/columns are always in
+// range (m0 + local_row + kTileM + 15 <= 4095, n0 + local_col + kTileN +
+// 3 + 12 <= 6139).
+// * Everything else stays byte-identical: tile, grid, waves, K stage,
+// single buffering, panel pack, A/B strides and vectorized staging, LDS
+// 14,848 B, kernel name / launch ABI, all shape guards, the generic
+// prefill arms and the scalar fallback.
+// * Falsifiable prediction: median_us <= 158.0229902267456 (expected
+// ~155.8-158.0 us) AND p90_us <= 158.49021911621094, with lds_instructions
+// per wave down by ~10 and lds_wait down by a similar amount while
+// arch_vgpr stays <= 64 (no spills) and the LDS footprint stays 14,848 B.
+// If the median does not improve, the register scales were not on the
+// critical path (the LDS broadcast reads were already cheap) and the
+// round-6 LDS slab is restored.
//
-// Round 5 (this session, glm_tp8_shared_down_proj_m4096 PACKING round): test
-// one weight packing/swizzle that makes each DUMMA B tile vector-loadable
-// and LDS-bank-safe, with the packing outside timing and the graph-stable
-// packed layout validated. The accepted round-3 arm stages its 128x64 B
-// tile from the plain n-major [N,K] packed weight: per block-stage the 512
-// int4 loads are 16 rows x 4 int4 per warp at 256-B row stride, i.e. every
-// warp instruction touches 16 DISTINCT 64-B segments (50% of each 128-B L2
-// line) - vector loads, but scattered at L2 granularity. The round changes
-// the one-time pack for (k,n) == (256,6144) ONLY to the panel layout
-// [N/128][K/64][128][64] (n-major 64-k-byte rows per panel; byte-exact index
-// relation packed[((nt*(k/64)+kt)*128+nn)*64+kk64] == raw[(nt*128+nn)*k +
-// kt*64+kk64], verified by simulation): each block's per-stage B tile is now
-// ONE contiguous 8 KiB region, so every warp instruction reads 1 KiB
-// contiguous (full sector/line utilization; per block-stage the B staging
-// collapses from 128 half-used 128-B lines to 64 fully-used lines). The new
-// exact-shape arm w8a8_dumma_prefill_tiled_kernel<64,128,64,true,false,
-// true,true,false,true> (new kBPanel template flag, default false) is
-// byte-identical in EVERYTHING else: same 64-K single-buffered stage, same
-// kBNMajor/kDirectFrag/kSkew8, same 64x128 tile / 512 thr / 8 waves / 32x32
-// quadrant, same grid (48,64) = 3072 blocks, same LDS A[64,88]+B[128,72] =
-// 14,848 B/block @ 2 blocks/CU, same ds_write2_b64 staging stores and same
-// conflict-free load_frag8 ds_read2_b64 fragment reads (LDS-bank-safe
-// preserved - only the GLOBAL byte order of the packed buffer changes,
-// never the staged LDS bytes), same k0-outer/kk-inner int32 order, same
-// fused epilogue -> mismatch 0 / max_abs_error 0.0 expected. The scalar
-// fallback decodes the panel layout for (k,n) == (256,6144) so M<128 tails
-// stay correct; gate_up (6144,512) keeps its plain [N,K] transpose pack and
-// its own arm; generic arms are untouched (kBPanel=false). Packing runs
-// out-of-timed/out-of-Graph in launch_pack_w8a8_weight, same byte count
-// (256*6144), same graph-stable buffer. Expected effect: B global staging
-// becomes perfectly coalesced 1 KiB-per-warp streams with half the L2 line
-// count; PMC vmem_read_instructions unchanged (~614,400 - same 512 int4
-// loads/block-stage), lds_* identical (LDS layout untouched), l2_hits/l2_
-// misses shift slightly in the direction of fewer miss sectors, valu ~flat
-// (address math per thread unchanged: one base + offset). FALSIFIABLE GATE:
-// median_us < 281.6507148742676 AND p90_us < 282.5235176086426 -> B-global
-// coalescing was binding on this shape and the panel pack wins; otherwise
-// the plain n-major pack is retained and this round records that B reads
-// are L2-resident (90.4% hit) and NOT binding on the LDS-bound critical
-// path (consistent with the lineage rule that pack layout stops paying once
-// fragment loads/strides are fixed).
+// Round 17 (fix_build, final round): restore a compiling candidate on the
+// frozen round-15 winning arm after round 16 produced no valid proposal (the
+// worker process exited on a rate-limit crash before any build ran). The only
+// landed correction is the round-16 dead-code cleanup of the exact arm's A
+// staging helper, re-applied with a consistent API so the definition and its
+// single call site agree:
+// * stage_proj_tiles is instantiated ONLY by the exact arm, whose host guard
+// requires m == kExactM and launches exactly grid(n/kBlockN, m/kBlockM), so
+// every staged row m0 + local_row is in range and the per-int4
+// `global_row < m` compare/select/zero-fill was dead. The helper drops its
+// now-unused `m` parameter (the call site drops the argument) and gains a
+// static_assert that the exact arm covers whole kBlockM row tiles; every
+// staged A byte is unchanged, so results stay bit-exact.
+// * Deliberately NOT included: the round-13/16 final-half post-store barrier
+// skip. That is a runtime change whose measured p90 regressed at the 158-us
+// baseline; this round is a build fix, not a latency change.
+// * Everything else is byte-identical to the accepted round-15 best: 64x128
+// macro-tile, 512 threads / 8 waves with a 32x32 quadrant per wave, 64-K
+// single-buffered two-barrier stages, round-5 panel pack
+// [N/128][K/64][128][64], round-15 register scales, round-6 LDS-staged
+// coalesced epilogue (both epilogue barriers kept), A[64,88] + B[128,72] =
+// 14,848 B LDS, grid (48,64) = 3072 blocks, 2 blocks/CU, kernel name /
+// launch ABI, every shape guard and the generic scalar fallback.
+// * Falsifiable prediction: build_success true, correctness exact (mismatch 0
+// / max_abs_error 0.0), graph capture true, median_us flat at the accepted
+// best 148.41579914093018 us (the removed compare/select was not on the
+// critical path), p90_us <= 148.8046169281006, arch_vgpr <= 64,
+// scratch_bytes 0 and LDS 14,848 B.
//
-// Header order is fixed by the control plane: hip_runtime first (du_mma.h
-// is not self-contained before it), hip_bfloat16 second, du_mma.h last.
+// Header order is fixed by the control plane: hip_runtime first (du_mma.h is
+// not self-contained before it), hip_bfloat16 second, du_mma.h last.
#include
#include
@@ -325,12 +278,32 @@ constexpr int kTileK = 32;
// gfx928 wavefront is 64 lanes; every blockDim must be a multiple of 64.
constexpr int kWaveSize = 64;
-// Large-prefill dispatch geometry. Each macro-tile config uses a 32x32
-// quadrant per wave, so waves per block = (BlockM/32) * (BlockN/32).
+// Large-prefill dispatch geometry. Each macro-tile config owns a 32x32
+// quadrant per wavefront, so waves per block = (BlockM/32) * (BlockN/32).
constexpr int kPrefillMinM = 128;
-constexpr int kPrefillStageK = 128;
constexpr int kDefaultBlockM = 64;
-constexpr int kDefaultBlockN = 64;
+constexpr int kDefaultBlockN = 128;
+constexpr int kDefaultStageK = 128;
+
+// Round 2 exact-shape arm: the assigned (4096, 6144, 256) shape only.
+constexpr int kExactM = 4096;
+constexpr int kExactN = 6144;
+constexpr int kExactK = 256;
+constexpr int kProjBlockM = 64;
+constexpr int kProjBlockN = 128;
+constexpr int kProjStageK = 64;
+// The exact-shape arm uses the bank-safe 8-mod-16 LDS row strides A = 88
+// (kProjStageK + 24) and B = 72 (kProjStageK + 8); see the round-2 header
+// comment. The stride scheme assumes kStageK == 64.
+
+// Round 6 (epilogue round) LDS staging geometry for the fused epilogue: the
+// 64x128 bf16 output tile is staged one 64-column half at a time (kEpiCols)
+// into the retired B staging buffer with the bank-safe 8-mod-16 element row
+// stride kEpiStride = 68 bf16 = 136 B (16 row starts land on 16 distinct LDS
+// banks and the stride stays 8-byte aligned for the int64 read-back). The
+// retired A staging buffer holds the 64 x_scale and 128 weight_scale values.
+constexpr int kEpiCols = 64;
+constexpr int kEpiStride = 68;
constexpr int kScalarThreads = 256;
@@ -338,8 +311,9 @@ using namespace du::dumma;
// Verified gfx928 INT8 DUMMA accumulator ownership:
// row = lane & 15, col_mod4 = lane >> 4, frag.x[i] -> columns
-// col_mod4 + 4*i. Scale by x_scale[row] and weight_scale[col] and store
-// bf16 directly from the fragment (no accumulator LDS round trip).
+// col_mod4 + 4*i.
+// Scale by x_scale[row] and weight_scale[col] and store bf16 directly from the
+// fragment (no accumulator LDS round trip). Tail rows (row >= m) are masked.
template
__device__ __forceinline__ void store_prefill_fragment(
const AccFragment& frag,
@@ -360,61 +334,37 @@ __device__ __forceinline__ void store_prefill_fragment(
#pragma unroll
for (int i = 0; i < 4; ++i) {
const int col = base_col + col_mod4 + 4 * i;
- const float scaled = static_cast(frag.x[i]) * xs * weight_scale[col];
- out[row * n + col] = __float2bfloat16(scaled);
+ const float scaled =
+ static_cast(frag.x[i]) * xs * weight_scale[col];
+ out[static_cast(row) * n + col] = __float2bfloat16(scaled);
}
}
-// Round 10 (fragment-load-form round): explicit 8-byte fragment loader
-// replacing the library du_load_matrix_sync for the exact-shape gate_up
-// kernel. Validated on the sibling TP8 down_proj accepted kernel (iter 8,
-// -25.9%: 230.91 -> 171.19 us) and the TP4 gate_up lineage (iter 13
-// load_fragment8, -21%: 280.05 -> 220.71 us). For BOTH the library row_major
-// matrix_a and col_major matrix_b m16n16k32 int8 loaders, lane l -> row
-// (l & 15), k-quarter (l >> 4), and the loader fetches the same eight
-// contiguous bytes at p[row*ldm + (l>>4)*8 .. +7] and stores them into the
-// same x[0..7] slots in the same little-endian order (the round-3 ISA shows
-// the compiler lowering this to ds_read2_b32 followed by a per-byte
-// v_and/v_or_b32_sdwa/v_or3 mask-OR "identity" reassembly - ~9.4M dead-ish
-// VALU instructions per replay, ~61% of the 15.43M VALU total). Loading the
-// same 8 bytes as one int64 forces the 8-byte ds_read_b64 form instead:
-// (a) the mask/OR reassembly VALU disappears, and (b) the LDS access
-// changes from ds_read2_b32 (32-lane dword phases - at the stride-80 row
-// pattern every phase groups rows 0-15 of k-quarter q with rows 0-15 of
-// k-quarter q+1 onto the same 16 bank phases, a 4-way conflict per phase)
-// to ds_read_b64 (16-lane phases: rows 0-15 land on 16 bank-pairs
-// {20*r mod 32} = 8 distinct pairs hit twice - a 2-way conflict per phase,
-// half of the current fragment-conflict factor). The element-to-slot
-// mapping is unchanged (f.x[0..7] = the same 8 bytes in the same
-// little-endian order), the v_mmac operands and the exact k0-outer/kk-inner
-// int32 accumulation sequence are untouched, and the staging stores, LDS
-// strides, barriers, tile, grid, occupancy and epilogue are byte-identical.
+// Round 2 exact-shape fragment loader. The library row_major matrix_a and
+// col_major matrix_b m16n16k32 int8 loaders both map lane l -> row (l & 15),
+// k-quarter (l >> 4), and fetch the eight contiguous bytes at
+// p[row*ldm + (l >> 4)*8 .. +7] into the fragment's x[0..7] slots in the same
+// little-endian order. Loading those 8 bytes as one int64 emits a single
+// 8-byte LDS read instead of the library's byte-wise loads plus per-byte
+// reassembly VALU; the element-to-slot mapping is identical, so the MMAC
+// operands and the exact int32 accumulation 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;
+__device__ __forceinline__ void load_frag8(Frag& f,
+ const int8_t* p,
+ unsigned ldm,
+ int lane) {
+ const unsigned row = static_cast