Repository navigation
[CUDA] Optimize HyperConnectionLowRankMix for S<=8 speculative decode and deterministic prefill - #33548
Open
Akshay Sonawane (apsonawane) wants to merge 15 commits into
Open
[CUDA] Optimize HyperConnectionLowRankMix for S<=8 speculative decode and deterministic prefill#33548Akshay Sonawane (apsonawane) wants to merge 15 commits into
Akshay Sonawane (apsonawane) wants to merge 15 commits into
Conversation
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
The environment-helper include order causes conflicting logging definitions in normal non-plugin CUDA-provider builds.
1 open issue
What changed in this PR
Adds opt-in single-row work groups for eligible four-row CUDA HyperConnectionLowRankMix decode, supporting the broader speculative-verification integration while retaining the default two-row grouping.
Changes:
- Adds
ORT_HC_DECODE_ROWSwith supported values 1 and 2. - Extends rank-320 reference coverage to four-row inputs.
- Documents tuning behavior and separate-process testing for both settings.
| File | Description |
|---|---|
| onnxruntime/test/contrib_ops/hyper_connection_low_rank_mix_op_test.cc | Adds four-row fused-decode reference cases. |
| onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu | Adds environment-controlled row grouping for eligible four-row decode. |
| docs/contrib_ops/cuda/hyper_connection_low_rank_mix.md | Documents supported settings and test configurations. |
Review effort: Balanced
…zation, and deterministic prefill post-gate Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Tianlei Wu (tianleiwu)
force-pushed
the
asonawane/dflash2-s4-hyperconnection
branch
from
October 23, 2026 23:01
60a9e2a to
1428ea7
Compare
Tianlei Wu (tianleiwu)
changed the base branch from
asonawane/main-ops-hyperconnection-tensor
to
main
October 23, 2026 23:01
…nectionLowRankMix Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…ction_low_rank_mix_impl.cu Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Accumulate K-split partial sums directly into va[r] inside HcUpPreKernel instead of staging through a float4 pq[8][ROWS] local array, avoiding dynamic register indexing and saving ~0.35 ms/cycle across 96 launches at S=8. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Tianlei Wu (tianleiwu)
added a commit
that referenced
this pull request
Oct 24, 2026
…2 QMoE, and PR cleanups - Extend FoldMid2/FoldTail to head_rows<=8 and pack tokens[4..7] into ring record words 1, 2, 3, and 5. - Add state_update_capacity (0..8) support to CausalConvWithState and enable state_update_capacity > 0 in QwenS2 GatedDeltaNetDecodeWarpKernel. - Raise kMaxProfiledExpandedRows to 80 in moe_gemv.h and kFusedRouteMaxRows to 8 in qmoe_fused_route.cuh; add fast INT3 and offline SM80 INT2/INT4 prepack QMoE support. - Eliminate pq[8][ROWS] local register array in HcUpPreKernel and sync PR cleanups (#33469, #33470, #33471, #33472, #33474, #33548, #33569). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… 5 blocks/SM - Parallelize HcUpPreKernel partial reduction across ROWS*nq (160) threads and move xv/inv/nw global loads to the epilogue (60->48 regs, 4->5 blocks/SM, -22.5% kernel time). - Alias red over wsm, stage kAsyncWeightSlots=9 in HcDownKernel shared memory (42.05 KB/block), and set __launch_bounds__(256, 5) + MaxShared carveout so both HcDownKernel and HcUpPreKernel run at 5 blocks/SM on Hopper. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Tianlei Wu (tianleiwu)
marked this pull request as draft
October 24, 2026 23:42
Tianlei Wu (tianleiwu)
marked this pull request as ready for review
October 25, 2026 05:38
…rows, and vectorize prefill - Replace runtime-indexed row_idx/row_ok local arrays with inline RowIdx/RowOk helpers, eliminating the 32-80B local-memory stack frame (STACK:0) across all HcDownKernel and HcUpPreKernel instantiations. - Default S=8 decode to LaunchFusedRows<T, 4, 2> (DOWN_ROWS=4, UP_ROWS=2), halving HcDownKernel grid from 672 CTAs (1.02 waves) to 336 CTAs (0.64 waves) and reducing 97-pair wall-clock span by 307 us/cycle on H200. - Add 16-byte vectorized PreMixVec8Kernel (1.68x faster) and 4-row/CTA FusedInjectPostGateB4Kernel (2.39x faster) for unfused prefill (-5.34 ms at L=2048). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…d HyperConnectionLowRankMix Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…case Drop the rows == 4 -> 1x1 dispatch (and its ORT_HC_DECODE_ROWS default of 1); four-row inputs now use the generic 2x2 fused dispatch. ORT_HC_DECODE_ROWS=4 and ORT_HC_S8_DECODE_ROWS are kept for benchmarking the 8-row configuration. Remove code that is unreachable afterwards: the ROWS >= 4 weight-slot ternary and the second cross-warp reduction branch of HcUpPreKernel (it only runs with at most 2 rows), and the unused `inject` scratch buffer/parameter of the unfused launcher. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… the unfused path FusedInjectPostGateKernel only accumulated eight branch columns into an smem[8][8] buffer, so num_branches > 8 read out of bounds and left the ninth gate unaccumulated. Tile the kernel over branch columns (grid.y) so any num_branches is computed correctly. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…ch and nine branches Add rows == 8 at (branches=4, hidden=2560, rank=320) to the rank-320 decode test (FP16/BF16, with and without prologue) so LaunchFusedRows<T, 4, 2> is exercised, and add a more-than-eight-branch regression test (rows=1, C=9, H=8, R=8 for FP32/FP16/BF16, plus C=17). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…nd env knobs Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…nob, validate ORT_HC_S8_DECODE_ROWS, tighten docs and decode test Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…s use 4x2 for 8 rows; document PDL benchmark Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
This branch has not been deployed
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.


Description
Optimizes$S \le 8$ speculative decoding (
com.microsoft::HyperConnectionLowRankMixon the CUDA Execution Provider forHcDownKernel+HcUpPreKernelfused path) and deterministic prefill (FusedInjectPostGateKernelunfused path):LaunchFusedRows<T, 4, 2>(4 rows per CTA inHcDownKernel, 2 rows per CTA inHcUpPreKernel) whenrows == 8,LaunchFusedRows<T, 2, 2>forrowsin2..7, andLaunchFusedRows<T, 1, 1>forrows == 1.HcUpPreKernel: Specializes shared-memory and loop bounds for(num_branches=4, rank=320)(HcUpPreKernel<T, ROWS, 4, 320>) and accumulates K-split partial sums directly intova[r]without afloat4 pq[8][ROWS]local register array.ORT_HYPER_CONNECTION_PDL=1default): EnablescudaLaunchAttributeProgrammaticStreamSerializationbetweenHcDownKernelandHcUpPreKernelon SM90+ (CUDA 12.3+) with__ldcggrid-synchronized reads when weight inputs are constant initializers.FusedInjectPostGateKernelwith Branch Tiling: Replaces the unfused split-K cuBLASinjectGEMM +PostGateKernelwith a deterministic fused inject-GEMV + post-gate kernel (FusedInjectPostGate4RowsKernelfor 16-bitnum_branches == 4with 8-byte alignedinject_weight, andFusedInjectPostGateKerneltiled over 8-branch column groups ingrid.yto support arbitrarynum_branches, includingnum_branches > 8).Summary of Changes
CUDA Kernel & Launcher
onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cuLaunchFusedRows<T, DOWN_ROWS, UP_ROWS>(1x1forrows == 1,2x2forrowsin2..7,4x2forrows == 8),HcUpPreKernel<T, ROWS, 4, 320>specialization without local register spills, Hopper PDL launch attributes (ORT_HYPER_CONNECTION_PDL=1), and deterministicFusedInjectPostGate4RowsKernel/ branch-tiledFusedInjectPostGateKernelon the unfused pathonnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix.{cc,impl.h}injectscratch buffer allocation from the unfused pathTests and Documentation
onnxruntime/test/contrib_ops/hyper_connection_low_rank_mix_op_test.cc: AddsFusedDecodeShapeRank320coveringrowsin{1, 2, 3, 4, 7, 8}at(branches=4, hidden=2560, rank=320)forBFloat16andMLFloat16(with and without fused prologue), andUnfusedMoreThanEightBranchescoveringnum_branches = 9(FP32,FP16,BF16) andnum_branches = 17(FP16).docs/contrib_ops/cuda/hyper_connection_low_rank_mix.md: Documents the fused and unfused paths, row-group dispatch table,ORT_HYPER_CONNECTION_PDLenvironment variable, and H200 DFlash2Performance
Measured on NVIDIA H200 (SM90, Qwen3.8-Flash-Next RC3 INT4 model, DFlash2$S=8$ decode ms/cycle across 5 repetitions):
4x2rows forORT_HYPER_CONNECTION_PDL=1)2x2rows forORT_HYPER_CONNECTION_PDL=0Testing
onnxruntime_provider_test --gtest_filter="HyperConnectionLowRankMixTest.*"qwen_flash_next_int4_int8_rc3_contrib_dflash2(Checklist