Repository navigation
feat(hipblaslt): run JIT GEMM solutions through the normal launch path - #13131
jaopaulolc wants to merge 5 commits into
Conversation
Codecov Report❌ Patch coverage is ❌ Your project check has failed because the head coverage (49.03%) 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 @@
## users/jolabega/jit-upstream-01-jit-interfaces #13131 +/- ##
=================================================================================
- Coverage 72.13% 72.13% -0.00%
=================================================================================
Files 2861 2861
Lines 472127 472165 +38
Branches 69811 69823 +12
=================================================================================
+ Hits 340551 340570 +19
- Misses 107831 107839 +8
- Partials 23745 23756 +11
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
8d1d41f to
6c4132a
Compare
3db7dd0 to
73e8602
Compare
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Jit runs one operation request for one device target through four stages: a backend generates library entries, the builder builds them, and a loader checks which of their solutions support the request and loads those, one bundle per solution. The requested count counts solutions, not entries. Each stage is an interface, a failure in one entry skips that entry and is recorded with its stage, and the scratch directory is kept after a failure. The opaque Backend, Request and Solution handles in hipblaslt-jit.hpp carry a Jit, an operation request and a compiled solution.
makeGemmRequest captures a GEMM from the existing descriptors, getJitAlgo generates, builds and loads a solution for it through a backend's Jit, and getGemmAlgo returns a hipblasLtMatmul algorithm that carries the solution's local index and a process-local token for it. get_library_and_adapter resolves the token to the bundle's library and adapter, so the algorithm runs through the same argument and launch path as a prebuilt one, including its synchronization storage. The TensileLite loader returns a TensileGemmBundle for each supported solution of a loaded entry, sharing its library and adapter. Its support check reuses ConstructTensileProblem and the solution's predicates and workspace. With HIPBLASLT_ENABLE_JIT off, the GEMM data keeps the selected algorithm and scale layouts, kernels are stored only once they are ready, a reinitialized handle keeps its synchronizer, running before initialization returns HIPBLAS_STATUS_NOT_INITIALIZED, and kernel-name lookups for unknown indices return an empty name.
HIPBLASLT_JIT_TESTING, off by default, compiles a replay backend into the library. It returns pre-generated source bundles with a solution whose predicates accept the device and the problem instead of running a generator, and the comgr builder still builds them. jit-end-to-end replays the two-solution bundle that jit-bundles writes for device 0's architecture. getJitAlgo builds and loads it and returns its first solution for K=512 and its second for K=256. Each runs an FP16 GEMM through hipblasLtMatmul and hipblaslt_ext::Gemm, and D is checked against a host reference. It also checks that a problem neither solution solves is not supported. jit-disabled checks that a build with JIT off installs no JIT header. The workflow turns on HIPBLASLT_JIT_TESTING and runs for the host files the GEMM path changes.
JIT.md describes the internal entry points, the Jit stages, how they count solutions and their failures, the TensileLite GEMM loader, HIPBLASLT_JIT_TESTING and the replay backend, and the lifetime of a JIT algorithm. The JIT test guide adds jit-end-to-end and jit-disabled.
The replay backend is the simplest Backend, and the shared build and load path is what every other backend uses. HIPBLASLT_JIT_TESTING no longer gates it.
73e8602 to
b37eeb1
Compare
030745a to
b35a2e1
Compare
|
Closing this as a review checkpoint. The branch stays, because #13202 is based on it. Solutions run only as JIT library indices. A query looks up the library first and publishes on a miss. The process-local in-memory algorithm path is not the intended way to run a JIT solution and will be removed on #13202. The pipeline and the replay backend in this branch remain, because a miss still has to build a code object before it can publish. |
JIRA ID : AIHPBLAS-4801
Motivation
hipBLASLt runs each matrix multiplication (GEMM) with a kernel from a library that is generated and tuned at build time, so a problem the library serves poorly has no better kernel at run time. Just-in-time (JIT) generation produces a kernel for a problem when it is needed. The previous PR builds generated kernel sources with comgr and loads the code object into a TensileLite library of one or more solutions. This PR runs those solutions through
hipblasLtMatmulandhipblaslt_ext::Gemmlike a prebuilt one, and it proves the path end to end on a GPU with the committed kernel.Technical Details
The feature is behind
HIPBLASLT_ENABLE_JIT, which is off by default. The commits each build on their own:Jitruns one request for one device through four stages, each an interface: Generate (a backend returns library entries and their kernel sources), Build, Support (each solution's predicates and workspace) and Load. It returns up to the requested count of supported solutions, best first. A failure in one entry skips it and records its stage.makeGemmRequest,getJitAlgoandgetGemmAlgoreturn ahipblasLtMatmulAlgo_tthat carries a process-local token and the solution's local index. The TensileLite loader intensile_host.cpploads an entry once and wraps each supported solution, with the shared library and adapter, in aTensileGemmBundle.get_library_and_adapterresolves the token, and the solution then runs through the normal launch path, including Stream-K's per-stream synchronization. A token from another process or device is rejected before any work is submitted. No installed header declares these entry points.Backendand returns pre-generated source bundles instead of running a generator, andJitbuilds and loads those entries with the same builder and loader as any other backend.JIT.mdand the JIT test guide. The workflow watches the GEMM host files.With JIT off, the shared host code still changes in small ways. The GEMM data keeps the selected algorithm and the scale layouts.
makeArgumentstores its kernels only after they are ready, and it restores the handle's synchronizer on reinitialization.runbeforeinitializereturnsNOT_INITIALIZED. Kernel-name lookups for an unknown index return an empty name.heuristicResult_copycopies the whole algorithm. Two internal status codes are added.Not in this PR: a generator backend and a solution store. Later PRs add them.
Test Plan
Host-only builds with
GPU_TARGETS=gfx950: JIT on with JIT testing, and JIT off. A third host build usesGPU_TARGETS=gfx942with JIT on. In each build, CTest runs the JIT tests inclients/tests/jit. The previous PR addsjit-bundles,jit-source-bundle,jit-builderandjit-loader. This PR adds:jit-end-to-endreplays the two-solutionplain-pairbundle thatjit-bundleswrites from the committed kernel assembly, builds and loads it, and runs M=256, N=128 throughhipblasLtMatmulandhipblaslt_ext::Gemmtwice, checking D against a host reference: K=512 runs the first solution and K=256 the second. It also checks that a transposed-A problem is not supported. It replays the bundle of device 0's architecture, so the gfx90a, gfx942 and gfx950 lanes each run their own kernel. CMake registers it in builds with JIT testing whenGPU_TARGETSinclude an architecture with committed bundles.jit-disabled, in the JIT-off build, checks that the JIT headers are not in the public include tree.A device-library build for gfx950, restricted to one FP16 logic file, ran the JIT tests and the
hipblaslt-testquick suite. The suite was compared with the same build ofdevelop.actionlintchecked the workflow.Test Result
jit-bundles,jit-source-bundle,jit-builder,jit-loader,jit-end-to-endjit-bundles,jit-source-bundle,jit-builderjit-source-bundle,jit-disabledjit-bundles,jit-source-bundle,jit-builder(-L jit-cpu; the build also registersjit-loaderandjit-end-to-end)hipblaslt-testquick suitedevelopand this branch, with identical test namesThe 785 failures report "NO solution found", because the build keeps one logic file. Each commit built with JIT on and passed the JIT tests it registers; the last commit also built with JIT off. There is no gfx90a or gfx942 GPU on the test host, so the workflow's lanes are the first native runs of
jit-loaderandjit-end-to-endon those architectures.actionlint1.7.12 reported no errors. The workflow's three lanes and the other automated GitHub checks are pending.Submission Checklist
Risk level
Low (2/5): the JIT sources build only with
HIPBLASLT_ENABLE_JIT=ON, which is off by default. With JIT off, the changes are the small shared host-code edits listed in Technical Details, and the quick suite matcheddevelop.