Conversation
Replant of the custom INT4 GEMV work onto develop. The matcher is lifted out of fuse_mlir.cpp into its own pass so it is available in non-MLIR builds; it runs immediately before fuse_mlir so MLIR does not claim the dot first. Opt-in via MIGRAPHX_ENABLE_INT4_GEMV=1.
…iene
Selection order for the M=1 INT4 GEMV launch geometry:
1. MIGRAPHX_INT4_GEMV_CONFIG override (global BS,TN or per-shape N:K=BS,TN)
2. measured config table, keyed on (gfx arch, N, K)
3. get_tuning_config -- compile_ops benchmarks the candidate grid, seeded
with the heuristic pick as solutions.front()
Also:
- trace output moved off std::cerr onto MIGRAPHX_INT4_GEMV_TRACE
- scalar int4_gemv_dot2 fallback behind MIGRAPHX_INT4_GEMV_HAS_FDOT2
- the three env vars documented in docs/reference/MIGraphX-dev-env-vars.rst
- clang-format (22.1.5, the version CI pins) clean on all touched files
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
test/gpu/fuse_int4_gemv.cpp, following the test/gpu/fuse_*.cpp convention -- the
glob in test/CMakeLists.txt picks it up as test_gpu_fuse_int4_gemv, so no build
file change is needed.
Two positive cases (asymmetric and symmetric, since the matcher branches on
zero-point presence) and four negative ones. The negative cases carry the
weight: a matcher that fired unconditionally would satisfy every positive
assertion, so M=8, a {4,1,K} batch, and a plain fp16 dot are each asserted to be
left for the default codegen path. The positive cases also assert the original
dot is gone, not just that the fused op appeared.
The file-scope environment initialiser is load-bearing rather than incidental.
The fusion is opt-in behind MIGRAPHX_ENABLE_INT4_GEMV, and migraphx::enabled()
memoises the lookup in a function-local `static const bool` -- frozen at the
first query in the process -- so a setenv() inside a TEST_CASE body would come
too late to matter. Setting it before main is what makes the assertions
meaningful; without it the pass early-returns and all six cases compare unfused
against unfused, passing green while testing nothing.
NOT YET COMPILED -- the build host is occupied. Committed so the work is not
lost; expect include-path and const-correctness fixes on first build.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
`batch` in apply() was assigned from a_shape and never read. MIGraphX CI compiles with -Werror -Weverything, so this one dead line failed every job that builds the GPU target -- All Targets Release, HIP Clang Release, Navi32, Navi4x, HIP Clang Static, HIP RTC Debug and MLIR Debug -- while the jobs that do not build the GPU target all passed. It does not reproduce in a Windows clang-cl build, which is how it got here. No behaviour change: the variable is a std::size_t read out of an already-constructed shape, never used, and -Wunused-variable can only fire on a variable nothing references. The identically-named `batch` in compute_shape is a different scope and is still used to build the output shape. Verified by compiling the translation unit with -Wunused-variable -Wunused-but-set-variable -Werror: fails on the line (error at 337:14), passes with it removed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
aditya-dl
force-pushed
the
int4-gemv-upstream
branch
from
September 17, 2026 20:58
c48dfae to
d1b53d2
Compare
test/gpu/fuse_int4_gemv.cpp was hand-written and never run through clang-format, which failed the `format` CI job. Two cosmetic fixes: a run_passes call that fits on one line, and default-argument alignment in make_int4_gemv_program. No behaviour change. The four source files in this PR were already clean; only the test was not. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The kernel's inner loop is __builtin_amdgcn_fdot2 (v_dot2_f32_f16). Previously a preprocessor guard selected a scalar multiply-add pair on targets lacking that instruction. That fallback was never compiled, and on such a target it would have been slower than the MLIR path this fusion replaces -- so it turned "no fusion" into "a fusion that loses", silently. The pass now takes a context and declines the match outright when the target has no fdot2, which is the same shape as the other five gates. The check is at pass level rather than per dot, since the architecture does not change between dots. Verified with __has_builtin on the device compilation pass: present on gfx1151, gfx1201 and gfx90a; absent on gfx900 and gfx1010. No change for any target we measure on. Removes MIGRAPHX_INT4_GEMV_HAS_FDOT2 and the untested branch it guarded. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
target.cpp now wraps the pass in enable_pass(enabled(MIGRAPHX_ENABLE_INT4_GEMV{}),
...), the same idiom already used for fuse_ck. Previously the pass was
registered unconditionally and returned early on the env check, so it was
constructed and run on every compile just to do nothing. With this it is not
constructed at all unless asked for.
The in-pass env check stays: the pass can be constructed directly, as the test
does, and the check documents the contract at the point it applies.
Also fixes test/gpu/fuse_int4_gemv.cpp, which constructed fuse_int4_gemv{} with
no context. Since the fdot2 gate added in the previous commit declines when
ctx is null, every positive expectation in that test would have passed
trivially by matching nothing. It now builds a real gpu::context, mirroring
test/gpu/fuse_gemm.cpp.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
aditya-dl
added a commit
to aditya-dl/onnxruntime-ep-amdgpu
that referenced
this pull request
Sep 18, 2026
MIGraphX gates its custom INT4 M=1 GEMV kernel behind MIGRAPHX_ENABLE_INT4_GEMV and leaves it off by default (ROCm/AMDMIGraphX#5272). Set it alongside the existing MIGRAPHX_MLIR_USE_SPECIFIC_OPS so the kernel is available to the EP, while MIGraphX used outside this process stays on its default path. Scope: SetEnvironmentVariable is process-global, so this enables the fusion for every MIGraphX consumer in the process, not only the OGA decode path. That is acceptable here because the EP is the MIGraphX consumer, but it is wider than "the decode path" and worth stating plainly. What actually limits it is the matcher, not this variable. MIGraphX only rewrites a dot when the weights are INT4 weight-only quantized, the output is M=1, shapes are static, and the target has fdot2. There is no "is this an LLM" check. In practice INT4 weight-only quantization is an LLM technique, so non-LLM models are not expected to reach the second gate -- but that is an expectation, not an enforced restriction. A batch-1 INT4 non-LLM model would match. If that ever needs to be prevented, the gate belongs at the OGA layer, which is the only layer that knows it is serving an LLM. Set in CreateEpFactories rather than later: MIGraphX reads the variable through enabled(), which memoises into a function-local static on first query, so a later write can be ignored. Windows only, matching the existing MIGRAPHX_MLIR_USE_SPECIFIC_OPS line in the same block. The kernel will therefore not fire through a Linux EP build. That is deliberate here -- following the file's convention rather than diverging from it -- and closing the gap is better done as one change covering both variables. Measured on STX-Halo (gfx1151), short prompt, 128 output tokens, 5 iterations after 2 warmups, against the DirectML backend on the same harness: Qwen2.5-Coder-0.5B RTN b32 380.83 tok/s (DML 321.70) Qwen2.5-0.5B RTN b32 336.50 tok/s (DML 298.93) Llama-3.2-1B AWQ b128 228.62 tok/s (DML 184.11) Mistral-7B-v0.2 SYM b128 57.04 tok/s (DML 50.67) Phi-4 14B RTN b32 27.76 tok/s (DML 6.08) Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
aditya-dl
added a commit
to aditya-dl/onnxruntime-ep-amdgpu
that referenced
this pull request
Sep 18, 2026
MIGraphX gates its custom INT4 M=1 GEMV kernel behind MIGRAPHX_ENABLE_INT4_GEMV and leaves it off by default (ROCm/AMDMIGraphX#5272). Set it from the umbrella factory so the kernel is available to this EP, while MIGraphX used outside this process stays on its default path. Set in gpu_factory.cc immediately before LoadDynamicLibrary(migraphxBackend), NOT in migraphx-backend.dll. MIGraphX reads the variable with std::getenv from src/env.cpp, which builds into migraphx.dll, and migraphx.dll is reached through a chain of static imports from migraphx-backend.dll -- so its CRT has already seeded its own copy of the environment before any code in that dll runs. The EP builds with a static CRT (/MT), so nothing written afterwards can reach it. amdgpu-ep.dll loads the backend dynamically, so code above that line runs before migraphx.dll is in the process at all. This was measured, not assumed: setting the variable from CreateEpFactories in migraphx-backend.dll does not work. A probe there confirmed the line executes and that GetEnvironmentVariableA reads back "1", while MIGraphX's own trace still reported the variable unset and the fusion never fired. Scope: SetEnvironmentVariable is process-global, so this enables the fusion for every MIGraphX consumer in the process. What narrows it is the matcher, which only rewrites INT4 weight-only M=1 dots with static shapes on a target with fdot2. There is no "is this an LLM" check. No off switch: enabled() treats any present value as on, so MIGRAPHX_ENABLE_INT4_GEMV=0 will not override this. An escape hatch would need a separate disable option on the EP side or a MIGRAPHX_DISABLE_* read upstream. Windows only, matching the existing MIGRAPHX_MLIR_USE_SPECIFIC_OPS line in mgx_factory.cc. Note that line has the same defect this change avoids and has therefore never taken effect; fixing it is a separate change. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
1 participant
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.
Draft — opening early for design feedback. Build/test gates are still being run; see Status below.
What
A standalone fused kernel for the M=1 INT4 GEMV that dominates LLM decode: one weight-only-quantized matrix-vector product per projection, per layer, per token. Off by default, behind
MIGRAPHX_ENABLE_INT4_GEMV=1.Two commits:
int4_gemv: standalone M=1 INT4 GEMV fusion pass, kernel and JIT compiler— the matcher, the kernel, the JIT.int4_gemv: three-tier config selection, autotune hookup, docs and hygiene— launch-geometry selection, tracing, docs.Config selection
The launch geometry is
(block_size, tile_n). Three tiers, in order:MIGRAPHX_INT4_GEMV_CONFIG— explicit override, either a globalBS,TNor per-shapeN:K=BS,TN.(gfx arch, N, K).get_tuning_config→compile_opsbenchmarks the candidate grid, seeded with the heuristic pick assolutions.front().Tiers 1 and 2 are trusted and short-circuit; tier 3 is what an unknown shape falls into, so a shape with no table row still gets a tuned config rather than a guess. The table short-circuit is
trusted and not exhaustive, soORT_MIGRAPHX_EXHAUSTIVE_TUNE=1forces a search even on a shape that has a row.Candidate grid: non-exhaustive
BS ∈ {64,128,256} × TN ∈ {2,4,8}(9 configs); exhaustiveBS ∈ {32,64,128,256,512} × TN ∈ {2,4,8,16}(20).Note on
--exhaustive-tune: it widens MLIR program-wide as well (we observed >1000 configs benchmarked pergpu::mlir_op), which makes it impractical on a whole model. It is not the recommended path here — the default autotune is.Kernel
__builtin_amdgcn_fdot2for the packed fp16 dot product, guarded byMIGRAPHX_INT4_GEMV_HAS_FDOT2with a scalar fallback for targets withoutv_dot2_f32_f16.dwordx4weight loads;_Float16activations; warp-shuffle reduction.Env vars
Documented in
docs/reference/MIGraphX-dev-env-vars.rst:MIGRAPHX_ENABLE_INT4_GEMVMIGRAPHX_INT4_GEMV_CONFIGMIGRAPHX_INT4_GEMV_TRACETrace output moved off
std::cerrontoMIGRAPHX_INT4_GEMV_TRACE.Status / what is still open
test/verify/case for the fused path.MIGRAPHX_INT4_GEMV_HAS_FDOT2=0fallback path — every build so far has been on a target withv_dot2_f32_f16, so that branch is untested.clang-format (22.1.5, the version CI pins) is clean on all touched files.
Feedback on the config-selection structure and on where the fusion pass belongs in the pipeline is especially welcome while this is still a draft.
🤖 Generated with Claude Code