Skip to content

Sparse MQA logits follow-ups for nv_dev: 16-byte paged KV stride, workspace reuse under graph capture #459

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We鈥檒l occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
RayWang96 merged 2 commits into from
Sep 24, 2026
Merged

Sparse MQA logits follow-ups for nv_dev: 16-byte paged KV stride, workspace reuse under graph capture #459

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
422debb
Reuse a warmed sparse MQA logits workspace during graph capture
lfr-0531 Sep 22, 2026
337b1f8
Accept 16-byte aligned KV pages in the paged sparse MQA logits
lfr-0531 Sep 22, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 13 additions & 3 deletions csrc/apis/attention.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -260,7 +260,16 @@ static const torch::Tensor& get_sparse_mqa_logits_workspace(const torch::TensorO
auto& workspace = workspaces[stream];
if (not workspace.defined()) {
// Warm up each stream before capture so one-time zeroing is not replayed with the graph.
DG_HOST_ASSERT(c10::cuda::currentStreamCaptureStatusMayInitCtx() == c10::cuda::CaptureStatus::None);
if (c10::cuda::currentStreamCaptureStatusMayInitCtx() != c10::cuda::CaptureStatus::None) {

@ds-review-bot ds-review-bot Sep 23, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

馃數 suggestion: The capture-time fallback scans the static workspaces map and returns the first defined workspace on the same device. This matches the single warmed stream per device used by TensorRT-LLM, but if more than one stream on the device has warmed a workspace, the choice is unordered_map-order dependent and may not be the stream the captured graph later replays on. Since the workspace is shared per-launch scratch, that could alias scratch across concurrently-executing streams. Consider documenting the 'one warmed workspace per device' assumption or tracking the warm-up stream so the fallback is deterministic. Not a blocker for the validated path.

馃 v4

@lfr-0531 lfr-0531 Sep 23, 2026 •
edited
Loading

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

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

Kept the simple scan (TensorRT-LLM warms exactly one stream per device); the comment now states that assumption.

// Capturing on a side stream: reuse a warmed workspace on this device instead of
// recording an allocation (assumes one warmed stream per device, whose eager
// launches are serialized with the replays).
for (const auto& [other_stream, other_workspace] : workspaces) {
if (other_workspace.defined() and other_workspace.device() == options.device())
return other_workspace;
}
DG_HOST_UNREACHABLE("Warm up the sparse MQA logits workspace on this device before capturing");
}
workspace = torch::zeros({kNumWorkspaceBytes}, options.dtype(torch::kByte));
}
return workspace;
Expand Down Expand Up @@ -434,15 +443,16 @@ static torch::Tensor fp8_fp4_paged_sparse_mqa_logits(const std::tuple<torch::Ten
head_dim_with_sf == (is_fp4 ? head_dim / 2 : head_dim) + static_cast<int>(sizeof(int)));
DG_HOST_ASSERT(fused_kv_cache.scalar_type() == torch::kUInt8 and fused_kv_cache.stride(1) == head_dim_with_sf and
fused_kv_cache.stride(3) == 1 and fused_kv_cache.stride(0) <= std::numeric_limits<int>::max() and
fused_kv_cache.stride(0) % 512 == 0);
fused_kv_cache.stride(0) % 16 == 0); // pages are read with 16-byte cp.async chunks

@ds-review-bot ds-review-bot Sep 23, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

馃數 suggestion: Relaxing the page stride to a multiple of 16 is correct for the 16-byte cp.async copies, but the paged entry point's 16-byte pointer-alignment loop covers {q_fp, q_sf, weights, metadata} and not fused_kv_cache. Page starts are base + page_idx * stride(0), so 16-byte KV access still requires the cache's data_ptr() to be 16-byte aligned. Torch allocations normally satisfy this, but an explicit assert or comment would make the contract complete and prevent a future strided/offset view from silently producing misaligned cp.async addresses.

馃 v4

@lfr-0531 lfr-0531 Sep 23, 2026 •
edited
Loading

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

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

Done: fused_kv_cache is now in the 16-byte pointer-alignment loop.

DG_HOST_ASSERT(fused_kv_cache.is_cuda() and fused_kv_cache.device() == q_fp.device());

const auto [_num_q_tokens_weights, _num_heads_weights] = get_shape<2>(weights);
DG_HOST_ASSERT(_num_q_tokens_weights == num_q_tokens and _num_heads_weights == num_heads);
DG_HOST_ASSERT(weights.scalar_type() == torch::kBFloat16 and weights.stride(1) == 1);
DG_HOST_ASSERT(metadata.dim() == 1 and metadata.scalar_type() == torch::kUInt8 and metadata.is_contiguous());
DG_HOST_ASSERT(metadata.numel() >= static_cast<int64_t>(sizeof(MetadataHeader)));
for (const auto& tensor: {q_fp, q_sf, weights, metadata}) {
// Pages start at `data_ptr() + page_idx * stride(0)`, so the base pointer needs 16-byte alignment too
for (const auto& tensor: {q_fp, q_sf, fused_kv_cache, weights, metadata}) {
DG_HOST_ASSERT(tensor.is_cuda() and tensor.device() == q_fp.device());
DG_HOST_ASSERT(reinterpret_cast<uintptr_t>(tensor.data_ptr()) % 16 == 0);
}
Expand Down
2 changes: 1 addition & 1 deletion deep_gemm/include/deep_gemm/impls/sm100_sparse_mqa_logits.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -616,7 +616,7 @@ void sm100_paged_sparse_mqa_logits(const uint32_t logits_stride, const uint32_t
// Keep the CUTLASS dtype out of the kernel template signature to avoid ptxas register spills
using qk_dtype_t = cute::conditional_t<kIsMXFP4, cutlass::float_e2m1_t, cutlass::float_e4m3_t>;
DG_STATIC_ASSERT(PAGE_KV % SPARSE_BLOCK_KV == 0, "Sparse KV blocks must not cross pages");
DG_DEVICE_ASSERT(kv_page_stride_bytes % 512 == 0);
DG_DEVICE_ASSERT(kv_page_stride_bytes % 16 == 0);
const auto kv_accessor = sparse_mqa_detail::PagedSparseKVAccessor<PAGE_KV, SPARSE_BLOCK_KV, qk_dtype_t>(
fused_kv_cache, kv_page_stride_bytes);
sm100_sparse_mqa_logits_core_impl<SPARSE_BLOCK_KV, kNumQStages, kNumKVStages, kNumTmemStages,
Expand Down