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
Extends com.microsoft::DynamicSparseAttention with INT8 per-channel KV cache support, cursor validation, and 12-warp $S \le 8$ grouped decode specialization, and extends com.microsoft::PackedSparseAttentionIndexer with FP32 past_key_state (T_KEY) support and fast $S \le 8$ QSA block scoring/TopK on CUDA.
Adds T_KEY type constraint (BFloat16 activations with floatpast_key_state), fused $S \le 8$ refresh/select kernels, and hierarchical/compact TopK selection while using mainline state_update (output 6) and state_status (output 8).
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>
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.
Extends$S \le 8$ grouped decode specialization, and extends $S \le 8$ QSA block scoring/TopK on CUDA.
com.microsoft::DynamicSparseAttentionwith INT8 per-channel KV cache support, cursor validation, and 12-warpcom.microsoft::PackedSparseAttentionIndexerwith FP32past_key_state(T_KEY) support and fastKey Changes
DynamicSparseAttentionINT8 KV Cache & 12-Warp Decodeonnxruntime/core/graph/contrib_ops/bert_defs.cc,onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention.{cc,h},dynamic_sparse_attention_impl.{h,cu}committed_length,key_scale,value_scaleinputs andvalidation_statusoutput, INT8 KV cache append/attention, and 12-warp (head_tiles=1) specialization inGroupedDecodeDynamicSparseAttentionKernel(kTargetGroupedQuantizedDecodeBlocks=512) forquery_heads_per_kv == 12(PackedSparseAttentionIndexerMixedT_KEY& Fast QSA Decodeonnxruntime/core/graph/contrib_ops/bert_defs.cc,onnxruntime/contrib_ops/cuda/sparse/packed_sparse_attention_indexer.{cc,h},packed_sparse_attention_indexer_impl.{h,cu},qsa_topk.cuhT_KEYtype constraint (BFloat16activations withfloatpast_key_state), fusedstate_update(output 6) andstate_status(output 8).onnxruntime/test/contrib_ops/dynamic_sparse_attention_op_test.cc,packed_sparse_attention_indexer_op_test.cc,docs/BFloat16+floatQSA state, and hierarchical TopK boundaries.Testing Notes
onnxruntime_provider_test --gtest_filter="DynamicSparseAttention*:PackedSparseAttentionIndexer*"qwen_flash_next_int4_int8_rc2_contrib_dflash2(state_update_capacity=7).