Kernel workflow: Add FlyDSL gfx950 skills - #386
Open
IzacharyI wants to merge 6 commits into
Open
Conversation
…ftware fp32 scale + XCD swizzle)
…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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Add FlyDSL gfx950 expert skills for GEAK
kernel_workflow.This PR covers two kernel operator families:
These skills are consumed when
use_expert_skills=true. They are advisory priors forkernel_workflowto reproduce and measure, not hard-coded replacements for correctness, parity, or A/B gates.Modifications
perf_knowledge/expert_skills/skills/flydsl_fp8_blockscale_gemm/skill.mdq_up_proj,down_proj, andqkv_projevidence. This is distinct fromperf_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.mdperf_knowledge/expert_skills/skills/flydsl_prefill_moe_stage2_fp8partial/skill.mdperf_knowledge/expert_skills/index.yamlkernel_workflow/roles/{tech_lead,engineer,deep_engineer}.mdperf_knowledge/operators/scaled_quant_gemm/backends/{ck,flydsl}.mdperf_knowledge/operators/dense_gemm/{overview,numerics}.mdperf_knowledge/operators/dense_gemm/backends/flydsl.mdperf_knowledge/operators/fused_moe_grouped_gemm/backends/flydsl.mdperf_knowledge/quantization/block_scaling_mxfp.mdSkill Organization
The submitted skills are kept separate by workflow layer, GPU generation, and kernel scenario.
flydsl_fp8_blockscale_gemmis separate from the existingflydsl_fp8_gemm_playbook: the existing playbook is a gfx942 e2e down-proj overlay focused on capture-safe integration, while the new skill is a gfx950kernel_workflowrecipe for CK/production blockscale GEMM -> FlyDSL candidates.The two Grouped-MoE skills are also kept separate:
flydsl_decode_moe_stage1_blkmaptargets decode / small-M stage-1 gate/up, whileflydsl_prefill_moe_stage2_fp8partialtargets prefill / large-M stage-2 down/reduce. They have different runtime signatures, validation scopes, and parity requirements.Workflow Impact
use_expert_skills=true.Validation
kernel_workflow/ Director-verified resultq_up_projflydsl_fp8_blockscale_gemm~1.226x-1.237xdown_projflydsl_fp8_blockscale_gemm1.4541xqkv_projflydsl_fp8_blockscale_gemm~1.05xflydsl_decode_moe_stage1_blkmap1.0778xflydsl_prefill_moe_stage2_fp8partial1.2019xChecks
index.yamlmatches skill frontmatter.git diff --checkpasses.perf_knowledge/expert_skills.