diff --git a/.github/workflows/ci-l0-checks.yml b/.github/workflows/ci-l0-checks.yml index 179deccc1..a150d6743 100644 --- a/.github/workflows/ci-l0-checks.yml +++ b/.github/workflows/ci-l0-checks.yml @@ -60,7 +60,7 @@ jobs: # path (typed scalars, folded multi-line EXTRA__ARGS). Without it those # tests would silently exercise the degraded line-scanner fallback instead. - run: pip install pytest pytest-cov coverage pyyaml - # Stdlib-only tests (import run_e2e via importlib; no claude_agent_sdk / GPU). + # CPU-only tests (no claude_agent_sdk / GPU); the expert-skill test also uses PyYAML. # Files are listed explicitly rather than discovered, so that adding a test # under examples/ (torch/triton + GPU) can never silently break this job. # @@ -98,8 +98,19 @@ jobs: e2e_workflow/scripts/tests/test_server_teardown.py \ e2e_workflow/scripts/tests/test_bench_e2e_teardown_lookup.py \ ci/node/test_run_model_material_audit.py \ + perf_knowledge/expert_skills/_contribute/test_flydsl_skill_portability.py \ geak/test_bootstrap.py + - name: expert skill static validation + run: | + for skill in \ + flydsl_decode_moe_stage1_blkmap \ + flydsl_prefill_moe_stage2_fp8partial \ + flydsl_fp8_blockscale_gemm + do + python perf_knowledge/expert_skills/_contribute/validate_skill.py "$skill" --static + done + # Standalone contract checks. NOT in the pytest list above: each runs its assertions at import # time and exits non-zero, which pytest collection cannot express. They were added with the # learned-KB work and were not wired into CI at all — a regression suite nothing runs is a diff --git a/e2e_workflow/roles/_fragments/expert_skills.md b/e2e_workflow/roles/_fragments/expert_skills.md index 99a9b64dc..14bdffa70 100644 --- a/e2e_workflow/roles/_fragments/expert_skills.md +++ b/e2e_workflow/roles/_fragments/expert_skills.md @@ -1,8 +1,9 @@ # Fragment: expert_skills (e2e layer) — ADVISORY, injected only when use_expert_skills is ON > This fragment is appended to a role's prompt by `e2e_workflow.js` **only when `use_expert_skills` -> is true (opt-in; default OFF)**. When OFF (the default), nothing is injected and behavior is -> byte-identical to a run without this feature. It is consumed by routing/integration roles (System +> is true (opt-in; default OFF)**. When OFF, this fragment is not injected; the base role prompt remains active, +> including general routing and target-backend guidance independent of this registry. It is +> consumed by routing/integration roles (System > Architect, Op Benchmarker, e2e > Integrator). It is **advisory**: a matched skill is a high-prior candidate to reproduce, never a > mandate, and never overrides your on-box A/B gate. @@ -18,7 +19,8 @@ the measurement wins (note it so the skill is later marked `stale`). 1. **Read the selector.** Open `EXPERT_SKILLS_DIR/index.yaml`. 2. **Match against the live bottleneck** you are routing/optimizing. A skill matches when ALL hold: - - `match.operator` == the bottleneck operator (same names as `capability_index.yaml`) + - `match.operator` is either a scalar equal to the bottleneck operator or a list containing that + operator (same names as `capability_index.yaml`) - the box `gen` ∈ `match.gens` - `env_report.model_arch_class` ∈ `match.arch_class` (or `match.arch_class` contains `'*'`) - if the skill is a migration skill (`from_backend`/`to_backend` set), the live path / your author diff --git a/kernel_workflow/roles/_fragments/expert_skills.md b/kernel_workflow/roles/_fragments/expert_skills.md index bb504c597..22b402e5f 100644 --- a/kernel_workflow/roles/_fragments/expert_skills.md +++ b/kernel_workflow/roles/_fragments/expert_skills.md @@ -1,7 +1,8 @@ # Fragment: expert_skills (kernel layer) — ADVISORY, injected only when use_expert_skills is ON > Appended to a role's prompt by `kernel_workflow.js` **only when `use_expert_skills` is true -> (opt-in; default OFF)**. When OFF (the default), nothing is injected and behavior is byte-identical. +> (opt-in; default OFF)**. When OFF, this fragment is not injected; the base role prompt remains active, +> including general target-backend and language guidance from `perf_knowledge`. > Consumed by the > tech_lead (planning) and author/engineer roles. **Advisory**: a matched skill is a high-prior > candidate to reproduce, never a mandate, and never overrides your isolated A/B vs the oracle. @@ -16,7 +17,8 @@ faster but can never reduce a result below your measured baseline. 1. Read `EXPERT_SKILLS_DIR/index.yaml`. 2. A skill matches the current op when ALL hold: - - `match.operator` == this op's operator (`KK_OPERATOR` / `op_spec.op_kind`) + - `match.operator` is either a scalar equal to this op's operator + (`KK_OPERATOR` / `op_spec.op_kind`) or a list containing that operator - box `gen` ∈ `match.gens`; `op_spec.dtype` ∈ `match.dtypes`; `op_spec.regime` ∈ `match.regimes` - migration skills: `from_backend`→`to_backend` fits this run's `mode`/`target_language` (e.g. authoring Triton from a TileLang source → a `tilelang→triton` skill applies) diff --git a/kernel_workflow/roles/tech_lead.md b/kernel_workflow/roles/tech_lead.md index f03c4d214..8f64eafc6 100644 --- a/kernel_workflow/roles/tech_lead.md +++ b/kernel_workflow/roles/tech_lead.md @@ -120,7 +120,7 @@ analysis below exactly as before.) programming model. (Dir names differ from the ids for the others: triton→`triton_amd`, hip→`hip_cpp`, ck→`composable_kernel`, asm→`asm_mfma`.) - **Cross-backend port / migration** (the TASK asks to rewrite the kernel into a DIFFERENT backend — - ANY `source→target`, e.g. `ck→flydsl`, `triton→tilelang`, `hip→ck`, `ck→ck_tile`): keep `kk_language` + ANY `source→target`, e.g. `ck→flydsl`, `triton→tilelang`, `hip→ck`): keep `kk_language` = the CURRENT editable source, but `kk_refs` MUST ALSO include (a) the TARGET backend card `operators//backends/.md` and (b) the TARGET language's authoring how-to under `languages//` — map the language id to its dir: triton→`triton_amd`, hip→`hip_cpp`, diff --git a/perf_knowledge/expert_skills/README.md b/perf_knowledge/expert_skills/README.md index e00aafc99..b7751e8dd 100644 --- a/perf_knowledge/expert_skills/README.md +++ b/perf_knowledge/expert_skills/README.md @@ -38,13 +38,14 @@ kernel, tuned config JSONs, a custom validation manifest — can carry those fil ## How a skill is selected by the workflows -Each skill's frontmatter has a `match:` block. Expert skills are **opt-in** — the workflows ignore this -directory entirely unless the run passes `use_expert_skills=true` (default OFF; when OFF the workflow -behaves byte-identically to a build without this feature). When enabled, the workflow filters -`index.yaml` by the current bottleneck: +Each skill's frontmatter has a `match:` block. Expert-skill selection is **opt-in** — the workflows do +not inject or query this registry unless the run passes `use_expert_skills=true` (default OFF). When +OFF, the base role prompt remains active, including general `perf_knowledge` guidance that is independent +of expert-skill selection. When enabled, the workflow filters `index.yaml` by the current bottleneck: ``` -match.operator == bottleneck.operator +(match.operator is scalar AND match.operator == bottleneck.operator) + OR (match.operator is a list AND bottleneck.operator ∈ match.operator) AND gen ∈ match.gens AND model_arch_class ∈ match.arch_class (or match.arch_class contains '*') AND (migration skills) from_backend/to_backend fit the live path diff --git a/perf_knowledge/expert_skills/_contribute/scaffold.py b/perf_knowledge/expert_skills/_contribute/scaffold.py index 83f45035c..dd3cb5fa2 100755 --- a/perf_knowledge/expert_skills/_contribute/scaffold.py +++ b/perf_knowledge/expert_skills/_contribute/scaffold.py @@ -42,6 +42,23 @@ def known_operators(): return {c["operator"] for c in (data.get("candidates") or []) if "operator" in c} +def skill_index_entry(sub, fm): + """Build one selector entry, including non-profile dependency skills.""" + scope = fm.get("scope", "kernel") + entry = { + "id": fm["id"], + "file": f"skills/{sub}/skill.md", + "scope": scope, + "match": fm.get("match", {}), + } + if scope == "dependency": + entry["validation_status"] = "n/a" + else: + entry["expects"] = fm.get("expects", {}) + entry["validation_status"] = (fm.get("validation") or {}).get("status", "draft") + return entry + + def reindex(): ops = known_operators() entries = [] @@ -57,14 +74,7 @@ def reindex(): if ops is not None and one != "*" and one not in ops: print(f" WARN: {sub}/skill.md: operator '{one}' not in capability_index.yaml", file=sys.stderr) - entries.append({ - "id": fm["id"], - "file": f"skills/{sub}/skill.md", - "scope": fm.get("scope", "kernel"), - "match": fm.get("match", {}), - "expects": fm.get("expects", {}), - "validation_status": (fm.get("validation") or {}).get("status", "draft"), - }) + entries.append(skill_index_entry(sub, fm)) header = ( "# index.yaml — expert_skills selector (AUTO-MAINTAINED by _contribute/scaffold.py + " "validate_skill.py).\n" diff --git a/perf_knowledge/expert_skills/_contribute/test_flydsl_skill_portability.py b/perf_knowledge/expert_skills/_contribute/test_flydsl_skill_portability.py index 31cb43e72..84dbcea2d 100644 --- a/perf_knowledge/expert_skills/_contribute/test_flydsl_skill_portability.py +++ b/perf_knowledge/expert_skills/_contribute/test_flydsl_skill_portability.py @@ -1,14 +1,29 @@ from pathlib import Path +import importlib.util +import subprocess +import sys import pytest import yaml SKILLS_ROOT = Path(__file__).resolve().parents[1] / "skills" +EXPERT_SKILLS_ROOT = SKILLS_ROOT.parent +GEAK_ROOT = EXPERT_SKILLS_ROOT.parents[1] MOE_SKILLS = [ "flydsl_decode_moe_stage1_blkmap", "flydsl_prefill_moe_stage2_fp8partial", ] +DENSE_SKILL = "flydsl_fp8_blockscale_gemm" +SUBMITTED_SKILLS = [*MOE_SKILLS, DENSE_SKILL] + + +def load_scaffold_module(): + path = EXPERT_SKILLS_ROOT / "_contribute" / "scaffold.py" + spec = importlib.util.spec_from_file_location("expert_skill_scaffold", path) + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module def load_skill(skill_id): @@ -35,3 +50,142 @@ def test_flydsl_skill_uses_portable_minimum_version_guidance(skill_id): evidence = yaml.safe_load(artifact.read_text()) assert evidence["skill_id"] == skill_id assert evidence["flydsl_version"] == "0.2.2" + + +@pytest.mark.parametrize("skill_id", MOE_SKILLS) +def test_moe_skills_match_both_taxonomy_ids_on_existing_flydsl(skill_id): + frontmatter, _ = load_skill(skill_id) + match = frontmatter["match"] + + assert match["operator"] == ["grouped_gemm_moe", "fused_moe_grouped_gemm"] + assert match["from_backend"] == "flydsl" + assert match["to_backend"] == "flydsl" + + +def test_dense_skill_matches_dense_and_scaled_quant_taxonomy(): + frontmatter, _ = load_skill(DENSE_SKILL) + + assert frontmatter["match"]["operator"] == ["dense_gemm", "scaled_quant_gemm"] + + +@pytest.mark.parametrize( + "fragment", + [ + GEAK_ROOT / "kernel_workflow" / "roles" / "_fragments" / "expert_skills.md", + GEAK_ROOT / "e2e_workflow" / "roles" / "_fragments" / "expert_skills.md", + ], +) +def test_selector_prompt_defines_list_operator_membership(fragment): + text = fragment.read_text() + + assert "a list containing" in text + + +@pytest.mark.parametrize("skill_id", MOE_SKILLS) +def test_emit_plan_uses_in_place_optimize_for_moe_skills(skill_id): + result = subprocess.run( + [ + sys.executable, + str(EXPERT_SKILLS_ROOT / "_contribute" / "validate_skill.py"), + skill_id, + "--emit-plan", + ], + check=True, + capture_output=True, + text=True, + ) + + assert "mode=optimize" in result.stdout + assert "target_language=triton" not in result.stdout + assert "mode=author" not in result.stdout + + +def test_emit_plan_lists_author_and_optimize_for_mixed_source_skill(): + result = subprocess.run( + [ + sys.executable, + str(EXPERT_SKILLS_ROOT / "_contribute" / "validate_skill.py"), + "gluon_authoring", + "--emit-plan", + ], + check=True, + capture_output=True, + text=True, + ) + + assert "mode=author target_language=gluon" in result.stdout + assert "mode=optimize" in result.stdout + + +@pytest.mark.parametrize("skill_id", SUBMITTED_SKILLS) +def test_kernel_skills_do_not_claim_an_unmeasured_e2e_gate(skill_id): + frontmatter, _ = load_skill(skill_id) + + assert "e2e_delta_min_pct" not in frontmatter["expects"] + assert frontmatter["validation"]["measured"]["e2e_pct"] == "" + + +@pytest.mark.parametrize( + "document", + [ + EXPERT_SKILLS_ROOT / "README.md", + GEAK_ROOT / "kernel_workflow" / "roles" / "_fragments" / "expert_skills.md", + GEAK_ROOT / "e2e_workflow" / "roles" / "_fragments" / "expert_skills.md", + ], +) +def test_disabled_wording_is_limited_to_the_skill_fragment(document): + text = document.read_text() + normalized = " ".join(text.split()) + + assert "base role prompt remains active" in normalized + + +def test_dependency_skill_reindexes_as_non_profile_matched(): + scaffold = load_scaffold_module() + entry = scaffold.skill_index_entry( + "ensure_flydsl", + { + "id": "ensure_flydsl", + "scope": "dependency", + "match": {"needs": "flydsl"}, + }, + ) + + assert entry["scope"] == "dependency" + assert "expects" not in entry + assert entry["validation_status"] == "n/a" + + +def test_cross_backend_examples_use_canonical_backend_ids(): + text = (GEAK_ROOT / "kernel_workflow" / "roles" / "tech_lead.md").read_text() + + assert "ck→ck_tile" not in text + + +def test_dense_skill_frontmatter_matches_archived_down_proj_evidence(): + frontmatter, body = load_skill(DENSE_SKILL) + artifact = EXPERT_SKILLS_ROOT / frontmatter["validation"]["artifact"] + evidence = yaml.safe_load(artifact.read_text()) + + assert "down_proj" in frontmatter["validation"]["measured"]["isolated"] + assert "tile_n=256/tile_k=128" in frontmatter["validation"]["measured"]["isolated"] + assert evidence["skill_id"] == DENSE_SKILL + assert evidence["provider"] == "standalone_flydsl" + assert evidence["method"] == ( + "CUDA-event same-session paired A/B, 3 archived on-box GEAK measurements" + ) + assert evidence["warmup_per_run"] == 10 + assert evidence["iterations_per_run"] == 100 + assert evidence["inner_repetitions"] == 3 + assert len(evidence["baseline_latency_ms"]) == 3 + assert len(evidence["skill_latency_ms"]) == 3 + assert evidence["speedup"] == pytest.approx( + evidence["baseline_median_ms"] / evidence["skill_median_ms"], rel=5e-5 + ) + assert evidence["parity"]["pass"] is True + assert evidence["conclusions"]["decisive_lever"] == "tile_n=256/tile_k=128" + assert evidence["conclusions"]["xcd_swizzle"] == "non-load-bearing for down_proj" + assert evidence["conclusions"]["eight_wave"] == "measured dead-end for down_proj" + assert evidence["raw_logs"].startswith("external archived GEAK artifacts") + assert "artifacts" not in evidence + assert "8-wave-blockscale port is also an open" not in body diff --git a/perf_knowledge/expert_skills/_contribute/validate_skill.py b/perf_knowledge/expert_skills/_contribute/validate_skill.py index 60d78738f..66999aa62 100755 --- a/perf_knowledge/expert_skills/_contribute/validate_skill.py +++ b/perf_knowledge/expert_skills/_contribute/validate_skill.py @@ -93,24 +93,25 @@ def emit_plan(skill_id, fm, args): print(f" model_path= use_expert_skills=true # expect |e2e delta| < noise band") else: match = fm.get("match") or {} - dst = match.get("to_backend") or "triton" + dst = match.get("to_backend") or "" srcs = match.get("from_backend") or [] srcs = [srcs] if isinstance(srcs, str) else list(srcs) + cross_backend_srcs = [src for src in srcs if src and dst and src != dst] print("# EFFICACY (kernel_workflow, isolated A/B vs the immutable oracle):") print(f"Workflow scriptPath={GEAK}/kernel_workflow/kernel_workflow.js args:") print(f" kernel_path= workflow_dir={GEAK}/kernel_workflow use_expert_skills=true") # kernel_workflow reads target_language ONLY on the mode=author branch, so a port has to ask for # that branch or the run silently measures the untouched source language instead of the migration. - if [s for s in srcs if s != dst]: + if cross_backend_srcs: + print(" # Cross-backend author arm:") print(f" mode=author target_language={dst}" - f" # port from {'|'.join(s for s in srcs if s != dst)};" + f" # port from {'|'.join(cross_backend_srcs)};" " target_language is inert without mode=author") - else: - print(f" target_language={dst}") + if not cross_backend_srcs or dst in srcs: + if cross_backend_srcs: + print(" # Existing target-backend arm (measure separately):") + print(" mode=optimize # in-place optimization; no target_language needed") print(f" task='reproduce expert_skill:{skill_id}; beat oracle, hold parity'") - if dst in srcs: - print(f"# The selector also matches an existing {dst} source ({dst}->{dst}); measure that entry") - print("# state separately in the default mode=optimize (no mode/target_language args).") print("# DO-NO-HARM (control op that does NOT match the selector must stay within noise band):") print(" kernel_path= use_expert_skills=true # expect no regression") print("\nThen stamp the result with: validate_skill.py", skill_id, diff --git a/perf_knowledge/expert_skills/index.yaml b/perf_knowledge/expert_skills/index.yaml index a8093a317..f5e0cfc8b 100644 --- a/perf_knowledge/expert_skills/index.yaml +++ b/perf_knowledge/expert_skills/index.yaml @@ -31,7 +31,7 @@ skills: validation_status: validated - id: ensure_flydsl file: skills/ensure_flydsl/skill.md - scope: dependency # build/dependency skill: makes flydsl importable; invoked BY the flydsl skills, not profile-matched + scope: dependency match: needs: flydsl arch_class: @@ -47,7 +47,9 @@ skills: file: skills/flydsl_decode_moe_stage1_blkmap/skill.md scope: kernel match: - operator: grouped_gemm_moe + operator: + - grouped_gemm_moe + - fused_moe_grouped_gemm arch_class: - '*' gens: @@ -59,8 +61,8 @@ skills: - fp8_e4m3_fnuz regimes: - decode - from_backend: '' - to_backend: '' + from_backend: flydsl + to_backend: flydsl profile_signature: op_name_regex: '' min_pct_gpu: 0.0 @@ -70,14 +72,15 @@ skills: this recipe adds. Do not include pre-existing sort/align/quant helpers in this isolated gate; report the full stage-1 window separately. The producer has an unrelated name, so a GEMM-only filter drops it and overstates the win (1.124x GEMM-only vs 1.104x with producer on FlyDSL 0.2.2). - e2e_delta_min_pct: 1.0 parity: required validation_status: validated - id: flydsl_fp8_blockscale_gemm file: skills/flydsl_fp8_blockscale_gemm/skill.md scope: kernel match: - operator: dense_gemm + operator: + - dense_gemm + - scaled_quant_gemm arch_class: - '*' gens: @@ -95,7 +98,6 @@ skills: min_pct_gpu: 0.0 expects: isolated_speedup_min: 1.05 - e2e_delta_min_pct: 1.0 parity: required validation_status: validated - id: flydsl_fp8_gemm_playbook @@ -124,7 +126,9 @@ skills: file: skills/flydsl_prefill_moe_stage2_fp8partial/skill.md scope: kernel match: - operator: grouped_gemm_moe + operator: + - grouped_gemm_moe + - fused_moe_grouped_gemm arch_class: - '*' gens: @@ -136,8 +140,8 @@ skills: - fp8_e4m3_fnuz regimes: - prefill - from_backend: '' - to_backend: '' + from_backend: flydsl + to_backend: flydsl profile_signature: op_name_regex: '' min_pct_gpu: 0.0 @@ -148,7 +152,6 @@ skills: in the reduce kernel, so a GEMM-only measurement tops out at a measured 1.053x and fails this gate even when the recipe is correctly reproduced. A per-kernel timing filter matched on the GEMM name will not match the reduce kernel name -- verify your filter catches both.' - e2e_delta_min_pct: 1.0 parity: relaxed validation_status: validated - id: flydsl_rewrite_quantized_moe diff --git a/perf_knowledge/expert_skills/skills/flydsl_decode_moe_stage1_blkmap/skill.md b/perf_knowledge/expert_skills/skills/flydsl_decode_moe_stage1_blkmap/skill.md index fac3a3a0b..1d01d2259 100644 --- a/perf_knowledge/expert_skills/skills/flydsl_decode_moe_stage1_blkmap/skill.md +++ b/perf_knowledge/expert_skills/skills/flydsl_decode_moe_stage1_blkmap/skill.md @@ -6,7 +6,9 @@ authors: - zhengy scope: kernel match: - operator: grouped_gemm_moe + operator: + - grouped_gemm_moe + - fused_moe_grouped_gemm arch_class: - '*' gens: @@ -18,8 +20,8 @@ match: - fp8_e4m3_fnuz regimes: - decode - from_backend: '' - to_backend: '' + from_backend: flydsl + to_backend: flydsl profile_signature: op_name_regex: '' min_pct_gpu: 0.0 @@ -29,7 +31,6 @@ expects: kernel this recipe adds. Do not include pre-existing sort/align/quant helpers in this isolated gate; report the full stage-1 window separately. The producer has an unrelated name, so a GEMM-only filter drops it and overstates the win (1.124x GEMM-only vs 1.104x with producer on FlyDSL 0.2.2).' - e2e_delta_min_pct: 1.0 parity: required validation: status: validated diff --git a/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.md b/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.md index 6e774e331..b5b942b86 100644 --- a/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.md +++ b/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.md @@ -1,12 +1,14 @@ --- id: flydsl_fp8_blockscale_gemm -title: Rewrite CK fp8 a8w8 blockscale GEMM to FlyDSL on gfx950 (software fp32 post-MFMA scale, NOT E8M0; shape-dependent perf lever — XCD swizzle on wide-N, 8-wave cluster on narrow-N) +title: Rewrite CK fp8 a8w8 blockscale GEMM to FlyDSL on gfx950 (software fp32 post-MFMA scale, NOT E8M0; shape-dependent perf lever — XCD locality on very-wide-N, tile geometry on K-light, 8-wave cluster on narrow-N) kind: expert_skill authors: - zhengy scope: kernel match: - operator: dense_gemm + operator: + - dense_gemm + - scaled_quant_gemm arch_class: - '*' gens: @@ -24,19 +26,18 @@ match: min_pct_gpu: 0.0 expects: isolated_speedup_min: 1.05 - e2e_delta_min_pct: 1.0 parity: required validation: status: validated - last_verified: '2026-07-16 (wide-N q_up), 2026-07-17 (narrow-N qkv)' + last_verified: '2026-07-16 (wide-N q_up), 2026-07-17 (narrow-N qkv and K-light down_proj)' gpu: 'MI355X / gfx950 (device 0x75a3)' model: 'claude (kernel_workflow expert-skill verify run)' measured: - isolated: 'wide-N q_up_proj M=4096 N=65536 K=1536 (4-wave preshuffle + xcd8): 1.226-1.237x same-session interleaved A/B vs production CK (gemm_a8w8_blockscale_bpreshuffle), xcd0 ablation 0.973x (loses to CK) confirming the XCD lever. narrow-N qkv_proj M=4096 N=2048 K=7168 (8-wave ping-pong cluster, BLOCK_M=128/BLOCK_N=256): director-verified 1.05x, accepted (interleaved A/B 1.0513/1.0500/1.0462, all >1.0); xcd_swizzle no-op and 4-wave -19% both confirmed for this shape' + isolated: 'wide-N q_up_proj M=4096 N=65536 K=1536 (4-wave preshuffle + xcd8): 1.226-1.237x same-session interleaved A/B vs production CK (gemm_a8w8_blockscale_bpreshuffle), xcd0 ablation 0.973x (loses to CK) confirming the XCD lever. narrow-N qkv_proj M=4096 N=2048 K=7168 (8-wave ping-pong cluster, BLOCK_M=128/BLOCK_N=256): director-verified 1.05x, accepted (interleaved A/B 1.0513/1.0500/1.0462, all >1.0); xcd_swizzle no-op and 4-wave -19% both confirmed for this shape. K-light down_proj M=16384 N=7168 K=768: 1.4506-1.4553x paired A/B; tile_n=256/tile_k=128 was decisive, XCD/fused-promote were non-load-bearing, and the completed 8-wave follow-up was a measured dead-end' e2e_pct: '' - parity: 'wide-N: err=0 / cos=1.0 seeds 0-3 vs fp32 dequant oracle (rtol=atol=1e-2). narrow-N: cos=1.0, maxabs_err=0.03125 (identical to CK), checkAllclose pass, 0 elements out of tol' - artifact: 'external validation artifacts for wide-N q_up manual verification and narrow-N qkv GEAK director verification (provenance only; exact local paths omitted)' - notes: 'Cross-machine reproduction, BOTH regimes now verified. wide-N/4-wave+xcd8 (q_up) and narrow-N/8-wave cluster (qkv) are both parity-clean and beat production CK. FlyDSL API drift confirmed as documented (kernels.mma -> kernels.common.mma; xcd_swizzle compile-arg is an uncommitted patch on mainline -> used the byte-frozen snapshot core via a module-alias shim). IMPORTANT install note: the blockscale software-promote cores live in a standalone FlyDSL checkout/build, NOT the aiter-embedded flydsl (aiter/aiter/ops/flydsl = rowscale/epilogue only: cos~0.999 but ~78% elements out of tol, a dead-end, not a blockscale drop-in). Point ports at a standalone build with software blockscale support.' + parity: 'wide-N: err=0 / cos=1.0 seeds 0-3 vs fp32 dequant oracle (rtol=atol=1e-2). narrow-N: cos=1.0, maxabs_err=0.03125 (identical to CK), checkAllclose pass, 0 elements out of tol. down_proj: FlyDSL maxabs_err=0.0078 / cos=1.0, at least as accurate as the CK oracle baseline' + artifact: skills/flydsl_fp8_blockscale_gemm/validation_gfx950.yaml + notes: 'Cross-machine reproduction across wide-N q_up, narrow-N qkv, and K-light down_proj. XCD is load-bearing only on the locality-limited q_up evidence; down_proj is grid-saturated and its measured lever is tile_n=256/tile_k=128, with XCD/fused-promote non-load-bearing and 8-wave closed as a dead-end. FlyDSL API drift confirmed as documented. IMPORTANT install note: the blockscale software-promote cores live in a standalone FlyDSL checkout/build, NOT the aiter-embedded flydsl (aiter/aiter/ops/flydsl = rowscale/epilogue only: cos~0.999 but ~78% elements out of tol, a dead-end, not a blockscale drop-in). Point ports at a standalone build with software blockscale support.' role: advisory_prior supersedes: [] --- @@ -65,10 +66,11 @@ which is why the tile re-sweep (not the mechanism) matters per shape: (…and any future / decode `-m`/`-nk` variant of the same operator). Two stages, and only the first is universal: **(1) the parity fix** — do the block-scale **in software** (never E8M0) — is the *same for -every shape* and is what unblocks the port; **(2) the perf lever is shape-dependent** — wide-N/large-C -shapes win with a 4-wave core + **XCD swizzle**, narrow-N/deep-K shapes win with an **8-wave "cluster"** -core where XCD is a no-op. Three shapes are already tuned end-to-end; use their configs as the per-regime -starting point (see **Per-shape recipes** below) rather than sweeping blind. +every shape* and is what unblocks the port; **(2) the perf lever is shape-dependent** — locality-limited +very-wide-N shapes need a 4-wave core + **XCD swizzle**, grid-saturated K-light shapes need the 4-wave +`tile_n=256/tile_k=128` geometry but not XCD, and narrow-N/deep-K shapes win with an **8-wave "cluster"** +core. Three shapes are already tuned end-to-end; use their configs as the per-regime starting point +(see **Per-shape recipes** below) rather than sweeping blind. ## Mechanism ### A. Universal — the parity fix (same for every shape; this is what unblocks the port) @@ -95,13 +97,14 @@ needed** — the fix is choosing a software-scale core over the native scaled-MF The correctness fix is universal; **which perf lever wins inverts with the shape**, because the two shape classes are bound by different things. Measured on the three tuned shapes: -- **Wide-N / large-C / cold-tail** (`q_up` N=65536/512 MiB C, `down_proj` N=7168/235 MiB C): bound by **L2 / - fabric locality**, so **XCD-aware grid rasterization is the lever**. MI355X has **8 XCDs**, each with its - own L2 slice; remapping CTA `(bx,by)` so CTAs sharing a B/N-panel land on one XCD lifts L2 hit-rate - (`74.8%→83.0%` on q_up) / cuts fabric read-latency (`1030→803 cyc` on down_proj). Worth **~1.23×** and - *load-bearing*: `xcd_swizzle=0` → **loses to CK (0.96×)**; `=8` → 1.11–1.36×. Here **8-wave spills** - (wide-N accumulator blows the VGPR file). GEAK's gate_up auto-win picked the right kernel but **left - `xcd_swizzle` off → only ~1.02×.** +- **Very-wide-N / locality-limited** (`q_up` N=65536/512 MiB C): **XCD-aware grid rasterization is the + lever**. MI355X has **8 XCDs**, each with its own L2 slice; remapping CTA `(bx,by)` so CTAs sharing a + B/N-panel land on one XCD lifts L2 hit-rate `74.8%→83.0%`. The `xcd0→xcd8` ablation is worth **~1.23×**; + `xcd0` loses to CK while `xcd8` wins. The 8-wave family spills on this accumulator footprint. +- **K-light / grid-saturated** (`down_proj` M=16384,N=7168,K=768): compute-bound with 7168 output tiles + over the device, so the decisive lever is the 4-wave **`tile_n=256/tile_k=128` geometry** that keeps + the short six-K-tile loop fed. XCD remapping and `fused_promote` are measured non-load-bearing on this + shape, and the 8-wave port is a measured dead-end (`147 us` ceiling vs the 4-wave winner near `129 us`). - **Narrow-N / deep-K / tiny-C** (`qkv` N=2048/16 MiB C/K=7168): grid+LDS-bound at **1 block/CU, 2 waves/SIMD**, no cold tail → **`xcd_swizzle` is a measured no-op** and the 4-wave core is ~19% slower. The win comes from the **8-wave ping-pong** kernel with the **`cluster` schedule** — move all four fp32 @@ -115,15 +118,18 @@ label** — match your `(M, N, K)` to the nearest row via the decision guide, ap only the shape-dependent knobs. The `seen as` column is provenance only and never a selector. **Decision guide (pick the kernel family purely from the shape):** -- **Wide-N (N ≳ 4096) or large C output or a cold-input tail** → **4-wave blockscale-preshuffle core**, - set **`xcd_swizzle=8`**, tile `t64×256×128 wpe2`. (XCD is the lever; 8-wave will spill.) +- **Very-wide-N with measured locality pressure** → **4-wave blockscale-preshuffle core**, tile + `t64×256×128 wpe2`, then sweep XCD; `xcd8` is the validated q_up winner. +- **K-light with enough tiles to saturate the grid** → the same 4-wave core at `t64×256×128 wpe2`; + lock the tile geometry first. On validated down_proj, XCD/fused-promote are non-load-bearing and the + 8-wave family cannot beat the 4-wave winner. - **Narrow-N (N ≲ 2048) + deep-K + tiny C output** → **8-wave ping-pong core with the `cluster` schedule**, `BLOCK_M=128, BLOCK_N=256`. (8-wave hides the promote; `xcd_swizzle` is a no-op; don't bother with it.) | shape M×N×K (drives selection) | bind | winning FlyDSL core | config | measured vs CK | decisive lever | what LOSES here | seen as | |---|---|---|---|---|---|---|---| | 4096×65536×1536 (wide-N, 512 MiB C) | L2-locality | 4-wave blockscale-preshuffle | `t64×256×128, wpe2, xcd8` | **1.11× hot / 1.15× cold / 1.18× ev** | XCD swizzle (`xcd0`→`xcd8` = 1.23×) | 8-wave (spills 0.14–0.72×), bigger tile/`wpe≥3` (spill), `fused_promote` (−1.6%) | `q_up_proj` | -| 16384×7168×768 (K-light, 235 MiB C) | cold-tail | 4-wave blockscale-preshuffle | `t64×256×128, wpe2, xcd8, fused_promote` | **1.4541× Director-verified** (`0.187491→0.128937 ms`) | XCD swizzle + `fused_promote`; second run `1.4568×` | `wpe≥3` (spill 4–5×); *8-wave is the top **open** lever here (~1.3× more), port unfinished* | `down_proj` | +| 16384×7168×768 (K-light, 235 MiB C) | compute/grid | 4-wave blockscale-preshuffle | `t64×256×128, wpe2` (`xcd`/`fused_promote` non-load-bearing) | **1.4541× Director-verified** (`0.187491→0.128937 ms`) | `tile_n=256/tile_k=128` | 8-wave measured no-win; `wpe≥3` spills | `down_proj` | | 4096×2048×7168 (narrow-N, deep-K, 16 MiB C) | LDS/grid | 8-wave ping-pong (`cluster`) | `BLOCK_M=128, BLOCK_N=256` (promotes after last MFMA) | **1.045× hot / 1.086–1.10× interleaved** | 8-wave + `cluster` promote-reorder (MFMA util 45.8→47.9%) | **XCD swizzle (no-op)**, 4-wave core (−19%), 16-wave (spill 1.35× slower), `BLOCK_M=256` (14× slower) | `qkv_proj` | Notes: the narrow-N/deep-K shape (`4096×2048×7168`) is the **marginal-speedup** case — it sits ~1.05–1.10× @@ -186,15 +192,15 @@ interleaved A/B only**, and correctness gates every candidate (a fast-but-wrong 3. **Scale = software fp32 post-MFMA (the parity key).** Both kernels already promote+scale after the MFMA. Never switch to `mfma_scale_*_f8f6f4` for arbitrary fp32 block scales. 4. **Apply the shape's perf lever (conditional — do not blind-apply q_up's).** - - 4-wave path (wide-N): set **`xcd_swizzle=8`** (MI355X = 8 XCDs) — the single most important knob here; - `0` = loses to CK. This is what the gate_up auto-win missed. Add `fused_promote=True` for K-light - shapes like `down_proj` (~1.03×). + - 4-wave, locality-limited path (`q_up`): set **`xcd_swizzle=8`** (MI355X = 8 XCDs); the measured + `xcd0` arm loses to CK. + - 4-wave, K-light/grid-saturated path (`down_proj`): lock `tile_n=256/tile_k=128`; do not credit + XCD or `fused_promote` without a fresh ablation because both were non-load-bearing in validation. - 8-wave path (narrow-N): ensure the **`cluster`** schedule (all four promotes after the last MFMA); **skip `xcd_swizzle`** (measured no-op) and do NOT raise occupancy to 16-wave (spills). 5. **Tile / config = start from the nearest recipe row, then re-sweep only shape-dependent knobs.** - 4-wave: `tile_m=64, tile_n=256, tile_k=128, waves_per_eu=2`; only these move the needle — - `tile_n∈{128,256}` (256 wins, better B reuse), `xcd_swizzle∈{4,8}` (8 wins), `waves_per_eu∈{0,1,2}` - (≈ equal — occupancy is NOT the lever); everything larger spills (see pitfalls). 8-wave: `BLOCK_M=128, + 4-wave: `tile_m=64, tile_n=256, tile_k=128, waves_per_eu=2`; re-sweep `tile_n∈{128,256}` and XCD only + when the shape shows locality pressure. Everything larger spills (see pitfalls). 8-wave: `BLOCK_M=128, BLOCK_N=256` is the only viable tile (`BLOCK_N` is hardwired to 256 by scale alignment; `BLOCK_M=256` idles half the CUs). 6. **Feed layout (must match the oracle).** Preshuffle B with the CK-matching 16×16 weight layout; flatten @@ -207,25 +213,25 @@ interleaved A/B only**, and correctness gates every candidate (a fast-but-wrong ## Knobs & pitfalls - **Only the *correctness* mechanism is universal — the perf answer is not "one kernel, one config".** The software fp32 scale transfers to every shape; the *kernel family, the lever, and the tile* do **not**. - Wide-N (`q_up`, `down_proj`) → 4-wave `t64×256×128 wpe2` + `xcd8`; narrow-N/deep-K (`qkv`, and - `gate_up` N=1536) → 8-wave `cluster` `BLOCK_M=128×256`. **Never assume q_up's answer transfers** — on - `qkv` the winning kernel, the tile, *and* the XCD verdict all flip. -- **`xcd_swizzle` is shape-conditional, NOT universal.** On wide-N/large-C it is the decisive lever: `8` - wins on MI355X (8 XCDs), `4` close, `0` = off (**loses to CK, 0.96×**), `16` too coarse (0.93×); bijective - tile remap → bit-exact, always safe to try. On narrow-N/tiny-C (`qkv`, 16 MiB C, LDS/grid-bound, no cold - tail) it is a **measured no-op** — don't waste a sweep axis on it there. + Very-wide/locality-limited `q_up` → 4-wave `t64×256×128 wpe2` + `xcd8`; K-light/grid-saturated + `down_proj` → the same 4-wave tile with XCD/fused-promote treated as non-load-bearing; narrow-N/deep-K + (`qkv`, and `gate_up` N=1536) → 8-wave `cluster` `BLOCK_M=128×256`. +- **`xcd_swizzle` is shape-conditional, NOT universal.** It is decisive on the validated q_up shape, but + a measured no-op on both grid-saturated down_proj and narrow-N/tiny-C qkv. Treat it as an ablation axis, + not a default inferred from N or output size alone. - **8-wave ping-pong WINS or LOSES depending on N — this is the sharpest shape trap.** - **Narrow-N/deep-K (`qkv`): 8-wave `cluster` is the winner** (4-wave is ~19% slower). Keep the `cluster` promote-reorder; do NOT go 16-wave (register wall → spill, 1.35× slower). - **Wide-N (`q_up`): 8-wave LOSES (0.14–0.72×)** — the 256×256 latency-hiding tile needs the unscaled MFMA partial live *simultaneously* with the promoted frag → blows the VGPR budget → MFMA serialises. Use the - 4-wave core instead. (`down_proj` is the exception where 8-wave is a promising *open* lever — its port - isn't finished, so its shipped winner is still 4-wave+xcd+`fused_promote`.) + 4-wave core instead. + - **K-light (`down_proj`): 8-wave is a measured dead-end** — its rowscale ceiling is about `147 us`, + already slower than the 4-wave blockscale winner near `129 us`. - **The 4-wave register-pressure wall is sharp.** Two-level accumulator + register-resident double-buffered B leave ~no VGPR headroom at the 2-wave budget. Anything bigger than `t64×256×128 wpe2` — larger `tile_m/tile_n/tile_k`, or `waves_per_eu≥3` — **spills to scratch → 3–10× slower** (measured - `t128×256`=3057µs, `wpe3`=3683µs). `fuse_promote` on wide-N is −1.6% (helps only K-light `down_proj`). - Don't chase bigger tiles / higher occupancy without HK-style explicit register allocation. + `t128×256`=3057µs, `wpe3`=3683µs). Do not credit `fused_promote` on down_proj without a fresh ablation; + it was non-load-bearing in the validated grid-saturated run. - **FlyDSL API drift**: installed FlyDSL builds sometimes change internal helpers (e.g. a `crd2idx` / `.ir_value()` signature), so a small compat shim in the harness may be needed. Prefer a source checkout matching the core you author against. @@ -241,16 +247,15 @@ interleaved A/B only**, and correctness gates every candidate (a fast-but-wrong K-light **1.4541×**, wide-N **~1.1–1.2×** (winner within ~1.7% of the promote-free floor), narrow-N/deep-K **~1.05–1.10×** (already near its promote-free *rowscale* ceiling). All are stall/locality-bound at ~25–36% of the 5 PFLOP fp8 peak, so the software promote is *not* the bottleneck. The one lever left for - every shape is **E2E epilog fusion** (cut the C round-trip — invisible in a standalone kernel; on the - K-light shape, an 8-wave-blockscale port is also an open kernel-level lever). Don't burn budget on + every shape is **E2E epilog fusion** (cut the C round-trip — invisible in a standalone kernel). Don't burn budget on tile/occupancy/prefetch/VGPR/numerics micro-opts (all measured <2%). - **Inert when not triggered.** Non-blockscale op, non-gfx950 box, or the FlyDSL blockscale entrypoint absent → fall back to the generic path, no regression. ## Sources -The evidence lives in external artifacts (manual FlyDSL-rewrite handoffs + a GEAK auto-run) and is cited for -**provenance only** — GEAK does not depend on those trees, and exact commits / file paths are intentionally -left out. The portable knowledge is the shapes, configs, and measured numbers below. +`validation_gfx950.yaml` records three archived on-box down_proj latency pairs, parity, the measured +conclusions, and the known provenance limitation. The underlying manual handoffs and GEAK run remain external; +GEAK does not depend on those trees. The portable knowledge is the shapes, configs, and measured numbers below. - **q_up_proj manual rewrite (primary evidence)** — M=4096,N=65536,K=1536; **err=0 / cos=1.0**; **1.114× hot / 1.148× cold / 1.179× CUDA-event median**; winner `t64×256×128, wpe2, xcd8`; XCD lever `xcd0`=0.958× vs @@ -260,9 +265,10 @@ left out. The portable knowledge is the shapes, configs, and measured numbers be software-scale core at `t64×256×128` (parity clean) but **omitted `xcd_swizzle`** → only ~1.02×. Shows the bypass works *and* why XCD swizzle is mandatory to actually win. - **down_proj workflow validation (Director accepted; artifact case `mlp_down_m16384_n7168_k768`)** — - **1.4541×** (`0.187491 -> 0.128937 ms`) with a second run at **1.4568×**. Correctness passed: - FlyDSL `maxabs_err=0.0078`, `cos=1.0`. Winner `t64×256×128, wpe2, xcd8, fused_promote`; - same 4-wave core as q_up. Open lever: 8-wave-blockscale port (~1.3× more, unfinished). + **1.4541×** (`0.187491 -> 0.128937 ms`). Correctness passed: + FlyDSL `maxabs_err=0.0078`, `cos=1.0`. Winner family is the same 4-wave core as q_up, but ablation + showed `tile_n=256/tile_k=128` was decisive while XCD/fused-promote were non-load-bearing; a completed + 8-wave follow-up could not beat the 4-wave winner. - **qkv_proj manual rewrite (narrow-N 8-wave `cluster` winner — where the levers invert)** — M=4096,N=2048,K=7168; **err=0 / cos=1.0** (24/24 seeds×shapes); **1.045× hot / 1.086–1.10× interleaved**. Winner = 8-wave ping-pong core with the **`cluster`** promote-reorder (all 4 promotes after last MFMA → diff --git a/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/validation_gfx950.yaml b/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/validation_gfx950.yaml new file mode 100644 index 000000000..a4d67902c --- /dev/null +++ b/perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/validation_gfx950.yaml @@ -0,0 +1,35 @@ +skill_id: flydsl_fp8_blockscale_gemm +date: '2026-07-17' +provider: standalone_flydsl +gpu: gfx950 / MI355X +scope: down_proj M=16384 N=7168 K=768, CK versus FlyDSL isolated kernel +method: CUDA-event same-session paired A/B, 3 archived on-box GEAK measurements +warmup_per_run: 10 +iterations_per_run: 100 +inner_repetitions: 3 +baseline_latency_ms: +- 0.187491 +- 0.187262 +- 0.187461 +skill_latency_ms: +- 0.128937 +- 0.129089 +- 0.128817 +baseline_median_ms: 0.187461 +skill_median_ms: 0.128937 +speedup: 1.4538960888 +latency_reduction_pct: 31.2192936131 +parity: + pass: true + oracle: fp32 block dequantization, matmul, then bf16 cast + baseline_maxabs_err_range: [0.535156, 0.586] + baseline_cosine_range: [0.99974823, 0.99975228] + skill_maxabs_err: 0.0078125 + skill_cosine: 1.0 +conclusions: + decisive_lever: tile_n=256/tile_k=128 + xcd_swizzle: non-load-bearing for down_proj + fused_promote: non-load-bearing for down_proj + eight_wave: measured dead-end for down_proj +raw_logs: external archived GEAK artifacts (director validation, reprofile, TechLead report) +notes: YAML was created after the measured run; exact FlyDSL/AITER commits were not captured diff --git a/perf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.md b/perf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.md index f67808eaa..9c5b41b7b 100644 --- a/perf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.md +++ b/perf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.md @@ -7,7 +7,9 @@ authors: - zhengy scope: kernel match: - operator: grouped_gemm_moe + operator: + - grouped_gemm_moe + - fused_moe_grouped_gemm arch_class: - '*' gens: @@ -19,8 +21,8 @@ match: - fp8_e4m3_fnuz regimes: - prefill - from_backend: '' - to_backend: '' + from_backend: flydsl + to_backend: flydsl profile_signature: op_name_regex: '' min_pct_gpu: 0.0 @@ -31,7 +33,6 @@ expects: in the reduce kernel, so a GEMM-only measurement tops out at a measured 1.053x and fails this gate even when the recipe is correctly reproduced. A per-kernel timing filter matched on the GEMM name will not match the reduce kernel name -- verify your filter catches both.' - e2e_delta_min_pct: 1.0 parity: relaxed validation: status: validated