Skip to content

Commit fe9cdac

Browse files
sufubaoniushengxiao
authored andcommitted
support glm52
1 parent cc74796 commit fe9cdac

41 files changed

Lines changed: 2446 additions & 201 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

docs/CN/source/models/supported_models.rst

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,8 @@ lightllm 支持大多数的主流的开源大语言模型以及多模态模型
5656
-
5757
* - `Qwen3-Moe <https://github.com/QwenLM/Qwen3>`_
5858
-
59+
* - `GLM-5.2 <https://huggingface.co/zai-org/GLM-5.2>`_
60+
- 支持 BF16/FP8 和 MTP。
5961

6062

6163
多模态模型
@@ -93,4 +95,4 @@ Reward模型
9395
* - `internLM-reward <https://huggingface.co/internlm/internlm2-1_8b-reward>`_
9496
- :code:`--use_reward_model`
9597
* - `Qwen2-Reward <https://huggingface.co/Qwen/Qwen2-Reward>`_
96-
- :code:`--use_reward_model`
98+
- :code:`--use_reward_model`

docs/EN/source/models/supported_models.rst

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,8 @@ Large Language Models
5656
-
5757
* - `DeepSeek-V3.2 `_
5858
-
59+
* - `GLM-5.2 <https://huggingface.co/zai-org/GLM-5.2>`_
60+
- Supports BF16/FP8 and MTP.
5961

6062
Multimodal Models
6163
^^^^^^^^^^^^^^^^^
@@ -94,4 +96,3 @@ Reward Models
9496
- :code:`--use_reward_model`
9597
* - `Qwen2-Reward <https://huggingface.co/Qwen/Qwen2-Reward>`_
9698
- :code:`--use_reward_model`
97-

lightllm/common/basemodel/attention/nsa/flashmla_sparse.py

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,7 @@
33

44
import dataclasses
55
import torch
6+
import torch.nn.functional as F
67
from typing import Tuple, TYPE_CHECKING
78

89
from ..base_att import BaseAttBackend, BasePrefillAttState, BaseDecodeAttState, AttControl
@@ -86,14 +87,20 @@ def _nsa_prefill_att(
8687
if topk_mem_indices.ndim == 2:
8788
topk_mem_indices = topk_mem_indices.unsqueeze(1)
8889

90+
real_head_num = q.shape[1]
91+
head_block_size = 64
92+
pad_head_num = (-real_head_num) % head_block_size
93+
if pad_head_num:
94+
q = F.pad(q, (0, 0, 0, pad_head_num))
95+
8996
mla_out, _, _ = flash_mla_sparse_fwd(
9097
q=q,
9198
kv=kv,
9299
indices=topk_mem_indices,
93100
sm_scale=softmax_scale,
94101
d_v=kv_lora_rank,
95102
)
96-
return mla_out
103+
return mla_out[:, :real_head_num, :]
97104

98105

99106
@dataclasses.dataclass

lightllm/common/basemodel/attention/nsa/fp8_flashmla_sparse.py

Lines changed: 30 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,6 @@
11
import dataclasses
22
import torch
3+
import torch.nn.functional as F
34
from typing import TYPE_CHECKING, Tuple
45

56
from ..base_att import AttControl, BaseAttBackend, BaseDecodeAttState, BasePrefillAttState
@@ -70,10 +71,9 @@ def _nsa_prefill_att(
7071
packed_kv: torch.Tensor,
7172
att_control: AttControl,
7273
) -> torch.Tensor:
73-
import flash_mla
74+
from flash_mla import flash_mla_sparse_fwd
7475

7576
nsa_dict = att_control.nsa_prefill_dict
76-
topk_indices = nsa_dict["topk_indices"]
7777
softmax_scale = nsa_dict["softmax_scale"]
7878
kv_lora_rank = nsa_dict["kv_lora_rank"]
7979
topk_mem_indices = nsa_dict["topk_mem_indices"]
@@ -91,18 +91,25 @@ def _nsa_prefill_att(
9191
)
9292
else:
9393
kv = prefill_cache_kv
94+
topk_indices = nsa_dict["topk_indices"]
9495

9596
if topk_indices.ndim == 2:
9697
topk_indices = topk_indices.unsqueeze(1)
9798

98-
mla_out, _, _ = flash_mla.flash_mla_sparse_fwd(
99+
real_head_num = q.shape[1]
100+
head_block_size = 64
101+
pad_head_num = (-real_head_num) % head_block_size
102+
if pad_head_num:
103+
q = F.pad(q, (0, 0, 0, pad_head_num))
104+
105+
mla_out, _, _ = flash_mla_sparse_fwd(
99106
q=q,
100107
kv=kv,
101108
indices=topk_indices,
102109
sm_scale=softmax_scale,
103110
d_v=kv_lora_rank,
104111
)
105-
return mla_out
112+
return mla_out[:, :real_head_num, :]
106113

107114

108115
@dataclasses.dataclass
@@ -141,9 +148,9 @@ def init_state(self):
141148
ragged_mem_index=self.ragged_mem_index,
142149
hold_req_idx=self.infer_state.req_manager.HOLD_REQUEST_ID,
143150
)
144-
import flash_mla
151+
from flash_mla import get_mla_metadata
145152

146-
self.flashmla_sched_meta, _ = flash_mla.get_mla_metadata()
153+
self.flashmla_sched_meta, _ = get_mla_metadata()
147154
return
148155

149156
def decode_att(
@@ -164,7 +171,7 @@ def _nsa_decode_att(
164171
packed_kv: torch.Tensor,
165172
att_control: AttControl,
166173
) -> torch.Tensor:
167-
import flash_mla
174+
from flash_mla import flash_mla_with_kvcache
168175

169176
nsa_dict = att_control.nsa_decode_dict
170177
topk_mem_indices = nsa_dict["topk_mem_indices"]
@@ -177,13 +184,26 @@ def _nsa_decode_att(
177184

178185
q_nope, q_rope = q
179186
q_all = torch.cat([q_nope, q_rope], dim=-1).unsqueeze(1).contiguous()
187+
188+
real_head_num = q_all.shape[2]
189+
if real_head_num <= 64:
190+
padded_head_num = 64
191+
elif real_head_num <= 128:
192+
padded_head_num = 128
193+
else:
194+
padded_head_num = real_head_num
195+
if padded_head_num != real_head_num:
196+
q_all = F.pad(q_all, (0, 0, 0, padded_head_num - real_head_num))
197+
198+
# Token-granular "pages": indices=topk_mem_indices are absolute KV-pool slots,
199+
# so the kv cache is viewed as (num_tokens, 1, 1, bytes) and no block_table is needed.
180200
kv = torch.as_strided(
181201
packed_kv,
182202
size=(packed_kv.shape[0], 1, 1, packed_kv.shape[-1]),
183203
stride=(packed_kv.stride(0), packed_kv.shape[-1], packed_kv.shape[-1], packed_kv.stride(-1)),
184204
)
185205

186-
o_tensor, _ = flash_mla.flash_mla_with_kvcache(
206+
o_tensor, _ = flash_mla_with_kvcache(
187207
q=q_all,
188208
k_cache=kv,
189209
block_table=None,
@@ -193,6 +213,7 @@ def _nsa_decode_att(
193213
softmax_scale=softmax_scale,
194214
causal=False,
195215
is_fp8_kvcache=True,
196-
indices=topk_mem_indices,
216+
indices=topk_mem_indices.to(dtype=torch.int32),
197217
)
218+
o_tensor = o_tensor[:, :, :real_head_num, :]
198219
return o_tensor[:, 0, :, :] # [b, 1, h, d] -> [b, h, d]

lightllm/common/basemodel/basemodel.py

Lines changed: 1 addition & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1177,12 +1177,7 @@ def _init_padded_req(self):
11771177
def _gen_special_model_input(self, token_num: int):
11781178
special_model_input = {}
11791179

1180-
is_mtp_draft_model = (
1181-
"Deepseek3MTPModel" in str(self.__class__)
1182-
or "Qwen3MOEMTPModel" in str(self.__class__)
1183-
or "MistralMTPModel" in str(self.__class__)
1184-
or "Glm4MoeLiteMTPModel" in str(self.__class__)
1185-
)
1180+
is_mtp_draft_model = getattr(self, "is_mtp_draft_model", False)
11861181
if is_mtp_draft_model:
11871182
special_model_input["mtp_draft_input_hiddens"] = torch.randn(
11881183
token_num, self.config["hidden_size"], dtype=self.data_type, device="cuda"

lightllm/common/basemodel/layer_weights/meta_weights/norm_weight.py

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,7 @@
33
from .base_weight import BaseWeightTpl
44
from lightllm.utils.dist_utils import get_current_device_id, get_current_rank_in_dp, get_dp_world_size
55
from lightllm.common.basemodel.triton_kernel.norm.rmsnorm import rmsnorm_forward
6+
from lightllm.common.basemodel.triton_kernel.norm.fused_add_rmsnorm import fused_add_rmsnorm_forward
67
from lightllm.common.basemodel.triton_kernel.norm.layernorm import layernorm_forward
78
from lightllm.common.basemodel.triton_kernel.norm.qk_norm import qk_rmsnorm_fused_forward
89
from lightllm.common.basemodel.triton_kernel.norm.gated_rmsnorm import gated_rmsnorm_forward
@@ -71,6 +72,21 @@ def __call__(
7172
) -> torch.Tensor:
7273
return self._forward(input=input, eps=eps, out=out, alloc_func=alloc_func)
7374

75+
def fused_add_forward(
76+
self,
77+
residual: torch.Tensor,
78+
x: torch.Tensor,
79+
eps: float,
80+
out: Optional[torch.Tensor] = None,
81+
alloc_func=torch.empty,
82+
) -> torch.Tensor:
83+
"""Fused residual-add + RMSNorm: ``residual <- residual + x`` (in place) and return
84+
``rmsnorm(residual) * weight``. Bit-identical to a plain ``residual.add_(x)`` followed
85+
by ``__call__`` but in a single Triton launch. CUDA/MUSA (Triton) only."""
86+
if out is None:
87+
out = alloc_func(residual.shape, dtype=residual.dtype, device=residual.device)
88+
return fused_add_rmsnorm_forward(residual=residual, x=x, weight=self.weight, eps=eps, out=out)
89+
7490

7591
class GatedRMSNormWeight(RMSNormWeight):
7692
def _triton_forward(

lightllm/common/basemodel/triton_kernel/fused_moe/grouped_fused_moe.py

Lines changed: 47 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@
1818
# See the License for the specific language governing permissions and
1919
# limitations under the License.
2020

21+
import os
2122
import torch
2223
import triton
2324
import triton.language as tl
@@ -26,6 +27,7 @@
2627
from lightllm.utils.vllm_utils import vllm_ops
2728
from lightllm.utils.device_utils import triton_support_tensor_descriptor
2829
from .moe_silu_and_mul import silu_and_mul_fwd
30+
from .moe_silu_and_mul_group_quant import silu_and_mul_group_quant_fwd
2931
from .moe_sum_reduce import moe_sum_reduce
3032
from lightllm.common.basemodel.triton_kernel.quantization.fp8act_quant_kernel import per_token_group_quant_fp8
3133
from lightllm.utils.torch_ops_utils import direct_register_custom_op
@@ -35,6 +37,11 @@
3537

3638
logger = init_logger(__name__)
3739

40+
# Fuse the down-projection's per-token-group fp8 quant into the silu_and_mul that produces its
41+
# input (eliminates a separate per_token_group_quant launch per MoE layer). Numerically matches
42+
# the unfused path; gate exists for A/B and rollback.
43+
_ENABLE_FUSED_SILU_QUANT = os.environ.get("LIGHTLLM_FUSED_SILU_QUANT", "1") == "1"
44+
3845

3946
@triton.jit
4047
def moe_align_kernel(
@@ -860,8 +867,12 @@ def grouped_matmul(
860867
assert BLOCK_SIZE_K == triton.next_power_of_2(BLOCK_SIZE_K)
861868

862869
if use_fp8_w8a8:
870+
if token_inputs.dtype == expert_weights.dtype:
871+
# token_inputs were already fp8-quantized by a fused producer (e.g. the fused
872+
# silu+mul+group-quant on the down projection); token_input_scale is its scale.
873+
assert token_input_scale is not None, "pre-quantized fp8 input requires its scale"
863874
# 当权重使用 block wise 量化时,激活也使用 per token, group size 量化
864-
if block_size_k == 0:
875+
elif block_size_k == 0:
865876
# input 使用 per token 量化
866877
token_inputs, token_input_scale = vllm_ops.scaled_fp8_quant(
867878
token_inputs, token_input_scale, use_per_token_if_dynamic=True
@@ -1082,18 +1093,44 @@ def fused_experts_impl(
10821093
bias=w1_bias,
10831094
)
10841095

1085-
silu_and_mul_fwd(
1086-
intermediate_cache1.view(-1, N),
1087-
intermediate_cache2.view(-1, N // 2),
1088-
limit=limit,
1089-
alpha=alpha,
1090-
layout=layout,
1091-
)
1096+
# Fuse the down-projection's per-token-group fp8 quant into silu_and_mul when the down
1097+
# weight is block-wise fp8 quantized: the silu output is emitted directly as fp8 + scales,
1098+
# so grouped_matmul below skips its internal per_token_group_quant launch.
1099+
use_fused_silu_quant = _ENABLE_FUSED_SILU_QUANT and use_fp8_w8a8 and w2_scale is not None and w2_scale.ndim == 3
1100+
if use_fused_silu_quant:
1101+
down_token_num = curr_topk_ids.numel()
1102+
down_k = N // 2
1103+
down_group_size = down_k // w2_scale.shape[2]
1104+
down_inputs = alloc_tensor_func(
1105+
(down_token_num, down_k), device=hidden_states.device, dtype=torch.float8_e4m3fn
1106+
)
1107+
down_input_scale = alloc_tensor_func(
1108+
(down_token_num, down_k // down_group_size), device=hidden_states.device, dtype=torch.float32
1109+
)
1110+
silu_and_mul_group_quant_fwd(
1111+
intermediate_cache1.view(-1, N),
1112+
down_inputs,
1113+
down_input_scale,
1114+
down_group_size,
1115+
layout=layout,
1116+
limit=limit,
1117+
alpha=alpha,
1118+
)
1119+
else:
1120+
silu_and_mul_fwd(
1121+
intermediate_cache1.view(-1, N),
1122+
intermediate_cache2.view(-1, N // 2),
1123+
limit=limit,
1124+
alpha=alpha,
1125+
layout=layout,
1126+
)
1127+
down_inputs = intermediate_cache2.view(-1, N // 2)
1128+
down_input_scale = a2_scale
10921129

10931130
grouped_matmul(
10941131
curr_topk_ids.numel(),
1095-
intermediate_cache2.view(-1, N // 2),
1096-
a2_scale,
1132+
down_inputs,
1133+
down_input_scale,
10971134
expert_to_token_num,
10981135
expert_to_tokens,
10991136
expert_to_weights=expert_to_weights,

0 commit comments

Comments
 (0)