04-model_executor_layers_sparse_attn_indexer.py.diff 8.4 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191
  1. --- baked-ours/model_executor/layers/sparse_attn_indexer.py 2026-09-10 20:13:20.000000000 +0200
  2. +++ baked/model_executor/layers/sparse_attn_indexer.py 2026-09-11 10:13:47.000000000 +0200
  3. @@ -17,6 +17,9 @@
  4. apply_candidate_mask as _apply_candidate_mask,
  5. )
  6. from vllm.model_executor.kernels.attention.dsa.candidate_blocks import (
  7. + _PrefillTimer as _DSV41PrefillTimer,
  8. +)
  9. +from vllm.model_executor.kernels.attention.dsa.candidate_blocks import (
  10. select_candidate_blocks as _select_candidate_blocks,
  11. )
  12. from vllm.model_executor.layers.quantization.utils.quant_utils import (
  13. @@ -325,6 +328,7 @@
  14. candidate_blocks: torch.Tensor | None = None,
  15. candidate_block_size: int = 0,
  16. candidate_write: bool = False,
  17. + compress_ratio: int = 1,
  18. ) -> torch.Tensor:
  19. # careful! this will be None in dummy run
  20. forward_context = get_forward_context()
  21. @@ -332,13 +336,25 @@
  22. fp8_dtype = current_platform.fp8_dtype()
  23. k_cache_prefix = _resolve_layer_name(k_cache_prefix)
  24. + # v4.1 two-level candidate filtering under DCP: the logits columns are this
  25. + # rank's local compressed states; with interleave 1 and world | block, local
  26. + # block b (size block // world) is exactly this rank's shard of global block
  27. + # b, so candidate ids need no translation. The producer MAX-reduces block
  28. + # scores across ranks and pins the global newest block (candidate_blocks.py).
  29. + local_cand_bs = candidate_block_size
  30. + dcp_group = None
  31. if candidate_blocks is not None:
  32. - # Candidate blocks are request-local; the DCP-sharded logits layout
  33. - # would need per-rank translation that is not implemented.
  34. - assert dcp_world_size == 1, (
  35. - "v4.1 two-level candidate filtering is not supported with DCP."
  36. - )
  37. assert candidate_block_size > 0
  38. + if dcp_world_size > 1:
  39. + assert cp_kv_cache_interleave_size == 1, (
  40. + "v4.1 candidate filtering under DCP needs cp_kv_cache_interleave_size=1"
  41. + )
  42. + assert candidate_block_size % dcp_world_size == 0, (
  43. + f"candidate_block_size ({candidate_block_size}) must be a multiple "
  44. + f"of dcp_world_size ({dcp_world_size})"
  45. + )
  46. + local_cand_bs = candidate_block_size // dcp_world_size
  47. + dcp_group = get_dcp_group()
  48. # assert isinstance(attn_metadata, dict)
  49. if not isinstance(attn_metadata, dict):
  50. @@ -494,6 +510,23 @@
  51. if chunk.local_total_seq_lens == 0:
  52. logits = q_slice.new_empty((q_slice.shape[0], 0), dtype=torch.float32)
  53. topk_indices.fill_(-1)
  54. + if candidate_blocks is not None and candidate_write and dcp_group is not None:
  55. + # Other ranks may hold states for this chunk: join the MAX
  56. + # reduce with an empty local contribution.
  57. + _select_candidate_blocks(
  58. + logits,
  59. + cu_seqlen_ks,
  60. + cu_seqlen_ke,
  61. + candidate_blocks.shape[1],
  62. + local_cand_bs,
  63. + candidate_blocks[chunk.token_start : chunk.token_end],
  64. + dcp_group=dcp_group,
  65. + global_row_ke=chunk.global_ctx_lens,
  66. + global_block_size=candidate_block_size,
  67. + max_global_blocks=triton.cdiv(
  68. + chunk.max_global_ctx, candidate_block_size
  69. + ),
  70. + )
  71. else:
  72. # DeepGEMM scalar-type tags (zero-copy): MXFP4 values → int8
  73. # (kPackedFP4), scales → int32 squeezed to 1-D kv_sf / 2-D q_sf.
  74. @@ -539,8 +572,16 @@
  75. cu_seqlen_ks,
  76. cu_seqlen_ke,
  77. chunk_candidates.shape[1],
  78. - candidate_block_size,
  79. + local_cand_bs,
  80. chunk_candidates,
  81. + dcp_group=dcp_group,
  82. + global_row_ke=chunk.global_ctx_lens,
  83. + global_block_size=candidate_block_size,
  84. + max_global_blocks=(
  85. + triton.cdiv(chunk.max_global_ctx, candidate_block_size)
  86. + if dcp_group is not None
  87. + else None
  88. + ),
  89. )
  90. else:
  91. _apply_candidate_mask(
  92. @@ -548,7 +589,7 @@
  93. cu_seqlen_ks,
  94. cu_seqlen_ke,
  95. chunk_candidates,
  96. - candidate_block_size,
  97. + local_cand_bs,
  98. )
  99. ops.top_k_per_row_prefill(
  100. logits,
  101. @@ -666,14 +707,44 @@
  102. vis = vis[:num_rows]
  103. decode_candidates = candidate_blocks[:num_rows]
  104. if candidate_write:
  105. + global_row_ke = None
  106. + max_global_blocks = None
  107. + if dcp_group is not None:
  108. + # Global (un-sharded) per-row context in STATES for the
  109. + # newest-block pin: per token when flattened / next_n == 1,
  110. + # per request (row_repeat) or expanded per row otherwise.
  111. + gl = decode_metadata.global_seq_lens
  112. + assert gl is not None
  113. + gl = gl.reshape(-1)
  114. + if row_repeat > 1:
  115. + gl = gl[:batch_size]
  116. + elif gl.numel() != num_rows:
  117. + # native spec decode: request lens -> per-row lens
  118. + gl = (
  119. + gl[:batch_size].unsqueeze(1)
  120. + - next_n
  121. + + 1
  122. + + torch.arange(next_n, device=gl.device, dtype=gl.dtype)
  123. + ).clamp_(min=0).reshape(-1)
  124. + else:
  125. + gl = gl[:num_rows]
  126. + global_row_ke = gl // compress_ratio
  127. + max_global_blocks = triton.cdiv(
  128. + attn_metadata_narrowed.max_seq_len // compress_ratio,
  129. + candidate_block_size,
  130. + )
  131. _select_candidate_blocks(
  132. logits,
  133. None,
  134. vis,
  135. decode_candidates.shape[1],
  136. - candidate_block_size,
  137. + local_cand_bs,
  138. decode_candidates,
  139. row_repeat,
  140. + dcp_group=dcp_group,
  141. + global_row_ke=global_row_ke,
  142. + global_block_size=candidate_block_size,
  143. + max_global_blocks=max_global_blocks,
  144. )
  145. else:
  146. _apply_candidate_mask(
  147. @@ -681,7 +752,7 @@
  148. None,
  149. vis,
  150. decode_candidates,
  151. - candidate_block_size,
  152. + local_cand_bs,
  153. row_repeat,
  154. )
  155. topk_indices = topk_indices_buffer[:num_padded_tokens, :topk_tokens]
  156. @@ -790,6 +861,7 @@
  157. candidate_blocks: torch.Tensor | None = None,
  158. candidate_block_size: int = 0,
  159. candidate_write: bool = False,
  160. + compress_ratio: int = 1,
  161. ) -> torch.Tensor:
  162. return topk_indices_buffer
  163. @@ -904,6 +976,17 @@
  164. k: torch.Tensor | None,
  165. weights: torch.Tensor,
  166. ):
  167. + # Stage 7: phase "indexer" of the prefill timing (inert unless enabled).
  168. + with _DSV41PrefillTimer("indexer", int(hidden_states.shape[0])):
  169. + return self._forward_cuda_impl(hidden_states, q_quant, k, weights)
  170. +
  171. + def _forward_cuda_impl(
  172. + self,
  173. + hidden_states: torch.Tensor,
  174. + q_quant: torch.Tensor | tuple[torch.Tensor, torch.Tensor],
  175. + k: torch.Tensor | None,
  176. + weights: torch.Tensor,
  177. + ):
  178. # FP8 path: single tensor (per-token scale is folded into `weights`).
  179. # FP4 path: (values, scales) tuple with scales required by the kernel.
  180. if isinstance(q_quant, tuple):
  181. @@ -935,6 +1018,7 @@
  182. candidate_blocks=self.candidate_blocks,
  183. candidate_block_size=self.candidate_block_size,
  184. candidate_write=self.candidate_write,
  185. + compress_ratio=self.compress_ratio,
  186. )
  187. def forward_xpu(