Skip to content

[CUDA] Optimize HyperConnectionLowRankMix for S<=8 speculative decode and deterministic prefill - #33548

Open
Akshay Sonawane (apsonawane) wants to merge 15 commits into
mainfrom
asonawane/dflash2-s4-hyperconnection
Open

Akshay Sonawane (apsonawane) wants to merge 15 commits into
mainfrom
asonawane/dflash2-s4-hyperconnection

Conversation

@apsonawane

@apsonawane Akshay Sonawane (apsonawane) commented Oct 22, 2026 •

Copy link
Copy Markdown
Contributor

Description

Optimizes com.microsoft::HyperConnectionLowRankMix on the CUDA Execution Provider for $S \le 8$ speculative decoding (HcDownKernel + HcUpPreKernel fused path) and deterministic prefill (FusedInjectPostGateKernel unfused path):

  • 4x2 Fused Row-Group Dispatch for 8-Row Speculative Decode: Uses LaunchFusedRows<T, 4, 2> (4 rows per CTA in HcDownKernel, 2 rows per CTA in HcUpPreKernel) when rows == 8, LaunchFusedRows<T, 2, 2> for rows in 2..7, and LaunchFusedRows<T, 1, 1> for rows == 1.
  • Rank-320 Specialization & Register-Spill-Free 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 into va[r] without a float4 pq[8][ROWS] local register array.
  • Hopper Programmatic Dependent Launch (ORT_HYPER_CONNECTION_PDL=1 default): Enables cudaLaunchAttributeProgrammaticStreamSerialization between HcDownKernel and HcUpPreKernel on SM90+ (CUDA 12.3+) with __ldcg grid-synchronized reads when weight inputs are constant initializers.
  • Deterministic Unfused FusedInjectPostGateKernel with Branch Tiling: Replaces the unfused split-K cuBLAS inject GEMM + PostGateKernel with a deterministic fused inject-GEMV + post-gate kernel (FusedInjectPostGate4RowsKernel for 16-bit num_branches == 4 with 8-byte aligned inject_weight, and FusedInjectPostGateKernel tiled over 8-branch column groups in grid.y to support arbitrary num_branches, including num_branches > 8).

Summary of Changes

CUDA Kernel & Launcher

File Change
onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu Adds LaunchFusedRows<T, DOWN_ROWS, UP_ROWS> (1x1 for rows == 1, 2x2 for rows in 2..7, 4x2 for rows == 8), HcUpPreKernel<T, ROWS, 4, 320> specialization without local register spills, Hopper PDL launch attributes (ORT_HYPER_CONNECTION_PDL=1), and deterministic FusedInjectPostGate4RowsKernel / branch-tiled FusedInjectPostGateKernel on the unfused path
onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix.{cc,impl.h} Removes the unused inject scratch buffer allocation from the unfused path

Tests and Documentation

  • onnxruntime/test/contrib_ops/hyper_connection_low_rank_mix_op_test.cc: Adds FusedDecodeShapeRank320 covering rows in {1, 2, 3, 4, 7, 8} at (branches=4, hidden=2560, rank=320) for BFloat16 and MLFloat16 (with and without fused prologue), and UnfusedMoreThanEightBranches covering num_branches = 9 (FP32, FP16, BF16) and num_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_PDL environment variable, and H200 DFlash2 $S=8$ benchmark results.

Performance

Measured on NVIDIA H200 (SM90, Qwen3.8-Flash-Next RC3 INT4 model, DFlash2 $S=8$ decode ms/cycle across 5 repetitions):

Configuration 2K 8K 32K
Default (4x2 rows for $S=8$, ORT_HYPER_CONNECTION_PDL=1) 14.93 ms 15.07 ms 15.30 ms
2x2 rows for $S=8$ 15.00 ms 15.11 ms 15.38 ms
ORT_HYPER_CONNECTION_PDL=0 15.09 ms 15.25 ms 15.48 ms

Testing

  • Run provider unit tests:
    onnxruntime_provider_test --gtest_filter="HyperConnectionLowRankMixTest.*"
  • End-to-end verification on qwen_flash_next_int4_int8_rc3_contrib_dflash2 ($S=8$ DFlash2 speculative decoding).

Checklist

  • Tests added/updated
  • No breaking changes (or documented in description)
  • Documentation updated (if applicable)

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 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_ROWS with 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

Comment thread onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu Outdated
…zation, and deterministic prefill post-gate

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
@tianleiwu
Tianlei Wu (tianleiwu) force-pushed the asonawane/dflash2-s4-hyperconnection branch from 60a9e2a to 1428ea7 Compare October 23, 2026 23:01
@tianleiwu
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>
@tianleiwu Tianlei Wu (tianleiwu) changed the title Tune hyperconnection row grouping for four-row decode [CUDA] Optimize HyperConnectionLowRankMix 4-row decode, rank-320 specialization, and deterministic prefill Oct 24, 2026
…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>

@github-actions github-actions Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You can commit the suggested changes from lintrunner.

Comment thread onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu Outdated
@tianleiwu
Tianlei Wu (tianleiwu) marked this pull request as draft October 24, 2026 23:42
@tianleiwu Tianlei Wu (tianleiwu) changed the title [CUDA] Optimize HyperConnectionLowRankMix 4-row decode, rank-320 specialization, and deterministic prefill [CUDA] Optimize HyperConnectionLowRankMix for S<=8 speculative decode and deterministic prefill Oct 25, 2026
@tianleiwu
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>

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Valid inputs with more than eight branches can trigger out-of-bounds shared-memory reads in the replacement post-gate kernel.

2 open issues
1 resolved since last review

Review effort: Balanced

Comment thread onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu Outdated
Comment thread onnxruntime/test/contrib_ops/hyper_connection_low_rank_mix_op_test.cc Outdated
…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

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants