feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1#9304
Open
jinchen62 wants to merge 60 commits into
Open
feat(hipblaslt): Gemm-From-Anywhere custom kernel framework V1#9304jinchen62 wants to merge 60 commits into
jinchen62 wants to merge 60 commits into
Conversation
Add code object V6. Fix some build/packaging issues with new custom kernel format. Add DebugPattern argument which puts easy to spot data in the argument buffer for debugging.
Update build scripts to allow subdirectories in CustomKernels folder to help organize files.
Update the way CustomKernel info gets stored in a solution to improve default build, and avoid having to come up with equivalent tuning arguments to describe external kernels.
- Add padding argument semantic
e.g. { type: uint32, semantic: Padding}
- Add padding param to the arg
e.g. { type: address, semantic: AddressD, padding: 8 }
AIGECORE-1985
Wave kernel is compiled from https://github.com/iree-org/wave Tensile log: https://gist.github.com/jinchen62/589e0b44525ccdf0662959169c349d6c
Fix `KeyError: 'MFMA_BF16_1K'` when building install.sh
Triton kernel is compiled from https://github.com/ROCm/aiter/tree/main/aiter/ops/triton/gemm/basic by the [script](https://gist.github.com/jinchen62/4316c7719b7a00bb901a965447c60586). Tensile log https://gist.github.com/jinchen62/7da14a7763e1e2b5594249d79715e777
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report❌ Patch coverage is ❌ Your project status has failed because the head coverage (76.84%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #9304 +/- ##
===========================================
- Coverage 65.36% 65.32% -0.04%
===========================================
Files 2685 2687 +2
Lines 422269 422979 +710
Branches 62742 62922 +180
===========================================
+ Hits 276012 276303 +291
- Misses 125466 125830 +364
- Partials 20791 20846 +55
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
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
4 participants
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.
Motivation
hipBLASLt/TensileLite currently only dispatches kernels generated by Tensile's own assembly writer. Several teams (AITER, Composable Kernel, rocRoller, Wave, Triton, and eventually hipKittens) independently write high-performance GEMM kernels for the same hardware, duplicating tuning effort and fragmenting the AMD GEMM kernel ecosystem instead of pooling it behind one selection/dispatch path. Gemm-From-Anywhere (GFA) lets hipBLASLt load and run kernels from any of these sources through the existing Tensile solution-selection and device-library loading path, so the fastest available kernel for a given problem shape can come from anywhere, not just Tensile.
This is "Version 1": the foundational metadata format, host-code dispatch path, and one working demo kernel per source generator (see Confluence: Hipblaslt-kernel-from-anywhere). Follow-on milestones (hipKittens, full production kernel sets, JIT generation) are tracked separately and are out of scope here.
Technical Details
.sfile embeds acustom.configYAML block in its.amdgpu_metadatasection, carrying provenance (Source.Origin/Repository/Version), aFeaturescapability map, and — for external (non-Tensile) kernels — the full Tensile-side interface (args,macrotile,threads,grid,ProblemType,MatrixInstruction). SeeTensile/CustomKernels/README.md.Tensile.AddCustomConfig(new): extracts a kernel's interface from a Tensile benchmark test YAML and mechanically injects thecustom.configblock into the.sfile.Tensile.ValidateMetadata(new): CI-gate script (--strict) that walksCustomKernels/and validates every kernel's embedded metadata; the same check is also available at build time viaGlobalParameters.ValidateMetadata/--validate-metadata.ContractionSolution/KernelArgumentsextended so the launch argument buffer is built generically from a kernel's declaredargs:list instead of requiring a hand-written C++ solution class per kernel.SolutionStructs/Solution.py,Naming.py, andContractions.pyextended withAssertFree0/1ElementMultiple/AssertSummationElementMultiple(now up to 256) so a custom kernel can declare exactly which problem shapes it supports.CustomKernels/split into per-origin subdirectories (tensile/,aiter/,ck/,rocroller/,wave/,triton/); loader behavior is unchanged, the split is organizational only.Tensile/Tests/custom/*.yaml) for: AITER bf16 GEMM, AITER MXFP4 GEMM (pre-shuffled scales), Composable Kernel bf16 GEMM, rocRoller fp16 GEMM, Wave bf16 GEMM, Triton MXFP4 GEMM, and a Tensile-native Stream-K custom kernel.roc::mxDataGeneratorwired intotensilelite's CMake for MX-type reference-data generation used by the MXFP4 demos.HandleCustomKernel.hasCustomKernel()for legacy-style logic-file entries, and an unhandled missing-file crash in the newToolchain.Component.Assembler._retargetAssemblySource.Test Plan
Tensile/Tests/unit/test_CustomKernelMetadata.py(AddCustomConfigextraction/injection,ValidateMetadata, predicate round-tripping) and extendedtest_custom_kernel_occupancy.py(occupancy over the new real gfx942/gfx950 custom kernels).Tensile/Tests/custom/build and dispatch each demo kernel end-to-end through the realCustomKernels:fork-parameter path withValidateMetadata: Trueenabled.python -m Tensile.ValidateMetadata --strictover the fullCustomKernels/tree as the metadata CI gate.pytest -m unit(the entire TensileLite Python characterization + unit suite) as the regression net for everything underTensile/.developbaseline checkout on the same gfx942/gfx950 hardware, running the same suites on both: the fullhipblaslt-testC++ client gtest suite, thetensilelite-testsC++ host-library gtest suite, and thecommon/gemm/*.yamlintegration tests — to separate any pre-existing issue from anything this PR introduces.Test Result
Ran locally against gfx942/gfx950 as part of preparing this PR, including a full side-by-side comparison against a
developbaseline checkout on the same hardware:python -m Tensile.ValidateMetadata --strict→ 126 kernels, 0 errors, 0 warnings.pytest -m unit→ 5368 passed, 0 failed, 222 skipped, confirmed byte-identical across two clean re-runs (cache cleared each time).hipblaslt-test(full C++ client gtest suite): 79102 run, 79096 passed, 6 skipped, 0 failed — identical ondevelopandGemmFromAnywhere, including the exact same 6 skipped tests.common/gemm/*.yaml(58 files): 52 PASS / 6 fail on both branches, same 6 files failing with "0 valid solutions" (f32_mac,f64_gfx90a,f64_mac,i8_gsu,gemm_bias_strided_batched[_broadcast]) — confirmed pre-existing and unrelated to this PR.custom/custom_{aiter,ck,rr,tensile_sk,triton,wave}.yaml(new, GFA-only, no equivalent exists ondevelop): all 7 PASS.tensilelite-tests(C++ host-library gtest suite): 349 passed, 1 failed, 3 skipped — identical on both branches, including the same pre-existing failure (RangeLibraryTest.SpecificSizes/0, a missingnonTemporalAkey inKernels_gfx950.yaml) and the same 3 skips.Risk level
Medium-High. Reasons:
develop, blowing through the Libraries PR Bot's size thresholds (50 files / 2000 changes / 700-per-file) — expect a non-blocking size warning. ~296K of the insertions are checked-in kernel.sfiles, not hand-written logic; the reviewable source surface is closer to 50 files.KernelWriter.py,KernelWriterAssembly.py,ContractionSolution.cpp) — per this repo'shipblaslt-pr-qualityskill these need 2 hipBLASLt code-owner approvals and a rebase+re-test ifdevelophas moved on them since this diverged (checked at PR-update time:develophad not moved since this branch's last merge).CustomKernel:-listed kernel), so existing Tensile-generated-kernel selection should be unaffected.Submission Checklist