You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Adds optional inputs committed_length ([1] or [batch_size]), key_scale, and value_scale ([num_kv_heads, head_size] or [1, num_kv_heads, 1, head_size]) plus optional Output(3, "validation_status") (int32[1], codes 0..10), and extends T_CACHE to include tensor(int8)
Adds T_KEY type constraint for past_key_state, present_key_state, and state_update so CUDA QSA can pair BFloat16 activations with float key state while preserving mainline output indices (output 6 = state_update, output 8 = state_status)
Strict validation reads the error flag back inside ValidateDynamicSparseAttentionOnDevice, before this cursor kernel runs. With batch_size=1, sequence_length=1, cache_capacity=2, seqlens_k={0}, and committed_length={1}, general validation succeeds, then the cursor writes 9 without any further host check. Run still succeeds and appends to the cache; omitting validation_status discards the error entirely. Check the cursor result before launching attention in strict mode, and add a regression without the status output.
Add QSA refresh snapshot and validation status coverage
The new QSA snapshot/status outputs lack refresh-mode numerical coverage: RunQsaPackedTest never requests outputs 8–10, and the only tests connecting them expect IndexShare reuse rejection. Add CUDA refresh tests comparing the first-token snapshot with a separate one-token transition, including a compression boundary and an empty request. Also verify validation_status values 0, 9, and 0x99 for valid requests and rejected overflow or inconsistent metadata. Existing state_update capture tests exercise output 6, not these outputs.
Pushed 9356555 with the review clarification and WebGPU pipeline fix.
Corrected mixed-state schema/generated docs to promise only CUDA QSA BFLOAT16 activations with FLOAT key state, not unsupported FLOAT16/FLOAT. WebGPU matching-type requirement is explicit.
Changed the binding-limit test to pass limit.c_str() to ConfigOptions::AddConfigEntry. Windows, Linux, macOS and WebGPU plugin logs all report this same std::string-to-const-char* compilation error; the static aggregate failure is downstream of compilation.
Android failed before compiling this PR while sdkmanager unpacked the emulator download: "Error on ZipFile unknown archive". This is an SDK download/infrastructure failure, not addressed by changing operator code. Fresh CI after this push will make another attempt; no unrelated workflow edits were made.
Validation: full-PR lintrunner -a and check-only lint, source-to-official-renderer documentation checks, and git diff --check passed. No local builds, compilation, native tests, or model execution were run as requested. Fresh CI is pending, not claimed green.
Added the scoped QSA and attention decode experiments on top of the current PR head, preserving its workspace-allocation, state-validation, and WebGPU review fixes.
Published commit: d21946e6cb37621b9a2e5f623ea56e635b3e7bdf.
Opt-in ORT_PACKED_SPARSE_INDEXER_COMPACT_TOPK=1: compact 2048-key chunks to 512 keys with device-side 512/1024/2048 sort probes, a fused merge/emit for at most five chunks, and existing pairwise merging above that. Preserves stable keys, causal tails, capacity stride/padding, and overflow handling. Rejects simultaneous radix and compact modes.
Opt-in ORT_DSA_INT8_VECTOR_STAGING=1: eight-byte grouped INT8 staging with independent per-channel key/value scales and the existing shape/resource guards.
Opt-in ORT_DSA_SHARED_SCALE_MERGE=1: share exponential factors across channels for grouped merges of at most 32 splits. Preserve nonzero-partial-sum activity and underflow/nonfinite accumulation semantics.
Added QSA chunk/fallback boundaries and distinct-channel signed INT8/auxiliary-cache/merge-boundary tests. Updated operator documentation.
No multi-snapshot state changes or dense-prefill normalization experiments were included.
Validation of the publishing sources: all changed CUDA translation units compile. Process-isolated native suites pass with all new switches both off and on: 48 attention tests and 36 packed-indexer tests per mode. This uses retained build dependencies, including build-compatible core headers and upstream test-utility definitions; it is not a clean full upstream build.
The earlier experimental tree had two malformed-state-length QSA/CSA failures with compact mode off and on. Those historical results must not be represented as a clean full suite. The latest publishing-version focused suites above pass, but the broader combined GDN run encountered an empty-request failure reproducible with the unmodified GDN PR-head kernel; clean-build CI remains required.
Historical Nsight measurements from experimental packages showed QSA summed kernel time 0.651 -> 0.409 ms/cycle, grouped attention compute 45.30 -> 31.02 us, and merge 9.28 -> 4.22 us. These are not new benchmark results for this transplanted PR tip. Full model accuracy remains deferred; there is no sustained-250-TPS-at-8K claim.
Fresh-build integration validation update (no source changes pushed to this ORT PR in this update):
Rebuilt ORT from scratch for H200 (sm90a), CUDA provider as a plugin, with FP4 QMoE, FP8/INT4 KV cache, GQA value layout, and FPA-INTB GEMM enabled. onnxruntime_QUICK_BUILD=OFF; no retained build objects were reused.
Rebuilt GenAI against the fresh ORT headers/libraries, and rebuilt the fold custom library.
Sparse/indexer/FP8 suite plus the isolated QSA capacity test: 140 additional passes.
Known limitations: QSA and CSA MalformedStateLengthsAreClamped tests still fail, including with specialized top-k disabled. These are not being claimed as passing. This is selected-suite validation, not a clean full-suite qualification.
Selected GenAI Qwen/executor/weight-pool tests: 69 passed, including the opt-in real-GPU mapped-weight test with proper CUDA plugin registration. One additional test is disabled.
The latest fused INT2 model completed one warmup plus three measured requests at each of 2K/8K/16K/32K, generating 512 tokens each. All 16 requests matched the previous-runtime token IDs exactly, with zero generation failures or standard fallbacks. Runtime-map inspection confirmed only the fresh ORT/GenAI libraries were used.
Prompt
Pooled decode TPS
Mean TTFT
2048
269.90
0.584 s
8192
250.17
2.448 s
16384
266.92
5.065 s
32768
264.12
10.578 s
Evidence scope: this is the full dirty experimental source stack containing this PR's optimizations plus other local work, not a standalone build/benchmark of this exact PR head. It is not an isolated speedup attribution, a fresh paired reference comparison, or proof of cross-process sustained 250 TPS. Full model accuracy qualification remains deferred.
The rebuild exposed local expert-cache function-pointer and INT2 direct-link test wiring issues, which were fixed locally. Those expert-cache extensions are not in the published PR and have deliberately not been mixed into this update.
Preserve state_status at output 8, append snapshot outputs at 9-11,
and retain BF16 prefill dispatch with an explicit quantized-cache guard.
Add coverage for snapshot outputs with and without state_status.
Validated using isolated merged-source provider/test links: 119 passed,
2 architecture-specific skips on GPU 1. Active worktrees are unchanged.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Resolved conflicts with main (1f4b965) in merge commit b42576e.
Preserved main’s established output 8 = state_status contract, including per-request invalid-metadata/capacity status semantics. Experimental snapshot outputs now occupy 9 = mid_kv_buffer, 10 = mid_state_lengths, 11 = validation_status. Earlier experimental graphs/exporters using slots 8–10 must insert an empty output at slot 8 before these three outputs. Mainline graph layouts are unchanged. Schema, kernel enum, inference, provider rejection paths, tests, and docs were updated consistently.
Kept mixed-type key-state support and main’s BF16 optimized prefill dispatch, explicitly excluding quantized caches from that BF16-only cache path.
Corrected raw-pointer/unique-pointer automatic-merge mismatches and preserved both branches’ tests. Added a regression requesting snapshot outputs both with and without state_status, including BF16 activations/FP32 state.
Compiled merged sources/tests and linked an isolated CUDA provider/test executable against existing integration-build dependencies. Combined suite for these two merges on H200 GPU 1: 119 passed, 2 architecture/image-specific skips, 0 failures.
This is merged-source validation using existing integration-build dependencies, not a full clean standalone-branch build or WebGPU qualification. GitHub currently reports MERGEABLE (no merge conflicts); required checks/reviews remain separate. No active model package or frozen GPU 3 LiveCodeBench snapshot was changed.
Follow-up validation: ran lintrunner -a on an isolated checkout of this PR against origin/main, reviewed and applied the formatter fixes, then reran both lintrunner -a and lintrunner successfully with no lint issues. Pushed formatting-only commit b71498b2660189a4d3638de2095cf1d74e691ce7 (3 files). Verified identical C++/CUDA token sequences before and after formatting, including preserved string/comment contents. No runtime behavior changes; active integration worktree and evaluation snapshot untouched.
…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>
… per KV head
- Specialize GroupedDecodeDynamicSparseAttentionKernel with 12 warps (384 threads,
head_tiles=1) when num_heads/kv_num_heads == 12, halving HBM sparse KV reads
and eliminating idle warps in tile 1 (-37.4% kernel time at S=8).
- Raise kTargetGroupedQuantizedDecodeBlocks to 512 and trim unused mma shared
memory when !(T==half && !kQuantizedCache).
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…dSparseAttentionIndexer
- Extend DynamicSparseAttentionPrefillKernel to INT8 KV caches using async
cp.async.ca.shared.global (CpAsync8) overlapped with Tensor Core MMA + online
softmax, alias sQ with sK (-8.45 KB smem), and set __launch_bounds__(128, 3)
(3 CTAs/SM, 5.3x lower long_scoreboard stall, 14.50 ms vs 34.66 ms at L=2048).
- Parallelize PackedSparseAttentionIndexer QsaUpdateState across compressed
blocks via QsaUpdateStateBlocksKernel and enable WarpHeads=true by default in
QsaRaggedPrefillScoreKernel.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…st coverage
- Stage all kQueryHeadsPerBlock heads into mma_query and populate the second 8-row MMA fragment/score slots in GroupedDecodeDynamicSparseAttentionKernel so warps 8-11 read computed scores for GQA ratio 12 on SM80+.
- Extend FloatingCacheGroupedReferenceParity_CUDA to sequence_length=9 to exercise GroupedPrefillDynamicSparseAttentionKernel, add GroupedDecodeRatio12ReferenceParity_CUDA, and add Int8CachePrefillBFloat16_CUDA.
- Remove unused legacy QsaPairedScoreTileTopKKernel and QsaCompactTileTopKKernel from packed_sparse_attention_indexer_impl.cu and sync packed_sparse_attention_indexer.md with current defaults and CSA-only position_ids.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Keep the BF16 prefill reference batch index in int64_t to match the sequence length and fix MSVC C4244 warnings treated as errors in Windows CUDA, CUDA plugin, and TensorRT CI.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…33627)
### Summary
Extends com.microsoft::DynamicSparseAttention (committed_length input
17, introduced in #33472) to accept either [1] or [batch_size] int32
elements and validates committed_length[b * committed_length_stride] per
batch item in ValidateInputsKernel.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Name validation code 9 (kDynamicSparseAttentionInvalidCommittedLength) and
return a cursor-specific error from the strict-validation host check instead
of treating it as success when no validation_status output is connected.
Add a regression test (strict env var, no status output) and clear
smooth_softmax in the INT8 prefill test so the reference matches and the
INT8 prefill kernel is reached.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
The INT8 vector-staging, shared-scale merge, compact/radix top-k and
warp-heads switches are cached in function-local statics, so add a small
helper that re-runs the relevant existing cases in a new process of the test
binary with each non-default combination set before the first eligible call.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…s codes
Give validation code 9 its documented name, state that the INT8 KV cache is
symmetric signed int8 like GroupQueryAttention/PagedAttention, list the exact
conditions and values for every CUDA dispatch switch, and sync
ContribOperators.md with the schema (committed_length is [1] or [batch];
validation_status code list).
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Benchmarks on the rc3 model (H200, DFlash2 S=8 decode) show
ORT_PACKED_SPARSE_INDEXER_RADIX_TOPK=1 is 0.15-0.34 ms/cycle slower than the
default. Remove the env var, the packed-key QsaPartialTopKKernel
specialization and the COMPACT/RADIX mutual-exclusion check; restore
QsaPartialTopKKernel to its non-templated form. Alternate-mode fresh-process
tests now cover only COMPACT_TOPK=0 and WARP_HEADS=0. Document the remaining
switches with benchmark deltas; schema and ContribOperators.md are unaffected.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
The reason will be displayed to describe this comment to others. Learn more.
From Copilot:
Completed an implementation pass and a second edge-case/regression pass at cd4e2e4. One new P1 alignment finding is attached inline. Checked INT8 scaling and cache handling, grouped decode/prefill, compact Top-K padding and causal visibility, optional outputs, and non-CUDA rejection paths. Existing review threads were checked to avoid duplicate comments.
Validation: a standalone CUDA probe reproduced the float4 misalignment failure on H200. This was not a build or session-level run of the unmodified ORT operators. Full PR-specific ORT suites, WebGPU execution, cross-architecture qualification, and benchmarks were not run.
…build
DynamicSparseAttention: the default-on fast paths issued 16-byte (uint4/float4), 8-byte (uint2, cp.async) and 4-byte (bfloat162) accesses on caller-provided buffers. Tensors bound through IOBinding can start at any element offset, so a misaligned scale, cache, auxiliary or output buffer faulted with cudaErrorMisalignedAddress (716) and poisoned the CUDA context.
- INT8 grouped decode: use vector staging only when the scales, auxiliary K/V and INT8 caches are aligned; otherwise use the scalar staging kernel.
- Floating grouped decode and grouped prefill: StageHeadRow copies a K/V row with one 16-byte chunk per lane only when the row is 16-byte aligned and element-wise otherwise, so the workspace layout is unchanged.
- Tensor-core prefill kernel: also require aligned caches, INT8 scales and output; otherwise the generic kernels run.
- Add IOBinding tests that offset each cache, scale, auxiliary and output buffer for INT8 and floating caches, for decode and prefill, and document the alignment requirements and fallbacks.
iOS CI: fresh_process_test_util.h uses std::system, which is unavailable on iOS. Include it and define QsaAlternateTopKModesInFreshProcess only under USE_CUDA.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
The iOS build failure is fixed in c8c7b77 (cc David Fan (@jiafatom)). The only compile error on the previous head was error: 'system' is unavailable: not available on iOS, from fresh_process_test_util.h, which packed_sparse_attention_indexer_op_test.cc included unconditionally. The include and QsaAlternateTopKModesInFreshProcess are now under USE_CUDA, since the helper is only needed by the CUDA kernel tests. The other jobs on the previous head were green; iOS_CI_on_Mac and iphone_simulator are re-running on the new commit.
This branch has not been deployed
No deployments
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
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
Extends$S \le 8$ grouped decode specialization, and extends $S \le 8$ QSA refresh/selection kernels, compact hierarchical TopK (
com.microsoft::DynamicSparseAttentionon the CUDA Execution Provider with symmetric signed INT8 per-channel KV cache support, per-batchcommitted_lengthcursor validation, and 12-warpcom.microsoft::PackedSparseAttentionIndexerwith mixed-precisionT_KEY(BFloat16activations withfloatpast_key_state), fusedORT_PACKED_SPARSE_INDEXER_COMPACT_TOPK=1), and warp-specialized ragged QSA prefill (ORT_PACKED_SPARSE_INDEXER_WARP_HEADS=1).Summary of Changes
com.microsoft::DynamicSparseAttentionINT8 KV Cache, Cursor Validation & 12-Warp Decodeonnxruntime/core/graph/contrib_ops/bert_defs.cccommitted_length([1]or[batch_size]),key_scale, andvalue_scale([num_kv_heads, head_size]or[1, num_kv_heads, 1, head_size]) plus optionalOutput(3, "validation_status")(int32[1], codes0..10), and extendsT_CACHEto includetensor(int8)onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention.{h,cc},dynamic_sparse_attention_impl.{h,cu},dynamic_sparse_attention_prefill.cuhORT_DSA_INT8_VECTOR_STAGING=1default), shared-scale split merge forsplit_count <= 32(ORT_DSA_SHARED_SCALE_MERGE=1default), 12-warp (head_tiles=1) specialization inGroupedDecodeDynamicSparseAttentionKernel(query_heads_per_kv == 12,9(kDynamicSparseAttentionInvalidCommittedLength)onnxruntime/contrib_ops/webgpu/bert/dynamic_sparse_attention.cccommitted_length/validation_statusextensions with clear status errorscom.microsoft::PackedSparseAttentionIndexerMixedT_KEY, Compact Hierarchical TopK & Ragged Prefillonnxruntime/core/graph/contrib_ops/bert_defs.ccT_KEYtype constraint forpast_key_state,present_key_state, andstate_updateso CUDA QSA can pairBFloat16activations withfloatkey state while preserving mainline output indices (output 6 = state_update,output 8 = state_status)onnxruntime/contrib_ops/cuda/sparse/packed_sparse_attention_indexer.{h,cc},packed_sparse_attention_indexer_impl.{h,cu},qsa_topk.cuhORT_PACKED_SPARSE_INDEXER_COMPACT_TOPK=1default: 2048-to-512 chunk compaction withQsaCompactTileTopKKerneland fusedQsaFinalMergeAndEmitKernelforQsaMergeTwoLevelsTopKKernelwhen=0), and warp-specialized ragged QSA prefill (ORT_PACKED_SPARSE_INDEXER_WARP_HEADS=1default)onnxruntime/contrib_ops/webgpu/bert/packed_sparse_attention_indexer.ccTests and Documentation
onnxruntime/test/contrib_ops/dynamic_sparse_attention_op_test.cc: Adds coverage for INT8 KV cache append/decode/prefill (withsmooth_softmax=0andhead_sinkcleared for prefill), 12-head grouped decode (FP16/BF16), scalar and per-batchcommitted_length, strict-validation cursor mismatch rejection withoutvalidation_statusoutput, and fresh-process tests (Int8CacheGroupedDecodeAlternateSwitchesInFreshProcess_CUDA) for non-defaultORT_DSA_INT8_VECTOR_STAGING/ORT_DSA_SHARED_SCALE_MERGEcombinations.onnxruntime/test/contrib_ops/packed_sparse_attention_indexer_op_test.cc,fresh_process_test_util.h: Adds coverage forBFloat16activations withfloatpast_key_state, compact/hierarchical TopK chunk boundaries, stable ties, partial merge groups, and fresh-process parity runs (QsaAlternateTopKModesInFreshProcess) withORT_PACKED_SPARSE_INDEXER_COMPACT_TOPK=0andORT_PACKED_SPARSE_INDEXER_WARP_HEADS=0.docs/ContribOperators.md,docs/OperatorKernels.md,docs/contrib_ops/cuda/dynamic_sparse_attention.md,docs/contrib_ops/packed_sparse_attention_indexer.md: Documents schema updates, INT8 symmetric quantization, validation status codes (0..10), dispatch switches, and H200 benchmark deltas.Performance
Measured on NVIDIA H200 (SM90, Qwen3.8-Flash-Next RC3 INT4 model, DFlash2$S=8$ decode across
2K / 8K / 32Kcontext):ORT_DSA_INT8_VECTOR_STAGING=1(default):~0.30 ms/cycle(~2%) faster than=0.ORT_DSA_SHARED_SCALE_MERGE=1(default):~0.05–0.10 ms/cycle(~0.3–0.5%) faster than=0.ORT_PACKED_SPARSE_INDEXER_COMPACT_TOPK=1(default):~0.20–0.27 ms/cycle(~1.5%) faster decode than=0.ORT_PACKED_SPARSE_INDEXER_WARP_HEADS=1(default): reduces prefill TTFT by3.7%at8Kand14%at32Kcompared to=0.Testing
onnxruntime_provider_test --gtest_filter="DynamicSparseAttention*:PackedSparseAttentionIndexer*"qwen_flash_next_int4_int8_rc3_contrib_dflash2(state_update_capacity=7).Checklist