kaizen9 commited on
Commit
7ced247
·
verified ·
1 Parent(s): ab887ff

Upload folder using huggingface_hub

Browse files
.gitattributes CHANGED
@@ -33,3 +33,4 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
33
  *.zip filter=lfs diff=lfs merge=lfs -text
34
  *.zst filter=lfs diff=lfs merge=lfs -text
35
  *tfevents* filter=lfs diff=lfs merge=lfs -text
 
 
33
  *.zip filter=lfs diff=lfs merge=lfs -text
34
  *.zst filter=lfs diff=lfs merge=lfs -text
35
  *tfevents* filter=lfs diff=lfs merge=lfs -text
36
+ tokenizer.json filter=lfs diff=lfs merge=lfs -text
added_tokens.json ADDED
@@ -0,0 +1,28 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "</think>": 151668,
3
+ "</tool_call>": 151658,
4
+ "</tool_response>": 151666,
5
+ "<think>": 151667,
6
+ "<tool_call>": 151657,
7
+ "<tool_response>": 151665,
8
+ "<|box_end|>": 151649,
9
+ "<|box_start|>": 151648,
10
+ "<|endoftext|>": 151643,
11
+ "<|file_sep|>": 151664,
12
+ "<|fim_middle|>": 151660,
13
+ "<|fim_pad|>": 151662,
14
+ "<|fim_prefix|>": 151659,
15
+ "<|fim_suffix|>": 151661,
16
+ "<|im_end|>": 151645,
17
+ "<|im_start|>": 151644,
18
+ "<|image_pad|>": 151655,
19
+ "<|object_ref_end|>": 151647,
20
+ "<|object_ref_start|>": 151646,
21
+ "<|quad_end|>": 151651,
22
+ "<|quad_start|>": 151650,
23
+ "<|repo_name|>": 151663,
24
+ "<|video_pad|>": 151656,
25
+ "<|vision_end|>": 151653,
26
+ "<|vision_pad|>": 151654,
27
+ "<|vision_start|>": 151652
28
+ }
chat_template.jinja ADDED
@@ -0,0 +1,88 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {%- if tools %}
2
+ {{- '<|im_start|>system\n' }}
3
+ {%- if messages[0].role == 'system' %}
4
+ {{- messages[0].content + '\n\n' }}
5
+ {%- endif %}
6
+ {{- '# Tools\n\nYou may call one or more functions to assist with the user query.\n\nYou are provided with function signatures within <tools></tools> XML tags:\n<tools>' }}
7
+ {%- for tool in tools %}
8
+ {{- '\n' }}
9
+ {{- tool | tojson }}
10
+ {%- endfor %}
11
+ {{- '\n</tools>\n\nFor each function call, return a json object with function name and arguments within <tool_call></tool_call> XML tags:\n<tool_call>\n{\"name\": <function-name>, \"arguments\": <args-json-object>}\n</tool_call><|im_end|>\n' }}
12
+ {%- else %}
13
+ {%- if messages[0].role == 'system' %}
14
+ {{- '<|im_start|>system\n' + messages[0].content + '<|im_end|>\n' }}
15
+ {%- endif %}
16
+ {%- endif %}
17
+
18
+
19
+ {%- for message in messages %}
20
+ {%- if message.content is string %}
21
+ {%- set content = message.content %}
22
+ {%- else %}
23
+ {%- set content = '' %}
24
+ {%- endif %}
25
+
26
+ {%- if message.role == 'user' or (message.role == 'system' and not loop.first) %}
27
+ {{- '<|im_start|>' + message.role + '\n' + content + '<|im_end|>\n' }}
28
+
29
+ {%- elif message.role == 'assistant' %}
30
+ {%- set reasoning_content = '' %}
31
+
32
+ {%- if message.reasoning_content is string %}
33
+ {%- set reasoning_content = message.reasoning_content %}
34
+ {%- elif '</think>' in content %}
35
+ {%- set reasoning_content = content.split('</think>')[0].rstrip('\n').split('<think>')[-1].lstrip('\n') %}
36
+ {%- set content = content.split('</think>')[-1].lstrip('\n') %}
37
+ {%- endif %}
38
+
39
+ {%- if reasoning_content %}
40
+ {{- '<|im_start|>' + message.role + '\n<think>\n' + reasoning_content.strip('\n') + '\n</think>\n\n' + content.lstrip('\n') }}
41
+ {%- else %}
42
+ {{- '<|im_start|>' + message.role + '\n' + content }}
43
+ {%- endif %}
44
+
45
+ {%- if message.tool_calls %}
46
+ {%- for tool_call in message.tool_calls %}
47
+ {%- if (loop.first and content) or not loop.first %}
48
+ {{- '\n' }}
49
+ {%- endif %}
50
+
51
+ {%- if tool_call.function %}
52
+ {%- set tool_call = tool_call.function %}
53
+ {%- endif %}
54
+
55
+ {{- '<tool_call>\n{\"name\": \"' }}
56
+ {{- tool_call.name }}
57
+ {{- '\", \"arguments\": ' }}
58
+
59
+ {%- if tool_call.arguments is string %}
60
+ {{- tool_call.arguments }}
61
+ {%- else %}
62
+ {{- tool_call.arguments | tojson }}
63
+ {%- endif %}
64
+
65
+ {{- '}\n</tool_call>' }}
66
+ {%- endfor %}
67
+ {%- endif %}
68
+
69
+ {{- '<|im_end|>\n' }}
70
+
71
+ {%- elif message.role == 'tool' %}
72
+ {%- if loop.first or messages[loop.index0 - 1].role != 'tool' %}
73
+ {{- '<|im_start|>user' }}
74
+ {%- endif %}
75
+
76
+ {{- '\n<tool_response>\n' }}
77
+ {{- content }}
78
+ {{- '\n</tool_response>' }}
79
+
80
+ {%- if loop.last or messages[loop.index0 + 1].role != 'tool' %}
81
+ {{- '<|im_end|>\n' }}
82
+ {%- endif %}
83
+ {%- endif %}
84
+ {%- endfor %}
85
+
86
+ {%- if add_generation_prompt %}
87
+ {{- '<|im_start|>assistant\n<think>\n' }}
88
+ {%- endif %}
config.json ADDED
@@ -0,0 +1,87 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "architectures": [
3
+ "MapleForCausalLM"
4
+ ],
5
+ "attention_dropout": 0.0,
6
+ "auto_map": {
7
+ "AutoConfig": "configuration_maple.MapleConfig",
8
+ "AutoModel": "modeling_maple.MapleModel",
9
+ "AutoModelForCausalLM": "modeling_maple.MapleForCausalLM"
10
+ },
11
+ "bos_token_id": 151643,
12
+ "dtype": "bfloat16",
13
+ "embedding_dropout": 0.0,
14
+ "eos_token_id": 151645,
15
+ "first_k_dense_replace": 0,
16
+ "hc": false,
17
+ "head_dim": 128,
18
+ "headwise_gate_swa": false,
19
+ "hidden_act": "silu",
20
+ "hidden_size": 2048,
21
+ "initializer_range": 0.02,
22
+ "intermediate_size": 5120,
23
+ "layer_types": [
24
+ "sliding_attention",
25
+ "sliding_attention",
26
+ "sliding_attention",
27
+ "full_attention",
28
+ "sliding_attention",
29
+ "sliding_attention",
30
+ "sliding_attention",
31
+ "full_attention",
32
+ "sliding_attention",
33
+ "sliding_attention",
34
+ "sliding_attention",
35
+ "full_attention",
36
+ "sliding_attention",
37
+ "sliding_attention",
38
+ "sliding_attention",
39
+ "full_attention",
40
+ "sliding_attention",
41
+ "sliding_attention",
42
+ "sliding_attention",
43
+ "full_attention",
44
+ "sliding_attention",
45
+ "sliding_attention",
46
+ "sliding_attention",
47
+ "full_attention"
48
+ ],
49
+ "max_position_embeddings": 140000,
50
+ "max_window_layers": 24,
51
+ "moe_intermediate_size": 512,
52
+ "moe_router_enable_expert_bias": false,
53
+ "moe_shared_expert_intermediate_size": 512,
54
+ "mtp_loss_scaling_factor": 0,
55
+ "n_group": 8,
56
+ "nope_on_global_attention": false,
57
+ "norm_topk_prob": true,
58
+ "num_attention_heads": 16,
59
+ "num_experts": 256,
60
+ "num_experts_per_tok": 8,
61
+ "num_hidden_layers": 24,
62
+ "num_key_value_heads": 4,
63
+ "num_nextn_predict_layers": 0,
64
+ "num_shared_experts": 0,
65
+ "output_dropout": 0.0,
66
+ "output_router_logits": false,
67
+ "pad_token_id": null,
68
+ "partial_rotary_factor": 0.5,
69
+ "preaffine": false,
70
+ "quantize": true,
71
+ "rms_norm_eps": 1e-06,
72
+ "rope_scaling": null,
73
+ "rope_theta": 10000,
74
+ "routed_scaling_factor": 2.5,
75
+ "router_dtype": "fp32",
76
+ "score_function": "sigmoid",
77
+ "sliding_window": 512,
78
+ "tie_word_embeddings": false,
79
+ "topk_group": 4,
80
+ "transformers_version": "4.57.1",
81
+ "use_bias": false,
82
+ "use_cache": true,
83
+ "use_qk_norm": true,
84
+ "use_qkv_bias": false,
85
+ "use_rmsnorm": true,
86
+ "vocab_size": 151936
87
+ }
configuration_maple.py ADDED
@@ -0,0 +1,84 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ """Bailing MoE V2 model configuration"""
2
+
3
+ from transformers.configuration_utils import PretrainedConfig
4
+
5
+
6
+ class MapleConfig(PretrainedConfig):
7
+
8
+ def __init__(
9
+ self,
10
+ vocab_size=157184,
11
+ hidden_size=2048,
12
+ intermediate_size=5120,
13
+ num_hidden_layers=20,
14
+ num_attention_heads=16,
15
+ num_key_value_heads=4,
16
+ hidden_act="silu",
17
+ use_qkv_bias=False, # bailing only
18
+ use_bias=False, # bailing only
19
+ rms_norm_eps=1e-06,
20
+ tie_word_embeddings=False, # PretrainedConfig key, here change default value.
21
+ embedding_dropout=0.0,
22
+ attention_dropout=0.0,
23
+ output_dropout=0.0,
24
+ initializer_range=0.02,
25
+ max_position_embeddings=32768,
26
+ rope_theta=600000.0,
27
+ use_cache=True,
28
+ max_window_layers=20,
29
+ rope_scaling=None,
30
+ pad_token_id=156892,
31
+ eos_token_id=156892,
32
+ num_experts=256,
33
+ num_shared_experts=1,
34
+ num_experts_per_tok=8,
35
+ n_group=8,
36
+ topk_group=4,
37
+ moe_intermediate_size=512,
38
+ first_k_dense_replace=1,
39
+ head_dim=128,
40
+ output_router_logits=False,
41
+ use_qk_norm=True,
42
+ num_nextn_predict_layers=0,
43
+ mtp_loss_scaling_factor=0,
44
+ moe_router_enable_expert_bias=True,
45
+ routed_scaling_factor=1.0,
46
+ **kwargs,
47
+ ):
48
+ self.num_hidden_layers = num_hidden_layers
49
+ self.vocab_size = vocab_size
50
+ self.hidden_size = hidden_size
51
+ self.intermediate_size = intermediate_size
52
+ self.num_attention_heads = num_attention_heads
53
+ self.num_key_value_heads = num_key_value_heads
54
+ self.hidden_act = hidden_act
55
+ self.use_qkv_bias = use_qkv_bias
56
+ self.use_bias = use_bias
57
+ self.rms_norm_eps = rms_norm_eps
58
+ self.embedding_dropout = embedding_dropout
59
+ self.attention_dropout = attention_dropout
60
+ self.output_dropout = output_dropout
61
+ self.num_nextn_predict_layers = num_nextn_predict_layers
62
+ self.mtp_loss_scaling_factor = mtp_loss_scaling_factor
63
+ self.initializer_range = initializer_range
64
+ self.max_position_embeddings = max_position_embeddings
65
+ self.rope_theta = rope_theta
66
+ self.use_cache = use_cache
67
+ self.max_window_layers = max_window_layers
68
+ self.head_dim = head_dim or self.hidden_size // self.num_attention_heads
69
+ self.rope_scaling = rope_scaling
70
+ self.use_qk_norm = use_qk_norm
71
+ self.moe_router_enable_expert_bias = moe_router_enable_expert_bias
72
+ self.routed_scaling_factor = routed_scaling_factor
73
+
74
+ # MoE configs
75
+ self.num_experts = num_experts
76
+ self.num_shared_experts = num_shared_experts
77
+ self.num_experts_per_tok = num_experts_per_tok
78
+ self.n_group = n_group
79
+ self.topk_group = topk_group
80
+ self.moe_intermediate_size = moe_intermediate_size
81
+ self.first_k_dense_replace = first_k_dense_replace
82
+ self.output_router_logits = output_router_logits
83
+
84
+ super().__init__(pad_token_id=pad_token_id, eos_token_id=eos_token_id, tie_word_embeddings=tie_word_embeddings, **kwargs)
fa3.py ADDED
@@ -0,0 +1,272 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ from typing import Optional, Tuple
2
+
3
+ import torch
4
+
5
+ import inspect
6
+ import os
7
+ from typing import Optional, TypedDict
8
+
9
+ import torch
10
+ import torch.nn.functional as F
11
+
12
+
13
+ try:
14
+ from flash_attn_interface import flash_attn_func, flash_attn_varlen_func
15
+ except:
16
+ from flash_attn import flash_attn_func, flash_attn_varlen_func
17
+ from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input
18
+
19
+
20
+ # Detect supported kwargs in FA3
21
+ _sig = inspect.signature(flash_attn_func)
22
+ _flash_supports_window_size = "window_size" in _sig.parameters
23
+ _flash_accepts_deterministic = "deterministic" in _sig.parameters
24
+ _flash_accepts_softcap = "softcap" in _sig.parameters
25
+
26
+
27
+ def _get_unpad_data(attention_mask: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor, int]:
28
+ seqlens_in_batch = attention_mask.sum(dim=-1, dtype=torch.int32)
29
+ indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
30
+ max_seqlen_in_batch = seqlens_in_batch.max().item()
31
+ cu_seqlens = F.pad(torch.cumsum(seqlens_in_batch, dim=0, dtype=torch.int32), (1, 0))
32
+ return indices, cu_seqlens, max_seqlen_in_batch
33
+
34
+
35
+ def _upad_input(
36
+ query_layer: torch.Tensor,
37
+ key_layer: torch.Tensor,
38
+ value_layer: torch.Tensor,
39
+ attention_mask: torch.Tensor,
40
+ query_length: int,
41
+ ):
42
+ indices_k, cu_seqlens_k, max_seqlen_in_batch_k = _get_unpad_data(attention_mask)
43
+ batch_size, kv_seq_len, num_key_value_heads, head_dim = key_layer.shape
44
+
45
+ key_layer = index_first_axis(key_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim), indices_k)
46
+ value_layer = index_first_axis(
47
+ value_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim), indices_k
48
+ )
49
+
50
+ if query_length == kv_seq_len:
51
+ query_layer = index_first_axis(query_layer.reshape(batch_size * kv_seq_len, -1, head_dim), indices_k)
52
+ cu_seqlens_q = cu_seqlens_k
53
+ max_seqlen_in_batch_q = max_seqlen_in_batch_k
54
+ indices_q = indices_k
55
+ elif query_length == 1:
56
+ max_seqlen_in_batch_q = 1
57
+ cu_seqlens_q = torch.arange(batch_size + 1, dtype=torch.int32, device=query_layer.device)
58
+ indices_q = cu_seqlens_q[:-1]
59
+ query_layer = query_layer.squeeze(1)
60
+ else:
61
+ attention_mask = attention_mask[:, -query_length:]
62
+ query_layer, indices_q, cu_seqlens_q, max_seqlen_in_batch_q, *_ = unpad_input(query_layer, attention_mask)
63
+
64
+ return (
65
+ query_layer,
66
+ key_layer,
67
+ value_layer,
68
+ indices_q,
69
+ (cu_seqlens_q, cu_seqlens_k),
70
+ (max_seqlen_in_batch_q, max_seqlen_in_batch_k),
71
+ )
72
+
73
+
74
+ def prepare_fa3_from_position_ids(query, key, value, position_ids):
75
+ query = query.view(-1, query.size(-2), query.size(-1))
76
+ key = key.contiguous().view(-1, key.size(-2), key.size(-1))
77
+ value = value.contiguous().view(-1, value.size(-2), value.size(-1))
78
+ position_ids = position_ids.flatten()
79
+ indices_q = torch.arange(position_ids.size(0), device=position_ids.device, dtype=torch.int32)
80
+
81
+ cu_seq_lens = torch.cat(
82
+ (
83
+ indices_q[position_ids == 0],
84
+ torch.tensor(position_ids.size(), device=position_ids.device, dtype=torch.int32),
85
+ )
86
+ )
87
+
88
+ max_length = position_ids.max() + 1
89
+ return query, key, value, indices_q, (cu_seq_lens, cu_seq_lens), (max_length, max_length)
90
+
91
+
92
+ def fa_peft_integration_check(
93
+ query: torch.Tensor,
94
+ key: torch.Tensor,
95
+ value: torch.Tensor,
96
+ target_dtype: Optional[torch.dtype] = None,
97
+ ):
98
+ if target_dtype is None:
99
+ return query, key, value
100
+ if query.dtype == torch.float32:
101
+ query = query.to(target_dtype)
102
+ key = key.to(target_dtype)
103
+ value = value.to(target_dtype)
104
+ return query, key, value
105
+
106
+
107
+ deterministic_g = os.environ.get("FLASH_ATTENTION_DETERMINISTIC", "0") == "1"
108
+
109
+
110
+ def _flash_attention_forward(
111
+ query_states: torch.Tensor,
112
+ key_states: torch.Tensor,
113
+ value_states: torch.Tensor,
114
+ attention_mask: Optional[torch.Tensor],
115
+ query_length: int,
116
+ is_causal: bool,
117
+ dropout: float = 0.0,
118
+ position_ids: Optional[torch.Tensor] = None,
119
+ softmax_scale: Optional[float] = None,
120
+ sliding_window: Optional[int] = None,
121
+ use_top_left_mask: bool = False,
122
+ softcap: Optional[float] = None,
123
+ deterministic: Optional[bool] = None,
124
+ cu_seq_lens_q: Optional[torch.LongTensor] = None,
125
+ cu_seq_lens_k: Optional[torch.LongTensor] = None,
126
+ max_length_q: Optional[int] = None,
127
+ max_length_k: Optional[int] = None,
128
+ target_dtype: Optional[torch.dtype] = None,
129
+ **kwargs,
130
+ ):
131
+ causal = is_causal if not use_top_left_mask else (is_causal and query_length != 1)
132
+
133
+ flash_kwargs = {}
134
+ if _flash_supports_window_size and sliding_window is not None and key_states.shape[1] > sliding_window:
135
+ flash_kwargs["window_size"] = (sliding_window, 0)
136
+ if _flash_accepts_deterministic:
137
+ if deterministic is None:
138
+ deterministic = deterministic_g
139
+ flash_kwargs["deterministic"] = deterministic
140
+
141
+ if attention_mask is not None:
142
+ batch_size = query_states.shape[0]
143
+ q_unpad, k_unpad, v_unpad, indices_q, (cu_seqlens_q, cu_seqlens_k), (max_q, max_k) = _upad_input(
144
+ query_states, key_states, value_states, attention_mask, query_length
145
+ )
146
+ attn_output_unpad = flash_attn_varlen_func(
147
+ q_unpad,
148
+ k_unpad,
149
+ v_unpad,
150
+ cu_seqlens_q=cu_seqlens_q,
151
+ cu_seqlens_k=cu_seqlens_k,
152
+ max_seqlen_q=max_q,
153
+ max_seqlen_k=max_k,
154
+ # dropout_p=dropout,
155
+ softmax_scale=softmax_scale,
156
+ causal=causal,
157
+ **flash_kwargs,
158
+ )
159
+ attn_output = pad_input(attn_output_unpad, indices_q, batch_size, query_length)
160
+
161
+ elif position_ids is not None and (
162
+ max_length_q is not None
163
+ # This fucks up compile
164
+ # or (query_length != 1 and not (torch.diff(position_ids, dim=-1) >= 0).all())
165
+ ):
166
+ batch_size = query_states.size(0)
167
+ if cu_seq_lens_q is None or cu_seq_lens_k is None:
168
+ q_unpad, k_unpad, v_unpad, indices_q, (cu_seq_lens_q, cu_seq_lens_k), (max_length_q, max_length_k) = (
169
+ prepare_fa3_from_position_ids(query_states, key_states, value_states, position_ids)
170
+ )
171
+ else:
172
+ q_unpad = query_states.reshape(-1, query_states.size(-2), query_states.size(-1))
173
+ k_unpad = key_states.reshape(-1, key_states.size(-2), key_states.size(-1))
174
+ v_unpad = value_states.reshape(-1, value_states.size(-2), value_states.size(-1))
175
+
176
+ attn_output = flash_attn_varlen_func(
177
+ q_unpad,
178
+ k_unpad,
179
+ v_unpad,
180
+ cu_seqlens_q=cu_seq_lens_q,
181
+ cu_seqlens_k=cu_seq_lens_k,
182
+ max_seqlen_q=max_length_q,
183
+ max_seqlen_k=max_length_k,
184
+ # dropout_p=dropout,
185
+ softmax_scale=softmax_scale,
186
+ causal=causal,
187
+ **flash_kwargs,
188
+ )
189
+ attn_output = attn_output.view(batch_size, -1, attn_output.size(-2), attn_output.size(-1))
190
+
191
+ else:
192
+ # print(f"scale {softmax_scale}")
193
+ attn_output = flash_attn_func(
194
+ query_states,
195
+ key_states,
196
+ value_states,
197
+ # dropout,
198
+ softmax_scale=softmax_scale,
199
+ causal=causal,
200
+ **flash_kwargs,
201
+ )
202
+
203
+ return attn_output
204
+
205
+
206
+ class FlashAttentionKwargs(TypedDict, total=False):
207
+ cu_seq_lens_q: Optional[torch.LongTensor]
208
+ cu_seq_lens_k: Optional[torch.LongTensor]
209
+ max_length_q: Optional[int]
210
+ max_length_k: Optional[int]
211
+
212
+
213
+ # _use_top_left_mask = flash_attn_supports_top_left_mask()
214
+
215
+ _use_top_left_mask = False
216
+
217
+
218
+ def flash_attention_forward(
219
+ module: torch.nn.Module,
220
+ query: torch.Tensor,
221
+ key: torch.Tensor,
222
+ value: torch.Tensor,
223
+ attention_mask: Optional[torch.Tensor],
224
+ dropout: float = 0.0,
225
+ scaling: Optional[float] = None,
226
+ sliding_window: Optional[int] = None,
227
+ softcap: Optional[float] = None,
228
+ **kwargs,
229
+ ) -> Tuple[torch.Tensor, None]:
230
+ # This is before the transpose
231
+ seq_len = query.shape[1]
232
+
233
+ # FA2 uses non-transposed inputs
234
+ query = query.transpose(1, 2)
235
+ key = key.transpose(1, 2)
236
+ value = value.transpose(1, 2)
237
+
238
+ # In PEFT, usually we cast the layer norms in float32 for training stability reasons
239
+ # therefore the input hidden states gets silently casted in float32. Hence, we need
240
+ # cast them back in the correct dtype just to be sure everything works as expected.
241
+ # This might slowdown training & inference so it is recommended to not cast the LayerNorms
242
+ # in fp32. (usually our RMSNorm modules handle it correctly)
243
+ target_dtype = None
244
+ if query.dtype == torch.float32:
245
+ if torch.is_autocast_enabled():
246
+ target_dtype = torch.get_autocast_gpu_dtype()
247
+ # Handle the case where the model is quantized
248
+ elif hasattr(module.config, "_pre_quantization_dtype"):
249
+ target_dtype = module.config._pre_quantization_dtype
250
+ else:
251
+ target_dtype = next(layer for layer in module.modules() if isinstance(layer, torch.nn.Linear)).weight.dtype
252
+
253
+ # FA2 always relies on the value set in the module, so remove it if present in kwargs to avoid passing it twice
254
+ kwargs.pop("is_causal", None)
255
+
256
+ attn_output = _flash_attention_forward(
257
+ query,
258
+ key,
259
+ value,
260
+ attention_mask,
261
+ query_length=seq_len,
262
+ is_causal=module.is_causal,
263
+ dropout=dropout,
264
+ softmax_scale=scaling,
265
+ sliding_window=sliding_window,
266
+ softcap=softcap,
267
+ use_top_left_mask=_use_top_left_mask,
268
+ target_dtype=target_dtype,
269
+ **kwargs,
270
+ )
271
+
272
+ return attn_output, None
fa3_utils.py ADDED
@@ -0,0 +1,206 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ import inspect
2
+ import os
3
+ from typing import Optional, TypedDict
4
+
5
+ import torch
6
+ import torch.nn.functional as F
7
+
8
+
9
+ try:
10
+ from flash_attn_interface import flash_attn_func, flash_attn_varlen_func
11
+ except:
12
+ from flash_attn import flash_attn_func, flash_attn_varlen_func
13
+ from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input
14
+
15
+
16
+ # Detect supported kwargs in FA3
17
+ _sig = inspect.signature(flash_attn_func)
18
+ _flash_supports_window_size = "window_size" in _sig.parameters
19
+ _flash_accepts_deterministic = "deterministic" in _sig.parameters
20
+ _flash_accepts_softcap = "softcap" in _sig.parameters
21
+
22
+
23
+ def _get_unpad_data(attention_mask: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor, int]:
24
+ seqlens_in_batch = attention_mask.sum(dim=-1, dtype=torch.int32)
25
+ indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
26
+ max_seqlen_in_batch = seqlens_in_batch.max().item()
27
+ cu_seqlens = F.pad(torch.cumsum(seqlens_in_batch, dim=0, dtype=torch.int32), (1, 0))
28
+ return indices, cu_seqlens, max_seqlen_in_batch
29
+
30
+
31
+ def _upad_input(
32
+ query_layer: torch.Tensor,
33
+ key_layer: torch.Tensor,
34
+ value_layer: torch.Tensor,
35
+ attention_mask: torch.Tensor,
36
+ query_length: int,
37
+ ):
38
+ indices_k, cu_seqlens_k, max_seqlen_in_batch_k = _get_unpad_data(attention_mask)
39
+ batch_size, kv_seq_len, num_key_value_heads, head_dim = key_layer.shape
40
+
41
+ key_layer = index_first_axis(key_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim), indices_k)
42
+ value_layer = index_first_axis(
43
+ value_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim), indices_k
44
+ )
45
+
46
+ if query_length == kv_seq_len:
47
+ query_layer = index_first_axis(query_layer.reshape(batch_size * kv_seq_len, -1, head_dim), indices_k)
48
+ cu_seqlens_q = cu_seqlens_k
49
+ max_seqlen_in_batch_q = max_seqlen_in_batch_k
50
+ indices_q = indices_k
51
+ elif query_length == 1:
52
+ max_seqlen_in_batch_q = 1
53
+ cu_seqlens_q = torch.arange(batch_size + 1, dtype=torch.int32, device=query_layer.device)
54
+ indices_q = cu_seqlens_q[:-1]
55
+ query_layer = query_layer.squeeze(1)
56
+ else:
57
+ attention_mask = attention_mask[:, -query_length:]
58
+ query_layer, indices_q, cu_seqlens_q, max_seqlen_in_batch_q, *_ = unpad_input(query_layer, attention_mask)
59
+
60
+ return (
61
+ query_layer,
62
+ key_layer,
63
+ value_layer,
64
+ indices_q,
65
+ (cu_seqlens_q, cu_seqlens_k),
66
+ (max_seqlen_in_batch_q, max_seqlen_in_batch_k),
67
+ )
68
+
69
+
70
+ def prepare_fa3_from_position_ids(query, key, value, position_ids):
71
+ query = query.view(-1, query.size(-2), query.size(-1))
72
+ key = key.contiguous().view(-1, key.size(-2), key.size(-1))
73
+ value = value.contiguous().view(-1, value.size(-2), value.size(-1))
74
+ position_ids = position_ids.flatten()
75
+ indices_q = torch.arange(position_ids.size(0), device=position_ids.device, dtype=torch.int32)
76
+
77
+ cu_seq_lens = torch.cat(
78
+ (
79
+ indices_q[position_ids == 0],
80
+ torch.tensor(position_ids.size(), device=position_ids.device, dtype=torch.int32),
81
+ )
82
+ )
83
+
84
+ max_length = position_ids.max() + 1
85
+ return query, key, value, indices_q, (cu_seq_lens, cu_seq_lens), (max_length, max_length)
86
+
87
+
88
+ def fa_peft_integration_check(
89
+ query: torch.Tensor,
90
+ key: torch.Tensor,
91
+ value: torch.Tensor,
92
+ target_dtype: Optional[torch.dtype] = None,
93
+ ):
94
+ if target_dtype is None:
95
+ return query, key, value
96
+ if query.dtype == torch.float32:
97
+ query = query.to(target_dtype)
98
+ key = key.to(target_dtype)
99
+ value = value.to(target_dtype)
100
+ return query, key, value
101
+
102
+
103
+ deterministic_g = os.environ.get("FLASH_ATTENTION_DETERMINISTIC", "0") == "1"
104
+
105
+
106
+ def _flash_attention_forward(
107
+ query_states: torch.Tensor,
108
+ key_states: torch.Tensor,
109
+ value_states: torch.Tensor,
110
+ attention_mask: Optional[torch.Tensor],
111
+ query_length: int,
112
+ is_causal: bool,
113
+ dropout: float = 0.0,
114
+ position_ids: Optional[torch.Tensor] = None,
115
+ softmax_scale: Optional[float] = None,
116
+ sliding_window: Optional[int] = None,
117
+ use_top_left_mask: bool = False,
118
+ softcap: Optional[float] = None,
119
+ deterministic: Optional[bool] = None,
120
+ cu_seq_lens_q: Optional[torch.LongTensor] = None,
121
+ cu_seq_lens_k: Optional[torch.LongTensor] = None,
122
+ max_length_q: Optional[int] = None,
123
+ max_length_k: Optional[int] = None,
124
+ target_dtype: Optional[torch.dtype] = None,
125
+ **kwargs,
126
+ ):
127
+ causal = is_causal if not use_top_left_mask else (is_causal and query_length != 1)
128
+
129
+ flash_kwargs = {}
130
+ if _flash_supports_window_size and sliding_window is not None and key_states.shape[1] > sliding_window:
131
+ flash_kwargs["window_size"] = (sliding_window, 0)
132
+ if _flash_accepts_deterministic:
133
+ if deterministic is None:
134
+ deterministic = deterministic_g
135
+ flash_kwargs["deterministic"] = deterministic
136
+
137
+ if attention_mask is not None:
138
+ batch_size = query_states.shape[0]
139
+ q_unpad, k_unpad, v_unpad, indices_q, (cu_seqlens_q, cu_seqlens_k), (max_q, max_k) = _upad_input(
140
+ query_states, key_states, value_states, attention_mask, query_length
141
+ )
142
+ attn_output_unpad = flash_attn_varlen_func(
143
+ q_unpad,
144
+ k_unpad,
145
+ v_unpad,
146
+ cu_seqlens_q=cu_seqlens_q,
147
+ cu_seqlens_k=cu_seqlens_k,
148
+ max_seqlen_q=max_q,
149
+ max_seqlen_k=max_k,
150
+ # dropout_p=dropout,
151
+ softmax_scale=softmax_scale,
152
+ causal=causal,
153
+ **flash_kwargs,
154
+ )
155
+ attn_output = pad_input(attn_output_unpad, indices_q, batch_size, query_length)
156
+
157
+ elif position_ids is not None and (
158
+ max_length_q is not None
159
+ # This fucks up compile
160
+ # or (query_length != 1 and not (torch.diff(position_ids, dim=-1) >= 0).all())
161
+ ):
162
+ batch_size = query_states.size(0)
163
+ if cu_seq_lens_q is None or cu_seq_lens_k is None:
164
+ q_unpad, k_unpad, v_unpad, indices_q, (cu_seq_lens_q, cu_seq_lens_k), (max_length_q, max_length_k) = (
165
+ prepare_fa3_from_position_ids(query_states, key_states, value_states, position_ids)
166
+ )
167
+ else:
168
+ q_unpad = query_states.reshape(-1, query_states.size(-2), query_states.size(-1))
169
+ k_unpad = key_states.reshape(-1, key_states.size(-2), key_states.size(-1))
170
+ v_unpad = value_states.reshape(-1, value_states.size(-2), value_states.size(-1))
171
+
172
+ attn_output = flash_attn_varlen_func(
173
+ q_unpad,
174
+ k_unpad,
175
+ v_unpad,
176
+ cu_seqlens_q=cu_seq_lens_q,
177
+ cu_seqlens_k=cu_seq_lens_k,
178
+ max_seqlen_q=max_length_q,
179
+ max_seqlen_k=max_length_k,
180
+ # dropout_p=dropout,
181
+ softmax_scale=softmax_scale,
182
+ causal=causal,
183
+ **flash_kwargs,
184
+ )
185
+ attn_output = attn_output.view(batch_size, -1, attn_output.size(-2), attn_output.size(-1))
186
+
187
+ else:
188
+ # print(f"scale {softmax_scale}")
189
+ attn_output = flash_attn_func(
190
+ query_states,
191
+ key_states,
192
+ value_states,
193
+ # dropout,
194
+ softmax_scale=softmax_scale,
195
+ causal=causal,
196
+ **flash_kwargs,
197
+ )
198
+
199
+ return attn_output
200
+
201
+
202
+ class FlashAttentionKwargs(TypedDict, total=False):
203
+ cu_seq_lens_q: Optional[torch.LongTensor]
204
+ cu_seq_lens_k: Optional[torch.LongTensor]
205
+ max_length_q: Optional[int]
206
+ max_length_k: Optional[int]
merges.txt ADDED
The diff for this file is too large to render. See raw diff
 
model-00001-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:1b2b2e6af85b11b894da849131bbe397ff0718a897bde75550725baede65aedb
3
+ size 4998990992
model-00002-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:b3d0a2d114b014efc61cba2a93616af4387c62a076519709616e3d5bd433ffde
3
+ size 4999904616
model-00003-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:5a7f06f47b52dc29b3decc192d417ce94fe528f581f05420facb5f64e6066ab0
3
+ size 4999904520
model-00004-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:aa658405f9bc7bb3f0a3b9fc50f43ba18a153681a76fcdfbf75b53a57f164134
3
+ size 4999905688
model-00005-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:d72fe5d8e5109c49dd49302f5aeb28c6f0cb187c7510ebb501438d6e871934e3
3
+ size 4999906952
model-00006-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:843bda728aba987e0b8b4ee67752700e2361c7c880f9826d947b86844ae824ec
3
+ size 4999906952
model-00007-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:e2b89a3474db730b334ede2a6b780ebdecb9c3a2a8d5a4dc88246a120d31a4f2
3
+ size 4999906952
model-00008-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:7e25757b077f840c267f2c084db1254d148d3c1c623a6bcf7a58f151ed46bd13
3
+ size 4999907000
model-00009-of-00009.safetensors ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:bb4469e9ee94df71acbe2d1fef70671f0bc762602dc10a7b2b76ae6329afa2a6
3
+ size 432038704
model.safetensors.index.json ADDED
The diff for this file is too large to render. See raw diff
 
modeling_maple.py ADDED
@@ -0,0 +1,932 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ # coding=utf-8
2
+ import math
3
+ from dataclasses import dataclass
4
+ from typing import Optional, Tuple, Union
5
+
6
+ import torch
7
+ import torch.nn.functional as F
8
+ import triton
9
+ import triton.language as tl
10
+ from torch import nn
11
+ from torch.library import triton_op, wrap_triton
12
+ from transformers.activations import ACT2FN
13
+ from transformers.cache_utils import Cache, DynamicCache
14
+ from transformers.generation.utils import GenerationMixin
15
+ from transformers.modeling_outputs import MoeModelOutputWithPast
16
+ from transformers.modeling_rope_utils import ROPE_INIT_FUNCTIONS, dynamic_rope_update
17
+ from transformers.modeling_utils import ALL_ATTENTION_FUNCTIONS, PreTrainedModel
18
+ from transformers.pytorch_utils import is_torch_greater_or_equal_than_1_13
19
+ from transformers.utils import (
20
+ ModelOutput,
21
+ add_start_docstrings,
22
+ add_start_docstrings_to_model_forward,
23
+ )
24
+ from transformers.utils import logging as hf_logging
25
+ from transformers.utils.import_utils import is_torch_fx_available
26
+
27
+ from .configuration_maple import MapleConfig
28
+ from .fa3 import flash_attention_forward
29
+
30
+ logger = hf_logging.get_logger(__name__)
31
+ _CONFIG_FOR_DOC = "MapleConfig"
32
+
33
+
34
+ # This makes `_prepare_4d_causal_attention_mask` a leaf function in the FX graph.
35
+ if is_torch_fx_available():
36
+ if not is_torch_greater_or_equal_than_1_13:
37
+ import torch.fx # noqa: F401
38
+
39
+ @dataclass
40
+ class MapleOutputWithPast(ModelOutput):
41
+ loss: Optional[torch.FloatTensor] = None
42
+ logits: Optional[torch.FloatTensor] = None
43
+ past_key_values: Optional[Cache] = None
44
+ hidden_states: Optional[tuple[torch.FloatTensor, ...]] = None
45
+ attentions: Optional[tuple[torch.FloatTensor, ...]] = None
46
+ z_loss: Optional[torch.FloatTensor] = None
47
+ aux_loss: Optional[torch.FloatTensor] = None
48
+ router_logits: Optional[tuple[torch.FloatTensor]] = None
49
+
50
+
51
+ class MoeV2ModelOutputWithPast(MoeModelOutputWithPast):
52
+ def __init__(self, aux_loss=0.0, **kwargs):
53
+ super().__init__(**kwargs)
54
+ self.aux_loss = aux_loss
55
+
56
+
57
+ def rmsnorm_ref(x: torch.Tensor, weight: torch.Tensor, eps: float):
58
+ # Match Linghe/Liger-llama RMSNorm: fp32 RMS reduction, then cast the
59
+ # normalized activation back to input dtype before the weight multiply.
60
+ dtype = x.dtype
61
+ x = x.float()
62
+ rstd = torch.rsqrt(x.pow(2).mean(dim=-1, keepdim=True) + eps)
63
+ x = (x * rstd).to(dtype)
64
+ return (x * weight).to(dtype)
65
+
66
+
67
+ def rotate_half_kernel_style(x: torch.Tensor):
68
+ """
69
+ Match the Triton kernel's rotate logic.
70
+
71
+ Kernel operates on only the first half of the full head dim.
72
+ For that half, it views as [2, d] and maps:
73
+ [a, b] -> [-b, a]
74
+ """
75
+ d2 = x.shape[-1]
76
+ assert d2 % 2 == 0
77
+ d = d2 // 2
78
+ x0 = x[..., :d]
79
+ x1 = x[..., d:]
80
+ return torch.cat((-x1, x0), dim=-1)
81
+
82
+
83
+ def half_rope_ref(x: torch.Tensor, freqs: torch.Tensor):
84
+ """
85
+ x: [B, L, H, D_full]
86
+ freqs: [L, D_full // 2]
87
+
88
+ Kernel applies RoPE to x[..., :D_full//2] and copies
89
+ x[..., D_full//2:] unchanged.
90
+ """
91
+ full_d = x.shape[-1]
92
+ half_d = full_d // 2
93
+ x0 = x[..., :half_d]
94
+ x1 = x[..., half_d:]
95
+
96
+ cos = freqs.cos()[None, :, None, :].to(torch.float32)
97
+ sin = freqs.sin()[None, :, None, :].to(torch.float32)
98
+
99
+ y0 = x0.float() * cos + rotate_half_kernel_style(x0.float()) * sin
100
+ y = torch.cat([y0.to(x.dtype), x1], dim=-1)
101
+ return y
102
+
103
+
104
+ def split_qkv_reference(
105
+ qkv: torch.Tensor,
106
+ q_norm_weight: torch.Tensor,
107
+ k_norm_weight: torch.Tensor,
108
+ freqs: torch.Tensor,
109
+ H: int,
110
+ h: int,
111
+ eps: float,
112
+ interleaved: bool,
113
+ silu: bool,
114
+ rope: bool,
115
+ ):
116
+ """
117
+ Reference matching the kernels' output layout:
118
+ qo: [B, L, H, D]
119
+ ko: [B, L, h, D]
120
+ vo: [B, L, h, D]
121
+
122
+ Input qkv is expected as [B, L, (H + 2h) * D] for non-transposed.
123
+ This function handles interleaved and non-interleaved packed QKV.
124
+ """
125
+ B, L, dim = qkv.shape
126
+ dtype = qkv.dtype
127
+ D = q_norm_weight.numel()
128
+ assert dim == (H + 2 * h) * D, (dim, H, h, D)
129
+
130
+ if not interleaved:
131
+ qkv_view = qkv.view(B, L, H + 2 * h, D)
132
+ q = qkv_view[:, :, :H, :]
133
+ k = qkv_view[:, :, H:H + h, :]
134
+ v = qkv_view[:, :, H + h:H + 2 * h, :]
135
+ else:
136
+ w = H // h
137
+ grouped = qkv.view(B, L, h, w + 2, D)
138
+ q = grouped[:, :, :, :w, :].reshape(B, L, H, D)
139
+ k = grouped[:, :, :, w, :]
140
+ v = grouped[:, :, :, w + 1, :]
141
+
142
+ if silu:
143
+ q = F.silu(q)
144
+ k = F.silu(k)
145
+ v = F.silu(v)
146
+
147
+ q = rmsnorm_ref(q, q_norm_weight, eps)
148
+ k = rmsnorm_ref(k, k_norm_weight, eps)
149
+ if rope:
150
+ q = half_rope_ref(q, freqs).to(dtype)
151
+ k = half_rope_ref(k, freqs).to(dtype)
152
+
153
+ return q, k, v
154
+
155
+ class MapleRotaryEmbedding(nn.Module):
156
+ def __init__(self, config: MapleConfig, device=None):
157
+ super().__init__()
158
+ if hasattr(config, "rope_scaling") and config.rope_scaling is not None:
159
+ self.rope_type = config.rope_scaling.get("rope_type", config.rope_scaling.get("type"))
160
+ else:
161
+ self.rope_type = "default"
162
+ self.max_seq_len_cached = config.max_position_embeddings
163
+ self.original_max_seq_len = config.max_position_embeddings
164
+
165
+ self.config = config
166
+ self.rope_init_fn = ROPE_INIT_FUNCTIONS[self.rope_type]
167
+
168
+ inv_freq, self.attention_scaling = self.rope_init_fn(self.config, device)
169
+ self.register_buffer("inv_freq", inv_freq, persistent=False)
170
+ self.original_inv_freq = self.inv_freq
171
+
172
+ @torch.no_grad()
173
+ @dynamic_rope_update
174
+ def forward(self, x, position_ids):
175
+ inv_freq_expanded = self.inv_freq[None, :, None].float().expand(position_ids.shape[0], -1, 1).to(x.device)
176
+ position_ids_expanded = position_ids[:, None, :].float()
177
+
178
+ device_type = x.device.type if isinstance(x.device.type, str) and x.device.type != "mps" else "cpu"
179
+ with torch.autocast(device_type=device_type, enabled=False):
180
+ freqs = (inv_freq_expanded.float() @ position_ids_expanded.float()).transpose(1, 2)
181
+ emb = torch.cat((freqs, freqs), dim=-1)
182
+ cos = emb.cos() * self.attention_scaling
183
+ sin = emb.sin() * self.attention_scaling
184
+ freqs = torch.cat([freqs, freqs], dim=-1)
185
+ return cos.to(dtype=x.dtype), sin.to(dtype=x.dtype), freqs.float()
186
+
187
+
188
+ def rotate_half(x):
189
+ x1 = x[..., : x.shape[-1] // 2]
190
+ x2 = x[..., x.shape[-1] // 2 :]
191
+ return torch.cat((-x2, x1), dim=-1)
192
+
193
+
194
+ def apply_rotary_pos_emb(q, k, cos, sin, unsqueeze_dim=1):
195
+ cos = cos.unsqueeze(unsqueeze_dim)
196
+ sin = sin.unsqueeze(unsqueeze_dim)
197
+
198
+ rotary_dim = cos.shape[-1]
199
+ q_rot, q_pass = q[..., :rotary_dim], q[..., rotary_dim:]
200
+ k_rot, k_pass = k[..., :rotary_dim], k[..., rotary_dim:]
201
+
202
+ q_embed = (q_rot * cos) + (rotate_half(q_rot) * sin)
203
+ k_embed = (k_rot * cos) + (rotate_half(k_rot) * sin)
204
+
205
+ q_embed = torch.cat([q_embed, q_pass], dim=-1)
206
+ k_embed = torch.cat([k_embed, k_pass], dim=-1)
207
+ return q_embed, k_embed
208
+
209
+
210
+ class MapleMLP(nn.Module):
211
+ def __init__(self, config: MapleConfig, intermediate_size: int):
212
+ super().__init__()
213
+ self.config = config
214
+ self.hidden_size = config.hidden_size
215
+ self.intermediate_size = intermediate_size
216
+
217
+ self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
218
+ self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
219
+ self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
220
+ self.act_fn = ACT2FN[config.hidden_act]
221
+
222
+ def forward(self, x, quantize_inplace_now: bool = False):
223
+ gate_weight, up_weight, down_weight = self.gate_proj.weight, self.up_proj.weight, self.down_proj.weight
224
+ return torch.nn.functional.linear(
225
+ self.act_fn(torch.clamp(torch.nn.functional.linear(x, gate_weight), max=7.0)) * torch.clamp(torch.nn.functional.linear(x, up_weight), min=-7.0, max=7.0),
226
+ down_weight,
227
+ )
228
+
229
+
230
+ class MapleRMSNorm(nn.Module):
231
+ def __init__(self, hidden_size, eps=1e-6):
232
+ super().__init__()
233
+ self.weight = nn.Parameter(torch.ones(hidden_size))
234
+ self.variance_epsilon = eps
235
+
236
+ def forward(self, hidden_states):
237
+ input_dtype = hidden_states.dtype
238
+ hidden_states = hidden_states.to(torch.float32)
239
+ variance = hidden_states.pow(2).mean(-1, keepdim=True)
240
+ hidden_states = hidden_states * torch.rsqrt(variance + self.variance_epsilon)
241
+ return self.weight * hidden_states.to(input_dtype)
242
+
243
+
244
+ try:
245
+ from liger_kernel.transformers.rms_norm import LigerRMSNorm
246
+
247
+ MapleRMSNorm = LigerRMSNorm
248
+ except:
249
+ print("no liger kernel")
250
+
251
+
252
+ class MapleGate(nn.Module):
253
+
254
+ def __init__(self, config: MapleConfig):
255
+ super().__init__()
256
+ self.config = config
257
+ self.top_k = config.num_experts_per_tok
258
+ self.num_experts = config.num_experts
259
+ self.gating_dim = config.hidden_size
260
+ self.weight = nn.Parameter(torch.empty((self.num_experts, self.gating_dim)))
261
+ self.reset_parameters()
262
+
263
+ def reset_parameters(self) -> None:
264
+ import torch.nn.init as init
265
+
266
+ init.kaiming_uniform_(self.weight, a=math.sqrt(5))
267
+
268
+ def forward(self, hidden_states: torch.Tensor):
269
+ hidden_states = hidden_states.view(-1, hidden_states.shape[-1])
270
+ logits = F.linear(hidden_states.type(torch.float32), self.weight.type(torch.float32))
271
+ routing_weights = F.softmax(logits, dim=1, dtype=torch.float)
272
+ scores, topk_idx = torch.topk(routing_weights, self.top_k, dim=-1)
273
+ scores = scores.type_as(logits)
274
+ topk_weight = scores / (scores.sum(dim=-1, keepdim=True) + 1e-20)
275
+
276
+ return topk_idx, topk_weight, logits
277
+
278
+
279
+ class MapleSparseMoeBlock(nn.Module):
280
+ """
281
+ Unfused MoE block matching HF layout (ModuleList experts).
282
+ """
283
+
284
+ def __init__(self, config) -> None:
285
+ super().__init__()
286
+ self.config = config
287
+ self.num_experts_per_tok = config.num_experts_per_tok
288
+ self._setup_experts()
289
+ self.gate = MapleGate(config)
290
+
291
+ def _setup_experts(self):
292
+ self.experts = nn.ModuleList(
293
+ [
294
+ MapleMLP(
295
+ config=self.config,
296
+ intermediate_size=self.config.moe_intermediate_size,
297
+ )
298
+ for _ in range(self.config.num_experts)
299
+ ]
300
+ )
301
+
302
+ def forward(
303
+ self,
304
+ hidden_states: torch.Tensor,
305
+ ) -> tuple[torch.Tensor, tuple[torch.Tensor, torch.Tensor]]:
306
+ original_shape = hidden_states.shape
307
+ identity = hidden_states
308
+
309
+ bsz, seq_len, h = hidden_states.shape
310
+ topk_idx, topk_weight, router_logits = self.gate(hidden_states)
311
+ hidden_states = hidden_states.view(-1, hidden_states.shape[-1])
312
+ flat_topk_idx = topk_idx.view(-1)
313
+
314
+ if self.training:
315
+ hidden_states = hidden_states.repeat_interleave(self.num_experts_per_tok, dim=0)
316
+ y = torch.empty_like(hidden_states)
317
+ for i, expert in enumerate(self.experts):
318
+ y[flat_topk_idx == i] = expert(hidden_states[flat_topk_idx == i])
319
+ y = (y.view(*topk_weight.shape, -1) * topk_weight.unsqueeze(-1)).sum(dim=1)
320
+ y = y.to(hidden_states.dtype).view(bsz, seq_len, h)
321
+ else:
322
+ y = self.moe_infer(hidden_states, topk_idx, topk_weight).view(bsz, seq_len, h)
323
+
324
+ return y
325
+
326
+ @torch.no_grad()
327
+ def moe_infer(self, x, topk_ids, topk_weight):
328
+ cnts = topk_ids.new_zeros((topk_ids.shape[0], len(self.experts)))
329
+ cnts.scatter_(1, topk_ids, 1)
330
+ tokens_per_expert = cnts.sum(dim=0)
331
+ idxs = topk_ids.view(-1).argsort()
332
+ sorted_tokens = x[idxs // topk_ids.shape[1]]
333
+ tokens_per_expert = tokens_per_expert.cpu().numpy()
334
+ outputs = []
335
+ start_idx = 0
336
+ for i, num_tokens in enumerate(tokens_per_expert):
337
+ end_idx = start_idx + num_tokens
338
+ if num_tokens == 0:
339
+ continue
340
+ expert = self.experts[i]
341
+ tokens_for_this_expert = sorted_tokens[start_idx:end_idx]
342
+ expert_out = expert(tokens_for_this_expert)
343
+ outputs.append(expert_out.to(x.device))
344
+ start_idx = end_idx
345
+
346
+ outs = torch.cat(outputs, dim=0) if len(outputs) else sorted_tokens.new_empty(0)
347
+ new_x = torch.empty_like(outs)
348
+ new_x[idxs] = outs
349
+ final_out = (
350
+ new_x.view(*topk_ids.shape, -1)
351
+ .type(topk_weight.dtype)
352
+ .mul_(topk_weight.unsqueeze(dim=-1))
353
+ .sum(dim=1)
354
+ .type(new_x.dtype)
355
+ )
356
+ return final_out
357
+
358
+
359
+ class MapleAttention(nn.Module):
360
+ """Fixed wiring for modern HF attention APIs: uses prepared causal_mask + cache_position + Cache.update()."""
361
+
362
+ def __init__(self, config: MapleConfig, layer_idx: Optional[int] = None):
363
+ super().__init__()
364
+ self.config = config
365
+ self.layer_idx = layer_idx
366
+ if layer_idx is None:
367
+ logger.warning_once(
368
+ f"Instantiating {self.__class__.__name__} without passing `layer_idx` is not recommended and will "
369
+ "lead to errors during the forward call if caching is used. Please pass `layer_idx`."
370
+ )
371
+
372
+ self.attention_dropout = config.attention_dropout
373
+ self.hidden_size = config.hidden_size
374
+ self.num_heads = config.num_attention_heads
375
+ self.head_dim = config.head_dim or self.hidden_size // self.num_heads
376
+ self.scaling = self.head_dim**-0.5
377
+
378
+ partial_rotary_factor = 0.5 # getattr(config, "partial_rotary_factor", 0.5)
379
+ self.rope_dim = int(self.head_dim * partial_rotary_factor)
380
+
381
+ self.num_key_value_heads = config.num_key_value_heads
382
+ self.num_key_value_groups = self.num_heads // self.num_key_value_heads
383
+ self.is_causal = True
384
+
385
+ self.sliding_window = None
386
+ self.layer_type = config.layer_types[layer_idx] if hasattr(config, "layer_types") else None
387
+ self.sliding_window = config.sliding_window if self.layer_type == "sliding_attention" else None
388
+
389
+ self.q_proj = nn.Linear(
390
+ config.hidden_size, config.num_attention_heads * self.head_dim, bias=False
391
+ )
392
+ self.k_proj = nn.Linear(
393
+ config.hidden_size, config.num_key_value_heads * self.head_dim, bias=False
394
+ )
395
+ self.v_proj = nn.Linear(
396
+ config.hidden_size, config.num_key_value_heads * self.head_dim, bias=False
397
+ )
398
+
399
+ self.q_norm = MapleRMSNorm(self.head_dim, eps=config.rms_norm_eps)
400
+ self.k_norm = MapleRMSNorm(self.head_dim, eps=config.rms_norm_eps)
401
+
402
+ self.o_proj = nn.Linear(self.num_heads * self.head_dim, self.hidden_size, bias=config.use_bias)
403
+
404
+ def forward(
405
+ self,
406
+ hidden_states: torch.Tensor,
407
+ attention_mask: Optional[torch.Tensor] = None, # IMPORTANT: pass prepared causal_mask here
408
+ position_ids: Optional[torch.LongTensor] = None,
409
+ past_key_value: Optional[Cache] = None,
410
+ output_attentions: bool = False,
411
+ use_cache: bool = False,
412
+ cache_position: Optional[torch.LongTensor] = None, # IMPORTANT: needed for modern cache update
413
+ position_embeddings: Optional[Tuple[torch.Tensor, torch.Tensor, torch.Tensor]] = None,
414
+ **kwargs,
415
+ ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Cache]]:
416
+ bsz, q_len, _ = hidden_states.size()
417
+ qkv_weight = torch.cat([self.q_proj.weight, self.k_proj.weight, self.v_proj.weight], dim=0)
418
+ out_qkv = torch.nn.functional.linear(hidden_states, qkv_weight)
419
+ cos, sin, _freqs = position_embeddings
420
+ # qkv = out_qkv.view(bsz, q_len, self.num_heads + 2 * self.num_key_value_heads, self.head_dim)
421
+
422
+ # query_states, key_states, value_states = qkv.split(
423
+ # [self.num_heads, self.num_key_value_heads, self.num_key_value_heads], dim=-2
424
+ # )
425
+ # query_states = query_states.transpose(1, 2)
426
+ # key_states = key_states.transpose(1, 2)
427
+ # value_states = value_states.transpose(1, 2)
428
+
429
+ # query_states = self.q_norm(query_states)
430
+ # key_states = self.k_norm(key_states)
431
+ # if self.sliding_window is not None:
432
+ # query_states, key_states = apply_rotary_pos_emb(query_states, key_states, cos, sin)
433
+ q, k, v = split_qkv_reference(
434
+ out_qkv,
435
+ self.query_layernorm.weight,
436
+ self.key_layernorm.weight,
437
+ _freqs.contiguous(),
438
+ self.num_heads,
439
+ self.config.num_key_value_heads,
440
+ eps=1e-6,
441
+ interleaved=False,
442
+ silu=False,
443
+ rope=True if self.sliding_window is not None else False,
444
+ )
445
+
446
+ # ---- Modern Cache update wiring (DynamicCache / StaticCache compatible) ----
447
+ if use_cache and past_key_value is not None:
448
+ cache_kwargs = {"sin": sin, "cos": cos, "cache_position": cache_position}
449
+ key_states, value_states = past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs)
450
+
451
+ # fa should transpose internally
452
+ attention_interface = ALL_ATTENTION_FUNCTIONS[self.config._attn_implementation]
453
+ attn_output, attn_weights = flash_attention_forward(
454
+ self,
455
+ query_states,
456
+ key_states,
457
+ value_states,
458
+ attention_mask, # prepared causal mask (or None for varlen flash path)
459
+ dropout=0.0,
460
+ position_ids=position_ids,
461
+ scaling=self.scaling,
462
+ sliding_window=self.sliding_window, # keep your prototype behavior
463
+ **kwargs,
464
+ )
465
+
466
+ attn_output = attn_output.reshape(bsz, q_len, -1).contiguous()
467
+ dense_weight = self.o_proj.weight
468
+ attn_output = torch.nn.functional.linear(attn_output, dense_weight)
469
+
470
+ if not output_attentions:
471
+ attn_weights = None
472
+
473
+ return attn_output, attn_weights, past_key_value
474
+
475
+
476
+ class MapleDecoderLayer(nn.Module):
477
+ def __init__(self, config: MapleConfig, layer_idx: int):
478
+ super().__init__()
479
+ self.hidden_size = config.hidden_size
480
+ self.layer_idx = layer_idx
481
+
482
+ self.self_attn = MapleAttention(config=config, layer_idx=layer_idx)
483
+
484
+ self.mlp = MapleSparseMoeBlock(config)
485
+
486
+ self.input_layernorm = MapleRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
487
+ self.post_attention_layernorm = MapleRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
488
+
489
+ def forward(
490
+ self,
491
+ hidden_states: torch.Tensor,
492
+ attention_mask: Optional[torch.Tensor] = None, # prepared causal mask
493
+ position_ids: Optional[torch.LongTensor] = None,
494
+ past_key_value: Optional[Cache] = None,
495
+ output_attentions: Optional[bool] = False,
496
+ output_router_logits: Optional[bool] = False, # your MOE doesn't return router logits; kept for API
497
+ use_cache: Optional[bool] = False,
498
+ cache_position: Optional[torch.LongTensor] = None,
499
+ position_embeddings: Optional[Tuple[torch.Tensor, torch.Tensor, torch.Tensor]] = None,
500
+ **kwargs,
501
+ ) -> Tuple[
502
+ torch.Tensor,
503
+ Optional[torch.Tensor],
504
+ Optional[Cache],
505
+ torch.Tensor,
506
+ Optional[torch.Tensor],
507
+ ]:
508
+ residual = hidden_states
509
+ hidden_states = self.input_layernorm(hidden_states)
510
+
511
+ attn_out, self_attn_weights, present_key_value = self.self_attn(
512
+ hidden_states=hidden_states,
513
+ attention_mask=attention_mask,
514
+ position_ids=position_ids,
515
+ past_key_value=past_key_value,
516
+ output_attentions=bool(output_attentions),
517
+ use_cache=bool(use_cache),
518
+ cache_position=cache_position,
519
+ position_embeddings=position_embeddings,
520
+ **kwargs,
521
+ )
522
+ hidden_states = residual + attn_out
523
+
524
+ residual = hidden_states
525
+ hidden_states = self.post_attention_layernorm(hidden_states)
526
+
527
+ mlp_out = self.mlp(hidden_states)
528
+ if isinstance(mlp_out, tuple):
529
+ hidden_states, aux_loss = mlp_out
530
+ else:
531
+ hidden_states, aux_loss = mlp_out, 0.0
532
+
533
+ hidden_states = residual + hidden_states.to(residual.device)
534
+ router_logits = None
535
+
536
+ return (
537
+ hidden_states,
538
+ self_attn_weights,
539
+ present_key_value,
540
+ aux_loss,
541
+ router_logits,
542
+ )
543
+
544
+
545
+ @add_start_docstrings(
546
+ "The bare Maple Model outputting raw hidden-states without any specific head on top.",
547
+ )
548
+ class MaplePreTrainedModel(PreTrainedModel):
549
+ config_class = MapleConfig
550
+ base_model_prefix = "model"
551
+ supports_gradient_checkpointing = True
552
+ _no_split_modules = ["MapleDecoderLayer"]
553
+ _skip_keys_device_placement = "past_key_values"
554
+ _supports_attention_backend = True
555
+ _supports_flash_attn_2 = True
556
+ _supports_sdpa = True
557
+ _supports_cache_class = True
558
+
559
+ def _init_weights(self, module):
560
+ std = self.config.initializer_range
561
+ if isinstance(module, nn.Linear):
562
+ module.weight.data.normal_(mean=0.0, std=std)
563
+ if module.bias is not None:
564
+ module.bias.data.zero_()
565
+ elif isinstance(module, nn.Embedding):
566
+ module.weight.data.normal_(mean=0.0, std=std)
567
+ if module.padding_idx is not None:
568
+ module.weight.data[module.padding_idx].zero_()
569
+
570
+
571
+ @add_start_docstrings(
572
+ "The bare Maple Model outputting raw hidden-states without any specific head on top.",
573
+ )
574
+ class MapleModel(MaplePreTrainedModel):
575
+ def __init__(self, config: MapleConfig):
576
+ super().__init__(config)
577
+ self.padding_idx = config.pad_token_id
578
+ self.vocab_size = config.vocab_size
579
+
580
+ self.word_embeddings = nn.Embedding(config.vocab_size, config.hidden_size, self.padding_idx)
581
+
582
+ layers = []
583
+ for layer_idx in range(config.num_hidden_layers):
584
+ layers.append(MapleDecoderLayer(config, layer_idx))
585
+ self.layers = nn.ModuleList(layers)
586
+ self.config = config
587
+
588
+ self.norm = MapleRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
589
+ self.rotary_emb = MapleRotaryEmbedding(config=config)
590
+ self.gradient_checkpointing = False
591
+ self._cache_debug_calls = 0
592
+ self.post_init()
593
+
594
+ def get_input_embeddings(self):
595
+ return self.word_embeddings
596
+
597
+ def set_input_embeddings(self, value):
598
+ self.word_embeddings = value
599
+
600
+ def prepare_fa2_from_position_ids(self, position_ids: torch.Tensor):
601
+ position_ids = position_ids.flatten()
602
+ T = position_ids.numel()
603
+ indices_q = torch.arange(T, device=position_ids.device, dtype=torch.int32)
604
+
605
+ starts = indices_q[position_ids == 0]
606
+
607
+ # If no segment-start markers exist (common in decoding where pos ids are offset),
608
+ # treat as a single sequence.
609
+ if starts.numel() == 0:
610
+ cu_seq_lens = torch.tensor([0, T], device=position_ids.device, dtype=torch.int32)
611
+ else:
612
+ # ensure boundaries valid
613
+ if starts[0].item() != 0:
614
+ starts = torch.cat([starts.new_zeros(1), starts], dim=0)
615
+ if starts[-1].item() != T:
616
+ starts = torch.cat([starts, starts.new_tensor([T])], dim=0)
617
+ cu_seq_lens = starts
618
+
619
+ max_length = (cu_seq_lens[1:] - cu_seq_lens[:-1]).max().item()
620
+ return (indices_q, (cu_seq_lens, cu_seq_lens), (max_length, max_length))
621
+
622
+ # def prepare_fa2_from_position_ids(self, position_ids: torch.Tensor):
623
+ # position_ids = position_ids.flatten()
624
+ # indices_q = torch.arange(position_ids.size(0), device=position_ids.device, dtype=torch.int32)
625
+
626
+ # cu_seq_lens = torch.cat(
627
+ # (
628
+ # indices_q[position_ids == 0],
629
+ # torch.tensor(position_ids.size(), device=position_ids.device, dtype=torch.int32),
630
+ # )
631
+ # )
632
+
633
+ # # max_length在不同的model里面type不同
634
+ # # modeling_qwen3_moe_foundation/modeling_qwen2_5_omni里为tensor
635
+ # # modeling_qwen2_vl的为int
636
+ # # 此处采用有.item()的写法,在decoder layers之前拿到int type的max_length
637
+ # # 否则在decoder里面仍然每一层都会触发.item()
638
+ # max_length = cu_seq_lens.diff().max().item()
639
+
640
+ # return (indices_q, (cu_seq_lens, cu_seq_lens), (max_length, max_length))
641
+
642
+ def forward(
643
+ self,
644
+ input_ids: torch.LongTensor = None,
645
+ attention_mask: Optional[torch.Tensor] = None, # 2D padding mask (B, S) coming in
646
+ position_ids: Optional[torch.LongTensor] = None,
647
+ past_key_values: Optional[Cache] = None,
648
+ inputs_embeds: Optional[torch.FloatTensor] = None,
649
+ use_cache: Optional[bool] = None,
650
+ output_attentions: Optional[bool] = None,
651
+ output_hidden_states: Optional[bool] = None,
652
+ output_router_logits: Optional[bool] = None,
653
+ return_dict: Optional[bool] = None,
654
+ cache_position: Optional[torch.LongTensor] = None,
655
+ **kwargs,
656
+ ) -> Union[Tuple, MoeV2ModelOutputWithPast]:
657
+ debug_cache = bool(kwargs.pop("debug_cache", False))
658
+ if debug_cache:
659
+ print(f"Debug cache enabled for call {self._cache_debug_calls}")
660
+ debug_call_id = self._cache_debug_calls
661
+ self._cache_debug_calls += 1
662
+
663
+ output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
664
+ output_router_logits = (
665
+ output_router_logits if output_router_logits is not None else self.config.output_router_logits
666
+ )
667
+ output_hidden_states = (
668
+ output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
669
+ )
670
+ use_cache = use_cache if use_cache is not None else self.config.use_cache
671
+ return_dict = return_dict if return_dict is not None else self.config.use_return_dict
672
+
673
+ # exactly one of input_ids / inputs_embeds
674
+ if (input_ids is None) == (inputs_embeds is None):
675
+ raise ValueError("You must specify exactly one of input_ids or inputs_embeds")
676
+
677
+ if self.gradient_checkpointing and self.training and use_cache:
678
+ logger.warning_once(
679
+ "`use_cache=True` is incompatible with gradient checkpointing. Setting `use_cache=False`."
680
+ )
681
+ use_cache = False
682
+
683
+ if use_cache and past_key_values is None:
684
+ past_key_values = DynamicCache()
685
+
686
+ if inputs_embeds is None:
687
+ inputs_embeds = self.word_embeddings(input_ids)
688
+
689
+ # SGLang transformers backend passes `forward_batch`; use it to identify
690
+ # decode mode (can have S>1 tokens due to token packing) and avoid
691
+ # decode-only dynamic metadata that harms CUDA graph capture.
692
+ forward_batch = kwargs.get("forward_batch", None)
693
+ is_decode_step = False
694
+ forward_mode = getattr(forward_batch, "forward_mode", None) if forward_batch is not None else None
695
+ if forward_mode is not None:
696
+ for mode_name in (
697
+ "is_decode",
698
+ "is_decode_or_idle",
699
+ "is_target_verify",
700
+ "is_draft_decode",
701
+ ):
702
+ mode_fn = getattr(forward_mode, mode_name, None)
703
+ if callable(mode_fn) and bool(mode_fn()):
704
+ is_decode_step = True
705
+ break
706
+
707
+ # Perform one-time in-place weight quantization during prefill (S > 1),
708
+
709
+ past_seen_tokens = past_key_values.get_seq_length() if past_key_values is not None else 0
710
+
711
+ if cache_position is None:
712
+ cache_position = torch.arange(
713
+ past_seen_tokens,
714
+ past_seen_tokens + inputs_embeds.shape[1],
715
+ device=inputs_embeds.device,
716
+ )
717
+
718
+ if position_ids is not None:
719
+ # For bsh cases, expand [1, S] position_ids to [B, S] before FA2 metadata prep.
720
+ batch_size = input_ids.shape[0] if input_ids is not None else inputs_embeds.shape[0]
721
+ if position_ids.shape[0] != batch_size:
722
+ position_ids = position_ids.expand(batch_size, -1)
723
+
724
+ # Decode does not need cu_seq_lens/max_length metadata and creating
725
+ # them every step hurts CUDA graph capture stability.
726
+ if (not is_decode_step) and inputs_embeds.shape[1] > 1:
727
+ _, (cu_seq_lens_q, cu_seq_lens_k), (max_length_q, max_length_k) = self.prepare_fa2_from_position_ids(
728
+ position_ids
729
+ )
730
+ kwargs["cu_seq_lens_q"] = cu_seq_lens_q
731
+ kwargs["cu_seq_lens_k"] = cu_seq_lens_k
732
+ kwargs["max_length_q"] = max_length_q
733
+ kwargs["max_length_k"] = max_length_k
734
+
735
+ if position_ids is None:
736
+ position_ids = cache_position.unsqueeze(0)
737
+
738
+ # IMPORTANT: build prepared causal_mask and pass it into layers (NOT raw attention_mask)
739
+ # mask_function = create_causal_mask # swap to create_sliding_window_causal_mask if you enable sliding window
740
+ # causal_mask = mask_function(
741
+ # config=self.config,
742
+ # input_embeds=inputs_embeds,
743
+ # attention_mask=attention_mask,
744
+ # cache_position=cache_position,
745
+ # past_key_values=past_key_values,
746
+ # position_ids=position_ids,
747
+ # )
748
+ # TODO: Im just disabling causal mask right now idk fix this later when we need SWA
749
+ causal_mask = None
750
+
751
+ hidden_states = inputs_embeds
752
+ position_embeddings = self.rotary_emb(hidden_states, position_ids)
753
+
754
+ all_hidden_states = () if output_hidden_states else None
755
+ all_self_attns = () if output_attentions else None
756
+ all_router_logits = () if output_router_logits else None
757
+
758
+ aux_loss_sum = 0.0
759
+
760
+ for decoder_layer in self.layers:
761
+ if output_hidden_states:
762
+ all_hidden_states += (hidden_states,)
763
+
764
+ if self.gradient_checkpointing and self.training:
765
+ layer_outputs = self._gradient_checkpointing_func(
766
+ decoder_layer.__call__,
767
+ hidden_states,
768
+ causal_mask,
769
+ position_ids,
770
+ past_key_values,
771
+ output_attentions,
772
+ output_router_logits,
773
+ use_cache,
774
+ cache_position,
775
+ position_embeddings,
776
+ **kwargs,
777
+ )
778
+ else:
779
+ layer_outputs = decoder_layer(
780
+ hidden_states,
781
+ attention_mask=causal_mask, # <-- FIXED
782
+ position_ids=position_ids,
783
+ past_key_value=past_key_values,
784
+ output_attentions=output_attentions,
785
+ output_router_logits=output_router_logits,
786
+ use_cache=use_cache,
787
+ cache_position=cache_position, # <-- FIXED
788
+ position_embeddings=position_embeddings,
789
+ **kwargs,
790
+ )
791
+
792
+ hidden_states = layer_outputs[0]
793
+
794
+ if output_attentions:
795
+ all_self_attns += (layer_outputs[1],)
796
+
797
+ # aux loss is at index 3 in our layer return
798
+ aux_loss_sum = aux_loss_sum + layer_outputs[3]
799
+
800
+ if output_router_logits:
801
+ all_router_logits += (layer_outputs[4],)
802
+
803
+ hidden_states = self.norm(hidden_states)
804
+
805
+ if debug_cache:
806
+ past_after_tokens = past_key_values.get_seq_length() if past_key_values is not None else 0
807
+ cache_start = int(cache_position[0].item()) if cache_position.numel() > 0 else -1
808
+ cache_end = int(cache_position[-1].item()) if cache_position.numel() > 0 else -1
809
+ cache_hit = bool(use_cache and inputs_embeds.shape[1] == 1 and past_seen_tokens > 0)
810
+ print(
811
+ "[cache-debug] "
812
+ f"call={debug_call_id} use_cache={use_cache} "
813
+ f"input_len={inputs_embeds.shape[1]} "
814
+ f"past_before={past_seen_tokens} past_after={past_after_tokens} "
815
+ f"cache_pos=[{cache_start},{cache_end}] "
816
+ f"cache_hit_expected={cache_hit}"
817
+ )
818
+
819
+ if output_hidden_states:
820
+ all_hidden_states += (hidden_states,)
821
+ moe_layer_count = len(self.layers) - 1
822
+ out = MoeV2ModelOutputWithPast(
823
+ last_hidden_state=hidden_states,
824
+ past_key_values=past_key_values if use_cache else None,
825
+ hidden_states=all_hidden_states,
826
+ attentions=all_self_attns,
827
+ router_logits=all_router_logits,
828
+ aux_loss=aux_loss_sum / moe_layer_count, # keeping your prototype behavior
829
+ )
830
+ return (
831
+ out
832
+ if return_dict
833
+ else (
834
+ out.last_hidden_state,
835
+ out.past_key_values,
836
+ out.hidden_states,
837
+ out.attentions,
838
+ )
839
+ )
840
+
841
+
842
+ class MapleForCausalLM(MaplePreTrainedModel, GenerationMixin):
843
+ _tied_weights_keys = ["lm_head.weight"]
844
+
845
+ def __init__(self, config: MapleConfig):
846
+ super().__init__(config)
847
+ self.model = MapleModel(config)
848
+ self.vocab_size = config.vocab_size
849
+
850
+ self.lm_head = nn.Linear(config.hidden_size, config.vocab_size, bias=False)
851
+ self.router_aux_loss_coef = 0.001
852
+ self.post_init()
853
+
854
+ def get_input_embeddings(self):
855
+ return self.model.word_embeddings
856
+
857
+ def set_input_embeddings(self, value):
858
+ self.model.word_embeddings = value
859
+
860
+ def get_output_embeddings(self):
861
+ return self.lm_head
862
+
863
+ def set_output_embeddings(self, new_embeddings):
864
+ self.lm_head = new_embeddings
865
+
866
+ def set_decoder(self, decoder):
867
+ self.model = decoder
868
+
869
+ def get_decoder(self):
870
+ return self.model
871
+
872
+ # @add_start_docstrings_to_model_forward(BAILINGMOEV2_INPUTS_DOCSTRING)
873
+ # @replace_return_docstrings(output_type=MoEV2CausalLMOutputWithPast, config_class=_CONFIG_FOR_DOC)
874
+ def forward(
875
+ self,
876
+ input_ids: torch.LongTensor = None,
877
+ attention_mask: Optional[torch.Tensor] = None,
878
+ position_ids: Optional[torch.LongTensor] = None,
879
+ past_key_values: Optional[Cache] = None,
880
+ inputs_embeds: Optional[torch.FloatTensor] = None,
881
+ labels: Optional[torch.Tensor] = None,
882
+ use_cache: Optional[bool] = None,
883
+ output_attentions: Optional[bool] = None,
884
+ output_hidden_states: Optional[bool] = None,
885
+ output_router_logits: Optional[bool] = None,
886
+ return_dict: Optional[bool] = None,
887
+ logits_to_keep: Union[int, torch.Tensor] = 0,
888
+ **kwargs,
889
+ ) -> Union[Tuple, MapleOutputWithPast]:
890
+ output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
891
+ output_hidden_states = (
892
+ output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
893
+ )
894
+ output_router_logits = (
895
+ output_router_logits if output_router_logits is not None else self.config.output_router_logits
896
+ )
897
+ return_dict = return_dict if return_dict is not None else self.config.use_return_dict
898
+
899
+ outputs = self.model(
900
+ input_ids=input_ids,
901
+ attention_mask=attention_mask,
902
+ position_ids=position_ids,
903
+ past_key_values=past_key_values,
904
+ inputs_embeds=inputs_embeds,
905
+ use_cache=use_cache,
906
+ output_attentions=output_attentions,
907
+ output_hidden_states=output_hidden_states,
908
+ output_router_logits=output_router_logits,
909
+ return_dict=True, # ensure attribute access
910
+ **kwargs,
911
+ )
912
+
913
+ hidden_states = outputs.last_hidden_state
914
+ assert isinstance(hidden_states, torch.Tensor)
915
+
916
+ # slice logits if requested
917
+ loss = None
918
+ logits = None
919
+ if labels is not None:
920
+ loss, logits = self.loss_function(hidden_states, self.lm_head.weight, labels)
921
+ else:
922
+ logits = self.lm_head(hidden_states)
923
+ out = MapleOutputWithPast(
924
+ loss=loss,
925
+ aux_loss=getattr(outputs, "aux_loss", 0.0),
926
+ logits=logits,
927
+ past_key_values=outputs.past_key_values if hasattr(outputs, "past_key_values") else None,
928
+ hidden_states=outputs.hidden_states if hasattr(outputs, "hidden_states") else None,
929
+ attentions=outputs.attentions if hasattr(outputs, "attentions") else None,
930
+ router_logits=outputs.router_logits if hasattr(outputs, "router_logits") else None,
931
+ )
932
+ return out
special_tokens_map.json ADDED
@@ -0,0 +1,31 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "additional_special_tokens": [
3
+ "<|im_start|>",
4
+ "<|im_end|>",
5
+ "<|object_ref_start|>",
6
+ "<|object_ref_end|>",
7
+ "<|box_start|>",
8
+ "<|box_end|>",
9
+ "<|quad_start|>",
10
+ "<|quad_end|>",
11
+ "<|vision_start|>",
12
+ "<|vision_end|>",
13
+ "<|vision_pad|>",
14
+ "<|image_pad|>",
15
+ "<|video_pad|>"
16
+ ],
17
+ "eos_token": {
18
+ "content": "<|im_end|>",
19
+ "lstrip": false,
20
+ "normalized": false,
21
+ "rstrip": false,
22
+ "single_word": false
23
+ },
24
+ "pad_token": {
25
+ "content": "<|endoftext|>",
26
+ "lstrip": false,
27
+ "normalized": false,
28
+ "rstrip": false,
29
+ "single_word": false
30
+ }
31
+ }
tokenizer.json ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:aeb13307a71acd8fe81861d94ad54ab689df773318809eed3cbe794b4492dae4
3
+ size 11422654
tokenizer_config.json ADDED
@@ -0,0 +1,240 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "add_bos_token": false,
3
+ "add_prefix_space": false,
4
+ "added_tokens_decoder": {
5
+ "151643": {
6
+ "content": "<|endoftext|>",
7
+ "lstrip": false,
8
+ "normalized": false,
9
+ "rstrip": false,
10
+ "single_word": false,
11
+ "special": true
12
+ },
13
+ "151644": {
14
+ "content": "<|im_start|>",
15
+ "lstrip": false,
16
+ "normalized": false,
17
+ "rstrip": false,
18
+ "single_word": false,
19
+ "special": true
20
+ },
21
+ "151645": {
22
+ "content": "<|im_end|>",
23
+ "lstrip": false,
24
+ "normalized": false,
25
+ "rstrip": false,
26
+ "single_word": false,
27
+ "special": true
28
+ },
29
+ "151646": {
30
+ "content": "<|object_ref_start|>",
31
+ "lstrip": false,
32
+ "normalized": false,
33
+ "rstrip": false,
34
+ "single_word": false,
35
+ "special": true
36
+ },
37
+ "151647": {
38
+ "content": "<|object_ref_end|>",
39
+ "lstrip": false,
40
+ "normalized": false,
41
+ "rstrip": false,
42
+ "single_word": false,
43
+ "special": true
44
+ },
45
+ "151648": {
46
+ "content": "<|box_start|>",
47
+ "lstrip": false,
48
+ "normalized": false,
49
+ "rstrip": false,
50
+ "single_word": false,
51
+ "special": true
52
+ },
53
+ "151649": {
54
+ "content": "<|box_end|>",
55
+ "lstrip": false,
56
+ "normalized": false,
57
+ "rstrip": false,
58
+ "single_word": false,
59
+ "special": true
60
+ },
61
+ "151650": {
62
+ "content": "<|quad_start|>",
63
+ "lstrip": false,
64
+ "normalized": false,
65
+ "rstrip": false,
66
+ "single_word": false,
67
+ "special": true
68
+ },
69
+ "151651": {
70
+ "content": "<|quad_end|>",
71
+ "lstrip": false,
72
+ "normalized": false,
73
+ "rstrip": false,
74
+ "single_word": false,
75
+ "special": true
76
+ },
77
+ "151652": {
78
+ "content": "<|vision_start|>",
79
+ "lstrip": false,
80
+ "normalized": false,
81
+ "rstrip": false,
82
+ "single_word": false,
83
+ "special": true
84
+ },
85
+ "151653": {
86
+ "content": "<|vision_end|>",
87
+ "lstrip": false,
88
+ "normalized": false,
89
+ "rstrip": false,
90
+ "single_word": false,
91
+ "special": true
92
+ },
93
+ "151654": {
94
+ "content": "<|vision_pad|>",
95
+ "lstrip": false,
96
+ "normalized": false,
97
+ "rstrip": false,
98
+ "single_word": false,
99
+ "special": true
100
+ },
101
+ "151655": {
102
+ "content": "<|image_pad|>",
103
+ "lstrip": false,
104
+ "normalized": false,
105
+ "rstrip": false,
106
+ "single_word": false,
107
+ "special": true
108
+ },
109
+ "151656": {
110
+ "content": "<|video_pad|>",
111
+ "lstrip": false,
112
+ "normalized": false,
113
+ "rstrip": false,
114
+ "single_word": false,
115
+ "special": true
116
+ },
117
+ "151657": {
118
+ "content": "<tool_call>",
119
+ "lstrip": false,
120
+ "normalized": false,
121
+ "rstrip": false,
122
+ "single_word": false,
123
+ "special": false
124
+ },
125
+ "151658": {
126
+ "content": "</tool_call>",
127
+ "lstrip": false,
128
+ "normalized": false,
129
+ "rstrip": false,
130
+ "single_word": false,
131
+ "special": false
132
+ },
133
+ "151659": {
134
+ "content": "<|fim_prefix|>",
135
+ "lstrip": false,
136
+ "normalized": false,
137
+ "rstrip": false,
138
+ "single_word": false,
139
+ "special": false
140
+ },
141
+ "151660": {
142
+ "content": "<|fim_middle|>",
143
+ "lstrip": false,
144
+ "normalized": false,
145
+ "rstrip": false,
146
+ "single_word": false,
147
+ "special": false
148
+ },
149
+ "151661": {
150
+ "content": "<|fim_suffix|>",
151
+ "lstrip": false,
152
+ "normalized": false,
153
+ "rstrip": false,
154
+ "single_word": false,
155
+ "special": false
156
+ },
157
+ "151662": {
158
+ "content": "<|fim_pad|>",
159
+ "lstrip": false,
160
+ "normalized": false,
161
+ "rstrip": false,
162
+ "single_word": false,
163
+ "special": false
164
+ },
165
+ "151663": {
166
+ "content": "<|repo_name|>",
167
+ "lstrip": false,
168
+ "normalized": false,
169
+ "rstrip": false,
170
+ "single_word": false,
171
+ "special": false
172
+ },
173
+ "151664": {
174
+ "content": "<|file_sep|>",
175
+ "lstrip": false,
176
+ "normalized": false,
177
+ "rstrip": false,
178
+ "single_word": false,
179
+ "special": false
180
+ },
181
+ "151665": {
182
+ "content": "<tool_response>",
183
+ "lstrip": false,
184
+ "normalized": false,
185
+ "rstrip": false,
186
+ "single_word": false,
187
+ "special": false
188
+ },
189
+ "151666": {
190
+ "content": "</tool_response>",
191
+ "lstrip": false,
192
+ "normalized": false,
193
+ "rstrip": false,
194
+ "single_word": false,
195
+ "special": false
196
+ },
197
+ "151667": {
198
+ "content": "<think>",
199
+ "lstrip": false,
200
+ "normalized": false,
201
+ "rstrip": false,
202
+ "single_word": false,
203
+ "special": false
204
+ },
205
+ "151668": {
206
+ "content": "</think>",
207
+ "lstrip": false,
208
+ "normalized": false,
209
+ "rstrip": false,
210
+ "single_word": false,
211
+ "special": false
212
+ }
213
+ },
214
+ "additional_special_tokens": [
215
+ "<|im_start|>",
216
+ "<|im_end|>",
217
+ "<|object_ref_start|>",
218
+ "<|object_ref_end|>",
219
+ "<|box_start|>",
220
+ "<|box_end|>",
221
+ "<|quad_start|>",
222
+ "<|quad_end|>",
223
+ "<|vision_start|>",
224
+ "<|vision_end|>",
225
+ "<|vision_pad|>",
226
+ "<|image_pad|>",
227
+ "<|video_pad|>"
228
+ ],
229
+ "bos_token": null,
230
+ "clean_up_tokenization_spaces": false,
231
+ "eos_token": "<|im_end|>",
232
+ "errors": "replace",
233
+ "extra_special_tokens": {},
234
+ "model_max_length": 1010000,
235
+ "pad_token": "<|endoftext|>",
236
+ "padding_side": "right",
237
+ "split_special_tokens": false,
238
+ "tokenizer_class": "Qwen2Tokenizer",
239
+ "unk_token": null
240
+ }
vocab.json ADDED
The diff for this file is too large to render. See raw diff