Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
adee029
[None][fix] DSpark FP8 producers: round UE8M0 scales up explicitly
dc3671 Nov 10, 2026
a4321c1
[None][fix] DSpark: one BF16 attention binder; V4.1 rejects fp8 proje…
dc3671 Nov 10, 2026
941e201
[None][chore] DSpark: check fused WKV width, document ring mapping, m…
dc3671 Nov 10, 2026
fbf3b55
[None][chore] DSpark: restore the dspark_attention_forward name and e…
dc3671 Nov 10, 2026
6fa9121
[None][fix] DSpark attention op: reject zero page stride, report reje…
dc3671 Nov 10, 2026
fce1445
[None][chore] DSpark draft-block prep: drop the per-device compile ca…
dc3671 Nov 10, 2026
1266c26
[None][chore] DSpark: remove the unwired fused_dsv4_dspark_attention_…
dc3671 Nov 10, 2026
a6aeac2
[None][chore] DFlash TRTLLM: drop context-kernel leftovers after the …
dc3671 Nov 10, 2026
54ccfb2
[None][chore] Revert no-op Helix V-up flag parametrization and autotu…
dc3671 Nov 10, 2026
af72cfa
[None][chore] CuteDSL FC12 MoE: drop the unreachable input-ready even…
dc3671 Nov 10, 2026
b84ed54
[None][fix] Rejection sampling: reject padding slots that alias the s…
dc3671 Nov 10, 2026
3f1e135
[None][chore] FP4 MLA packed plan test: pin the default rule plan as …
dc3671 Nov 10, 2026
ef653dd
[None][chore] FP4 MLA mufu16: delete the dead score-slot row-sum prot…
dc3671 Nov 10, 2026
793558a
[None][chore] DSpark attention op: one source for the shape contract,…
xguannv Nov 10, 2026
1efaab7
[None][chore] DSpark FP8 producers: reuse the dense GEMM UE8M0 helper…
xguannv Nov 10, 2026
1b38a68
[None][fix] DSpark: post_load_weights(weights) is the only attention-…
xguannv Nov 10, 2026
a020441
[None][chore] Annotate the rejection-guard padding test and the Helix…
xguannv Nov 10, 2026
6e14cb3
[None][chore] FP4 MLA packed plan test: drop the duplicated default-p…
xguannv Nov 10, 2026
7385f85
[None][chore] FP4 MLA mufu16: delete code left unreferenced by the de…
xguannv Nov 10, 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
4 changes: 3 additions & 1 deletion docs/source/features/speculative-decoding.md
Original file line number Diff line number Diff line change
Expand Up @@ -360,7 +360,9 @@ use FP8 on Rubin SM107 with TP=1 or attention DP, and BF16 otherwise. Set
bf16 to use dequantized BF16 draft weights. An explicit fp8 selection
requires TP=1 or attention DP; unsupported tensor-parallel mappings are
rejected during initialization. Non-Rubin FP8 is opt-in and requires workload
accuracy and acceptance-length validation.
accuracy and acceptance-length validation. The variable applies only to
DeepSeek-V4 drafters: DeepSeek-V4.1 DSpark always uses BF16 projections and
rejects fp8.

### Packed FP4 MLA wave planning

Expand Down

Large diffs are not rendered by default.

Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@

"""Blackwell DSpark attention specialization and persistent scheduler."""

from typing import Optional, Tuple, Type
from typing import Tuple, Type

import cutlass
import cutlass.cute as cute
Expand Down Expand Up @@ -205,18 +205,18 @@ class DSparkAttention(DSparkAttentionKernel):
pv_tiler_mn = (128, 256)
qk_tiler_k = 128
page_size_draft = 8
page_size_win = 128
window_len = 128

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.

[Nice to have] The rename leaves two class attributes for the same 128-row window. window_size (:200) is used for fixed_cache_seq_len. window_len is used only by the age mask (attention_kernel.py:1073-1074) and is validated only against itself. Suggest using window_size for the mask (or checking that the two match) and replacing the hard-coded 128s in the ring arithmetic. mO (attention_kernel.py:590) can also be typed cute.Tensor now that the output is always BF16.


@staticmethod
def _compute_grid(
o: cute.Tensor,
q_latent: cute.Tensor,
cluster_shape_mnk: cute.Shape,
max_active_clusters: int,
) -> Tuple[DSparkPersistentTileSchedulerParams, cute.Shape]:
"""Build scheduler parameters and cap the persistent DSpark grid."""
tile_sched_params = DSparkPersistentTileSchedulerParams(
cute.size(o.shape[3]),
cute.size(o.shape[2]),
cute.size(q_latent.shape[3]),
cute.size(q_latent.shape[2]),
cluster_shape_mnk,
)
grid = DSparkPersistentTileScheduler.get_grid_shape(tile_sched_params, max_active_clusters)
Expand All @@ -237,14 +237,13 @@ def __init__(
mma_pv_tiler_mn: Tuple[int, int],
max_active_clusters: int,
page_size_draft: int,
page_size_win: int,
window_len: int,
skip_correction_threshold: float,
*,
arch_str: str,
seq_len_q: int = block_size,
mma_qk_tiler_k: int = qk_tiler_k,
inverse_rope_dim: int = 0,
output_groups: int | None = None,
history_page_size: int = 32,
num_heads: int = num_heads,
):
Expand All @@ -256,7 +255,7 @@ def __init__(
"mma_qk_tiler_mn": (mma_qk_tiler_mn, self.qk_tiler_mn),
"mma_pv_tiler_mn": (mma_pv_tiler_mn, self.pv_tiler_mn),
"page_size_draft": (page_size_draft, self.page_size_draft),
"page_size_win": (page_size_win, self.page_size_win),
"window_len": (window_len, self.window_len),
"mma_qk_tiler_k": (mma_qk_tiler_k, self.qk_tiler_k),
}
mismatches = [
Expand All @@ -269,23 +268,19 @@ def __init__(
if seq_len_q not in (5, 6):
raise ValueError(f"DSpark block size must be 5 or 6, got {seq_len_q}")
if history_page_size not in (16, 32, 64, 128, 256):
raise ValueError("DSpark history page size must be 16, 32, 64, 128, or 256")
if inverse_rope_dim not in (0, 64):
raise ValueError(f"DSpark inverse_rope_dim must be 0 or 64, got {inverse_rope_dim}")
if output_groups is not None and (
output_groups <= 0 or self.num_heads % output_groups != 0
):
raise ValueError(
f"DSpark output groups must divide {self.num_heads}, got {output_groups}"
f"DSpark history page size must be 16, 32, 64, 128, or 256, got {history_page_size}"
)
if inverse_rope_dim not in (0, 64):
raise ValueError(f"DSpark inverse_rope_dim must be 0 or 64, got {inverse_rope_dim}")

super().__init__(
acc_dtype,
mma_qk_tiler_mn,
mma_pv_tiler_mn,
max_active_clusters,
page_size_draft,
page_size_win,
window_len,
skip_correction_threshold,
arch_str=arch_str,
seq_len_q=seq_len_q,
Expand All @@ -297,7 +292,6 @@ def __init__(
self.tma_page_size_draft = self.qk_tiler_mn[1]
self.fixed_cache_seq_len = self.window_size + seq_len_q
self.inverse_rope_dim = inverse_rope_dim
self.output_groups = output_groups
self.history_page_size = history_page_size
# Subdivide the 64-row loads at physical-page and cyclic-window boundaries.
self.history_tile_sizes = tuple(sorted({1, 4, 16, min(history_page_size, 64)}))
Expand All @@ -320,7 +314,6 @@ def wrapper(
attn_sink_ptr: cute.Pointer,
inverse_rope_freqs_ptr: cute.Pointer,
o_tensor: cute.Tensor,
o_scale_tensor: Optional[cute.Tensor],
softmax_scale: cutlass.Float32,
stream: cuda.CUstream,
):
Expand Down Expand Up @@ -362,10 +355,7 @@ def wrapper(
page_table_ptr,
layout=cute.make_layout((table_width, batch_size), stride=(1, table_stride)),
)
if cutlass.const_expr(self.output_groups is None):
output = cute.make_tensor(o_tensor.iterator, layout=q_layout)
else:
output = o_tensor
output = cute.make_tensor(o_tensor.iterator, layout=q_layout)
cache_seqs = cute.make_tensor(
cache_seqs_ptr,
layout=cute.make_layout((batch_size,), stride=(1,)),
Expand Down Expand Up @@ -395,7 +385,6 @@ def wrapper(
draft,
page_table,
output,
o_scale_tensor,
cache_seqs,
valid_len,
capacities,
Expand Down
Loading
Loading