Skip to content
Merged
Show file tree
Hide file tree
Changes from 4 commits
Commits
Show all changes
31 commits
Select commit Hold shift + click to select a range
2c7f33e
Initial draft version
kunlunl Nov 6, 2025
501ad65
Run through single GPU version
kunlunl Nov 10, 2025
c3754a7
Fix the usage of attention_mask
kunlunl Nov 10, 2025
451a36b
Update mask & indexer loss
kunlunl Nov 10, 2025
186dcdc
Minor changes about code style
kunlunl Nov 11, 2025
5271406
Fix TP
kunlunl Nov 17, 2025
6e6fb50
Fix attn mask and norm type
kunlunl Nov 18, 2025
237164e
Format
kunlunl Nov 18, 2025
72ba916
Resolve minor comments
kunlunl Nov 18, 2025
0e77355
Add unit test
kunlunl Nov 18, 2025
8851f4e
Fix UT
kunlunl Nov 20, 2025
7cdedf1
Use hidden_size to initialize indexer when q_lora_rank is None
kunlunl Nov 24, 2025
c915e4a
Add fused tilelang kernels
kunlunl Nov 25, 2025
de64232
Add indexer loss tracker
kunlunl Nov 25, 2025
1d6f4d3
Fix sparse indexer loss
kunlunl Nov 26, 2025
8495e03
Address minor comments
kunlunl Nov 27, 2025
ad04d58
Temporarily delete fused kernels
kunlunl Nov 27, 2025
9dae1ce
Fix lint error
kunlunl Nov 27, 2025
9cc697b
Make variable/class names more specific (SparseAttention -> DSA)
kunlunl Nov 27, 2025
b8f5c46
Merge branch 'dev' into kunlunl/deepseek_v3.2
kunlunl Nov 27, 2025
0d7e1d1
Fix lint error
kunlunl Nov 27, 2025
6630c7f
Merge linear-attention-type and sparse-attention-type to experimental…
kunlunl Nov 28, 2025
4b7f3c5
Merge branch 'dev' into kunlunl/deepseek_v3.2
kunlunl Nov 28, 2025
b1e309e
Rename test_sparse_attention to test_attention_variant_dsa
kunlunl Nov 28, 2025
6fe2f35
Minor fixes
kunlunl Nov 28, 2025
33684d2
Fix qk norm spec in dsa spec
kunlunl Nov 28, 2025
a0b6fd9
Add fast-hadamard-transform to pyproject.toml
kunlunl Nov 28, 2025
f22344a
Remove fast-hadamard-transform
kunlunl Dec 1, 2025
ce99e9f
Minor fix for args
kunlunl Dec 1, 2025
23ba310
Merge branch 'dev' into kunlunl/deepseek_v3.2
kunlunl Dec 1, 2025
cbfa053
Add mock hadamard_transformer
kunlunl Dec 1, 2025
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
2 changes: 2 additions & 0 deletions gpt_builders.py
Original file line number Diff line number Diff line change
Expand Up @@ -112,6 +112,7 @@ def _get_transformer_layer_spec(use_te, config):
args.moe_grouped_gemm,
args.qk_layernorm,
args.multi_latent_attention,
args.sparse_attention_type,
args.linear_attention_type,
moe_use_legacy_grouped_gemm=args.moe_use_legacy_grouped_gemm,
qk_l2_norm=args.qk_l2_norm,
Expand All @@ -123,6 +124,7 @@ def _get_transformer_layer_spec(use_te, config):
args.moe_grouped_gemm,
args.qk_layernorm,
args.multi_latent_attention,
args.sparse_attention_type,
args.linear_attention_type,
moe_use_legacy_grouped_gemm=args.moe_use_legacy_grouped_gemm,
normalization=args.normalization,
Expand Down
30 changes: 28 additions & 2 deletions megatron/core/models/gpt/gpt_layer_specs.py
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,9 @@
get_linear_attention_module_spec_for_backend,
)
from megatron.core.models.gpt.moe_module_specs import get_moe_module_spec_for_backend
from megatron.core.models.gpt.sparse_attention_module_specs import (
get_sparse_attention_module_spec_for_backend,
)
from megatron.core.transformer.attention import SelfAttention, SelfAttentionSubmodules
from megatron.core.transformer.enums import AttnMaskType, LayerType
from megatron.core.transformer.identity_op import IdentityOp
Expand Down Expand Up @@ -77,6 +80,7 @@ def get_gpt_layer_with_transformer_engine_spec(
moe_grouped_gemm: Optional[bool] = False,
qk_layernorm: Optional[bool] = False,
multi_latent_attention: Optional[bool] = False,
sparse_attention_type: Optional[str] = None,
linear_attention_type: Optional[str] = None,
fp8: Optional[str] = None, # pylint: disable=unused-argument
moe_use_legacy_grouped_gemm: Optional[bool] = False,
Expand All @@ -94,6 +98,7 @@ def get_gpt_layer_with_transformer_engine_spec(
moe_grouped_gemm (bool, optional): To use Grouped GEMM. Defaults to False.
qk_layernorm (bool, optional): To use layernorm for queries/keys. Defaults to False.
multi_latent_attention (bool, optional): To use multi-latent attention. Defaults to False.
sparse_attention_type (str, optional): The type of sparse attention. Defaults to None.
linear_attention_type (str, optional): The type of linear attention. Defaults to None.
fp8 (str, optional): Deprecated. For temporary Nemo compatibility.
moe_use_legacy_grouped_gemm (bool, optional): Force use the legacy GroupedMLP.
Expand Down Expand Up @@ -129,6 +134,7 @@ def get_gpt_layer_with_transformer_engine_spec(
attention = get_attention_module_spec_for_backend(
backend=backend,
sharded_state_dict_keys_map=sharded_state_dict_keys_map,
sparse_attention_type=sparse_attention_type,
linear_attention_type=linear_attention_type,
qk_layernorm=qk_layernorm,
qk_l2_norm=qk_l2_norm,
Expand Down Expand Up @@ -161,6 +167,7 @@ def get_gpt_layer_local_spec(
moe_grouped_gemm: Optional[bool] = False,
qk_layernorm: Optional[bool] = False,
multi_latent_attention: Optional[bool] = False,
sparse_attention_type: Optional[str] = None,
linear_attention_type: Optional[str] = None,
fp8: Optional[str] = None, # pylint: disable=unused-argument
moe_use_legacy_grouped_gemm: Optional[bool] = False,
Expand All @@ -176,6 +183,7 @@ def get_gpt_layer_local_spec(
moe_grouped_gemm (bool, optional): To use Grouped GEMM. Defaults to False.
qk_layernorm (bool, optional): To use layernorm for queries/keys. Defaults to False.
multi_latent_attention (bool, optional): To use multi-latent attention. Defaults to False.
sparse_attention_type (str, optional): The type of sparse attention. Defaults to None.
linear_attention_type (str, optional): The type of linear attention. Defaults to None.
fp8 (str, optional): Deprecated. For temporary Nemo compatibility.
moe_use_legacy_grouped_gemm (bool, optional): Force use the legacy GroupedMLP.
Expand All @@ -200,6 +208,9 @@ def get_gpt_layer_local_spec(
" and will be removed soon. Please update your code accordingly."
)

if sparse_attention_type is not None:
raise NotImplementedError("Sparse attention is not supported with local spec yet.")

if linear_attention_type is not None:
raise NotImplementedError("Linear attention is not supported with local spec yet.")

Expand All @@ -208,6 +219,7 @@ def get_gpt_layer_local_spec(
attention = get_attention_module_spec_for_backend(
backend=backend,
sharded_state_dict_keys_map=sharded_state_dict_keys_map,
sparse_attention_type=sparse_attention_type,
linear_attention_type=linear_attention_type,
qk_layernorm=qk_layernorm,
qk_l2_norm=qk_l2_norm,
Expand Down Expand Up @@ -272,6 +284,7 @@ def get_transformer_layer_spec_for_backend(
def get_attention_module_spec_for_backend(
backend: BackendSpecProvider,
sharded_state_dict_keys_map: dict,
sparse_attention_type: Optional[str] = None,
linear_attention_type: Optional[str] = None,
qk_layernorm: Optional[bool] = False,
qk_l2_norm: Optional[bool] = False,
Expand All @@ -292,6 +305,16 @@ def get_attention_module_spec_for_backend(
rms_norm = normalization == "RMSNorm"
qk_norm = backend.layer_norm(rms_norm=rms_norm, for_qk=True)

if sparse_attention_type is not None:
assert multi_latent_attention, "Currently only MLA supports sparse attention."
core_attention = get_sparse_attention_module_spec_for_backend(
backend=backend,
sparse_attention_type=sparse_attention_type,
normalization=normalization,
)
else:
core_attention = backend.core_attention()

if multi_latent_attention:
assert qk_l2_norm is False, "qk_l2_norm is not supported with MLA."
linear_q_down_proj = (
Expand Down Expand Up @@ -328,7 +351,7 @@ def get_attention_module_spec_for_backend(
linear_q_up_proj=linear_q_up_proj,
linear_kv_down_proj=linear_kv_down_proj,
linear_kv_up_proj=linear_kv_up_proj,
core_attention=backend.core_attention(),
core_attention=core_attention,
linear_proj=backend.row_parallel_linear(),
q_layernorm=qk_norm,
kv_layernorm=qk_norm,
Expand All @@ -352,7 +375,7 @@ def get_attention_module_spec_for_backend(
params={"attn_mask_type": AttnMaskType.causal},
submodules=SelfAttentionSubmodules(
linear_qkv=linear_qkv,
core_attention=backend.core_attention(),
core_attention=core_attention,
linear_proj=backend.row_parallel_linear(),
q_layernorm=qk_norm,
k_layernorm=qk_norm,
Expand Down Expand Up @@ -522,15 +545,18 @@ def get_gpt_decoder_block_spec(
continue
linear_attention_type = config.linear_attention_type
multi_latent_attention = None
sparse_attention_type = None
else:
linear_attention_type = None
multi_latent_attention = config.multi_latent_attention
sparse_attention_type = config.sparse_attention_type

layer_spec_key = f"{mlp_type}_{attention_type}"
layer_spec_dict[layer_spec_key] = get_layer_spec_fn(
num_experts=num_experts,
moe_grouped_gemm=moe_grouped_gemm,
multi_latent_attention=multi_latent_attention,
sparse_attention_type=sparse_attention_type,
linear_attention_type=linear_attention_type,
**get_layer_spec_kwargs,
)
Expand Down
62 changes: 62 additions & 0 deletions megatron/core/models/gpt/sparse_attention_module_specs.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,62 @@
# Copyright (c) 2025, NVIDIA CORPORATION. All rights reserved.

from typing import Optional

from megatron.core.models.backends import BackendSpecProvider
from megatron.core.transformer.sparse_attention import (
Indexer,
IndexerSubmodules,
SparseAttention,
SparseAttentionSubmodules,
)
from megatron.core.transformer.spec_utils import ModuleSpec


def get_indexer_spec_for_backend(
backend: BackendSpecProvider, normalization: Optional[str] = None
) -> ModuleSpec:
"""Helper function to get Indexer module spec for a given backend.

Args:
backend: Backend specification provider (TE or Local).
normalization: Normalization type ("RMSNorm" or None for LayerNorm).

Returns:
ModuleSpec for Indexer with appropriate submodules.
"""
rms_norm = normalization == "RMSNorm"
return ModuleSpec(
module=Indexer,
submodules=IndexerSubmodules(
linear_wq_b=backend.linear(),
linear_wk=backend.linear(),
k_norm=backend.layer_norm(rms_norm=rms_norm, for_qk=True),
linear_weights_proj=backend.linear(),
),
)


def get_sparse_attention_module_spec_for_backend(
backend: BackendSpecProvider, sparse_attention_type: str, normalization: Optional[str] = None
) -> ModuleSpec:
"""Helper function to get module spec for Sparse Attention.

Args:
backend: Backend specification provider (TE or Local).
sparse_attention_type: Type of sparse attention.
normalization: Normalization type ("RMSNorm" or None for LayerNorm).

Returns:
ModuleSpec for the sparse attention implementation with appropriate submodules.
"""
if sparse_attention_type == "dsa":
# Because TransformerEngine does not support sparse attention yet, we use local
# implementation whether the backend is TransformerEngine or not.
return ModuleSpec(
module=SparseAttention,
submodules=SparseAttentionSubmodules(
indexer=get_indexer_spec_for_backend(backend, normalization=normalization),
),
)
else:
raise ValueError(f"Invalid sparse attention type: {sparse_attention_type}")
2 changes: 2 additions & 0 deletions megatron/core/transformer/attention.py
Original file line number Diff line number Diff line change
Expand Up @@ -172,6 +172,7 @@ def __init__(
self.key_hidden_size = self.hidden_size_per_attention_head
self.val_hidden_size = self.hidden_size_per_attention_head

# TODO: This is built twice when using MLA, should be refactored.
self.core_attention = build_module(
submodules.core_attention,
config=self.config,
Expand All @@ -188,6 +189,7 @@ def __init__(
and "core_attn" in self.config.recompute_modules
)

# TODO: This is built twice when using MLA, should be refactored.
# Output.
self.linear_proj = build_module(
submodules.linear_proj,
Expand Down
77 changes: 54 additions & 23 deletions megatron/core/transformer/multi_latent_attention.py
Original file line number Diff line number Diff line change
Expand Up @@ -233,13 +233,22 @@ def forward(
# Get the query, key and value tensors based on the type of attention -
# self or cross attn.
# query: [96, 1, 16, 128], key:[96, 1, 16, 128], value:[96, 1, 16, 128]
query, key, value = self.get_query_key_value_tensors(
hidden_states,
key_value_states,
position_ids,
packed_seq_params,
inference_context=inference_context,
)
if self.config.sparse_attention_type is None:
query, key, value = self.get_query_key_value_tensors(
hidden_states,
key_value_states,
position_ids,
packed_seq_params,
inference_context=inference_context,
)
else:
query, key, value, q_compressed, _ = self.get_query_key_value_and_compressed_tensors(
hidden_states,
key_value_states,
position_ids,
packed_seq_params,
inference_context=inference_context,
)

# ===================================================
# Adjust key, value for inference
Expand Down Expand Up @@ -267,14 +276,28 @@ def forward(
)
else:
if inference_context is None or inference_context.is_static_batching():
core_attn_out = self.core_attention(
query,
key,
value,
attention_mask,
packed_seq_params=packed_seq_params,
attn_mask_type=attn_mask_type,
)
if self.config.sparse_attention_type is None:
core_attn_out = self.core_attention(
query,
key,
value,
attention_mask,
packed_seq_params=packed_seq_params,
attn_mask_type=attn_mask_type,
)
else:
# For sparse attention, use a specialized forward.
core_attn_out = self.core_attention(
query,
key,
value,
x=hidden_states,
qr=q_compressed,
attention_mask=attention_mask,
attn_mask_type=attn_mask_type,
attention_bias=None,
packed_seq_params=packed_seq_params,
)
elif self.cache_mla_latents:
# Dynamic batching attention kernel.
q, k, v = (query, key, value)
Expand Down Expand Up @@ -460,7 +483,7 @@ def __init__(
eps=self.config.layernorm_epsilon,
)

def get_query_key_value_tensors(
def get_query_key_value_and_compressed_tensors(
self,
hidden_states,
key_value_states=None,
Expand Down Expand Up @@ -578,6 +601,16 @@ def get_query_key_value_tensors(
kv_compressed = kv_compressed.squeeze(1)
k_pos_emb = k_pos_emb.squeeze(1)

# =========================================
# Apply norm
# =========================================

if self.config.q_lora_rank is not None:
# q_compressed: [num_tokens, q_lora_rank]
q_compressed = self.q_layernorm(q_compressed)

kv_compressed = self.kv_layernorm(kv_compressed)
Comment thread
yuzhongw-nvidia marked this conversation as resolved.

# =========================================
# QKV up projection and RoPE apply
# =========================================
Expand All @@ -588,7 +621,6 @@ def qkv_up_proj_and_rope_apply_for_cached_latent_kv(
if self.config.q_lora_rank is not None:
# q_compressed: [num_tokens, q_lora_rank]
# q: [num_tokens, n * (qk_head_dim + qk_pos_emb_head_dim)]
q_compressed = self.q_layernorm(q_compressed)
q, _ = self.linear_q_up_proj(q_compressed)
else:
# q_compressed: [num_tokens, hidden_size]
Expand All @@ -598,8 +630,6 @@ def qkv_up_proj_and_rope_apply_for_cached_latent_kv(
# q: [num_tokens, n, q_head_dim]
q = q.view(*q.size()[:-1], self.num_attention_heads_per_partition, self.q_head_dim)

kv_compressed = self.kv_layernorm(kv_compressed)

# [num_tokens, qk_pos_emb_head_dim] -> [num_tokens, 1, qk_pos_emb_head_dim]
k_pos_emb = torch.unsqueeze(k_pos_emb, -2)

Expand Down Expand Up @@ -663,7 +693,6 @@ def qkv_up_proj_and_rope_apply(q_compressed, kv_compressed, k_pos_emb, rotary_po
if self.config.q_lora_rank is not None:
# q_compressed: [num_tokens, q_lora_rank]
# q: [num_tokens, n * (qk_head_dim + qk_pos_emb_head_dim)]
q_compressed = self.q_layernorm(q_compressed)
q, _ = self.linear_q_up_proj(q_compressed)
else:
# q_compressed: [num_tokens, hidden_size]
Expand All @@ -673,8 +702,6 @@ def qkv_up_proj_and_rope_apply(q_compressed, kv_compressed, k_pos_emb, rotary_po
# q: [num_tokens, n, q_head_dim]
q = q.view(*q.size()[:-1], self.num_attention_heads_per_partition, self.q_head_dim)

kv_compressed = self.kv_layernorm(kv_compressed)

# kv: [num_tokens, n * (qk_head_dim + v_head_dim)]
kv, _ = self.linear_kv_up_proj(kv_compressed)

Expand Down Expand Up @@ -798,7 +825,11 @@ def qkv_up_proj_and_rope_apply(q_compressed, kv_compressed, k_pos_emb, rotary_po
q_compressed, kv_compressed, k_pos_emb, rotary_pos_emb
)

return query, key, value
return query, key, value, q_compressed, kv_compressed

def get_query_key_value_tensors(self, *args, **kwargs):
# Only return query, key and value.
return get_query_key_value_and_compressed_tensors(self, *args, **kwargs)[:3]
Comment thread
kunlunl marked this conversation as resolved.
Outdated

def uncompress_kv_from_cache(self, kv_cached):
"""
Expand Down
Loading