Skip to content

Kernel workflow: Add FlyDSL gfx950 skills - #386

Open
IzacharyI wants to merge 6 commits into
AMD-AGI:mainfrom
IzacharyI:upstream_dsv4pro_workload
Open

Kernel workflow: Add FlyDSL gfx950 skills#386
IzacharyI wants to merge 6 commits into
AMD-AGI:mainfrom
IzacharyI:upstream_dsv4pro_workload

Conversation

@IzacharyI

@IzacharyI IzacharyI commented Jul 31, 2026

Copy link
Copy Markdown

Summary

Add FlyDSL gfx950 expert skills for GEAK kernel_workflow.

This PR covers two kernel operator families:

  1. Dense fp8 a8w8 blockscale GEMM
  2. Grouped-MoE GEMM

These skills are consumed when use_expert_skills=true. They are advisory priors for kernel_workflow to reproduce and measure, not hard-coded replacements for correctness, parity, or A/B gates.

Modifications

Path Change
perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.md Adds the gfx950 fp8 a8w8 blockscale GEMM kernel skill with validated q_up_proj, down_proj, and qkv_proj evidence. This is distinct from perf_knowledge/expert_skills/skills/flydsl_fp8_gemm_playbook/skill.md, which is a gfx942 e2e down-proj overlay playbook.
perf_knowledge/expert_skills/skills/flydsl_decode_moe_stage1_blkmap/skill.md Adds the decode stage-1 grouped-MoE blkmap skill with phase/stage guardrails, descriptor timing scope, and runtime-signature checks.
perf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.md Adds the prefill stage-2 grouped-MoE fp8-partial skill with phase/stage guardrails, scale-consistency checks, and relaxed-parity guidance.
perf_knowledge/expert_skills/index.yaml Regenerates expert-skill selector metadata for the submitted validated skills.
kernel_workflow/roles/{tech_lead,engineer,deep_engineer}.md Surfaces target backend/language authoring knowledge for cross-backend CK -> FlyDSL ports.
perf_knowledge/operators/scaled_quant_gemm/backends/{ck,flydsl}.md Updates CK/FlyDSL scaled-quant GEMM guidance and expert-skill references for blockscale GEMM authoring.
perf_knowledge/operators/dense_gemm/{overview,numerics}.md Updates dense GEMM overview and blockscale numerics references for CK -> FlyDSL workflow knowledge.
perf_knowledge/operators/dense_gemm/backends/flydsl.md Updates dense GEMM FlyDSL backend guidance and expert-skill references.
perf_knowledge/operators/fused_moe_grouped_gemm/backends/flydsl.md Updates FlyDSL grouped-MoE backend knowledge surfaced during workflow authoring.
perf_knowledge/quantization/block_scaling_mxfp.md Updates block-scaling guidance to reference the current blockscale GEMM expert skill.

Skill Organization

The submitted skills are kept separate by workflow layer, GPU generation, and kernel scenario.

flydsl_fp8_blockscale_gemm is separate from the existing flydsl_fp8_gemm_playbook: the existing playbook is a gfx942 e2e down-proj overlay focused on capture-safe integration, while the new skill is a gfx950 kernel_workflow recipe for CK/production blockscale GEMM -> FlyDSL candidates.

The two Grouped-MoE skills are also kept separate: flydsl_decode_moe_stage1_blkmap targets decode / small-M stage-1 gate/up, while flydsl_prefill_moe_stage2_fp8partial targets prefill / large-M stage-2 down/reduce. They have different runtime signatures, validation scopes, and parity requirements.

Workflow Impact

  • Expert skills are only consulted when use_expert_skills=true.
  • Matched skills are advisory candidates for kernel workflow reproduction.
  • Final acceptance still depends on on-box A/B measurement and parity validation.
  • Default behavior with expert skills disabled is unchanged.
  • CK -> FlyDSL ports now surface target backend/language knowledge earlier, so workflow agents can reuse validated GEMM recipes instead of rediscovering the porting rules.

Validation

Operator / Scenario Skill kernel_workflow / Director-verified result Correctness
Dense fp8 blockscale GEMM / wide-N q_up_proj flydsl_fp8_blockscale_gemm ~1.226x-1.237x Pass
Dense fp8 blockscale GEMM / K-light down_proj flydsl_fp8_blockscale_gemm 1.4541x Pass
Dense fp8 blockscale GEMM / narrow-N qkv_proj flydsl_fp8_blockscale_gemm ~1.05x Pass
Grouped-MoE GEMM / decode stage-1 gate/up flydsl_decode_moe_stage1_blkmap 1.0778x Pass
Grouped-MoE GEMM / prefill stage-2 down/reduce flydsl_prefill_moe_stage2_fp8partial 1.2019x pass

Checks

  • Static validation passes for all submitted skills.
  • index.yaml matches skill frontmatter.
  • git diff --check passes.
  • No stale expert-skill links remain under perf_knowledge/expert_skills.
  • No local validation artifact paths remain in the submitted skill files.
  • No co-author trailers are present in the submitted commits.

…xcd8, 1.23x same-session A/B vs production CK, parity err=0/cos=1.0); flip skill status draft->validated; note FlyDSL API drift + xcd_swizzle byte-frozen snapshot core
…d 1.05x, parity clean); note blockscale core lives in standalone /sgl-wor
…nowledge on cross-backend ports

When a TASK rewrites a kernel into a different backend (any source->target, e.g.
ck->flydsl, triton->tilelang, hip->ck), the source backend card alone doesn't teach
how to author the target. Now:
- tech_lead analyze: kk_refs also includes the target backend card + target
  language authoring how-to under languages/<dir>/ (with id->dir map).
- engineer / deep_engineer: also read/mine the target backend card + target
  language authoring docs on cross-backend rewrites.
- operator cards (scaled_quant_gemm ck+flydsl, fused_moe_grouped_gemm flydsl,
  dense_gemm flydsl): cross-link the FlyDSL GEMM authoring docs.

Fixes: ck->flydsl (and any a2b) migration runs previously only saw the source
backend card and never surfaced target-language general-GEMM authoring knowledge.
Teach the workflow to match and reproduce decode block-pair reuse and prefill FP8-partial traffic reduction on gfx950.
@IzacharyI IzacharyI changed the title Add FlyDSL gfx950 kernel workflow skills Kernel workflow: Add FlyDSL gfx950 skills Jul 31, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant