[CUDA] Add persistent fpA_intB MatMulNBits tactic autotune cache by tianleiwu · Pull Request #29588 · microsoft/onnxruntime

Tianlei Wu (tianleiwu) · 2026-07-07T03:50:24Z

Description

Adds an opt-in, hardware-keyed on-disk tactic cache for the fpA_intB weight-only MatMulNBits path, plus a reduced first-time profiling sweep with lazy fallback. Persistence is off by default: with no cache location configured, behavior is unchanged (in-process only).

Cache subsystem (gemm_tactic_cache.{h,cc})

HardwareSignature (GPU name, SM, CUDA runtime, ORT version; commit/build config/driver/MP count recorded for diagnostics) names and guards cache files, so tactics are never reused across incompatible GPUs/toolkits/ORT versions. A strict tactic_selection_version header field rejects caches tuned under an older selection policy, even when the ORT version is unchanged.
TSV format: percent-encoded free-text fields, a column-name header (appended columns are backward/forward compatible), file locking + atomic write through a per-writer temp file and rename (MoveFileExW on Windows; paths are UTF-8), in-memory merge on flush so concurrent sessions tuning different shapes don't lose updates, and creation of a missing cache directory on first write.
MatMulNBitsTacticCache keyed by (N16b, K, act dtype, weight type, bits, block_size, zero-points, gemv_enabled, packing_sm). packing_sm keeps the SM80-compat and native-SM90 tactics for the same shape distinct on one device (e.g. H200). 2-, 4-, and 8-bit weights are keyed separately.

Profiler (gemm_profiler.h, fpA_intB_gemm_profiler.{h,cc})

Reduced initial sweep to a small M bucket set ({1,2,4,…,2048} default; override with ep.cuda.fpa_intb_profile_m / ORT_FPA_INTB_PROFILE_M). getBestConfigOrProfile lazily profiles any missing bucket on first use.
Persistent-cache load/store hooks are no-ops unless a cache is attached.
CUTLASS tactics loaded from disk are validated against the runner's current getConfigs() and dropped (re-profiled) if incompatible.

Wiring (matmul_nbits.{h,cc}, cuda_execution_provider.cc, plugin cuda_ep.cc)

Cache location: session config (ep.cuda.gemm_tactic_cache_prefix / _dir) takes precedence over env (ORT_CUDA_GEMM_TACTIC_CACHE_PREFIX / _DIR); a prefix wins over a dir.
Process-global caches are keyed by resolved file path; the file is loaded once per location (not once per node), and the signature is built from the kernel's cached device properties.
The process-global in-memory profile map (GemmIdCore) is keyed by the same quantization fields as the persistent key (bits, block size, zero points, GEMV support), so kernels with different quantization never share tactics.
Construction-time sweeps and lazily profiled buckets are staged in memory as soon as they are profiled, and all disk writes are batched at CUDA EP teardown (~CUDAExecutionProvider, or ~CudaEp for the plugin EP), so there is no disk I/O on the inference path.
Lazy profiling is skipped while the compute stream is capturing a CUDA graph; run a warmup inference before capture.

Offline tuning tool: onnxruntime/python/tools/fpa_intb_tune.py (uses session config entries; releases the session before summarizing so lazily tuned buckets are flushed).

Testing

gemm_tactic_cache_test.cc (built into onnxruntime_providers_cuda_ut with onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON, run via --gtest_filter=CUDA_EP_Unittest.All): TSV round-trip and malformed escapes, SM80/SM90 config serialize/parse, null tactic, store/load, signature-mismatch rejection, non-strict commit/build config, selection-version rejection, appended-column tolerance, concurrent flush-merge, missing-directory creation, file path resolution, and GemmIdCore separation of quantization variants.
*MatMulNBits* CUDA integration tests.

Copilot

Pull request overview

This PR adds an opt-in, hardware-keyed persistent on-disk tactic autotune cache for the CUDA fpA_intB MatMulNBits path, and updates the profiler to reduce the initial M-bucket sweep with a lazy single-bucket fallback (with CUDA graph capture safety). It also includes a Python offline tuning tool, unit tests, and documentation for the cache format and configuration.

Changes:

Introduce gemm_tactic_cache utilities (hardware signature, TSV encoding/decoding, config serialization, atomic/locked merge flush).
Update GemmPluginProfiler / fpA_intB profiler to support reduced initial profiling + lazy per-bucket profiling and persistent-cache hooks.
Wire cache resolution into MatMulNBits and add an offline tuner + docs/tests.

Reviewed changes

Copilot reviewed 11 out of 11 changed files in this pull request and generated 3 comments.

Show a summary per file

File	Description
onnxruntime/test/contrib_ops/cuda_kernels/gemm_tactic_cache_test.cc	New pure-logic unit tests for TSV/config round-trip and cache merge/signature behavior.
onnxruntime/python/tools/fpa_intb_tune.py	New offline tuning CLI tool that generates a persistent cache by constructing a CUDA session and optionally running dummy inferences.
onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.h	Use new profiler max-M logic for the initial sweep.
onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.cc	Add process-global persistent cache wiring + CUDA graph capture guard for lazy profiling.
onnxruntime/contrib_ops/cuda/llm/gemm_tactic_cache.h	New cache API: hardware signature, TSV helpers, MatMulNBits cache key and disk-backed cache wrapper.
onnxruntime/contrib_ops/cuda/llm/gemm_tactic_cache.cc	New cache implementation: TSV IO, signature checks, file locking, merge + atomic replace.
onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h	Add reduced-bucket sweep support, persistent-cache hooks, and lazy `getBestConfigOrProfile`.
onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemm_profiler.h	Add env-controlled M-bucket override, persistent-cache attachment, and overrides for new hook points.
onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemm_profiler.cc	Implement reduced bucket set, persistent cache load/store, and stale-tactic validation.
docs/contrib_ops/cuda/matmul_nbits.md	Document cache enablement, reduced sweep, lazy profiling, CUDA graph capture constraints, and tooling/tests.
docs/contrib_ops/cuda/gemm_profiler_cache.md	New design/spec doc describing the cache format, signature strictness, and phased plan.

…peed up MatMulNBits GEMV (#32876) ## Description Speculative decoding verifies a small batch of rows per step (8 for DFlash2 width 7), which lands two CUDA decode paths on poor kernels. Which kernel is best for these small-M shapes depends on the GPU: on SM 8.x cuBLAS picks a serial split-K kernel that is 10x+ slower than the small-N GEMV, on SM 12.0 the GEMV only helps at M = 2..8, and on SM 9.0 cuBLAS is as fast or faster. A fixed architecture policy regressed some of these (e.g. RTX 3060 at M = 9..64, N = 1024, up to 4.1x slower), so this PR selects the kernel per shape and per device by measuring it. 1. **GEMM auto-tuning for fp16/bf16 `MatMul` (opt-in).** With session config `ep.cuda.enable_gemm_auto_tune=1` (or env `ORT_CUDA_GEMM_AUTO_TUNE=1`), the first run of each eligible shape times cuBLAS and the small-N GEMV on the current device and caches the faster one for the process. When disabled (default), MatMul uses cuBLAS as on `main`. 2. **bf16 small-N GEMV.** The GEMV is templated on fp16/bf16 (fp32 accumulation, deterministic split-K). 3. **MatMulNBits fpA_intB GEMV at M >= 4.** The M x StepK activation tile plus the CtaM x CtaN accumulators push the 4/8-bit kernel to about 250 registers; halving CtaN frees registers and doubles the block count. The existing fpA_intB tactic profiler still chooses GEMV vs CUTLASS per M bucket. ## Summary of Changes ### GEMM auto-tuner | File | Change | |------|--------| | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | New key `ep.cuda.enable_gemm_auto_tune` (`kOrtSessionOptionsCudaEnableGemmAutoTune`). | | `onnxruntime/core/providers/cuda/math/gemm_auto_tuner.{h,cc}`, `gemm_auto_tuner_impl.cu` | Dispatch policy resolution, candidate timing, and a process-wide cache keyed by (device UUID, dtype, M, N, K, operand-alignment class). | | `onnxruntime/core/providers/cuda/math/matmul.{h,cc}` | `MatMul<MLFloat16/BFloat16>` resolves the policy at construction and, for eligible single GEMMs, looks up / tunes the kernel. Ineligible shapes (M > 64, N > 1024, K < 128, transposes, alpha != 1, batched) go straight to cuBLAS with no overhead. | Dispatch policy, in precedence order: | Setting | Behavior | |---|---| | `ORT_ENABLE_SMALL_N_GEMV=1` / `0` | Forces the small-N GEMV (when eligible) / cuBLAS; bypasses tuning. Kept for A/B testing and back-compat with the opt-in from #31478. | | session config `ep.cuda.enable_gemm_auto_tune` | `1` auto-tunes, `0` uses cuBLAS. | | env `ORT_CUDA_GEMM_AUTO_TUNE` | Fallback when the session config is unset. | | none | cuBLAS (same as `main`). | Tuning details: - **CUDA graph safe.** Tuning synchronizes the stream, so it is skipped while the stream is capturing; an untuned shape then uses cuBLAS and is not cached. ORT's warm-up runs before capture do the tuning. - **Representative timing.** Each timed run starts from a decode-like L2 state: a read-only pass over a scratch buffer (2x L2, capped at 256 MiB) evicts the weights without leaving dirty lines to write back, then the activation is read back in, since in decode it was just produced by the previous op (with a cold activation the H200 decisions disagreed with CUDA graph replay). A GPU delay kernel is queued first so host launch gaps of multi-launch candidates are hidden, candidates are interleaved, the fixed event-pair overhead (measured with an empty region) is subtracted, and the median of 10 runs is used. - **Stable choice.** cuBLAS is replaced only if the candidate is at least 5% faster. The cache is process-global (first insertion wins, tuning serialized by a mutex), so every session in a process runs the same kernel for the same shape. Each decision is logged at VERBOSE. - Cost: about 3 ms once per eligible (M, N, K) shape, M <= 64. ### Small-N GEMV | File | Change | |------|--------| | `onnxruntime/core/providers/cuda/math/matmul_small_n_gemv.cu`, `.h` | Covers M <= 64 in ordered 8-row chunks. Adds `SmallNGemvVecSplitKKernel` for even N, K % 8 == 0 and 16-byte aligned A. Each lane owns two columns and eight consecutive K rows, so one 16-byte broadcast A load feeds 16 FMAs; warps reduce through shared memory, then a deterministic last-block split-K reduction runs over `st.cg`/`ld.cg` partials. Completion counters are cleared by a one-block kernel instead of `cudaMemsetAsync` (a memset node costs several microseconds inside a CUDA graph). Templated on fp16/bf16. Every eligible vectorized shape splits K at least twice, so the unreachable single-slice store path was removed. | ### MatMulNBits fpA_intB GEMV | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` | For M >= 4, 4/8-bit weights (StepK < 64) use `CtaNLargeM = CtaN / 2` (8 -> 4 without zero points). M = 1..3 and the 2-bit layout are unchanged. | ## Testing - CUDA internal tests (`onnxruntime_providers_cuda_ut`, H200): - `GemmAutoTunerTest.*`: policy precedence and parsing, candidate selection margin, first-insertion-wins cache, timing/selection with synthetic GPU-delay candidates, cached keys not re-timed, CUDA graph capture detection. - `MatMulSmallNGemvTest.*`: fp16 and bf16 for M = 1..8, 17, vectorized/scalar variants, column-tile boundaries, stale counters cleared by the launcher, vectorized-kernel selection. - `MatMulSmallNGemvOpTest.*`: fp16 and bf16 MatMul through the CUDA EP, both forced and auto-tuned, for eligible shapes (M = 1..64, N = 1..1024) and ineligible fallbacks (K = 127, M = 65, N = 1025). - `onnxruntime_provider_test --gtest_filter='*MatMul*:*Gemm*'` (1003 tests) passes with and without `ORT_CUDA_GEMM_AUTO_TUNE=1`. - Python `test_matmul_gemm_auto_tune_cuda_graph_replay`: fp16 MatMuls tuned during warm-up, then captured and replayed with changing inputs. - Compiled for sm_86 and sm_90 in both the in-tree and the plugin CUDA EP builds. - H200 (SM 9.0), CUDA graph replay of enough MatMul copies that weights stream from DRAM, per-MatMul us, K = 5120, best of two interleaved sessions per mode. Auto-tune picks the GEMV only where it wins, and is within 1% of the faster kernel everywhere except one near-tie inside the 5% margin (bf16 4x48). Decisions were identical across five separate processes. | dtype | M | N | cuBLAS | forced GEMV | auto-tune (choice) | |---|---:|---:|---:|---:|---| | fp16 | 1 | 48 | 6.25 | 6.23 | 6.23 (GEMV) | | fp16 | 2 | 48 | 6.95 | 6.21 | 6.20 (GEMV) | | fp16 | 4 | 48 | 6.94 | 6.72 | 6.72 (GEMV) | | fp16 | 8 | 48 | 7.00 | 7.57 | 6.99 (cuBLAS) | | fp16 | 16 | 48 | 7.31 | 13.03 | 7.31 (cuBLAS) | | fp16 | 8 | 1024 | 8.17 | 13.60 | 8.19 (cuBLAS) | | fp16 | 16 | 1024 | 8.33 | 22.62 | 8.29 (cuBLAS) | | bf16 | 2 | 48 | 6.95 | 6.31 | 6.31 (GEMV) | | bf16 | 4 | 48 | 6.94 | 6.72 | 6.96 (cuBLAS) | | bf16 | 8 | 1024 | 8.28 | 13.89 | 8.17 (cuBLAS) | Tuning adds about 3 ms to the first run of a session per eligible shape. - Previous results for the kernels themselves (RTX 4090, sm_89, CUDA 13.3, nsys, weights streamed from DRAM), which auto-tuning now reaches when it selects the GEMV: | Shape (M x N x K) | cuBLAS / before | GEMV / after | |---|---:|---:| | fp16 MatMul 8 x 48 x 5120 | 78.6 us | 4.2 us | | fp16 MatMul 8 x 1024 x 5120 | 8.4 us | 7.3 us | | MatMulNBits GEMV 8 x 34816 x 5120 | 117.1 us | 112.0 us | | MatMulNBits GEMV 8 x 10240 x 5120 | 40.4 us | 35.7 us | | MatMulNBits GEMV 8 x 5120 x 17408 | 60.0 us | 57.3 us | | MatMulNBits GEMV 8 x 1024 x 5120 | 10.4 us | 8.9 us | End-to-end on RTX 4090 (ORT GenAI, CUDA plugin EP, Qwen3.8-27B INT4 + DFlash2 width 7, greedy, 1,024 tokens) the GEMV + MatMulNBits changes gave x1.28-x1.31 decode throughput (8/8 prompts per context) with unchanged TTFT and MMLU-Pro (88/100). With this revision, set `ep.cuda.enable_gemm_auto_tune=1` in the model's session options to get the MatMul part of that gain. ## Motivation and Context - The small-N GEMV (#31478, extended to 64 rows in #32289) was opt-in because on SM 9.0 cuBLAS was faster in the real model, while on SM 8.x cuBLAS is 10x+ slower at M = 2..8. An architecture default (the first revision of this PR) regressed RTX 3060 above M = 8 (4.1x at M = 64, N = 1024) and was neutral on RTX 5060 Ti; per-shape measurement on the actual device avoids hand-maintained thresholds. - A dedicated tuner is used instead of TunableOp: TunableOp is not built into the CUDA plugin EP, has no CUDA graph capture guard, times with a hot L2, sleeps 50 ms per tuned shape, and is enabled EP-wide. - Candidates are a table, so follow-ups can add kernels (e.g. cuBLASLt heuristic algorithms, CUTLASS, TRT-LLM `cudaCoreGemm`/tinygemm2, FlashInfer TGV on Blackwell) and a persistent TSV tuning cache like #29588. - The follow-up #32884 removes the remaining memset nodes from the decode graph (fpA_intB split-K, paged XQA, causal conv) and makes the fpA_intB tactic profiler prefer the GEMV on near ties at small M. ## Checklist - [x] Tests added/updated - [x] No breaking changes. Without the new option MatMul uses cuBLAS as on `main`; `ORT_ENABLE_SMALL_N_GEMV` keeps its meaning. - [x] Documentation updated (session option key comment) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… on SM 9.0+ (#32901) ## Description Adds TensorRT-LLM's `tinygemm2` kernel as a third candidate for the CUDA fp16/bf16 `MatMul` auto-tuner introduced in #32876 (`ep.cuda.enable_gemm_auto_tune=1`). On H200 it makes decode-sized projections up to 2.2x faster than cuBLAS under CUDA graph replay, and the tuner keeps cuBLAS wherever tinygemm2 is slower. Without the session option, behavior is unchanged. **Stacked on #32876** (base branch `tlwu/cuda-small-m-decode-gemv`); only the last commit belongs to this PR. ## Summary of Changes | File | Change | |------|--------| | `onnxruntime/core/providers/cuda/math/tinygemm2_kernel.cuh` | Kernel adapted from TensorRT-LLM `cpp/tensorrt_llm/kernels/tinygemm2` (Apache 2.0, header retained). A CTA computes 16 output columns x 8 rows; 4 compute warps run `mma.sync.m16n8k16` with fp32 accumulation, 4 warps TMA-load B and 4 load A through a 16-stage mbarrier pipeline. | | `onnxruntime/core/providers/cuda/math/tinygemm2.{h,cu}` | Eligibility, per-device capability probe, tensor-map encoding and launch with programmatic dependent launch. | | `onnxruntime/core/providers/cuda/math/matmul.{h,cc}` | Eligible candidates (small-N GEMV, tinygemm2) are collected into a mask; the tuner times cuBLAS against every eligible candidate. `ORT_ENABLE_SMALL_N_GEMV` keeps forcing only the small-N GEMV. | | `onnxruntime/core/providers/cuda/math/gemm_auto_tuner.{h,cc}` | `GemmKernel::kTinyGemm2`; the tuning key includes the eligible-candidate mask, since operand alignment can change it. | | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | Documents the new candidate. | Changes from the TensorRT-LLM kernel: - **B is read in the ONNX `MatMul` `[K, N]` layout in place.** TensorRT-LLM expects `[N, K]` weights. The B tile is loaded as 64 K-rows x 16 columns with a 32B TMA swizzle and read with `ldmatrix.x4.trans`, so no transposed weight copy (and no extra memory) is needed. - fp16 in addition to bf16 (only the MMA instruction and output conversion differ). - No bias; raw `mbarrier` PTX instead of `cuda::barrier`; the compute-warp reduction uses a named barrier for its 128 threads instead of a `__syncthreads()` that the loader warps never reach. Eligibility: `1 <= M <= 64`, `N * K <= 32M` elements, `N % 8 == 0` and `K % 8 == 0` (TMA row pitch), 16-byte aligned A and B, no transposes, `alpha == 1`. Device support is probed once per device: the kernel's PTX version must be >= 9.0 (a binary without SM 9.0+ code, or JIT from older PTX, has an empty kernel body) and its ~195 KiB of shared memory must fit `sharedMemPerBlockOptin` (so e.g. SM 12.0 consumer Blackwell is excluded). The kernel is launched with programmatic dependent launch (CUDA >= 12.3): the weight loads overlap the preceding kernel, and only the activation loads wait for it via `cudaGridDependencySynchronize()`, which also orders the output stores. ## Testing - Standalone check of the ported kernel against a float64 reference for fp16 and bf16, M in {1,3,8,9,17,64} x N in {8,16,24,48,136,1000} x K in {8,64,136,720,1032,5120}, plus 4x2880x2880 and 64x4096x5120: all within output rounding. `compute-sanitizer` memcheck, synccheck and racecheck report 0 errors. - CUDA internal tests (`onnxruntime_providers_cuda_ut`, H200): new `TinyGemm2Test.Eligibility` and `TinyGemm2Test.MatchesReference` (partial row/column tiles, K below one tile and not a multiple of the 1024-element K loop, NaN-poisoned output), plus the existing `GemmAutoTunerTest.*` / `MatMulSmallNGemv*` suites with tinygemm2-eligible operator shapes (N = 2880 and 1032) added: 17 passed. - `onnxruntime_provider_test --gtest_filter='*MatMul*:*Gemm*'` (1003 tests) passes with and without `ORT_CUDA_GEMM_AUTO_TUNE=1`; Python CUDA graph tests pass. - CUDA plugin EP build: auto-tune selects tinygemm2 and CUDA graph capture/replay (with the programmatic launch edges) produces correct results. - Compiled for sm_86 and sm_90a in both the in-tree and plugin builds. H200, CUDA graph replay of enough MatMul copies that weights stream from DRAM, K = 5120, per-MatMul us, best of two sessions per mode: | dtype | M | N | cuBLAS | auto-tune | choice | speedup | |---|---:|---:|---:|---:|---|---:| | bf16 | 1 | 48 | 6.24 | 3.35 | tinygemm2 | 1.87x | | bf16 | 8 | 48 | 6.99 | 3.38 | tinygemm2 | 2.07x | | bf16 | 64 | 48 | 7.05 | 3.41 | tinygemm2 | 2.07x | | bf16 | 8 | 1024 | 8.06 | 4.73 | tinygemm2 | 1.70x | | bf16 | 16 | 1024 | 8.16 | 8.17 | cuBLAS | 1.00x | | bf16 | 8 | 2048 | 10.99 | 8.09 | tinygemm2 | 1.36x | | bf16 | 8 | 5120 | 19.23 | 19.30 | cuBLAS | 1.00x | | fp16 | 16 | 48 | 7.29 | 3.35 | tinygemm2 | 2.18x | | fp16 | 16 | 1024 | 8.10 | 6.94 | tinygemm2 | 1.17x | | fp16 | 8 | 2048 | 10.75 | 7.82 | tinygemm2 | 1.37x | | fp16 | 64 | 1024 | 8.51 | 8.51 | cuBLAS | 1.00x | Auto-tune was never slower than cuBLAS across M in {1,4,8,16,64} x N in {48,1024,2048,5120} for both dtypes. The isolated tuner timing does not capture the programmatic-launch overlap, so near ties (e.g. bf16 16x1024) conservatively stay on cuBLAS. ## Motivation and Context - Stage 2 of the GEMM auto-tuner in #32876, which was designed as a candidate table so kernels like this can be added without changing the dispatch policy. - tinygemm2 targets exactly the latency-bound decode/speculative-verify shapes (a handful of rows against a small-to-medium weight) where cuBLAS spends most of its time on launch/ramp and tile under-utilization. - Follow-ups: FlashInfer TGV GEMM for SM 10.x, cuBLASLt heuristic algorithms as candidates, and a persistent tuning cache (#29588). ## Checklist - [x] Tests added/updated - [x] No breaking changes (only active with `ep.cuda.enable_gemm_auto_tune=1` on SM 9.0+) - [x] Documentation updated (session option key comment) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Copilot

Copilot review overview

🟡 Changes recommended

Hardware isolation, rejected-entry replacement, unsupported GEMV validation, and lock-failure handling contain unresolved defects.

Review effort: Balanced
Findings: 3 · 2 · 1

Open (6)

Reject CUDA kernel tactics without runner support · New
Do not write persistent cache without the advisory lock · New
Partition profiler state by actual GPU identity · New
Replace rejected cached tactics during re-profiling · New
Include bias presence in tactic identity · New
Fix duplicate subsection numbering and cross-references · New

Resolved since last review (3)

WriteAllLocked uses std::rename() to do an atomic temp->final replace. On Windows, rename/MoveFile…
Dummy inference for bf16 models will likely fail because _numpy_dtype_for maps "tensor(bfloat16)"…
The process-global cache map is keyed by the raw session-config strings (and uses a newline…

Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com>

Baiju Meswani (baijumeswani)

Thanks for adding persistent tactic caching. I consolidated the independent reviews into the inline comments below. The main concerns are saving temporary profiling failures permanently, replacing a cache after a read failure, validating loaded tactics, and keeping per-session cache files in sync with the shared profile map. I also included performance, tooling, build-guard, documentation, and test suggestions, clearly marked where they are design suggestions rather than confirmed regressions.

Reviewed commit 0a6966b. Several failure paths were reproduced with isolated host checks using the PR code and mocked runtime dependencies. The file-permission and failed-write checks used real filesystem behavior. The BF16 input-binding check used an installed ORT wheel. No full CUDA build or GPU inference was run for this review.

Copilot

Copilot review overview

🟡 Changes recommended

The tuning tool rejects caches on nightly/local wheels, and construction flushing contradicts the stated PR contract.

Review effort: Balanced
Findings: 2

Open (2)

Initial buckets are lost without destructor-triggered flush · New
Use native runtime version for cache compatibility · New

Resolved since last review (6)

Partition profiler state by actual GPU identity
Do not write persistent cache without the advisory lock
Reject CUDA kernel tactics without runner support
Include bias presence in tactic identity
Replace rejected cached tactics during re-profiling
Fix duplicate subsection numbering and cross-references

Adds an opt-in, hardware-keyed on-disk tactic cache for the fpA_intB MatMulNBits path, a reduced first-time M-bucket sweep with lazy single-bucket fallback, CUDA EP teardown flush of lazily profiled tactics, an offline tuning tool, docs, and unit tests. Squashed and rebased onto origin/main.

- Key 2-bit weights as uint2b_t (2-bit fpA_intB landed on main). - Session-config cache location now fully overrides env vars, matching docs. - Resolve the cache path before loading so each file is parsed once per location instead of once per MatMulNBits node; build the signature from the kernel's cached device properties and refuse to share one file across different GPUs in a process. - Harden TSV hex decoding and integer parsing. - Tuning tool: use session config entries, drop the obsolete GEMV option bits, make ml_dtypes optional, and release the session before summarizing so lazily profiled buckets are flushed. - Fix stale docs/comments (teardown flush, strict signature fields, test invocation, sample TSV header).

- Restore setL2CacheBytes/mL2CacheBytes and the FORCE_CHUNKED doc text dropped by the rebase. - Key the in-process profile map by bits/group size/zero points/GEMV support so kernels with different quantization do not share tactics or persist them under the wrong cache key. - Stage lazily profiled buckets right after profiling instead of in kernel destructors, so any flush point sees them regardless of teardown order. - Flush from ~CudaEp: the plugin build does not compile cuda_execution_provider.cc. - Log flush failures at EP teardown.

- Add a strict tactic_selection_version header field so caches tuned under an older selection policy are rejected even when ORT_VERSION is unchanged. - Use a per-writer temp file so writers never share it if file locking fails. - Open cache, lock, and temp files through UTF-8 std::filesystem paths and MoveFileExW so non-ASCII paths work on Windows. - Create a missing cache directory on first write.

- Reject cached CUDA-GEMV tactics when the runner has no GEMV kernel. - Overwrite cache rows rejected on load with the freshly profiled tactic. - Fail Flush when the advisory file lock cannot be acquired, instead of writing unlocked and risking lost rows. - Add has_bias to GemmIdCore and the persistent key/TSV schema, and bump tactic_selection_version. - Add the GPU name to GemmIdCore so the process-global profile map is not shared between devices that have the same packing SM. - Renumber the tactic cache doc section to 6.3 and fix its references.

Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com>

- Import onnx only via `from onnx import ...`. - Move platform-specific file locking into helpers so fcntl is always initialized where it is used.

Wheel dev/local suffixes in ort.__version__ made caches written by the C++ side (base ORT_VERSION) look inapplicable.

Baiju Meswani (baijumeswani)

Thanks for the updates. I still see a few issues with loading the cache and checking whether tuning succeeded. Please also connect the new tests to CI. The inline comments include suggestions for saving shared results and reducing startup work.

Bound cached tactic validation by each session's launch limit, stage shared capture-time hits, and require session-specific tuning evidence. Wire regression tests into Python and CUDA CI and document the tuning-tool package limitation and measured startup costs. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Baiju Meswani (baijumeswani)

Thanks for the fixes. I still see three issues below. Please check the saved settings, fix the duplicate plugin registration, and make sure the built-in CUDA job runs the new Python test.

CUDA plugin wheels register their bundled provider during import. Check registered EP devices before explicitly registering the worker's CUDA plugin so fresh-process cache tests also work with plugin wheels. Add regression coverage for existing registration, explicit registration, non-CUDA devices, and plugin-loading failures. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Merge origin/main without rewriting the published PR history. Preserve paired-K and wave-aware GEMV choices in profile and persistent keys, serialize the selected GEMV variant, and reject older selection-policy caches. Verify the tuning tool's selected full configurations reached disk, cover a failed main-cache flush with unit and CUDA integration regressions, and explicitly enable pybind in the built-in CUDA test job. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Validate plugin cache selections through per-session tactic reports instead of relying on import-time filtered INFO logs. Initialize CLI arguments before environment-dependent checks, provision CPU Torch for Linux CUDA packer tests, and recognize the current compact prepacked contract diagnostic. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Resolve the profiler and MatMulNBits conflicts while preserving persistent-cache identity, validation, paired-mode behavior and teardown flushing. Integrate streaming-L2 scratch sizing and INT4 decode variants, serialize their geometry choices, and advance the tactic-selection signature to version 4. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Keep the standalone weight-packer oracle independent of compute-kernel exclusions. Undefine EXCLUDE_SM_90 only for that target so force_arch=90 cannot silently produce the SM80 column-interleaved byte layout.\n\nRetain all native/Torch byte-parity cases and add a regression guard against SM90-to-SM80 layout substitution. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Tianlei Wu (tianleiwu) requested a review from Copilot July 7, 2026 07:07

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) July 7, 2026 07:07 View session

Copilot AI reviewed Jul 7, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.cc Outdated

Comment thread onnxruntime/contrib_ops/cuda/llm/gemm_tactic_cache.cc Outdated

Comment thread onnxruntime/python/tools/fpa_intb_tune.py Outdated

Tianlei Wu (tianleiwu) marked this pull request as draft July 7, 2026 17:00

Base automatically changed from tlwu/20260707/fpa_intb_sm90_bs32_bias to main July 7, 2026 20:58

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch from 5a2ea51 to c43f96e Compare September 23, 2026 10:26

This was referenced Sep 29, 2026

[CUDA] Auto-tune small-M decode MatMul (cuBLAS vs small-N GEMV) and speed up MatMulNBits GEMV #32876

Merged

[CUDA] Add tinygemm2 as a GEMM auto-tune candidate for small-M MatMul on SM 9.0+ #32901

Merged

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch from c43f96e to 8661a3f Compare October 3, 2026 00:56

Tianlei Wu (tianleiwu) mentioned this pull request Oct 3, 2026

[CUDA] Load fpA_intB tactic cache on session create and save on close #33078

Open

Tianlei Wu (tianleiwu) marked this pull request as ready for review October 3, 2026 01:25

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch from 8198c22 to 50b9740 Compare October 5, 2026 04:40

Tianlei Wu (tianleiwu) added this pull request to stack #33093 October 5, 2026 04:41

Tianlei Wu (tianleiwu) requested a balanced review from Copilot October 5, 2026 04:51

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) October 5, 2026 04:52 View session

Copilot AI reviewed Oct 5, 2026

View reviewed changes

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch 3 times, most recently from efa1df1 to 91d0f38 Compare October 5, 2026 21:32

Copilot AI mentioned this pull request Oct 5, 2026

Fix CUDA tactic-cache environment lookup linkage #33119

Closed

Copilot AI added a commit that referenced this pull request Oct 5, 2026

Merge PR #29588 tactic-cache source for focused follow-up

715dbe3

Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com>

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch from 91d0f38 to 0a6966b Compare October 5, 2026 23:47

Baiju Meswani (baijumeswani) reviewed Oct 6, 2026

View reviewed changes

github-advanced-security Bot found potential problems Oct 7, 2026

View reviewed changes

Comment thread onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py Fixed

Comment thread onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py Fixed

Tianlei Wu (tianleiwu) requested a balanced review from Copilot October 7, 2026 08:14

Copilot started reviewing on behalf of Tianlei Wu (tianleiwu) October 7, 2026 08:14 View session

Copilot AI reviewed Oct 7, 2026

View reviewed changes

Comment thread onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h

Comment thread onnxruntime/python/tools/fpa_intb_tune.py Outdated

Tianlei Wu (tianleiwu) and others added 9 commits October 7, 2026 17:03

Fix CUDA tactic-cache environment lookups across provider boundaries

f61b65a

Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com>

review feedbacks

c8c62f3

Fix CodeQL alerts in fpA_intB tune test

cc7624a

- Import onnx only via `from onnx import ...`. - Move platform-specific file locking into helpers so fcntl is always initialized where it is used.

Use native ORT version in fpa_intb_tune cache signature

94a9245

Wheel dev/local suffixes in ort.__version__ made caches written by the C++ side (base ORT_VERSION) look inapplicable.

Tianlei Wu (tianleiwu) force-pushed the tlwu/20260707/fpa_intb_autotune_cache branch from c5e7855 to 94a9245 Compare October 7, 2026 17:06

Baiju Meswani (baijumeswani) reviewed Oct 7, 2026

View reviewed changes

Baiju Meswani (baijumeswani) reviewed Oct 8, 2026

View reviewed changes

Comment thread onnxruntime/python/tools/fpa_intb_tune.py

Comment thread onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py

Comment thread .github/workflows/linux_cuda_ci.yml

Baiju Meswani (baijumeswani) previously approved these changes Oct 8, 2026

View reviewed changes

Tianlei Wu (tianleiwu) dismissed Baiju Meswani (baijumeswani)’s stale review via c235647 October 8, 2026 00:29

github-advanced-security Bot found potential problems Oct 8, 2026

View reviewed changes

Comment thread onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py Fixed

Tianlei Wu (tianleiwu) and others added 3 commits October 9, 2026 21:03

Conversation

Tianlei Wu (tianleiwu) commented Jul 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Testing

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Baiju Meswani (baijumeswani) left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Baiju Meswani (baijumeswani) left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Baiju Meswani (baijumeswani) left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Tianlei Wu (tianleiwu) commented Jul 7, 2026 •

edited

Loading