Upload checkpoint-1550

added_tokens.json

@@ -0,0 +1,30 @@
+{
+  "</think>": 151668,
+  "</tool_call>": 151658,
+  "</tool_response>": 151666,
+  "<EOB>": 151670,
+  "<think>": 151667,
+  "<tool_call>": 151657,
+  "<tool_response>": 151665,
+  "<|MASK|>": 151669,
+  "<|box_end|>": 151649,
+  "<|box_start|>": 151648,
+  "<|endoftext|>": 151643,
+  "<|file_sep|>": 151664,
+  "<|fim_middle|>": 151660,
+  "<|fim_pad|>": 151662,
+  "<|fim_prefix|>": 151659,
+  "<|fim_suffix|>": 151661,
+  "<|im_end|>": 151645,
+  "<|im_start|>": 151644,
+  "<|image_pad|>": 151655,
+  "<|object_ref_end|>": 151647,
+  "<|object_ref_start|>": 151646,
+  "<|quad_end|>": 151651,
+  "<|quad_start|>": 151650,
+  "<|repo_name|>": 151663,
+  "<|video_pad|>": 151656,
+  "<|vision_end|>": 151653,
+  "<|vision_pad|>": 151654,
+  "<|vision_start|>": 151652
+}

chat_template.jinja

@@ -0,0 +1,85 @@
+{%- if tools %}
+    {{- '<|im_start|>system\n' }}
+    {%- if messages[0].role == 'system' %}
+        {{- messages[0].content + '\n\n' }}
+    {%- endif %}
+    {{- "# 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>" }}
+    {%- for tool in tools %}
+        {{- "\n" }}
+        {{- tool | tojson }}
+    {%- endfor %}
+    {{- "\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" }}
+{%- else %}
+    {%- if messages[0].role == 'system' %}
+        {{- '<|im_start|>system\n' + messages[0].content + '<|im_end|>\n' }}
+    {%- endif %}
+{%- endif %}
+{%- set ns = namespace(multi_step_tool=true, last_query_index=messages|length - 1) %}
+{%- for message in messages[::-1] %}
+    {%- set index = (messages|length - 1) - loop.index0 %}
+    {%- if ns.multi_step_tool and message.role == "user" and not(message.content.startswith('<tool_response>') and message.content.endswith('</tool_response>')) %}
+        {%- set ns.multi_step_tool = false %}
+        {%- set ns.last_query_index = index %}
+    {%- endif %}
+{%- endfor %}
+{%- for message in messages %}
+    {%- if (message.role == "user") or (message.role == "system" and not loop.first) %}
+        {{- '<|im_start|>' + message.role + '\n' + message.content + '<|im_end|>' + '\n' }}
+    {%- elif message.role == "assistant" %}
+        {%- set content = message.content %}
+        {%- set reasoning_content = '' %}
+        {%- if message.reasoning_content is defined and message.reasoning_content is not none %}
+            {%- set reasoning_content = message.reasoning_content %}
+        {%- else %}
+            {%- if '</think>' in message.content %}
+                {%- set content = message.content.split('</think>')[-1].lstrip('\n') %}
+                {%- set reasoning_content = message.content.split('</think>')[0].rstrip('\n').split('<think>')[-1].lstrip('\n') %}
+            {%- endif %}
+        {%- endif %}
+        {%- if loop.index0 > ns.last_query_index %}
+            {%- if loop.last or (not loop.last and reasoning_content) %}
+                {{- '<|im_start|>' + message.role + '\n<think>\n' + reasoning_content.strip('\n') + '\n</think>\n\n' + content.lstrip('\n') }}
+            {%- else %}
+                {{- '<|im_start|>' + message.role + '\n' + content }}
+            {%- endif %}
+        {%- else %}
+            {{- '<|im_start|>' + message.role + '\n' + content }}
+        {%- endif %}
+        {%- if message.tool_calls %}
+            {%- for tool_call in message.tool_calls %}
+                {%- if (loop.first and content) or (not loop.first) %}
+                    {{- '\n' }}
+                {%- endif %}
+                {%- if tool_call.function %}
+                    {%- set tool_call = tool_call.function %}
+                {%- endif %}
+                {{- '<tool_call>\n{"name": "' }}
+                {{- tool_call.name }}
+                {{- '", "arguments": ' }}
+                {%- if tool_call.arguments is string %}
+                    {{- tool_call.arguments }}
+                {%- else %}
+                    {{- tool_call.arguments | tojson }}
+                {%- endif %}
+                {{- '}\n</tool_call>' }}
+            {%- endfor %}
+        {%- endif %}
+        {{- '<|im_end|>\n' }}
+    {%- elif message.role == "tool" %}
+        {%- if loop.first or (messages[loop.index0 - 1].role != "tool") %}
+            {{- '<|im_start|>user' }}
+        {%- endif %}
+        {{- '\n<tool_response>\n' }}
+        {{- message.content }}
+        {{- '\n</tool_response>' }}
+        {%- if loop.last or (messages[loop.index0 + 1].role != "tool") %}
+            {{- '<|im_end|>\n' }}
+        {%- endif %}
+    {%- endif %}
+{%- endfor %}
+{%- if add_generation_prompt %}
+    {{- '<|im_start|>assistant\n' }}
+    {%- if enable_thinking is defined and enable_thinking is false %}
+        {{- '<think>\n\n</think>\n\n' }}
+    {%- endif %}
+{%- endif %}

config.json

@@ -0,0 +1,45 @@
+{
+  "architectures": [
+    "SDARForCausalLM"
+  ],
+  "attention_bias": false,
+  "attention_dropout": 0.0,
+  "auto_map": {
+    "AutoConfig": "configuration_sdar.SDARConfig",
+    "AutoModel": "modeling_sdar.SDARForCausalLM",
+    "AutoModelForCausalLM": "modeling_sdar.SDARForCausalLM"
+  },
+  "block_size": 4,
+  "bos_token_id": 151643,
+  "debug": false,
+  "dynamic_blocks": false,
+  "eob_token_id": 151670,
+  "eos_token_id": 151643,
+  "ep_size": 1,
+  "fuse_cross_entropy": true,
+  "head_dim": 128,
+  "hidden_act": "silu",
+  "hidden_size": 2560,
+  "initializer_range": 0.02,
+  "intermediate_size": 9728,
+  "mask_token_id": 151669,
+  "max_position_embeddings": 32768,
+  "max_window_layers": 36,
+  "micro_forward": false,
+  "model_type": "sdar",
+  "num_attention_heads": 32,
+  "num_hidden_layers": 36,
+  "num_key_value_heads": 8,
+  "rms_norm_eps": 1e-06,
+  "rope_scaling": null,
+  "rope_theta": 1000000,
+  "skip_checkpoint": false,
+  "sliding_window": null,
+  "tie_word_embeddings": false,
+  "torch_dtype": "bfloat16",
+  "transformers_version": "4.52.4",
+  "use_cache": false,
+  "use_deepep": false,
+  "use_sliding_window": false,
+  "vocab_size": 151936
+}

configuration_sdar.py

@@ -0,0 +1,212 @@
+# coding=utf-8
+# Copyright 2024 The Qwen team, Alibaba Group and the HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""SDAR model configuration"""
+
+from transformers.configuration_utils import PretrainedConfig
+from transformers.modeling_rope_utils import rope_config_validation
+from transformers.utils import logging
+
+
+logger = logging.get_logger(__name__)
+
+
+class SDARConfig(PretrainedConfig):
+    r"""
+    This is the configuration class to store the configuration of a [`SDARModel`]. It is used to instantiate a
+    SDAR model according to the specified arguments, defining the model architecture. Instantiating a configuration
+    with the defaults will yield a similar configuration to that of
+    SDAR-1.7B [DiffuOpen/SDAR-1.7B-Chat](https://huggingface.co/DiffuOpen/SDAR-1.7B-Chat/).
+
+    Configuration objects inherit from [`PretrainedConfig`] and can be used to control the model outputs. Read the
+    documentation from [`PretrainedConfig`] for more information.
+
+
+    Args:
+        vocab_size (`int`, *optional*, defaults to 151936):
+            Vocabulary size of the SDAR model. Defines the number of different tokens that can be represented by the
+            `inputs_ids` passed when calling [`SDARModel`]
+        hidden_size (`int`, *optional*, defaults to 4096):
+            Dimension of the hidden representations.
+        intermediate_size (`int`, *optional*, defaults to 22016):
+            Dimension of the MLP representations.
+        num_hidden_layers (`int`, *optional*, defaults to 32):
+            Number of hidden layers in the Transformer encoder.
+        num_attention_heads (`int`, *optional*, defaults to 32):
+            Number of attention heads for each attention layer in the Transformer encoder.
+        num_key_value_heads (`int`, *optional*, defaults to 32):
+            This is the number of key_value heads that should be used to implement Grouped Query Attention. If
+            `num_key_value_heads=num_attention_heads`, the model will use Multi Head Attention (MHA), if
+            `num_key_value_heads=1` the model will use Multi Query Attention (MQA) otherwise GQA is used. When
+            converting a multi-head checkpoint to a GQA checkpoint, each group key and value head should be constructed
+            by meanpooling all the original heads within that group. For more details checkout [this
+            paper](https://arxiv.org/pdf/2305.13245.pdf). If it is not specified, will default to `32`.
+        head_dim (`int`, *optional*, defaults to 128):
+            The attention head dimension.
+        hidden_act (`str` or `function`, *optional*, defaults to `"silu"`):
+            The non-linear activation function (function or string) in the decoder.
+        max_position_embeddings (`int`, *optional*, defaults to 32768):
+            The maximum sequence length that this model might ever be used with.
+        initializer_range (`float`, *optional*, defaults to 0.02):
+            The standard deviation of the truncated_normal_initializer for initializing all weight matrices.
+        rms_norm_eps (`float`, *optional*, defaults to 1e-06):
+            The epsilon used by the rms normalization layers.
+        use_cache (`bool`, *optional*, defaults to `True`):
+            Whether or not the model should return the last key/values attentions (not used by all models). Only
+            relevant if `config.is_decoder=True`.
+        tie_word_embeddings (`bool`, *optional*, defaults to `False`):
+            Whether the model's input and output word embeddings should be tied.
+        rope_theta (`float`, *optional*, defaults to 10000.0):
+            The base period of the RoPE embeddings.
+        rope_scaling (`Dict`, *optional*):
+            Dictionary containing the scaling configuration for the RoPE embeddings. NOTE: if you apply new rope type
+            and you expect the model to work on longer `max_position_embeddings`, we recommend you to update this value
+            accordingly.
+            Expected contents:
+                `rope_type` (`str`):
+                    The sub-variant of RoPE to use. Can be one of ['default', 'linear', 'dynamic', 'yarn', 'longrope',
+                    'llama3'], with 'default' being the original RoPE implementation.
+                `factor` (`float`, *optional*):
+                    Used with all rope types except 'default'. The scaling factor to apply to the RoPE embeddings. In
+                    most scaling types, a `factor` of x will enable the model to handle sequences of length x *
+                    original maximum pre-trained length.
+                `original_max_position_embeddings` (`int`, *optional*):
+                    Used with 'dynamic', 'longrope' and 'llama3'. The original max position embeddings used during
+                    pretraining.
+                `attention_factor` (`float`, *optional*):
+                    Used with 'yarn' and 'longrope'. The scaling factor to be applied on the attention
+                    computation. If unspecified, it defaults to value recommended by the implementation, using the
+                    `factor` field to infer the suggested value.
+                `beta_fast` (`float`, *optional*):
+                    Only used with 'yarn'. Parameter to set the boundary for extrapolation (only) in the linear
+                    ramp function. If unspecified, it defaults to 32.
+                `beta_slow` (`float`, *optional*):
+                    Only used with 'yarn'. Parameter to set the boundary for interpolation (only) in the linear
+                    ramp function. If unspecified, it defaults to 1.
+                `short_factor` (`List[float]`, *optional*):
+                    Only used with 'longrope'. The scaling factor to be applied to short contexts (<
+                    `original_max_position_embeddings`). Must be a list of numbers with the same length as the hidden
+                    size divided by the number of attention heads divided by 2
+                `long_factor` (`List[float]`, *optional*):
+                    Only used with 'longrope'. The scaling factor to be applied to long contexts (<
+                    `original_max_position_embeddings`). Must be a list of numbers with the same length as the hidden
+                    size divided by the number of attention heads divided by 2
+                `low_freq_factor` (`float`, *optional*):
+                    Only used with 'llama3'. Scaling factor applied to low frequency components of the RoPE
+                `high_freq_factor` (`float`, *optional*):
+                    Only used with 'llama3'. Scaling factor applied to high frequency components of the RoPE
+        attention_bias (`bool`, defaults to `False`, *optional*, defaults to `False`):
+            Whether to use a bias in the query, key, value and output projection layers during self-attention.
+        use_sliding_window (`bool`, *optional*, defaults to `False`):
+            Whether to use sliding window attention.
+        sliding_window (`int`, *optional*, defaults to 4096):
+            Sliding window attention (SWA) window size. If not specified, will default to `4096`.
+        max_window_layers (`int`, *optional*, defaults to 28):
+            The number of layers that use SWA (Sliding Window Attention). The bottom layers use SWA while the top use full attention.
+        attention_dropout (`float`, *optional*, defaults to 0.0):
+            The dropout ratio for the attention probabilities.
+
+    ```python
+    >>> from transformers import SDARModel, SDARConfig
+
+    >>> # Initializing a SDAR style configuration
+    >>> configuration = SDARConfig()
+
+    >>> # Initializing a model from the SDAR-8B style configuration
+    >>> model = SDARModel(configuration)
+
+    >>> # Accessing the model configuration
+    >>> configuration = model.config
+    ```"""
+
+    model_type = "sdar"
+    keys_to_ignore_at_inference = ["past_key_values"]
+
+    # Default tensor parallel plan for base model `SDAR`
+    base_model_tp_plan = {
+        "layers.*.self_attn.q_proj": "colwise",
+        "layers.*.self_attn.k_proj": "colwise",
+        "layers.*.self_attn.v_proj": "colwise",
+        "layers.*.self_attn.o_proj": "rowwise",
+        "layers.*.mlp.gate_proj": "colwise",
+        "layers.*.mlp.up_proj": "colwise",
+        "layers.*.mlp.down_proj": "rowwise",
+    }
+    base_model_pp_plan = {
+        "embed_tokens": (["input_ids"], ["inputs_embeds"]),
+        "layers": (["hidden_states", "attention_mask"], ["hidden_states"]),
+        "norm": (["hidden_states"], ["hidden_states"]),
+    }
+
+    def __init__(
+        self,
+        vocab_size=151936,
+        hidden_size=4096,
+        intermediate_size=22016,
+        num_hidden_layers=32,
+        num_attention_heads=32,
+        num_key_value_heads=32,
+        head_dim=128,
+        hidden_act="silu",
+        max_position_embeddings=32768,
+        initializer_range=0.02,
+        rms_norm_eps=1e-6,
+        use_cache=True,
+        tie_word_embeddings=False,
+        rope_theta=10000.0,
+        rope_scaling=None,
+        attention_bias=False,
+        use_sliding_window=False,
+        sliding_window=4096,
+        max_window_layers=28,
+        attention_dropout=0.0,
+        **kwargs,
+    ):
+        self.vocab_size = vocab_size
+        self.max_position_embeddings = max_position_embeddings
+        self.hidden_size = hidden_size
+        self.intermediate_size = intermediate_size
+        self.num_hidden_layers = num_hidden_layers
+        self.num_attention_heads = num_attention_heads
+        self.use_sliding_window = use_sliding_window
+        self.sliding_window = sliding_window  # we check `use_sliding_window` in the modeling code
+        self.max_window_layers = max_window_layers
+
+        # for backward compatibility
+        if num_key_value_heads is None:
+            num_key_value_heads = num_attention_heads
+
+        self.num_key_value_heads = num_key_value_heads
+        self.head_dim = head_dim
+        self.hidden_act = hidden_act
+        self.initializer_range = initializer_range
+        self.rms_norm_eps = rms_norm_eps
+        self.use_cache = use_cache
+        self.rope_theta = rope_theta
+        self.rope_scaling = rope_scaling
+        self.attention_bias = attention_bias
+        self.attention_dropout = attention_dropout
+        # Validate the correctness of rotary position embeddings parameters
+        # BC: if there is a 'type' field, move it to 'rope_type'.
+        if self.rope_scaling is not None and "type" in self.rope_scaling:
+            self.rope_scaling["rope_type"] = self.rope_scaling["type"]
+        rope_config_validation(self)
+
+        super().__init__(
+            tie_word_embeddings=tie_word_embeddings,
+            **kwargs,
+        )
+
+
+__all__ = ["SDARConfig"]

dynamic_blocks_utils.py

@@ -0,0 +1,223 @@
+"""
+Utility functions for dynamic block training with variable-length blocks.
+
+This module provides functions to extract block boundaries from <EOB> tokens
+and generate attention masks for variable-length blocks.
+"""
+
+import torch
+from typing import List, Tuple
+
+
+def calculate_block_nums_from_eob(
+    input_ids: torch.Tensor,
+    num_tokens_list: List[List[int]],
+    eob_token_id: int
+) -> List[List[torch.Tensor]]:
+    """
+    Extract variable block lengths from <EOB> token positions, respecting packed sample boundaries.
+    
+    Args:
+        input_ids: Token IDs tensor, shape (batch_size, seq_len)
+        num_tokens_list: List of lists, where each inner list contains sequence lengths for a batch item.
+                        (Output from calculate_token_nums)
+        eob_token_id: Token ID for <EOB>
+        
+    Returns:
+        List of lists of tensors. Outer list is batch. Inner list is samples. 
+        Each tensor contains block lengths for that sample.
+    """
+    batch_size, seq_len = input_ids.shape
+    all_batch_block_lengths = []
+    
+    for i in range(batch_size):
+        current_ids = input_ids[i]
+        sample_lengths = num_tokens_list[i] # List of integers
+        
+        current_sample_block_lengths = []
+        start_idx = 0
+        
+        for length in sample_lengths:
+            # Handle tensor or int
+            if isinstance(length, torch.Tensor):
+                length = length.item()
+                
+            # Extract sample tokens
+            end_idx = start_idx + length
+            # Ensure we don't go out of bounds (e.g. if sum(lengths) != seq_len due to padding logic differences)
+            # But typically sum(lengths) == seq_len for packed data + padding
+            end_idx = min(end_idx, seq_len)
+            
+            if start_idx >= seq_len:
+                break
+                
+            sample_ids = current_ids[start_idx:end_idx]
+            
+            # Find positions of <EOB> tokens in this sample
+            eob_positions = torch.nonzero(sample_ids == eob_token_id).flatten()
+            
+            # Calculate block lengths for this sample
+            if len(eob_positions) == 0:
+                # No EOB tokens, treat entire sample as one block
+                block_lengths = torch.tensor([length], device=input_ids.device)
+            else:
+                # Add start and end positions
+                # EOB is included in its block (boundary marker)
+                boundaries = torch.cat([
+                    torch.tensor([0], device=input_ids.device),
+                    eob_positions + 1,  # +1 to include EOB token in block
+                    torch.tensor([length], device=input_ids.device)
+                ])
+                block_lengths = torch.diff(boundaries)
+            # Filter out 0-length blocks (happens when EOB is at the end of the sample)
+            block_lengths = block_lengths[block_lengths > 0]
+            
+            current_sample_block_lengths.append(block_lengths)
+            start_idx = end_idx
+            
+        all_batch_block_lengths.append(current_sample_block_lengths)
+    
+    return all_batch_block_lengths
+
+
+def block_diff_mask_dynamic(b, h, q_idx, kv_idx, block_boundaries=None, n=None):
+    """
+    Dynamic block diffusion mask using precomputed block boundaries.
+    
+    This replaces the fixed block_size arithmetic with torch.searchsorted
+    to support variable-length blocks.
+    
+    Args:
+        b: Batch index (unused in mask logic)
+        h: Head index (unused in mask logic)
+        q_idx: Query indices tensor
+        kv_idx: Key-value indices tensor
+        block_boundaries: Cumulative sum of block lengths, e.g., [0, 4, 12, 16]
+                         This maps: tokens 0-3 → block 0, 4-11 → block 1, 12-15 → block 2
+        n: Number of denoised (clean) tokens
+    
+    Returns:
+        Boolean attention mask (True = can attend)
+        
+    The mask combines three types:
+        - M_BD (Block Diagonal): Self-attention within noised blocks
+        - M_OBC (Offset Block Causal): Cross-attention from noised to conditional context
+        - M_BC (Block Causal): Attention to denoised blocks
+    """
+    # Map indices to block IDs (handling both Noisy 0..n-1 and Clean n..2n-1)
+    # We use modulo n to map Clean tokens back to their relative position
+    q_mod = q_idx % n
+    kv_mod = kv_idx % n
+    
+    # Use searchsorted to find which block each index belongs to
+    # right=True ensures that [0, 4] maps 0,1,2,3 to the first interval
+    # We subtract 1 to get 0-based block indices
+    q_block_id = torch.searchsorted(block_boundaries, q_mod, right=True) - 1
+    kv_block_id = torch.searchsorted(block_boundaries, kv_mod, right=True) - 1
+    
+    # Clamp to handle edge cases
+    q_block_id = torch.clamp(q_block_id, 0, len(block_boundaries) - 2)
+    kv_block_id = torch.clamp(kv_block_id, 0, len(block_boundaries) - 2)
+    
+    # Identify Noisy vs Clean
+    # Noisy: < n (x0_flag = False)
+    # Clean: >= n (x0_flag = True)
+    is_clean_q = q_idx >= n
+    is_clean_kv = kv_idx >= n
+    
+    # **1. Block Diagonal Mask (M_BD) **
+    # Self-attention within blocks (Noisy->Noisy, Clean->Clean)
+    M_BD = (q_block_id == kv_block_id) & (is_clean_q == is_clean_kv)
+
+    # **2. Offset Block-Causal Mask (M_OBC) **
+    # Noisy i attends to Clean j < i
+    # (Original code: block_q > block_kv & clean_kv & noisy_q)
+    M_OBC = (q_block_id > kv_block_id) & (is_clean_kv) & (~is_clean_q)
+
+    # **3. Block-Causal Mask (M_BC) **
+    # Clean i attends to Clean j <= i
+    M_BC = (q_block_id >= kv_block_id) & (is_clean_kv) & (is_clean_q)
+
+    # **4. Combine Masks **
+    return M_BD | M_OBC | M_BC
+
+
+def block_attn_mask_dynamic(
+    nested_block_lengths_list: List[List[torch.Tensor]],
+    device: torch.device
+) -> torch.Tensor:
+    """
+    Construct attention masks for variable-length blocks, handling packed sequences.
+    
+    Args:
+        nested_block_lengths_list: List (batch) of Lists (samples) of Tensors (block lengths).
+        device: Device to create tensors on
+        
+    Returns:
+        Attention mask tensor, shape (batch_size, total_seq_len*2, total_seq_len*2)
+    """
+    masks = []
+    
+    for sample_block_lengths_list in nested_block_lengths_list:
+        sample_masks = []
+        
+        for block_lengths in sample_block_lengths_list:
+            # Calculate total sequence length for this sample
+            total_len = block_lengths.sum().item()
+            if total_len == 0:
+                continue
+                
+            n = total_len  # Number of clean tokens
+            
+            # Create block boundaries (cumulative sum)
+            block_boundaries = torch.cat([
+                torch.tensor([0], device=device),
+                torch.cumsum(block_lengths, dim=0)
+            ])
+            
+            # Create index tensors for the full 2n x 2n mask
+            seq_len_doubled = total_len * 2
+            q_idx = torch.arange(seq_len_doubled, device=device)[:, None]
+            kv_idx = torch.arange(seq_len_doubled, device=device)[None, :]
+            
+            # Generate mask using dynamic block boundaries
+            mask = block_diff_mask_dynamic(
+                b=None,
+                h=None,
+                q_idx=q_idx,
+                kv_idx=kv_idx,
+                block_boundaries=block_boundaries,
+                n=n
+            )
+            
+            sample_masks.append(mask)
+        
+        # Combine sample masks into a single block-diagonal mask for the batch item
+        if sample_masks:
+            row_mask = torch.block_diag(*sample_masks)
+        else:
+            # Should not happen if input is valid
+            row_mask = torch.zeros((0, 0), device=device, dtype=torch.bool)
+            
+        masks.append(row_mask)
+    
+    # Stack into batch
+    # We assume all row_masks have the same size (2 * seq_len)
+    # If not (due to padding differences?), we might need to pad them.
+    # But calculate_token_nums usually covers the whole seq_len including padding.
+    
+    # Check sizes
+    sizes = [m.shape[0] for m in masks]
+    max_size = max(sizes)
+    
+    padded_masks = []
+    for m in masks:
+        if m.shape[0] < max_size:
+            # Pad with False (no attention)
+            pad_size = max_size - m.shape[0]
+            m = torch.nn.functional.pad(m, (0, pad_size, 0, pad_size), value=False)
+        padded_masks.append(m)
+        
+    masks = torch.stack(padded_masks, dim=0)
+    return masks
+

fused_linear_diffusion_cross_entropy.py

@@ -0,0 +1,723 @@
+# -*- coding: utf-8 -*-
+
+# Code adapted from
+# https://github.com/fla-org/flash-linear-attention/blob/main/fla/modules/fused_linear_cross_entropy.py
+# Implementation of element-wise division of cross entropy loss
+
+
+# Code adapted from
+# https://github.com/linkedin/Liger-Kernel/blob/main/src/liger_kernel/ops/fused_linear_cross_entropy.py
+
+from functools import partial
+from typing import Optional, Tuple
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+import triton
+import triton.language as tl
+from torch.distributed import DeviceMesh
+from torch.distributed.tensor import DTensor, Replicate, Shard, distribute_module
+from torch.distributed.tensor.parallel import ParallelStyle
+
+# The hard limit of TRITON_MAX_TENSOR_NUMEL is 1048576
+# https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/language/core.py#L19
+# However, setting limit as 65536 as in LayerNorm tutorial is faster because of less register spilling
+# The optimal maximum block size depends on your hardware, your kernel, and your dtype
+MAX_FUSED_SIZE = 65536 // 2
+
+
+@triton.heuristics({
+    'HAS_SCALE': lambda args: args['scale'] is not None
+})
+@triton.autotune(
+    configs=[
+        triton.Config({}, num_warps=num_warps)
+        for num_warps in [1, 2, 4, 8, 16, 32]
+    ],
+    key=['D']
+)
+@triton.jit
+def logsumexp_fwd_kernel(
+    x,
+    z,
+    scale,
+    D: tl.constexpr,
+    B: tl.constexpr,
+    HAS_SCALE: tl.constexpr
+):
+    i_n, i_d = tl.program_id(0).to(tl.int64), tl.program_id(1).to(tl.int64)
+    o_d = i_d * B + tl.arange(0, B)
+    m_d = o_d < D
+
+    b_x = tl.load(x + i_n * D + o_d, mask=m_d, other=-float('inf'))
+    if HAS_SCALE:
+        b_x = b_x * scale
+    b_m = tl.max(b_x, 0)
+    b_z = tl.log(tl.sum(tl.exp(b_x - b_m), 0)) + b_m
+    tl.store(z + i_n * tl.cdiv(D, B) + i_d, b_z)
+
+
+def logsumexp_fwd(
+    x,
+    scale: Optional[float] = None,
+    dtype: Optional[torch.dtype] = None
+):
+    r"""
+    Compute the logsumexp of the input tensor over the last dimension.
+
+    Args:
+        x (Tensor):
+            The input tensor of any shape.
+        scale (Optional[float]):
+            The scale applied to the input tensor. Default: `None`.
+        dtype (Optional[torch.dtype]):
+            The data type of the output tensor. Default: `None`.
+    Returns:
+        Tensor: The logsumexp of the input tensor.
+    """
+
+    shape = x.shape
+    x = x.view(-1, shape[-1])
+    N, D = x.shape
+    B = min(triton.next_power_of_2(D), 64 * 1024)
+    ND = triton.cdiv(D, B)
+
+    z = x.new_empty(N, ND, dtype=torch.float)
+    logsumexp_fwd_kernel[(N, ND)](
+        x=x,
+        z=z,
+        scale=scale,
+        D=D,
+        B=B
+    )
+    z = z.logsumexp(-1).view(*shape[:-1])
+    if dtype is not None and dtype != torch.float:
+        z = z.to(dtype)
+    return z
+
+@triton.jit
+def cross_entropy_kernel(
+    logits,
+    lse,
+    target,
+    p_mask,
+    loss,
+    total,
+    ignore_index,
+    label_smoothing: tl.constexpr,
+    logit_scale: tl.constexpr,
+    reduction: tl.constexpr,
+    V: tl.constexpr,
+    BV: tl.constexpr
+):
+    """
+    This kernel computes both cross entropy loss and the gradient of the input.
+    We only consider hard label + mean reduction for now.
+    Please refer to https://pytorch.org/docs/stable/generated/torch.nn.CrossEntropyLoss.html for the math.
+
+    Args:
+        logits:
+            Pointer to logits tensor.
+        lse:
+            Pointer to logsumexp tensor.
+        target: Pointer to target tensor.
+        loss:
+            Pointer to tensor to store the loss.
+        V (int):
+            The number of columns in the input tensor.
+        total (int):
+            The number of non-ignored classes.
+        ignore_index (int):
+            The index to ignore in the target.
+        label_smoothing (float):
+            The amount of smoothing when computing the loss, where 0.0 means no smoothing.
+        reduction (str):
+            The string for the reduction to apply
+        BV (int):
+            The block size for vocab.
+    """
+
+    # https://github.com/triton-lang/triton/issues/1058
+    # If B*T*V is too large, i_n * stride will overflow out of int32, so we convert to int64
+    i_n = tl.program_id(0).to(tl.int64)
+    NV = tl.cdiv(V, BV)
+
+    # 1. Load target first because if the target is ignore_index, we can return right away
+    b_y = tl.load(target + i_n)
+    # load p_mask
+    b_p_mask = tl.load(p_mask + i_n)
+
+    # 2. locate the start index
+    logits += i_n * V
+
+    if b_y == ignore_index:
+        # set all x as 0
+        for i in range(0, V, BV):
+            o_v = i + tl.arange(0, BV)
+            tl.store(logits + o_v, 0.0, mask=o_v < V)
+        return
+
+    # Online softmax: 2 loads + 1 store (compared with 3 loads + 1 store for the safe softmax)
+    # Refer to Algorithm 3 in the paper: https://arxiv.org/pdf/1805.02867
+
+    # 3. [Online softmax] first pass: compute logsumexp
+    # we did this in anouter kernel
+    b_l = tl.load(logits + b_y) * logit_scale
+    b_lse = tl.load(lse + i_n)
+
+    # 4. Calculate the loss
+    # loss = lse - logits_l
+    # celoss = -log(q_y) = -log(softmax(x_y))
+    b_loss = (b_lse - b_l) / b_p_mask  # Diffusion Scaled '1/t'
+
+    # Label smoothing is a general case of normal cross entropy
+    # See the full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issue-2503665310
+    b_z = 0.0
+    eps = label_smoothing / V
+
+    # We need tl.debug_barrier() as mentioned in
+    # https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/ops/cross_entropy.py#L34
+    tl.debug_barrier()
+
+    # 5. [Online Softmax] Second pass: compute gradients
+    # For 'mean' reduction, gradients are normalized by number of non-ignored elements
+    # dx_y = (softmax(x_y) - 1) / N
+    # dx_i = softmax(x_i) / N, i != y
+    # For label smoothing:
+    # dx_i = (softmax(x_y) - label_smoothing / V) / N, i != y
+    # dx_y = (softmax(x_y) - label_smoothing / V - (1 - label_smoothing)) / N
+    #      = dx_i - (1 - label_smoothing) / N
+    for iv in range(0, NV):
+        o_v = iv * BV + tl.arange(0, BV)
+        b_logits = tl.load(logits + o_v, mask=o_v < V, other=float('-inf')) * logit_scale
+        if label_smoothing > 0:
+            # scale X beforehand to avoid overflow
+            b_z += tl.sum(tl.where(o_v < V, -eps * b_logits, 0.0))
+        b_p = (tl.exp(b_logits - b_lse) - eps) * logit_scale
+        b_p /= b_p_mask # 修改
+        if reduction == "mean":
+            b_p = b_p / total
+        tl.store(logits + o_v, b_p, mask=o_v < V)
+
+        tl.debug_barrier()
+
+    # Orginal loss = H(q, p),  with label smoothing regularization = H(q', p) and (label_smoothing / V) = eps
+    # H(q', p) = (1 - label_smoothing) * H(q, p) + label_smoothing * H(u, p)
+    #          = (1 - label_smoothing) * H(q, p) + eps * sum(logsoftmax(x_i))
+    # By using m (global max of xi) and d (sum of e^(xi-m)), we can simplify as:
+    #          = (1 - label_smoothing) * H(q, p) + (-sum(x_i * eps) + label_smoothing * (m + logd))
+    # Refer to H(q', p) in section 7 of the paper:
+    # https://arxiv.org/pdf/1512.00567
+    # pytorch:
+    # https://github.com/pytorch/pytorch/blob/2981534f54d49fa3a9755c9b0855e7929c2527f0/aten/src/ATen/native/LossNLL.cpp#L516
+    # See full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issuecomment-2333753087
+    if label_smoothing > 0:
+        b_loss = b_loss * (1 - label_smoothing) + (b_z + label_smoothing * b_lse)
+
+    # 6. Specially handle the i==y case where `dx_y = (softmax(x_y) - (1 - label_smoothing) / N`
+    b_l = tl.load(logits + b_y)
+
+    # Normalize the loss by the number of non-ignored elements if reduction is "mean"
+    if reduction == 'mean':
+        b_loss = b_loss / total
+        # b_l += (label_smoothing - 1) / total * logit_scale
+        # b_l has already been divided by b_p_mask and total
+        b_l += (label_smoothing - 1) / b_p_mask / total * logit_scale
+    else:
+        # b_l += (label_smoothing - 1) * logit_scale
+        b_l += (label_smoothing - 1) / b_p_mask * logit_scale
+
+    tl.store(loss + i_n, b_loss)
+    tl.store(logits + b_y, b_l)
+
+
+@triton.jit
+def elementwise_mul_kernel(
+    x,
+    g,
+    N: tl.constexpr,
+    B: tl.constexpr
+):
+    """
+    This function multiplies each element of the tensor pointed by x with the value pointed by g.
+    The multiplication is performed in-place on the tensor pointed by x.
+
+    Parameters:
+    x:
+        Pointer to the input tensor.
+    g:
+        Pointer to the gradient output value.
+    N (int):
+        The number of columns in the input tensor.
+    B (int):
+        The block size for Triton operations.
+    """
+
+    # Get the program ID and convert it to int64 to avoid overflow
+    i_x = tl.program_id(0).to(tl.int64)
+    o_x = i_x * B + tl.arange(0, B)
+
+    # Load the gradient output value
+    b_g = tl.load(g)
+    b_x = tl.load(x + o_x, mask=o_x < N)
+    tl.store(x + o_x, b_x * b_g, mask=o_x < N)
+
+
+def fused_linear_cross_entropy_forward(
+    x: torch.Tensor,
+    target: torch.LongTensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor = None,
+    p_mask: torch.Tensor = None,
+    ignore_index: int = -100,
+    label_smoothing: float = 0.0,
+    logit_scale: float = 1.0,
+    num_chunks: int = 8,
+    reduction: str = "mean"
+):
+    device = x.device
+    # inputs have shape: [N, H]
+    # materialized activations will have shape: [N, V]
+    # the increase in memory = [N, V]
+    # reduction can be achieved by partitioning the number of tokens N into smaller chunks.
+
+    # ideally, we would like to achieve the same memory consumption as [N, H],
+    # so the expected chunk size should be:
+    # NC = ceil(V / H)
+    # C = ceil(N / NC)
+    # for ex: N = 4096*4, V = 32000, H = 4096 ==> NC = 8, C = ceil(N / NC) = 2048
+    N, H, V = *x.shape, weight.shape[0]
+    BV = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
+    # TODO: in real cases, we may need to limit the number of chunks NC to
+    # ensure the precisions of accumulated gradients
+    NC = min(num_chunks, triton.cdiv(V, H))
+    C = triton.next_power_of_2(triton.cdiv(N, NC))
+    NC = triton.cdiv(N, C)
+
+    # [N, H]
+    dx = torch.zeros_like(x, device=device)
+    # [V, H]
+    dw = torch.zeros_like(weight, device=device, dtype=torch.float) if weight is not None else None
+    # [V]
+    db = torch.zeros_like(bias, device=device, dtype=torch.float) if bias is not None else None
+    # [N]
+    loss = torch.zeros(N, device=device, dtype=torch.float)
+
+    total = target.ne(ignore_index).sum().item()
+
+    for ic in range(NC):
+        start, end = ic * C, min((ic + 1) * C, N)
+        # [C, N]
+        c_x = x[start:end]
+        # when doing matmul, use the original precision
+        # [C, V]
+        c_logits = F.linear(c_x, weight, bias)
+        c_target = target[start:end]
+        c_p_mask = p_mask[start:end]
+        # [C]
+        # keep lse in fp32 to maintain precision
+        c_lse = logsumexp_fwd(c_logits, scale=logit_scale, dtype=torch.float)
+
+        # unreduced loss
+        c_loss = loss[start:end]
+
+        # Here we calculate the gradient of c_logits in place so we can save memory.
+        cross_entropy_kernel[(c_logits.shape[0],)](
+            logits=c_logits,
+            lse=c_lse,
+            target=c_target,
+            p_mask=c_p_mask,
+            loss=c_loss,
+            total=total,
+            ignore_index=ignore_index,
+            label_smoothing=label_smoothing,
+            logit_scale=logit_scale,
+            reduction=reduction,
+            V=V,
+            BV=BV,
+            num_warps=32
+        )
+
+        # gradient of logits is computed in-place by the above triton kernel and is of shape: C x V
+        # thus dx should be of shape: C x H
+        dx[start:end] = torch.mm(c_logits, weight)
+
+        # keep dw in fp32 to maintain precision
+        if weight is not None:
+            dw += c_logits.t() @ c_x
+
+        if bias is not None:
+            torch.add(input=db, other=c_logits.sum(0), out=db)
+
+    loss = loss.sum()
+    if dw is not None:
+        dw = dw.to(weight)
+    if db is not None:
+        db = db.to(bias)
+    return loss, dx, dw, db
+
+
+def fused_linear_cross_entropy_backward(
+    do: torch.Tensor,
+    dx: torch.Tensor,
+    dw: torch.Tensor,
+    db: torch.Tensor
+):
+    # If cross entropy is the last layer, do is 1.0. Skip the mul to save time
+    if torch.ne(do, torch.tensor(1.0, device=do.device)):
+        # We use a Triton kernel instead of a PyTorch operation because modifying inputs in-place
+        # for gradient storage and backward multiple times causes anomalies with PyTorch but not with Triton.
+        N, H = dx.shape
+        B = min(MAX_FUSED_SIZE, triton.next_power_of_2(H))
+
+        elementwise_mul_kernel[(triton.cdiv(N * H, B),)](
+            x=dx,
+            g=do,
+            N=N*H,
+            B=B,
+            num_warps=32,
+        )
+
+        # handle dw
+        if dw is not None:
+            V, H = dw.shape
+            elementwise_mul_kernel[(triton.cdiv(V * H, B),)](
+                x=dw,
+                g=do,
+                N=V*H,
+                B=B,
+                num_warps=32,
+            )
+
+        if db is not None:
+            V = db.shape[0]
+            elementwise_mul_kernel[(triton.cdiv(V, B),)](
+                x=db,
+                g=do,
+                N=V,
+                B=B,
+                num_warps=32,
+            )
+    return dx, dw, db
+
+
+class FusedLinearCrossEntropyFunction(torch.autograd.Function):
+
+    @staticmethod
+    def forward(
+        ctx,
+        x: torch.Tensor,
+        target: torch.LongTensor,
+        weight: torch.Tensor,
+        bias: torch.Tensor = None,
+        p_mask: torch.Tensor = None,
+        ignore_index: int = -100,
+        label_smoothing: float = 0.0,
+        logit_scale: float = 1.0,
+        num_chunks: int = 8,
+        reduction: str = "mean"
+    ):
+        """
+        Fusing the last linear layer with cross-entropy loss
+            Reference: https://github.com/mgmalek/efficient_cross_entropy
+
+        Handle the forward and backward pass of the final linear layer via cross-entropy loss by avoiding
+        the materialization of the large logits tensor. Since Cross Entropy Loss is the last layer, we can
+        compute the gradient at the forward pass. By doing so, we don't have to store the x and target
+        for the backward pass.
+
+        x (torch.Tensor): [batch_size * seq_len, hidden_size]
+        target (torch.LongTensor): [batch_size * seq_len]
+            where each value is in [0, vocab_size).
+        weight (torch.Tensor): [vocab_size, hidden_size]
+            where `vocab_size` is the number of classes.
+        bias (Optional[torch.Tensor]): [vocab_size]
+            where `vocab_size` is the number of classes.
+        p_mask(torch.Tensor): [batch_size * seq_len]
+            Its shape should be same as target. 
+        ignore_index:
+            the index to ignore in the target.
+        label_smoothing:
+            the amount of smoothing when computing the loss, where 0.0 means no smoothing.
+        logit_scale: float = 1.0,
+            A scaling factor applied to the logits. Default: 1.0
+        num_chunks: int
+            The number of chunks to split the input tensor into for processing.
+            This can help optimize memory usage and computation speed.
+            Default: 8
+        reduction:
+            Specifies the reduction to apply to the output: 'mean' | 'sum'.
+            'mean': the weighted mean of the output is taken,
+            'sum': the output will be summed.
+            Default: 'mean'.
+        """
+        loss, dx, dw, db = fused_linear_cross_entropy_forward(
+            x,
+            target,
+            weight,
+            bias,
+            p_mask,
+            ignore_index,
+            label_smoothing,
+            logit_scale,
+            num_chunks,
+            reduction
+        )
+        # downcast to dtype and store for backward
+        ctx.save_for_backward(
+            dx.detach(),
+            dw.detach() if weight is not None else None,
+            db.detach() if bias is not None else None,
+        )
+        return loss
+
+    @staticmethod
+    def backward(ctx, do):
+        dx, dw, db = ctx.saved_tensors
+        dx, dw, db = fused_linear_cross_entropy_backward(do, dx, dw, db)
+        # 10 gradients should be returned, with `p_mask` having no grads
+        # Check the number of arguments in the `forward` method
+        return dx, None, dw, db, None, None, None, None, None, None
+
+
+def fused_linear_cross_entropy_loss(
+    x: torch.Tensor,
+    target: torch.LongTensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor = None,
+    p_mask: torch.Tensor = None,
+    ignore_index: int = -100,
+    label_smoothing: float = 0.0,
+    logit_scale: float = 1.0,
+    num_chunks: int = 8,
+    reduction: str = "mean"
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    """
+    Args:
+        x (torch.Tensor): [batch_size * seq_len, hidden_size]
+        target (torch.LongTensor): [batch_size * seq_len]
+            where each value is in [0, vocab_size).
+        weight (torch.Tensor): [vocab_size, hidden_size]
+            where `vocab_size` is the number of classes.
+        bias (Optional[torch.Tensor]): [vocab_size]
+            where `vocab_size` is the number of classes.
+        p_mask(torch.Tensor): [batch_size * seq_len]
+            Its shape should be same as target. 
+        ignore_index: int.
+            If target == ignore_index, the loss is set to 0.0.
+        label_smoothing: float
+        logit_scale: float
+            A scaling factor applied to the logits. Default: 1.0
+        num_chunks: int
+            The number of chunks to split the input tensor into for processing.
+            This can help optimize memory usage and computation speed.
+            Default: 8
+        reduction:
+            Specifies the reduction to apply to the output: 'mean' | 'sum'.
+            'mean': the weighted mean of the output is taken,
+            'sum': the output will be summed.
+            Default: 'mean'.
+    Returns:
+        losses: [batch,], float
+    """
+    return FusedLinearCrossEntropyFunction.apply(
+        x,
+        target,
+        weight,
+        bias,
+        p_mask,
+        ignore_index,
+        label_smoothing,
+        logit_scale,
+        num_chunks,
+        reduction
+    )
+
+
+class FusedLinearDiffusionCrossEntropyLoss(nn.Module):
+
+    def __init__(
+        self,
+        ignore_index: int = -100,
+        label_smoothing: float = 0.0,
+        logit_scale: float = 1.0,
+        num_chunks: int = 8,
+        reduction: str = "mean"
+    ):
+        """
+        Args:
+            ignore_index: int.
+                If target == ignore_index, the loss is set to 0.0.
+            label_smoothing: float
+            logit_scale: float
+                A scaling factor applied to the logits. Default: 1.0
+            num_chunks: int
+                The number of chunks to split the input tensor into for processing.
+                This can help optimize memory usage and computation speed.
+                Default: 8
+            reduction:
+                Specifies the reduction to apply to the output: 'mean' | 'sum'.
+                'mean': the weighted mean of the output is taken,
+                'sum': the output will be summed.
+                Default: 'mean'.
+        """
+        super().__init__()
+
+        assert reduction in ["mean", "sum"], f"reduction: {reduction} is not supported"
+
+        self.ignore_index = ignore_index
+        self.label_smoothing = label_smoothing
+        self.logit_scale = logit_scale
+        self.num_chunks = num_chunks
+        self.reduction = reduction
+
+    @torch.compiler.disable
+    def forward(
+        self,
+        x: torch.Tensor,
+        target: torch.LongTensor,
+        weight: torch.Tensor,
+        bias: Optional[torch.Tensor] = None,
+        p_mask: torch.Tensor = None,
+        eob_token_id: Optional[int] = None,
+        eob_weight: float = 1.0
+    ):
+        """
+        Args:
+            x (torch.Tensor): [batch_size, seq_len, hidden_size]
+            target (torch.LongTensor): [batch_size, seq_len]
+                where each value is in [0, V).
+            weight (torch.Tensor): [vocab_size, hidden_size]
+                where `vocab_size` is the number of classes.
+            bias (Optional[torch.Tensor]): [vocab_size]
+                where `vocab_size` is the number of classes.
+            p_mask(torch.Tensor): [batch_size, seq_len]
+                Its shape is same as target. 
+                Shape: (1, packed_length) when varlen attn is used.
+        Returns:
+            loss
+            
+        TODO:
+            follow https://github.com/ML-GSAI/LLaDA/blob/main/GUIDELINES.md#pre-training
+            ```py
+            unreduced_loss /= p_mask
+            ```
+            Scale the values of `unreduced_loss at different positions
+        """
+        if p_mask is None:
+            p_mask = torch.ones_like(target, dtype=torch.float, device=x.device)
+        
+        # Apply EOB weight if provided
+        if eob_token_id is not None and eob_weight != 1.0:
+            # We want to multiply the loss of EOB tokens by eob_weight.
+            # The kernel calculates loss = unreduced_loss / p_mask.
+            # So to get loss * eob_weight, we need to divide p_mask by eob_weight.
+            # p_mask_new = p_mask / eob_weight
+            # loss_new = unreduced_loss / (p_mask / eob_weight) = (unreduced_loss / p_mask) * eob_weight
+            
+            # Create a mask for EOB tokens
+            is_eob = (target == eob_token_id)
+            if is_eob.any():
+                # We need to modify p_mask. Since p_mask might be reused or is a view, let's clone it if needed.
+                # However, p_mask is usually created fresh in forward_add_noise_packed.
+                # But to be safe and avoid side effects if it's used elsewhere (unlikely), we can modify in place if it's not a leaf.
+                # p_mask is likely a tensor from the graph.
+                
+                # Let's modify it. Note: p_mask shape matches target shape here.
+                # We use a float mask to avoid in-place modification issues if possible, or just modify.
+                # p_mask = p_mask.clone() # Safer
+                
+                # Actually, we can just do:
+                # p_mask[is_eob] = p_mask[is_eob] / eob_weight
+                # But p_mask might be on a different device or flattened?
+                # target is flattened below. Let's do it before flattening or after?
+                # The code flattens target and p_mask below.
+                pass
+
+        x = x.contiguous().view(-1, x.shape[-1])
+        target = target.contiguous().view(-1)
+        weight = weight.contiguous()
+        bias = bias.contiguous() if bias else None
+        p_mask = p_mask.contiguous().view(-1)
+        
+        # Apply EOB weight (after flattening to be safe and consistent)
+        if eob_token_id is not None and eob_weight != 1.0:
+            is_eob = (target == eob_token_id)
+            if is_eob.any():
+                # We divide p_mask by eob_weight to effectively multiply loss by eob_weight
+                # We use a multiplier tensor to avoid in-place ops on p_mask if it causes issues, 
+                # but modifying p_mask is the most direct way for the kernel.
+                # We need to ensure p_mask is floating point.
+                p_mask = p_mask.clone() # Clone to avoid modifying the input tensor
+                p_mask[is_eob] = p_mask[is_eob] / eob_weight
+
+        l, d = x.shape
+        assert l == target.shape[0] == p_mask.shape[0], f"{x.shape=}, {target.shape=}, {p_mask.shape=}"
+            
+        loss = fused_linear_cross_entropy_loss(
+            x,
+            target,
+            weight=weight,
+            bias=bias,
+            p_mask=p_mask,
+            ignore_index=self.ignore_index,
+            label_smoothing=self.label_smoothing,
+            logit_scale=self.logit_scale,
+            num_chunks=self.num_chunks,
+            reduction=self.reduction
+        )
+        return loss
+
+
+class LinearLossParallel(ParallelStyle):
+    def __init__(
+        self,
+        *,
+        sequence_dim: int = 1,
+        use_local_output: bool = False,
+    ):
+        super().__init__()
+
+        self.sequence_sharding = (Shard(sequence_dim),)
+        self.use_local_output = use_local_output
+
+    @staticmethod
+    def _prepare_input_fn(sequence_sharding, mod, inputs, device_mesh):
+        x, target, weight, bias = inputs
+
+        if not isinstance(x, DTensor):
+            # assume the input passed in already sharded on the sequence dim and create the DTensor
+            x = DTensor.from_local(x, device_mesh, sequence_sharding)
+        if x.placements != sequence_sharding:
+            x = x.redistribute(placements=sequence_sharding, async_op=True)
+        if not isinstance(target, DTensor):
+            target = DTensor.from_local(target, device_mesh, [Replicate()])
+        if target.placements != sequence_sharding:
+            target = target.redistribute(placements=sequence_sharding, async_op=True)
+
+        if not isinstance(weight, DTensor):
+            weight = DTensor.from_local(weight, device_mesh, [Replicate()])
+        if weight.placements != [Replicate()]:
+            # we replicate the weight/bias in FLCE
+            weight = weight.redistribute(placements=[Replicate()], async_op=True)
+
+        if bias is not None and not isinstance(bias, DTensor):
+            bias = DTensor.from_local(bias, device_mesh, [Replicate()])
+        if bias is not None and bias.placements != [Replicate()]:
+            bias = bias.redistribute(placements=[Replicate()], async_op=True)
+
+        return x.to_local(), target.to_local(), weight.to_local(), bias.to_local() if bias is not None else bias
+
+    @staticmethod
+    def _prepare_output_fn(use_local_output, mod, outputs, device_mesh):
+        return outputs.to_local() if use_local_output else outputs
+
+    def _apply(self, module: nn.Module, device_mesh: DeviceMesh) -> nn.Module:
+        return distribute_module(
+            module,
+            device_mesh,
+            partition_fn=None,
+            input_fn=partial(self._prepare_input_fn, self.sequence_sharding),
+            output_fn=partial(self._prepare_output_fn, self.use_local_output)
+        )

generation_config.json

@@ -0,0 +1,13 @@
+{
+  "bos_token_id": 151643,
+  "do_sample": true,
+  "eos_token_id": [
+    151645,
+    151643
+  ],
+  "pad_token_id": 151643,
+  "temperature": 0.6,
+  "top_k": 20,
+  "top_p": 0.95,
+  "transformers_version": "4.52.4"
+}

latest

@@ -0,0 +1 @@
+global_step1549

model-00001-of-00002.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:785fe86ef4cad56e5e4c4fdb04046db53e3f94e3c0eabf07fa63c8bb8b4bc6c9
+size 4967215360

model-00002-of-00002.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:929f19e3802619c326f0d972f3824328bea9b978ca8918f8b86fcd9799e8d62d
+size 3855679144

model.safetensors.index.json

@@ -0,0 +1,406 @@
+{
+  "metadata": {
+    "total_size": 8822848512
+  },
+  "weight_map": {
+    "lm_head.weight": "model-00002-of-00002.safetensors",
+    "model.embed_tokens.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.16.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.17.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.19.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.20.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.20.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.20.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.21.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.21.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.22.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.25.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.27.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.3.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.30.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.32.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.k_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.q_norm.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.4.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.5.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.k_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.q_norm.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "model.norm.weight": "model-00002-of-00002.safetensors"
+  }
+}

modeling_sdar.py

@@ -0,0 +1,1245 @@
+# This file is modified based on https://github.com/huggingface/transformers/blob/v4.52.4/src/transformers/models/qwen3/modeling_qwen3.py.
+#
+#                🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨
+#           This file was automatically generated from src/transformers/models/qwen3/modular_qwen3.py.
+#               Do NOT edit this file manually as any edits will be overwritten by the generation of
+#             the file from the modular. If any change should be done, please apply the change to the
+#                          modular_qwen3.py file directly. One of our CI enforces this.
+#                🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨🚨
+# coding=utf-8
+# Copyright 2025 The Qwen team, Alibaba Group and the HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+from typing import Optional, Tuple, Union, List
+
+import torch
+from torch import nn
+from einops import rearrange
+
+from transformers.activations import ACT2FN
+from transformers.cache_utils import Cache, DynamicCache, SlidingWindowCache, StaticCache
+from transformers.generation import GenerationMixin
+from transformers.integrations import use_kernel_forward_from_hub
+from transformers.modeling_attn_mask_utils import AttentionMaskConverter
+from transformers.modeling_flash_attention_utils import FlashAttentionKwargs
+from transformers.modeling_layers import GradientCheckpointingLayer
+from transformers.modeling_outputs import (
+    BaseModelOutputWithPast,
+    CausalLMOutputWithPast,
+)
+from transformers.modeling_rope_utils import ROPE_INIT_FUNCTIONS, dynamic_rope_update
+from transformers.modeling_utils import  PreTrainedModel
+from transformers.processing_utils import Unpack
+from transformers.utils import LossKwargs, auto_docstring, can_return_tuple, is_torch_flex_attn_available, logging
+from .configuration_sdar import SDARConfig
+from .fused_linear_diffusion_cross_entropy import FusedLinearDiffusionCrossEntropyLoss
+from .dynamic_blocks_utils import calculate_block_nums_from_eob, block_attn_mask_dynamic
+
+from flash_attn.ops.triton.layer_norm import rms_norm_fn as flash_rms_norm
+
+import torch.nn.functional as F
+try:
+    from flash_attn import flash_attn_func, flash_attn_varlen_func
+    from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input
+except:
+    pass
+
+try:
+    from liger_kernel.ops.swiglu import LigerSiLUMulFunction  # noqa: F401
+    liger_kernel_is_available = True
+except ImportError:
+    liger_kernel_is_available = False
+
+
+if is_torch_flex_attn_available():
+    from torch.nn.attention.flex_attention import BlockMask, create_block_mask, flex_attention
+    from transformers.integrations.flex_attention import make_flex_block_causal_mask
+
+
+logger = logging.get_logger(__name__)
+
+
+def modify_padded_position_ids_2d(position_ids: torch.LongTensor) -> torch.LongTensor:
+    """
+    This function uses fully vectorized PyTorch operations to modify the packed position_ids of a batch.
+    It assumes that the input is a 2D Tensor, shape (batch_size, sequence_length).
+    It will independently process each row in the batch.
+
+    Args:
+        position_ids: a 2D Tensor, shape (batch_size, sequence_length).
+
+    Returns:
+        the modified position_ids Tensor, shape (batch_size, sequence_length).
+    """
+    if position_ids.dim() != 2:
+        raise ValueError(f"Input tensor must be 2D, but got {position_ids.dim()} dimensions.")
+        
+    batch_size, seq_len = position_ids.shape
+    device = position_ids.device
+
+    col_indices = torch.arange(seq_len, device=device, dtype=position_ids.dtype).expand(batch_size, -1)
+    mask = (position_ids != 0)
+
+    masked_indices = col_indices * mask
+    last_nonzero_idx = torch.max(masked_indices, dim=1).values
+    has_nonzero = torch.any(mask, dim=1)
+    pad_start_idx = torch.where(has_nonzero, last_nonzero_idx + 1, torch.tensor(0, device=device, dtype=position_ids.dtype))
+
+    padding_mask = col_indices >= pad_start_idx.unsqueeze(1)
+    new_pad_values = col_indices - pad_start_idx.unsqueeze(1)
+    position_ids = torch.where(padding_mask, new_pad_values, position_ids)
+
+    return position_ids
+
+
+def calculate_token_nums(position_ids: torch.Tensor):
+    """
+    This function uses PyTorch to efficiently calculate the length of each packed sequence in a batch.
+
+    Args:
+        position_ids (torch.Tensor): a 2D Tensor, shape (batch_size, sequence_length).
+                                      For example: tensor([[0,1,2,3,4,0,1,2,3,4,5,0,1,2,3,0,0,0]])
+    Returns:
+        list[list[int]]: a nested list, containing the length of each sequence in each batch item.
+                          For example: [[5, 6, 4, 1, 1, 1]]
+    """
+    if position_ids.dim() != 2:
+        raise ValueError(f"The input must be a 2D Tensor, but got {position_ids.dim()}D")
+
+    all_lengths = []
+    
+    # we process the batch by batch item by batch item. Because the number of sequence lengths in each row is different (ragged),
+    # so loop is the most efficient and clear  
+    # the op in loop is fully vectorize
+    for pids_row in position_ids:
+        # get the total length of the current row
+        seq_len = pids_row.shape[0]
+        
+        # 1. find the indices of all elements that are equal to 0   
+        # pids_row == 0  Tensor: [True, False, ..., True, ...]
+        # torch.nonzero will return index of these zero
+        # .flatten() will change the shape from (N, 1) to (N,)
+        zero_indices = torch.nonzero(pids_row == 0).flatten()
+        
+
+        # it is very important for calculate last seq length 
+        # note : same device (cpu/cuda)
+        split_points = torch.cat([
+            zero_indices,
+            torch.tensor([seq_len], device=pids_row.device, dtype=zero_indices.dtype)
+        ])
+        
+        # 3. compute difference , get length
+        # torch.diff([a, b, c, d]) will return  [b-a, c-b, d-c]
+        lengths = torch.diff(split_points)
+
+        all_lengths.append(lengths)
+
+    return all_lengths
+
+
+def forward_add_noise_packed(
+    inputs_ids: torch.Tensor,
+    num_tokens_list: List[torch.Tensor],
+    prompt_mask: torch.Tensor,
+    mask_id: int,
+    eob_token_id: Optional[int] = None,
+    eps: float = 1e-3,
+    max_tries: int = 10,
+) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+    """
+    This function adds noise to the token IDs of a batch of packed sequences.
+
+    This function keeps the logic of generating independent random noise rates for each logical sample (concatenated within each batch item).
+    It will randomly replace some token IDs with mask_id.
+    This process will avoid the positions marked by prompt_mask.
+
+    Args:
+        inputs_ids (torch.Tensor): 
+            the token ID tensor, shape (bsz, total_tokens), the token IDs of the packed sequences.
+        num_tokens_list (List[torch.Tensor]): 
+            a list of tensors, length is bsz. Each tensor records the length of each logical sample in the corresponding batch item. For example: [tensor([len1, len2]), tensor([len3, len4, len5])].
+        prompt_mask (torch.Tensor): 
+            a boolean tensor, shape (bsz, total_tokens), True positions represent prompt, should not add noise.
+        mask_id (int): 
+            the ID of the mask token to replace.
+        eob_token_id (int, optional):
+            the ID of the EOB token. If provided, EOB tokens will ALWAYS be masked.
+        eps (float): 
+            a small value, used to prevent the noise rate t from being exactly 0, ensure p_mask > 0.
+        max_tries (int): 
+            the maximum number of attempts to ensure at least one non-prompt token is masked for each batch item.
+
+    Returns:
+        Tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
+        - noisy_input_ids (torch.Tensor): 
+            the token ID tensor, shape (bsz, total_tokens).
+        - final_masked_indices (torch.Tensor): 
+            a boolean tensor, shape (bsz, total_tokens), True positions represent the positions that are actually masked.
+        - p_masks (torch.Tensor): 
+            a one-dimensional tensor, containing the actual noise rates of the tokens that are masked.
+    """
+    # 1. validate and get shape
+    bsz, total_tokens = inputs_ids.shape
+    device = inputs_ids.device
+
+    # validate the consistency of the input
+    assert len(num_tokens_list) == bsz, f"num_tokens_list 的长度 ({len(num_tokens_list)}) 必须等于 bsz ({bsz})"
+    assert prompt_mask.shape == (bsz, total_tokens), f"prompt_mask 形状不匹配, 期望 {(bsz, total_tokens)}, 得到 {prompt_mask.shape}"
+
+    noisy_ids_list = []
+    final_masked_indices_list = []
+    p_masks_per_token_list = []
+
+    # 2.  iterate  with loop. it is efficient because length is different
+    for i in range(bsz):
+        # get the data of the current batch item
+        current_ids = inputs_ids[i:i+1] # shape: (1, total_tokens)
+        current_num_tokens = num_tokens_list[i]
+        current_prompt_mask = prompt_mask[i:i+1] # shape: (1, total_tokens)
+        
+        num_samples_in_item = len(current_num_tokens)
+        # validate the consistency of the token number in the current batch item
+        assert total_tokens == torch.sum(current_num_tokens), \
+            f"the sum of num_tokens in batch item {i} ({torch.sum(current_num_tokens)}) does not match total_tokens ({total_tokens})"
+
+        eligible_for_masking = ~current_prompt_mask
+
+        # if no token can be masked, use the original input and set p_mask to eps
+        if not eligible_for_masking.any():
+            noisy_ids_list.append(current_ids)
+            final_masked_indices_list.append(torch.zeros_like(current_prompt_mask, dtype=torch.bool))
+            # the shape of p_mask_per_token should be (1, total_tokens) for subsequent concatenation
+            p_masks_per_token_list.append(torch.full((1, total_tokens), eps, device=device, dtype=torch.float))
+            continue
+
+        # --- try to generate mask, ensure at least one token is masked ---
+        final_masked_indices_item = torch.zeros_like(current_prompt_mask, dtype=torch.bool)
+        p_mask_per_token = None
+        
+        for _ in range(max_tries):
+            # generate a independent noise rate t for each logical sample
+            t = torch.rand(num_samples_in_item, device=device)
+            p_mask_per_sample = (1 - eps) * t + eps
+
+            # extend the noise rate of each sample to all tokens
+            p_mask_per_token_1d = torch.repeat_interleave(p_mask_per_sample, current_num_tokens)
+            p_mask_per_token = p_mask_per_token_1d.unsqueeze(0) # shape: (1, total_tokens)
+
+            # generate random mask based on the noise rate
+            masked_indices = torch.rand_like(p_mask_per_token) < p_mask_per_token
+            
+            # Note: We do NOT force EOB tokens to always be masked.
+            # The eob_weight parameter in the loss function (default 0.1) handles 
+            # reducing the loss contribution of EOB tokens.
+            # Allowing probabilistic masking lets the model learn natural EOB patterns.
+            
+            #  apply  prompt mask，ensure prompt is not  mask
+            final_masked_indices_item = masked_indices & eligible_for_masking
+
+            # if at least one token is masked, break the loop   
+            if final_masked_indices_item.any():
+                break
+        
+        # if  max_tries , still not mask any token ( very low propobility)，force mask one token
+        if not final_masked_indices_item.any():
+            eligible_indices = torch.nonzero(eligible_for_masking.squeeze(0), as_tuple=True)[0]
+            if len(eligible_indices) > 0:
+                # random choose one to mask
+                random_choice = torch.randint(0, len(eligible_indices), (1,)).item()
+                force_mask_idx = eligible_indices[random_choice]
+                final_masked_indices_item[0, force_mask_idx] = True
+
+
+        # generate noisy IDs based on the final mask
+        noisy_ids_item = torch.where(
+            final_masked_indices_item,
+            mask_id,
+            current_ids
+        )
+        
+        # save the result of the current batch item
+        noisy_ids_list.append(noisy_ids_item)
+        final_masked_indices_list.append(final_masked_indices_item)
+        p_masks_per_token_list.append(p_mask_per_token)
+
+    # 3. stack the results in the list into the final batch tensor
+    noisy_input_ids = torch.cat(noisy_ids_list, dim=0)
+    final_masked_indices = torch.cat(final_masked_indices_list, dim=0)
+    p_mask_full = torch.cat(p_masks_per_token_list, dim=0)
+    
+    # 4. extract the noise rate corresponding to the masked positions
+    p_masks = p_mask_full[final_masked_indices]
+    
+    return noisy_input_ids, final_masked_indices, p_masks
+
+
+def block_diff_mask(b, h, q_idx, kv_idx, block_size=None, n=None):
+    """
+    Constructs the specialized block diffusion attention mask for training
+    composed of three masks:
+    - **Block Diagonal Mask (M_BD)**: Self-attention within noised blocks
+    - **Offset Block Causal Mask (M_OBC)**: Cross-attention for conditional context
+    - **Block Causal Mask (M_BC)**: Attention to update x0
+
+    Args:
+        b, h: Batch and head indices (ignored for mask logic).
+        q_idx, kv_idx: Query and Key indices.
+        seq_len: Total sequence length.
+        block_size: Defines the block structure.
+
+    Returns:
+        A boolean attention mask.
+    """
+
+    # Indicate whether token belongs to xt or x0
+    x0_flag_q = q_idx >= n
+    x0_flag_kv = kv_idx >= n
+
+    # Compute block indices
+    block_q = torch.where(
+        x0_flag_q == 1, (q_idx - n) // block_size, q_idx // block_size
+    )
+    block_kv = torch.where(
+        x0_flag_kv == 1, (kv_idx - n) // block_size, kv_idx // block_size
+    )
+
+    # **1. Block Diagonal Mask (M_BD) **
+    block_diagonal = (block_q == block_kv) & (x0_flag_q == x0_flag_kv)
+
+    # **2. Offset Block-Causal Mask (M_OBC) **
+    offset_block_causal = (block_q > block_kv) & (
+        x0_flag_kv == 1) & (x0_flag_q == 0)
+
+    # **3. Block-Causal Mask (M_BC) **
+    block_causal = (block_q >= block_kv) & (x0_flag_kv == 1) & (x0_flag_q == 1)
+
+    # **4. Combine Masks **
+    return block_diagonal | offset_block_causal | block_causal
+
+
+def block_attn_mask(num_tokens, block_size, device):
+    masks = []
+    for i in range(len(num_tokens)):
+        cur_masks = []
+        for num in num_tokens[i]:
+            # return n*n instead of 2n*2n
+            single_mask = block_diff_mask(
+                b=None,
+                h=None,
+                q_idx=torch.arange(num * 2, device=device)[:, None],
+                kv_idx=torch.arange(num * 2, device=device)[None, :],
+                block_size=block_size,
+                n=num,
+            )
+            cur_masks.append(single_mask)
+        masks.append(torch.block_diag(*cur_masks))
+    masks = torch.stack(masks, dim=0)
+    return masks
+
+
+@torch.compile(fullgraph=True, mode="max-autotune-no-cudagraphs")
+def fused_flex_attention(query, key, value, attention_mask, **kwargs):
+    return flex_attention(query, key, value, block_mask=attention_mask, **kwargs)
+
+
+@use_kernel_forward_from_hub("RMSNorm")
+class SDARRMSNorm(nn.Module):
+    def __init__(self, hidden_size, eps=1e-6):
+        """
+        SDARRMSNorm is equivalent to T5LayerNorm
+        """
+        super().__init__()
+        self.weight = nn.Parameter(torch.ones(hidden_size))
+        self.variance_epsilon = eps
+
+    def forward(self, hidden_states):
+        return flash_rms_norm(
+            hidden_states, weight=self.weight, bias=None, eps=self.variance_epsilon)
+        '''
+        input_dtype = hidden_states.dtype
+        hidden_states = hidden_states.to(torch.float32)
+        variance = hidden_states.pow(2).mean(-1, keepdim=True)
+        hidden_states = hidden_states * \
+            torch.rsqrt(variance + self.variance_epsilon)
+        return self.weight * hidden_states.to(input_dtype)
+        '''
+
+    def extra_repr(self):
+        return f"{tuple(self.weight.shape)}, eps={self.variance_epsilon}"
+
+
+class SDARMLP(nn.Module):
+    def __init__(self, config):
+        super().__init__()
+        self.config = config
+        self.hidden_size = config.hidden_size
+        self.intermediate_size = config.intermediate_size
+        self.gate_proj = nn.Linear(
+            self.hidden_size, self.intermediate_size, bias=False)
+        self.up_proj = nn.Linear(
+            self.hidden_size, self.intermediate_size, bias=False)
+        self.down_proj = nn.Linear(
+            self.intermediate_size, self.hidden_size, bias=False)
+        self.act_fn = ACT2FN[config.hidden_act]
+
+    def forward(self, x):
+        if liger_kernel_is_available:
+            return self.down_proj(LigerSiLUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
+        else:
+            down_proj = self.down_proj(self.act_fn(
+                self.gate_proj(x)) * self.up_proj(x))
+            return down_proj
+
+
+def rotate_half(x):
+    """Rotates half the hidden dims of the input."""
+    x1 = x[..., : x.shape[-1] // 2]
+    x2 = x[..., x.shape[-1] // 2:]
+    return torch.cat((-x2, x1), dim=-1)
+
+
+def apply_rotary_pos_emb(q, k, cos, sin, position_ids=None, unsqueeze_dim=1):
+    """Applies Rotary Position Embedding to the query and key tensors.
+
+    Args:
+        q (`torch.Tensor`): The query tensor.
+        k (`torch.Tensor`): The key tensor.
+        cos (`torch.Tensor`): The cosine part of the rotary embedding.
+        sin (`torch.Tensor`): The sine part of the rotary embedding.
+        position_ids (`torch.Tensor`, *optional*):
+            Deprecated and unused.
+        unsqueeze_dim (`int`, *optional*, defaults to 1):
+            The 'unsqueeze_dim' argument specifies the dimension along which to unsqueeze cos[position_ids] and
+            sin[position_ids] so that they can be properly broadcasted to the dimensions of q and k. For example, note
+            that cos[position_ids] and sin[position_ids] have the shape [batch_size, seq_len, head_dim]. Then, if q and
+            k have the shape [batch_size, heads, seq_len, head_dim], then setting unsqueeze_dim=1 makes
+            cos[position_ids] and sin[position_ids] broadcastable to the shapes of q and k. Similarly, if q and k have
+            the shape [batch_size, seq_len, heads, head_dim], then set unsqueeze_dim=2.
+    Returns:
+        `tuple(torch.Tensor)` comprising of the query and key tensors rotated using the Rotary Position Embedding.
+    """
+    cos = cos.unsqueeze(unsqueeze_dim)
+    sin = sin.unsqueeze(unsqueeze_dim)
+    q_embed = (q * cos) + (rotate_half(q) * sin)
+    k_embed = (k * cos) + (rotate_half(k) * sin)
+    return q_embed, k_embed
+
+
+def repeat_kv(hidden_states: torch.Tensor, n_rep: int) -> torch.Tensor:
+    """
+    This is the equivalent of torch.repeat_interleave(x, dim=1, repeats=n_rep). The hidden states go from (batch,
+    num_key_value_heads, seqlen, head_dim) to (batch, num_attention_heads, seqlen, head_dim)
+    """
+    batch, num_key_value_heads, slen, head_dim = hidden_states.shape
+    if n_rep == 1:
+        return hidden_states
+    hidden_states = hidden_states[:, :, None, :, :].expand(
+        batch, num_key_value_heads, n_rep, slen, head_dim)
+    return hidden_states.reshape(batch, num_key_value_heads * n_rep, slen, head_dim)
+
+
+def eager_attention_forward(
+    module: nn.Module,
+    query: torch.Tensor,
+    key: torch.Tensor,
+    value: torch.Tensor,
+    attention_mask: Optional[torch.Tensor],
+    scaling: float,
+    dropout: float = 0.0,
+    **kwargs,
+):
+    key_states = repeat_kv(key, module.num_key_value_groups)
+    value_states = repeat_kv(value, module.num_key_value_groups)
+
+    attn_weights = torch.matmul(query, key_states.transpose(2, 3)) * scaling
+    if attention_mask is not None:
+        causal_mask = attention_mask[:, :, :, : key_states.shape[-2]]
+        attn_weights = attn_weights + causal_mask
+
+    attn_weights = nn.functional.softmax(
+        attn_weights, dim=-1, dtype=torch.float32).to(query.dtype)
+    attn_weights = nn.functional.dropout(
+        attn_weights, p=dropout, training=module.training)
+    attn_output = torch.matmul(attn_weights, value_states)
+    attn_output = attn_output.transpose(1, 2).contiguous()
+
+    return attn_output, attn_weights
+
+
+class SDARAttention(nn.Module):
+    """Multi-headed attention from 'Attention Is All You Need' paper"""
+
+    def __init__(self, config: SDARConfig, layer_idx: int):
+        super().__init__()
+        self.config = config
+        self.layer_idx = layer_idx
+        self.head_dim = getattr(
+            config, "head_dim", config.hidden_size // config.num_attention_heads)
+        self.num_key_value_groups = config.num_attention_heads // config.num_key_value_heads
+        self.scaling = self.head_dim**-0.5
+        self.attention_dropout = config.attention_dropout
+        self.is_causal = True
+
+        self.hidden_size = config.hidden_size
+        self.num_attention_heads = config.num_attention_heads
+        self.num_key_value_heads = config.num_key_value_heads
+
+        self.q_proj = nn.Linear(
+            config.hidden_size, config.num_attention_heads * self.head_dim, bias=config.attention_bias
+        )
+        self.k_proj = nn.Linear(
+            config.hidden_size, config.num_key_value_heads * self.head_dim, bias=config.attention_bias
+        )
+        self.v_proj = nn.Linear(
+            config.hidden_size, config.num_key_value_heads * self.head_dim, bias=config.attention_bias
+        )
+        self.o_proj = nn.Linear(
+            config.num_attention_heads * self.head_dim, config.hidden_size, bias=config.attention_bias
+        )
+        # unlike olmo, only on the head dim!
+        self.q_norm = SDARRMSNorm(self.head_dim, eps=config.rms_norm_eps)
+        # thus post q_norm does not need reshape
+        self.k_norm = SDARRMSNorm(self.head_dim, eps=config.rms_norm_eps)
+        self.sliding_window = config.sliding_window
+        if not (
+            self.config.use_sliding_window
+            and getattr(self.config, "sliding_window", None) is not None
+            and self.layer_idx >= self.config.max_window_layers
+        ):
+            self.sliding_window = None
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        position_embeddings: Tuple[torch.Tensor, torch.Tensor],
+        attention_mask: Optional[torch.Tensor],
+        past_key_value: Optional[Cache] = None,
+        cache_position: Optional[torch.LongTensor] = None,
+        **kwargs: Unpack[FlashAttentionKwargs],
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
+        input_shape = hidden_states.shape[:-1]
+        bsz, q_len = input_shape
+        hidden_shape = (*input_shape, -1, self.head_dim)
+
+        query_states = self.q_norm(self.q_proj(
+            hidden_states).view(hidden_shape)).transpose(1, 2)
+        key_states = self.k_norm(self.k_proj(
+            hidden_states).view(hidden_shape)).transpose(1, 2)
+        value_states = self.v_proj(hidden_states).view(
+            hidden_shape).transpose(1, 2)
+
+        cos, sin = position_embeddings
+        query_states, key_states = apply_rotary_pos_emb(
+            query_states, key_states, cos, sin)
+
+        if past_key_value is not None and kwargs.get("store_kv", False):
+            # sin and cos are specific to RoPE models; cache_position needed for the static cache
+            key_states, value_states = past_key_value.update(
+                key_states, value_states, self.layer_idx)
+        elif past_key_value is not None and not kwargs.get("store_kv", False) and len(past_key_value) > self.layer_idx:
+            # only retrive, do not store kv
+            past_key_states, past_value_states = past_key_value[self.layer_idx]
+            key_states = torch.cat(
+                [past_key_states, key_states], dim=-2)
+            value_states = torch.cat(
+                [past_value_states, value_states], dim=-2)
+
+        if self.training:
+            attn_output, attn_weights = fused_flex_attention(
+                query=query_states,
+                key=key_states,
+                value=value_states,
+                attention_mask=attention_mask,
+                enable_gqa=True,
+                scale=self.scaling,
+                return_lse=True
+            )
+            attn_weights = attn_weights.to(
+                value_states.dtype) if attn_weights is not None else None
+            attn_output = rearrange(attn_output, 'b h l d -> b l (h d)')
+        else:
+            attention_mask = attention_mask.bool() if attention_mask is not None else None
+            attn_weights = None
+            if torch.all(attention_mask):  # decoding
+                query_states = query_states.transpose(1, 2)
+                key_states = key_states.transpose(1, 2)
+                value_states = value_states.transpose(1, 2)
+                attn_output = flash_attn_func(
+                    query_states,
+                    key_states,
+                    value_states,
+                    causal=False,
+                    softmax_scale=self.scaling
+                )
+                attn_output = rearrange(attn_output, 'b l h d -> b l (h d)')
+            else:  # prefilling
+                attn_output = F.scaled_dot_product_attention(
+                    query=query_states,
+                    key=key_states,
+                    value=value_states,
+                    attn_mask=attention_mask,
+                    is_causal=False,
+                    scale=self.scaling,
+                    enable_gqa=True
+                )
+                attn_output = rearrange(attn_output, 'b h l d -> b l (h d)')
+        attn_output = self.o_proj(attn_output)
+        return attn_output, attn_weights  # , attn_weights
+
+
+class SDARDecoderLayer(GradientCheckpointingLayer):
+    def __init__(self, config: SDARConfig, layer_idx: int):
+        super().__init__()
+        self.hidden_size = config.hidden_size
+        self.self_attn = SDARAttention(config=config, layer_idx=layer_idx)
+        self.mlp = SDARMLP(config)
+        self.input_layernorm = SDARRMSNorm(
+            config.hidden_size, eps=config.rms_norm_eps)
+        self.post_attention_layernorm = SDARRMSNorm(
+            config.hidden_size, eps=config.rms_norm_eps)
+        if (
+            config.sliding_window and config._attn_implementation != "flash_attention_2"
+        ):  # diff with Llama is this warning
+            logger.warning_once(
+                f"Sliding Window Attention is enabled but not implemented for `{config._attn_implementation}`; "
+                "unexpected results may be encountered."
+            )
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: Optional[bool] = False,
+        use_cache: Optional[bool] = False,
+        store_kv: Optional[bool] = False,
+        cache_position: Optional[torch.LongTensor] = None,
+        # necessary, but kept here for BC
+        position_embeddings: Optional[Tuple[torch.Tensor,
+                                            torch.Tensor]] = None,
+        **kwargs: Unpack[FlashAttentionKwargs],
+    ) -> Tuple[torch.FloatTensor, Optional[Tuple[torch.FloatTensor, torch.FloatTensor]]]:
+        residual = hidden_states
+        hidden_states = self.input_layernorm(hidden_states)
+
+        # Self Attention
+        hidden_states, self_attn_weights = self.self_attn(
+            hidden_states=hidden_states,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_value=past_key_value,
+            output_attentions=output_attentions,
+            use_cache=use_cache,
+            store_kv=store_kv,
+            cache_position=cache_position,
+            position_embeddings=position_embeddings,
+            **kwargs,
+        )
+        hidden_states = residual + hidden_states
+
+        # Fully Connected
+        residual = hidden_states
+        hidden_states = self.post_attention_layernorm(hidden_states)
+        hidden_states = self.mlp(hidden_states)
+        hidden_states = residual + hidden_states
+
+        outputs = (hidden_states,)
+        if output_attentions:
+            outputs += (self_attn_weights,)
+
+        return outputs
+
+
+@auto_docstring
+class SDARPreTrainedModel(PreTrainedModel):
+    config_class = SDARConfig
+    base_model_prefix = "model"
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["SDARDecoderLayer"]
+    _skip_keys_device_placement = ["past_key_values"]
+    _supports_flash_attn_2 = True
+    _supports_sdpa = True
+    _supports_flex_attn = True
+    _supports_cache_class = True
+    _supports_quantized_cache = True
+    _supports_static_cache = True
+    _supports_attention_backend = True
+
+    def _init_weights(self, module):
+        std = self.config.initializer_range
+        if isinstance(module, nn.Linear):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.bias is not None:
+                module.bias.data.zero_()
+        elif isinstance(module, nn.Embedding):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.padding_idx is not None:
+                module.weight.data[module.padding_idx].zero_()
+        elif isinstance(module, SDARRMSNorm):
+            module.weight.data.fill_(1.0)
+
+
+class SDARRotaryEmbedding(nn.Module):
+    def __init__(self, config: SDARConfig, device=None):
+        super().__init__()
+        # BC: "rope_type" was originally "type"
+        if hasattr(config, "rope_scaling") and config.rope_scaling is not None:
+            self.rope_type = config.rope_scaling.get(
+                "rope_type", config.rope_scaling.get("type"))
+        else:
+            self.rope_type = "default"
+        self.max_seq_len_cached = config.max_position_embeddings
+        self.original_max_seq_len = config.max_position_embeddings
+
+        self.config = config
+        self.rope_init_fn = ROPE_INIT_FUNCTIONS[self.rope_type]
+
+        inv_freq, self.attention_scaling = self.rope_init_fn(
+            self.config, device)
+        self.register_buffer("inv_freq", inv_freq, persistent=False)
+        self.original_inv_freq = self.inv_freq
+
+    @torch.no_grad()
+    # power user: used with advanced RoPE types (e.g. dynamic rope)
+    @dynamic_rope_update
+    def forward(self, x, position_ids):
+        inv_freq_expanded = self.inv_freq[None, :, None].float().expand(
+            position_ids.shape[0], -1, 1).to(x.device)
+        position_ids_expanded = position_ids[:, None, :].float()
+
+        device_type = x.device.type if isinstance(
+            x.device.type, str) and x.device.type != "mps" else "cpu"
+        with torch.autocast(device_type=device_type, enabled=False):  # Force float32
+            freqs = (inv_freq_expanded.float() @
+                     position_ids_expanded.float()).transpose(1, 2)
+            emb = torch.cat((freqs, freqs), dim=-1)
+            cos = emb.cos() * self.attention_scaling
+            sin = emb.sin() * self.attention_scaling
+
+        return cos.to(dtype=x.dtype), sin.to(dtype=x.dtype)
+
+
+@auto_docstring
+class SDARModel(SDARPreTrainedModel):
+    def __init__(self, config: SDARConfig):
+        super().__init__(config)
+        self.padding_idx = config.pad_token_id
+        self.vocab_size = config.vocab_size
+
+        self.embed_tokens = nn.Embedding(
+            config.vocab_size, config.hidden_size, self.padding_idx)
+        self.layers = nn.ModuleList(
+            [SDARDecoderLayer(config, layer_idx)
+             for layer_idx in range(config.num_hidden_layers)]
+        )
+        self.norm = SDARRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+        self.rotary_emb = SDARRotaryEmbedding(config=config)
+        self.gradient_checkpointing = False
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    def get_input_embeddings(self):
+        return self.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.embed_tokens = value
+
+    @can_return_tuple
+    @auto_docstring
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Cache] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        use_cache: Optional[bool] = None,
+        store_kv: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        cache_position: Optional[torch.LongTensor] = None,
+        **flash_attn_kwargs: Unpack[FlashAttentionKwargs],
+    ) -> BaseModelOutputWithPast:
+        output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
+        output_hidden_states = (
+            output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
+        )
+        use_cache = use_cache if use_cache is not None else self.config.use_cache
+
+        if (input_ids is None) ^ (inputs_embeds is not None):
+            raise ValueError(
+                "You must specify exactly one of input_ids or inputs_embeds")
+
+        if self.gradient_checkpointing and self.training and use_cache:
+            logger.warning_once(
+                "`use_cache=True` is incompatible with gradient checkpointing. Setting `use_cache=False`."
+            )
+            use_cache = False
+
+        # TODO (joao): remove this exception in v4.56 -- it exists for users that try to pass a legacy cache
+        if not isinstance(past_key_values, (type(None), Cache)):
+            raise ValueError(
+                "The `past_key_values` should be either a `Cache` object or `None`.")
+
+        if inputs_embeds is None:
+            inputs_embeds = self.embed_tokens(input_ids)
+
+        if use_cache and past_key_values is None:
+            past_key_values = DynamicCache()
+
+        if cache_position is None:
+            past_seen_tokens = past_key_values.get_seq_length(
+            ) if past_key_values is not None else 0
+            cache_position = torch.arange(
+                past_seen_tokens, past_seen_tokens + inputs_embeds.shape[1], device=inputs_embeds.device
+            )
+
+        if position_ids is None:
+            position_ids = cache_position.unsqueeze(0)
+
+        # causal_mask = self._update_causal_mask(
+        #     attention_mask, inputs_embeds, cache_position, past_key_values, output_attentions
+        # )
+
+        hidden_states = inputs_embeds
+
+        # create position embeddings to be shared across the decoder layers
+        position_embeddings = self.rotary_emb(hidden_states, position_ids)
+
+        # decoder layers
+        all_hidden_states = () if output_hidden_states else None
+        all_self_attns = () if output_attentions else None
+
+        for decoder_layer in self.layers[: self.config.num_hidden_layers]:
+            if output_hidden_states:
+                all_hidden_states += (hidden_states,)
+
+            layer_outputs = decoder_layer(
+                hidden_states,
+                attention_mask=attention_mask,
+                position_ids=position_ids,
+                past_key_value=past_key_values,
+                output_attentions=output_attentions,
+                use_cache=use_cache,
+                store_kv=store_kv,
+                cache_position=cache_position,
+                position_embeddings=position_embeddings,
+                **flash_attn_kwargs,
+            )
+
+            hidden_states = layer_outputs[0]
+
+            if output_attentions:
+                all_self_attns += (layer_outputs[1],)
+
+        hidden_states = self.norm(hidden_states)
+
+        # add hidden states from the last decoder layer
+        if output_hidden_states:
+            all_hidden_states += (hidden_states,)
+
+        return BaseModelOutputWithPast(
+            last_hidden_state=hidden_states,
+            past_key_values=past_key_values if use_cache else None,
+            hidden_states=all_hidden_states,
+            attentions=all_self_attns,
+        )
+
+    def _update_causal_mask(
+        self,
+        attention_mask: Union[torch.Tensor, "BlockMask"],
+        input_tensor: torch.Tensor,
+        cache_position: torch.Tensor,
+        past_key_values: Cache,
+        output_attentions: bool = False,
+    ):
+        if self.config._attn_implementation == "flash_attention_2":
+            if attention_mask is not None and past_key_values is not None:
+                is_padding_right = attention_mask[:, -
+                                                  1].sum().item() != input_tensor.size()[0]
+                if is_padding_right:
+                    raise ValueError(
+                        "You are attempting to perform batched generation with padding_side='right'"
+                        " this may lead to unexpected behaviour for Flash Attention version of Qwen3. Make sure to "
+                        " call `tokenizer.padding_side  = 'left'` before tokenizing the input. "
+                    )
+            if attention_mask is not None and 0.0 in attention_mask:
+                return attention_mask
+            return None
+        if self.config._attn_implementation == "flex_attention":
+            if isinstance(attention_mask, torch.Tensor):
+                seq_len_q, seq_len_kv = attention_mask.shape
+                assert seq_len_q == seq_len_kv, f"got {attention_mask.shape=}"
+                attention_mask = create_block_mask(
+                    # 2d bool tensor, shape: [2*seqlen, 2*seqlen]
+                    lambda b, h, q_idx, kv_idx: attention_mask[q_idx, kv_idx],
+                    B=None, H=None, Q_LEN=seq_len_q, KV_LEN=seq_len_kv,
+                )
+            else:
+                # Here we pass in flex mask computed externally
+                assert isinstance(attention_mask, BlockMask)
+            return attention_mask
+
+        # For SDPA, when possible, we will rely on its `is_causal` argument instead of its `attn_mask` argument, in
+        # order to dispatch on Flash Attention 2. This feature is not compatible with static cache, as SDPA will fail
+        # to infer the attention mask.
+        past_seen_tokens = past_key_values.get_seq_length(
+        ) if past_key_values is not None else 0
+        using_static_cache = isinstance(past_key_values, StaticCache)
+        using_sliding_window_cache = isinstance(
+            past_key_values, SlidingWindowCache)
+
+        # When output attentions is True, sdpa implementation's forward method calls the eager implementation's forward
+        if (
+            self.config._attn_implementation == "sdpa"
+            and not (using_static_cache or using_sliding_window_cache)
+            and not output_attentions
+        ):
+            if AttentionMaskConverter._ignore_causal_mask_sdpa(
+                attention_mask,
+                inputs_embeds=input_tensor,
+                past_key_values_length=past_seen_tokens,
+                sliding_window=self.config.sliding_window,
+                is_training=self.training,
+            ):
+                return None
+
+        dtype = input_tensor.dtype
+        min_dtype = torch.finfo(dtype).min
+        sequence_length = input_tensor.shape[1]
+        # SlidingWindowCache or StaticCache
+        if using_sliding_window_cache or using_static_cache:
+            target_length = past_key_values.get_max_cache_shape()
+        # DynamicCache or no cache
+        else:
+            target_length = (
+                attention_mask.shape[-1]
+                if isinstance(attention_mask, torch.Tensor)
+                else past_seen_tokens + sequence_length + 1
+            )
+
+        # In case the provided `attention` mask is 2D, we generate a causal mask here (4D).
+        causal_mask = self._prepare_4d_causal_attention_mask_with_cache_position(
+            attention_mask,
+            sequence_length=sequence_length,
+            target_length=target_length,
+            dtype=dtype,
+            cache_position=cache_position,
+            batch_size=input_tensor.shape[0],
+            config=self.config,
+            past_key_values=past_key_values,
+        )
+
+        if (
+            self.config._attn_implementation == "sdpa"
+            and attention_mask is not None
+            and attention_mask.device.type in ["cuda", "xpu", "npu"]
+            and not output_attentions
+        ):
+            # Attend to all tokens in fully masked rows in the causal_mask, for example the relevant first rows when
+            # using left padding. This is required by F.scaled_dot_product_attention memory-efficient attention path.
+            # Details: https://github.com/pytorch/pytorch/issues/110213
+            causal_mask = AttentionMaskConverter._unmask_unattended(
+                causal_mask, min_dtype)
+
+        return causal_mask
+
+    @staticmethod
+    def _prepare_4d_causal_attention_mask_with_cache_position(
+        attention_mask: torch.Tensor,
+        sequence_length: int,
+        target_length: int,
+        dtype: torch.dtype,
+        cache_position: torch.Tensor,
+        batch_size: int,
+        config: SDARConfig,
+        past_key_values: Cache,
+    ):
+        """
+        Creates a causal 4D mask of shape `(batch_size, 1, query_length, key_value_length)` from a 2D mask of shape
+        `(batch_size, key_value_length)`, or if the input `attention_mask` is already 4D, do nothing.
+
+        Args:
+            attention_mask (`torch.Tensor`):
+                A 2D attention mask of shape `(batch_size, key_value_length)` or a 4D attention mask of shape `(batch_size, 1, query_length, key_value_length)`.
+            sequence_length (`int`):
+                The sequence length being processed.
+            target_length (`int`):
+                The target length: when generating with static cache, the mask should be as long as the static cache, to account for the 0 padding, the part of the cache that is not filled yet.
+            dtype (`torch.dtype`):
+                The dtype to use for the 4D attention mask.
+            cache_position (`torch.Tensor`):
+                Indices depicting the position of the input sequence tokens in the sequence.
+            batch_size (`torch.Tensor`):
+                Batch size.
+            config (`SDARConfig`):
+                The model's configuration class
+            past_key_values (`Cache`):
+                The cache class that is being used currently to generate
+        """
+        if attention_mask is not None and attention_mask.dim() == 4:
+            # In this case we assume that the mask comes already in inverted form and requires no inversion or slicing.
+            causal_mask = attention_mask
+        else:
+            min_dtype = torch.finfo(dtype).min
+            causal_mask = torch.full(
+                (sequence_length, target_length), fill_value=min_dtype, dtype=dtype, device=cache_position.device
+            )
+            diagonal_attend_mask = torch.arange(target_length, device=cache_position.device) > cache_position.reshape(
+                -1, 1
+            )
+            text_config = config.get_text_config()
+            if getattr(text_config, "use_sliding_window", True) and text_config.sliding_window is not None:
+                # if we have sliding window, we should not attend to tokens beyond sliding window length, so we mask them out also
+                # the check is needed to verify is current checkpoint was trained with sliding window or not
+                if not isinstance(past_key_values, SlidingWindowCache) or sequence_length > target_length:
+                    sliding_attend_mask = torch.arange(target_length, device=cache_position.device) <= (
+                        cache_position.reshape(-1, 1) -
+                        text_config.sliding_window
+                    )
+                    diagonal_attend_mask.bitwise_or_(sliding_attend_mask)
+            causal_mask *= diagonal_attend_mask
+            causal_mask = causal_mask[None, None,
+                                      :, :].expand(batch_size, 1, -1, -1)
+            if attention_mask is not None:
+                causal_mask = causal_mask.clone()  # copy to contiguous memory for in-place edit
+                if attention_mask.shape[-1] > target_length:
+                    attention_mask = attention_mask[:, :target_length]
+                mask_length = attention_mask.shape[-1]
+                padding_mask = causal_mask[:, :, :, :mask_length] + attention_mask[:, None, None, :].to(
+                    causal_mask.device
+                )
+                padding_mask = padding_mask == 0
+                causal_mask[:, :, :, :mask_length] = causal_mask[:, :, :, :mask_length].masked_fill(
+                    padding_mask, min_dtype
+                )
+        return causal_mask
+
+
+class KwargsForCausalLM(FlashAttentionKwargs, LossKwargs):
+    ...
+
+
+@auto_docstring
+class SDARForCausalLM(SDARPreTrainedModel, GenerationMixin):
+    _tied_weights_keys = ["lm_head.weight"]
+    _tp_plan = {"lm_head": "colwise_rep"}
+    _pp_plan = {"lm_head": (["hidden_states"], ["logits"])}
+
+    def __init__(self, config):
+        super().__init__(config)
+        self.model = SDARModel(config)
+        self.vocab_size = config.vocab_size
+        self.lm_head = nn.Linear(
+            config.hidden_size, config.vocab_size, bias=False)
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    def get_input_embeddings(self):
+        return self.model.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.model.embed_tokens = value
+
+    def get_output_embeddings(self):
+        return self.lm_head
+
+    def set_output_embeddings(self, new_embeddings):
+        self.lm_head = new_embeddings
+
+    def set_decoder(self, decoder):
+        self.model = decoder
+
+    def get_decoder(self):
+        return self.model
+
+    def prepare_for_bd_training(self, inputs_ids, position_ids, prompt_mask):
+        bsz, seq_len = inputs_ids.shape
+        num_tokens = calculate_token_nums(position_ids) # List[torch.Tensor]
+        noisy_inputs_ids, logits_to_keep_half, p_mask = forward_add_noise_packed(
+            inputs_ids=inputs_ids,
+            num_tokens_list=num_tokens,
+            prompt_mask=prompt_mask,
+            mask_id=self.config.mask_token_id,
+            eob_token_id=getattr(self.config, "eob_token_id", None),
+        )
+        router_noisy_part_list = []
+        for i in range(bsz):
+            cur_router_noisy_part = (torch.arange(num_tokens[i].shape[0] *2) % 2 == 0).to(inputs_ids.device)
+            cur_router_noisy_part = cur_router_noisy_part.repeat_interleave(num_tokens[i].repeat_interleave(2))
+            router_noisy_part_list.append(cur_router_noisy_part)
+        router_noisy_part = torch.stack(router_noisy_part_list, dim=0)
+
+        # concated inputs_ids: (bzs, seq_len x 2)
+        concat_inputs_ids = inputs_ids.repeat(1, 2)
+        # concated logits_to_keep: (bsz, seq_len x 2)
+        logits_to_keep = torch.zeros(
+                    bsz, 2 * seq_len, dtype=torch.bool, device=inputs_ids.device)
+        # concated position_ids: (bsz, seq_len x 2)
+        concat_position_ids = torch.zeros(
+                    bsz, 2 * seq_len, dtype=position_ids.dtype, device=position_ids.device)
+        for i in range(bsz):
+            concat_inputs_ids[i][router_noisy_part[i]] = noisy_inputs_ids[i]
+            concat_inputs_ids[i][~router_noisy_part[i]] = inputs_ids[i]
+
+            logits_to_keep[i][router_noisy_part[i]] = logits_to_keep_half[i]
+
+            concat_position_ids[i][router_noisy_part[i]] = position_ids[i]
+            concat_position_ids[i][~router_noisy_part[i]] = position_ids[i]
+
+        # create flex_attention mask
+        if getattr(self.config, "dynamic_blocks", False) and getattr(self.config, "eob_token_id", None) is not None:
+            # Dynamic blocks based on EOB tokens
+            block_lengths_list = calculate_block_nums_from_eob(inputs_ids, num_tokens, self.config.eob_token_id)
+            attention_mask = block_attn_mask_dynamic(block_lengths_list, inputs_ids.device)
+        else:
+            # Fixed blocks
+            attention_mask = block_attn_mask(num_tokens, self.config.block_size, inputs_ids.device)
+            
+        flex_attention_mask_3d = create_block_mask(
+                            lambda b, h, q_idx, kv_idx: attention_mask[b, q_idx, kv_idx],
+                            B=attention_mask.size(0), H=None,
+                            Q_LEN=attention_mask.size(1), KV_LEN=attention_mask.size(2),
+        )
+
+        return concat_inputs_ids, concat_position_ids, flex_attention_mask_3d, logits_to_keep_half, logits_to_keep, p_mask
+
+    @can_return_tuple
+    @auto_docstring
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Cache] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        labels: Optional[torch.LongTensor] = None,
+        use_cache: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        cache_position: Optional[torch.LongTensor] = None,
+        logits_to_keep: Union[int, torch.Tensor] = 0,
+        **kwargs: Unpack[KwargsForCausalLM],
+    ) -> CausalLMOutputWithPast:
+        r"""
+        labels (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Labels for computing the masked language modeling loss. Indices should either be in `[0, ...,
+            config.vocab_size]` or -100 (see `input_ids` docstring). Tokens with indices set to `-100` are ignored
+            (masked), the loss is only computed for the tokens with labels in `[0, ..., config.vocab_size]`.
+
+        Example:
+
+        ```python
+        >>> from transformers import AutoTokenizer, SDARForCausalLM
+
+        >>> model = SDARForCausalLM.from_pretrained("DiffuOpen/SDAR-1.7B-Chat")
+        >>> tokenizer = AutoTokenizer.from_pretrained("DiffuOpen/SDAR-1.7B-Chat")
+
+        >>> prompt = "Hey, are you conscious? Can you talk to me?"
+        >>> inputs = tokenizer(prompt, return_tensors="pt")
+
+        >>> # Generate
+        >>> generate_ids = model.generate(inputs.input_ids, max_length=30)
+        >>> tokenizer.batch_decode(generate_ids, skip_special_tokens=True, clean_up_tokenization_spaces=False)[0]
+        "Hey, are you conscious? Can you talk to me?\nI'm not conscious, but I can talk to you."
+        ```"""
+        output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
+        output_hidden_states = (
+            output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
+        )
+        if self.training:
+            assert inputs_embeds is None, "only support input_ids during training"
+            prompt_mask = (labels == -100) if labels is not None else None
+            position_ids = modify_padded_position_ids_2d(position_ids)
+            concat_inputs_ids, concat_position_ids, flex_attention_mask_3d, logits_to_keep_half, logits_to_keep, p_mask = self.prepare_for_bd_training(input_ids, position_ids, prompt_mask)
+            outputs = self.model(
+                input_ids=concat_inputs_ids,
+                attention_mask=flex_attention_mask_3d,
+                position_ids=concat_position_ids,
+                output_attentions=output_attentions,
+                output_hidden_states=output_hidden_states,
+                return_dict=True,
+                cache_position=cache_position,
+                **kwargs,
+            )
+            hidden_states = outputs.last_hidden_state
+            hidden_states = hidden_states[logits_to_keep].contiguous()
+            assert labels is not None, "Labels must be provided for training."
+            labels = labels[logits_to_keep_half].contiguous()
+            loss_fct = FusedLinearDiffusionCrossEntropyLoss(reduction='sum')
+            loss = loss_fct(  # it will return (sum_loss, unreduced_loss)
+                    # conduct `view(-1, V)` inside the function
+                    x=hidden_states,
+                    target=labels,
+                    weight=self.lm_head.weight,
+                    bias=self.lm_head.bias,
+                    p_mask=p_mask,
+                    eob_token_id=getattr(self.config, "eob_token_id", None),
+                    eob_weight=getattr(self.config, "eob_weight", 1.0),
+                )
+            loss = loss / labels.numel()
+            logits = None
+        else:
+            # decoder outputs consists of (dec_features, layer_state, dec_hidden, dec_attn)
+            outputs: BaseModelOutputWithPast = self.model(
+                input_ids=input_ids,
+                attention_mask=attention_mask,
+                position_ids=position_ids,
+                past_key_values=past_key_values,
+                inputs_embeds=inputs_embeds,
+                use_cache=use_cache,
+                output_attentions=output_attentions,
+                output_hidden_states=output_hidden_states,
+                cache_position=cache_position,
+                **kwargs,
+            )
+
+            hidden_states = outputs.last_hidden_state
+            # Only compute necessary logits, and do not upcast them to float if we are not computing the loss
+            slice_indices = slice(-logits_to_keep,
+                                None) if isinstance(logits_to_keep, int) else logits_to_keep
+            hidden_states = hidden_states[:, slice_indices, :].contiguous()
+            fuse_linear_and_cross_entropy = self.config.fuse_cross_entropy and self.training
+            if fuse_linear_and_cross_entropy:
+                # When using fused_linear_ce_loss, we do not compute the whole logits on HBM
+                logits = None
+            else:
+                logits = self.lm_head(hidden_states)
+
+            loss = None
+            if labels is not None:
+                # FusedLinearCrossEntropyLoss will be implemented by monkey patch when training
+                # We don't use it when inferencing
+                loss_fct = nn.CrossEntropyLoss()  # nn.CE
+                loss = loss_fct(
+                    logits.view(-1, self.config.vocab_size), labels.view(-1))
+
+        return CausalLMOutputWithPast(
+            loss=loss,
+            logits=logits,
+            past_key_values=outputs.past_key_values,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )
+
+
+__all__ = [
+    "SDARForCausalLM",
+    "SDARModel",
+    "SDARPreTrainedModel",
+]

rng_state_0.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:25c62d345be8e040c196fb33f0e649d52b41ea74e0f2ba1fb163bd3eb2abe1a7
+size 16133

rng_state_1.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9f842350c63e633f32f33db4a1fcd5da03a73c43d86665c6e935919c09efda43
+size 16133

rng_state_2.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:b3c6b1564d12078d05ae00b585e01af634598a41e7a5166a489bf7e00024943d
+size 16133

rng_state_3.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:1a6da4163f04f24e920a83f2647a1ac25e4ef5b4ee0c8cbe5e94570a84ffb029
+size 16133

rng_state_4.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:7f404a929daa54e7845502018f60a891cdbda19558e0b8d18a6452d57c61a8e0
+size 16133

rng_state_5.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:cbc0b9bb52448a6f9dedccf684c93906a535f1d5ba3b9851bfd098d15f291df9
+size 16133

rng_state_6.pth

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:99b132e59a3199b46987dc184d1675f4aef42e3dcb2d13bbd5d888152aa0afb4
+size 16133

scheduler.pt

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:f14e5100356e29b1d39819aba28314a603c8339a038828cf64e69f89bd67e7a5
+size 1465

special_tokens_map.json

@@ -0,0 +1,40 @@
+{
+  "additional_special_tokens": [
+    "<|im_start|>",
+    "<|im_end|>",
+    "<|object_ref_start|>",
+    "<|object_ref_end|>",
+    "<|box_start|>",
+    "<|box_end|>",
+    "<|quad_start|>",
+    "<|quad_end|>",
+    "<|vision_start|>",
+    "<|vision_end|>",
+    "<|vision_pad|>",
+    "<|image_pad|>",
+    "<|video_pad|>",
+    "<|MASK|>",
+    "<EOB>"
+  ],
+  "eos_token": {
+    "content": "<|im_end|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  },
+  "mask_token": {
+    "content": "<|MASK|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  },
+  "pad_token": {
+    "content": "<|endoftext|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  }
+}

tokenization_qwen2.py

@@ -0,0 +1,342 @@
+# coding=utf-8
+# Copyright 2024 The Qwen team, Alibaba Group and The HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""Tokenization classes for Qwen2."""
+
+import json
+import os
+import unicodedata
+from functools import lru_cache
+from typing import Optional, Tuple
+
+import regex as re
+
+from transformers.tokenization_utils import AddedToken, PreTrainedTokenizer
+from transformers.utils import logging
+
+
+logger = logging.get_logger(__name__)
+
+VOCAB_FILES_NAMES = {
+    "vocab_file": "vocab.json",
+    "merges_file": "merges.txt",
+}
+
+
+MAX_MODEL_INPUT_SIZES = {"qwen/qwen-tokenizer": 32768}
+
+PRETOKENIZE_REGEX = r"""(?i:'s|'t|'re|'ve|'m|'ll|'d)|[^\r\n\p{L}\p{N}]?\p{L}+|\p{N}| ?[^\s\p{L}\p{N}]+[\r\n]*|\s*[\r\n]+|\s+(?!\S)|\s+"""
+
+
+@lru_cache()
+# Copied from transformers.models.gpt2.tokenization_gpt2.bytes_to_unicode
+def bytes_to_unicode():
+    """
+    Returns list of utf-8 byte and a mapping to unicode strings. We specifically avoids mapping to whitespace/control
+    characters the bpe code barfs on.
+
+    The reversible bpe codes work on unicode strings. This means you need a large # of unicode characters in your vocab
+    if you want to avoid UNKs. When you're at something like a 10B token dataset you end up needing around 5K for
+    decent coverage. This is a significant percentage of your normal, say, 32K bpe vocab. To avoid that, we want lookup
+    tables between utf-8 bytes and unicode strings.
+    """
+    bs = (
+        list(range(ord("!"), ord("~") + 1)) + list(range(ord("¡"), ord("¬") + 1)) + list(range(ord("®"), ord("ÿ") + 1))
+    )
+    cs = bs[:]
+    n = 0
+    for b in range(2**8):
+        if b not in bs:
+            bs.append(b)
+            cs.append(2**8 + n)
+            n += 1
+    cs = [chr(n) for n in cs]
+    return dict(zip(bs, cs))
+
+
+# Copied from transformers.models.gpt2.tokenization_gpt2.get_pairs
+def get_pairs(word):
+    """
+    Return set of symbol pairs in a word.
+
+    Word is represented as tuple of symbols (symbols being variable-length strings).
+    """
+    pairs = set()
+    prev_char = word[0]
+    for char in word[1:]:
+        pairs.add((prev_char, char))
+        prev_char = char
+    return pairs
+
+
+class Qwen2Tokenizer(PreTrainedTokenizer):
+    """
+    Construct a Qwen2 tokenizer. Based on byte-level Byte-Pair-Encoding.
+
+    Same with GPT2Tokenizer, this tokenizer has been trained to treat spaces like parts of the tokens so a word will
+    be encoded differently whether it is at the beginning of the sentence (without space) or not:
+
+    ```python
+    >>> from transformers import Qwen2Tokenizer
+
+    >>> tokenizer = Qwen2Tokenizer.from_pretrained("Qwen/Qwen-tokenizer")
+    >>> tokenizer("Hello world")["input_ids"]
+    [9707, 1879]
+
+    >>> tokenizer(" Hello world")["input_ids"]
+    [21927, 1879]
+    ```
+    This is expected.
+
+    You should not use GPT2Tokenizer instead, because of the different pretokenization rules.
+
+    This tokenizer inherits from [`PreTrainedTokenizer`] which contains most of the main methods. Users should refer to
+    this superclass for more information regarding those methods.
+
+    Args:
+        vocab_file (`str`):
+            Path to the vocabulary file.
+        merges_file (`str`):
+            Path to the merges file.
+        errors (`str`, *optional*, defaults to `"replace"`):
+            Paradigm to follow when decoding bytes to UTF-8. See
+            [bytes.decode](https://docs.python.org/3/library/stdtypes.html#bytes.decode) for more information.
+        unk_token (`str`, *optional*, defaults to `"<|endoftext|>"`):
+            The unknown token. A token that is not in the vocabulary cannot be converted to an ID and is set to be this
+            token instead.
+        bos_token (`str`, *optional*):
+            The beginning of sequence token. Not applicable for this tokenizer.
+        eos_token (`str`, *optional*, defaults to `"<|endoftext|>"`):
+            The end of sequence token.
+        pad_token (`str`, *optional*, defaults to `"<|endoftext|>"`):
+            The token used for padding, for example when batching sequences of different lengths.
+        clean_up_tokenization_spaces (`bool`, *optional*, defaults to `False`):
+            Whether or not the model should cleanup the spaces that were added when splitting the input text during the
+            tokenization process. Not applicable to this tokenizer, since tokenization does not add spaces.
+        split_special_tokens (`bool`, *optional*, defaults to `False`):
+            Whether or not the special tokens should be split during the tokenization process. The default behavior is
+            to not split special tokens. This means that if `<|endoftext|>` is the `eos_token`, then `tokenizer.tokenize("<|endoftext|>") =
+            ['<|endoftext|>`]. Otherwise, if `split_special_tokens=True`, then `tokenizer.tokenize("<|endoftext|>")` will be give `['<',
+            '|', 'endo', 'ft', 'ext', '|', '>']`. This argument is only supported for `slow` tokenizers for the moment.
+    """
+
+    vocab_files_names = VOCAB_FILES_NAMES
+    model_input_names = ["input_ids", "attention_mask"]
+
+    def __init__(
+        self,
+        vocab_file,
+        merges_file,
+        errors="replace",
+        unk_token="<|endoftext|>",
+        bos_token=None,
+        eos_token="<|endoftext|>",
+        pad_token="<|endoftext|>",
+        clean_up_tokenization_spaces=False,
+        split_special_tokens=False,
+        **kwargs,
+    ):
+        # Qwen vocab does not contain control tokens; added tokens need to be special
+        bos_token = (
+            AddedToken(bos_token, lstrip=False, rstrip=False, special=True, normalized=False)
+            if isinstance(bos_token, str)
+            else bos_token
+        )
+        eos_token = (
+            AddedToken(eos_token, lstrip=False, rstrip=False, special=True, normalized=False)
+            if isinstance(eos_token, str)
+            else eos_token
+        )
+        unk_token = (
+            AddedToken(unk_token, lstrip=False, rstrip=False, special=True, normalized=False)
+            if isinstance(unk_token, str)
+            else unk_token
+        )
+        pad_token = (
+            AddedToken(pad_token, lstrip=False, rstrip=False, special=True, normalized=False)
+            if isinstance(pad_token, str)
+            else pad_token
+        )
+
+        with open(vocab_file, encoding="utf-8") as vocab_handle:
+            self.encoder = json.load(vocab_handle)
+        self.decoder = {v: k for k, v in self.encoder.items()}
+        self.errors = errors  # how to handle errors in decoding
+        self.byte_encoder = bytes_to_unicode()
+        self.byte_decoder = {v: k for k, v in self.byte_encoder.items()}
+        bpe_merges = []
+        with open(merges_file, encoding="utf-8") as merges_handle:
+            for i, line in enumerate(merges_handle):
+                line = line.strip()
+                if (i == 0 and line.startswith("#version:")) or not line:
+                    continue
+                bpe_merges.append(tuple(line.split()))
+        self.bpe_ranks = dict(zip(bpe_merges, range(len(bpe_merges))))
+        # NOTE: the cache can grow without bound and will get really large for long running processes
+        # (esp. for texts of language that do not use space between word, e.g. Chinese); technically
+        # not a memory leak but appears as one.
+        # GPT2Tokenizer has the same problem, so let's be consistent.
+        self.cache = {}
+
+        self.pat = re.compile(PRETOKENIZE_REGEX)
+
+        if kwargs.get("add_prefix_space", False):
+            logger.warning_once(
+                f"{self.__class__.__name} does not support `add_prefix_space`, setting it to True has no effect."
+            )
+
+        super().__init__(
+            errors=errors,
+            bos_token=bos_token,
+            eos_token=eos_token,
+            pad_token=pad_token,
+            unk_token=unk_token,
+            clean_up_tokenization_spaces=clean_up_tokenization_spaces,
+            split_special_tokens=split_special_tokens,
+            **kwargs,
+        )
+
+    @property
+    def vocab_size(self) -> int:
+        return len(self.encoder)
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer.get_vocab
+    def get_vocab(self):
+        return dict(self.encoder, **self.added_tokens_encoder)
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer.bpe
+    def bpe(self, token):
+        if token in self.cache:
+            return self.cache[token]
+        word = tuple(token)
+        pairs = get_pairs(word)
+
+        if not pairs:
+            return token
+
+        while True:
+            bigram = min(pairs, key=lambda pair: self.bpe_ranks.get(pair, float("inf")))
+            if bigram not in self.bpe_ranks:
+                break
+            first, second = bigram
+            new_word = []
+            i = 0
+            while i < len(word):
+                try:
+                    j = word.index(first, i)
+                except ValueError:
+                    new_word.extend(word[i:])
+                    break
+                else:
+                    new_word.extend(word[i:j])
+                    i = j
+
+                if word[i] == first and i < len(word) - 1 and word[i + 1] == second:
+                    new_word.append(first + second)
+                    i += 2
+                else:
+                    new_word.append(word[i])
+                    i += 1
+            new_word = tuple(new_word)
+            word = new_word
+            if len(word) == 1:
+                break
+            else:
+                pairs = get_pairs(word)
+        word = " ".join(word)
+        self.cache[token] = word
+        return word
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer._tokenize
+    def _tokenize(self, text):
+        """Tokenize a string."""
+        bpe_tokens = []
+        for token in re.findall(self.pat, text):
+            token = "".join(
+                self.byte_encoder[b] for b in token.encode("utf-8")
+            )  # Maps all our bytes to unicode strings, avoiding control tokens of the BPE (spaces in our case)
+            bpe_tokens.extend(bpe_token for bpe_token in self.bpe(token).split(" "))
+        return bpe_tokens
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer._convert_token_to_id
+    def _convert_token_to_id(self, token):
+        """Converts a token (str) in an id using the vocab."""
+        return self.encoder.get(token, self.encoder.get(self.unk_token))
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer._convert_id_to_token
+    def _convert_id_to_token(self, index):
+        """Converts an index (integer) in a token (str) using the vocab."""
+        return self.decoder.get(index)
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer.convert_tokens_to_string
+    def convert_tokens_to_string(self, tokens):
+        """Converts a sequence of tokens (string) in a single string."""
+        text = "".join(tokens)
+        text = bytearray([self.byte_decoder[c] for c in text]).decode("utf-8", errors=self.errors)
+        return text
+
+    def decode(
+        self,
+        token_ids,
+        skip_special_tokens: bool = False,
+        clean_up_tokenization_spaces: Optional[bool] = False,
+        spaces_between_special_tokens: bool = False,
+        **kwargs,
+    ) -> str:
+        # `spaces_between_special_tokens` defaults to True for _decode in slow tokenizers
+        # and cannot be configured elsewhere, but it should default to False for Qwen2Tokenizer
+        return super().decode(
+            token_ids,
+            skip_special_tokens=skip_special_tokens,
+            clean_up_tokenization_spaces=clean_up_tokenization_spaces,
+            spaces_between_special_tokens=spaces_between_special_tokens,
+            **kwargs,
+        )
+
+    # Copied from transformers.models.gpt2.tokenization_gpt2.GPT2Tokenizer.save_vocabulary
+    def save_vocabulary(self, save_directory: str, filename_prefix: Optional[str] = None) -> Tuple[str]:
+        if not os.path.isdir(save_directory):
+            logger.error(f"Vocabulary path ({save_directory}) should be a directory")
+            return
+        vocab_file = os.path.join(
+            save_directory, (filename_prefix + "-" if filename_prefix else "") + VOCAB_FILES_NAMES["vocab_file"]
+        )
+        merge_file = os.path.join(
+            save_directory, (filename_prefix + "-" if filename_prefix else "") + VOCAB_FILES_NAMES["merges_file"]
+        )
+
+        with open(vocab_file, "w", encoding="utf-8") as f:
+            f.write(json.dumps(self.encoder, indent=2, sort_keys=True, ensure_ascii=False) + "\n")
+
+        index = 0
+        with open(merge_file, "w", encoding="utf-8") as writer:
+            writer.write("#version: 0.2\n")
+            for bpe_tokens, token_index in sorted(self.bpe_ranks.items(), key=lambda kv: kv[1]):
+                if index != token_index:
+                    logger.warning(
+                        f"Saving vocabulary to {merge_file}: BPE merge indices are not consecutive."
+                        " Please check that the tokenizer is not corrupted!"
+                    )
+                    index = token_index
+                writer.write(" ".join(bpe_tokens) + "\n")
+                index += 1
+
+        return vocab_file, merge_file
+
+    def prepare_for_tokenization(self, text, **kwargs):
+        text = unicodedata.normalize("NFC", text)
+        return (text, kwargs)
+
+
+__all__ = ["Qwen2Tokenizer"]

tokenizer_config.json

@@ -0,0 +1,265 @@
+{
+  "add_bos_token": false,
+  "add_prefix_space": false,
+  "added_tokens_decoder": {
+    "151643": {
+      "content": "<|endoftext|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151644": {
+      "content": "<|im_start|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151645": {
+      "content": "<|im_end|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151646": {
+      "content": "<|object_ref_start|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151647": {
+      "content": "<|object_ref_end|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151648": {
+      "content": "<|box_start|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151649": {
+      "content": "<|box_end|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151650": {
+      "content": "<|quad_start|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151651": {
+      "content": "<|quad_end|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151652": {
+      "content": "<|vision_start|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151653": {
+      "content": "<|vision_end|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151654": {
+      "content": "<|vision_pad|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151655": {
+      "content": "<|image_pad|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151656": {
+      "content": "<|video_pad|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151657": {
+      "content": "<tool_call>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151658": {
+      "content": "</tool_call>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151659": {
+      "content": "<|fim_prefix|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151660": {
+      "content": "<|fim_middle|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151661": {
+      "content": "<|fim_suffix|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151662": {
+      "content": "<|fim_pad|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151663": {
+      "content": "<|repo_name|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151664": {
+      "content": "<|file_sep|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151665": {
+      "content": "<tool_response>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151666": {
+      "content": "</tool_response>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151667": {
+      "content": "<think>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151668": {
+      "content": "</think>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": false
+    },
+    "151669": {
+      "content": "<|MASK|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    },
+    "151670": {
+      "content": "<EOB>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    }
+  },
+  "additional_special_tokens": [
+    "<|im_start|>",
+    "<|im_end|>",
+    "<|object_ref_start|>",
+    "<|object_ref_end|>",
+    "<|box_start|>",
+    "<|box_end|>",
+    "<|quad_start|>",
+    "<|quad_end|>",
+    "<|vision_start|>",
+    "<|vision_end|>",
+    "<|vision_pad|>",
+    "<|image_pad|>",
+    "<|video_pad|>",
+    "<|MASK|>",
+    "<EOB>"
+  ],
+  "auto_map": {
+    "AutoTokenizer": [
+      "tokenization_qwen2.Qwen2Tokenizer",
+      null
+    ]
+  },
+  "bos_token": null,
+  "clean_up_tokenization_spaces": false,
+  "eos_token": "<|im_end|>",
+  "errors": "replace",
+  "extra_special_tokens": {},
+  "mask_token": "<|MASK|>",
+  "model_max_length": 131072,
+  "pad_token": "<|endoftext|>",
+  "padding_side": "right",
+  "split_special_tokens": false,
+  "tokenizer_class": "Qwen2Tokenizer",
+  "unk_token": null
+}

trainer_state.json

@@ -0,0 +1,2204 @@
+{
+  "best_global_step": null,
+  "best_metric": null,
+  "best_model_checkpoint": null,
+  "epoch": 2.094658553076403,
+  "eval_steps": 500,
+  "global_step": 1550,
+  "is_hyper_param_search": false,
+  "is_local_process_zero": true,
+  "is_world_process_zero": true,
+  "log_history": [
+    {
+      "epoch": 0.0067613252197430695,
+      "grad_norm": 1272.255859375,
+      "learning_rate": 5.970149253731343e-07,
+      "loss": 37.2004,
+      "step": 5
+    },
+    {
+      "epoch": 0.013522650439486139,
+      "grad_norm": 538.208251953125,
+      "learning_rate": 1.3432835820895524e-06,
+      "loss": 34.1505,
+      "step": 10
+    },
+    {
+      "epoch": 0.02028397565922921,
+      "grad_norm": 241.84115600585938,
+      "learning_rate": 2.08955223880597e-06,
+      "loss": 27.7417,
+      "step": 15
+    },
+    {
+      "epoch": 0.027045300878972278,
+      "grad_norm": 149.00253295898438,
+      "learning_rate": 2.835820895522388e-06,
+      "loss": 22.3086,
+      "step": 20
+    },
+    {
+      "epoch": 0.03380662609871535,
+      "grad_norm": 137.87466430664062,
+      "learning_rate": 3.582089552238806e-06,
+      "loss": 19.628,
+      "step": 25
+    },
+    {
+      "epoch": 0.04056795131845842,
+      "grad_norm": 166.65280151367188,
+      "learning_rate": 4.3283582089552236e-06,
+      "loss": 17.6907,
+      "step": 30
+    },
+    {
+      "epoch": 0.04732927653820149,
+      "grad_norm": 215.9780731201172,
+      "learning_rate": 5.074626865671642e-06,
+      "loss": 16.4974,
+      "step": 35
+    },
+    {
+      "epoch": 0.054090601757944556,
+      "grad_norm": 223.51185607910156,
+      "learning_rate": 5.820895522388061e-06,
+      "loss": 15.188,
+      "step": 40
+    },
+    {
+      "epoch": 0.060851926977687626,
+      "grad_norm": 225.0447998046875,
+      "learning_rate": 6.567164179104478e-06,
+      "loss": 14.3741,
+      "step": 45
+    },
+    {
+      "epoch": 0.0676132521974307,
+      "grad_norm": 224.58688354492188,
+      "learning_rate": 7.313432835820896e-06,
+      "loss": 13.2594,
+      "step": 50
+    },
+    {
+      "epoch": 0.07437457741717377,
+      "grad_norm": 227.6384735107422,
+      "learning_rate": 8.059701492537314e-06,
+      "loss": 12.5621,
+      "step": 55
+    },
+    {
+      "epoch": 0.08113590263691683,
+      "grad_norm": 221.85934448242188,
+      "learning_rate": 8.805970149253732e-06,
+      "loss": 11.5079,
+      "step": 60
+    },
+    {
+      "epoch": 0.0878972278566599,
+      "grad_norm": 216.83883666992188,
+      "learning_rate": 9.552238805970149e-06,
+      "loss": 10.5443,
+      "step": 65
+    },
+    {
+      "epoch": 0.09465855307640297,
+      "grad_norm": 207.60470581054688,
+      "learning_rate": 9.999978708249054e-06,
+      "loss": 9.3134,
+      "step": 70
+    },
+    {
+      "epoch": 0.10141987829614604,
+      "grad_norm": 200.57904052734375,
+      "learning_rate": 9.999739178133424e-06,
+      "loss": 8.1168,
+      "step": 75
+    },
+    {
+      "epoch": 0.10818120351588911,
+      "grad_norm": 187.943603515625,
+      "learning_rate": 9.999233516005972e-06,
+      "loss": 7.0911,
+      "step": 80
+    },
+    {
+      "epoch": 0.11494252873563218,
+      "grad_norm": 162.5326385498047,
+      "learning_rate": 9.99846174878268e-06,
+      "loss": 6.1073,
+      "step": 85
+    },
+    {
+      "epoch": 0.12170385395537525,
+      "grad_norm": 120.48332214355469,
+      "learning_rate": 9.99742391754408e-06,
+      "loss": 5.2959,
+      "step": 90
+    },
+    {
+      "epoch": 0.12846517917511832,
+      "grad_norm": 75.04917907714844,
+      "learning_rate": 9.99612007753308e-06,
+      "loss": 4.682,
+      "step": 95
+    },
+    {
+      "epoch": 0.1352265043948614,
+      "grad_norm": 41.34440612792969,
+      "learning_rate": 9.994550298152012e-06,
+      "loss": 4.2617,
+      "step": 100
+    },
+    {
+      "epoch": 0.14198782961460446,
+      "grad_norm": 25.1535587310791,
+      "learning_rate": 9.992714662958944e-06,
+      "loss": 3.962,
+      "step": 105
+    },
+    {
+      "epoch": 0.14874915483434753,
+      "grad_norm": 15.239022254943848,
+      "learning_rate": 9.990613269663231e-06,
+      "loss": 3.789,
+      "step": 110
+    },
+    {
+      "epoch": 0.1555104800540906,
+      "grad_norm": 17.221874237060547,
+      "learning_rate": 9.988246230120311e-06,
+      "loss": 3.6555,
+      "step": 115
+    },
+    {
+      "epoch": 0.16227180527383367,
+      "grad_norm": 10.792862892150879,
+      "learning_rate": 9.985613670325759e-06,
+      "loss": 3.4862,
+      "step": 120
+    },
+    {
+      "epoch": 0.16903313049357674,
+      "grad_norm": 7.69727897644043,
+      "learning_rate": 9.982715730408568e-06,
+      "loss": 3.3919,
+      "step": 125
+    },
+    {
+      "epoch": 0.1757944557133198,
+      "grad_norm": 7.8089118003845215,
+      "learning_rate": 9.979552564623703e-06,
+      "loss": 3.3046,
+      "step": 130
+    },
+    {
+      "epoch": 0.18255578093306288,
+      "grad_norm": 8.153103828430176,
+      "learning_rate": 9.97612434134388e-06,
+      "loss": 3.2778,
+      "step": 135
+    },
+    {
+      "epoch": 0.18931710615280595,
+      "grad_norm": 6.2102484703063965,
+      "learning_rate": 9.972431243050611e-06,
+      "loss": 3.1411,
+      "step": 140
+    },
+    {
+      "epoch": 0.19607843137254902,
+      "grad_norm": 4.64615535736084,
+      "learning_rate": 9.968473466324486e-06,
+      "loss": 3.1302,
+      "step": 145
+    },
+    {
+      "epoch": 0.2028397565922921,
+      "grad_norm": 6.153476238250732,
+      "learning_rate": 9.964251221834712e-06,
+      "loss": 3.1216,
+      "step": 150
+    },
+    {
+      "epoch": 0.20960108181203516,
+      "grad_norm": 4.58518648147583,
+      "learning_rate": 9.959764734327894e-06,
+      "loss": 3.0453,
+      "step": 155
+    },
+    {
+      "epoch": 0.21636240703177823,
+      "grad_norm": 5.034542560577393,
+      "learning_rate": 9.955014242616083e-06,
+      "loss": 3.0575,
+      "step": 160
+    },
+    {
+      "epoch": 0.2231237322515213,
+      "grad_norm": 5.907400608062744,
+      "learning_rate": 9.949999999564049e-06,
+      "loss": 3.045,
+      "step": 165
+    },
+    {
+      "epoch": 0.22988505747126436,
+      "grad_norm": 5.186819553375244,
+      "learning_rate": 9.944722272075834e-06,
+      "loss": 3.0121,
+      "step": 170
+    },
+    {
+      "epoch": 0.23664638269100743,
+      "grad_norm": 5.540721416473389,
+      "learning_rate": 9.93918134108054e-06,
+      "loss": 2.9834,
+      "step": 175
+    },
+    {
+      "epoch": 0.2434077079107505,
+      "grad_norm": 5.014519214630127,
+      "learning_rate": 9.933377501517375e-06,
+      "loss": 2.9621,
+      "step": 180
+    },
+    {
+      "epoch": 0.2501690331304936,
+      "grad_norm": 5.306411266326904,
+      "learning_rate": 9.927311062319953e-06,
+      "loss": 2.9895,
+      "step": 185
+    },
+    {
+      "epoch": 0.25693035835023664,
+      "grad_norm": 3.703613042831421,
+      "learning_rate": 9.920982346399854e-06,
+      "loss": 2.9428,
+      "step": 190
+    },
+    {
+      "epoch": 0.26369168356997974,
+      "grad_norm": 4.594870567321777,
+      "learning_rate": 9.91439169062943e-06,
+      "loss": 2.9494,
+      "step": 195
+    },
+    {
+      "epoch": 0.2704530087897228,
+      "grad_norm": 3.66827130317688,
+      "learning_rate": 9.907539445823877e-06,
+      "loss": 2.9214,
+      "step": 200
+    },
+    {
+      "epoch": 0.2772143340094659,
+      "grad_norm": 4.361055850982666,
+      "learning_rate": 9.90042597672256e-06,
+      "loss": 2.9404,
+      "step": 205
+    },
+    {
+      "epoch": 0.2839756592292089,
+      "grad_norm": 4.3119916915893555,
+      "learning_rate": 9.893051661969599e-06,
+      "loss": 2.9174,
+      "step": 210
+    },
+    {
+      "epoch": 0.290736984448952,
+      "grad_norm": 3.507157325744629,
+      "learning_rate": 9.885416894093714e-06,
+      "loss": 2.9036,
+      "step": 215
+    },
+    {
+      "epoch": 0.29749830966869506,
+      "grad_norm": 4.4753522872924805,
+      "learning_rate": 9.877522079487334e-06,
+      "loss": 2.905,
+      "step": 220
+    },
+    {
+      "epoch": 0.30425963488843816,
+      "grad_norm": 3.661130666732788,
+      "learning_rate": 9.869367638384955e-06,
+      "loss": 2.904,
+      "step": 225
+    },
+    {
+      "epoch": 0.3110209601081812,
+      "grad_norm": 3.0701863765716553,
+      "learning_rate": 9.860954004840787e-06,
+      "loss": 2.8509,
+      "step": 230
+    },
+    {
+      "epoch": 0.3177822853279243,
+      "grad_norm": 3.502415895462036,
+      "learning_rate": 9.85228162670563e-06,
+      "loss": 2.8465,
+      "step": 235
+    },
+    {
+      "epoch": 0.32454361054766734,
+      "grad_norm": 3.181335210800171,
+      "learning_rate": 9.843350965603058e-06,
+      "loss": 2.859,
+      "step": 240
+    },
+    {
+      "epoch": 0.33130493576741044,
+      "grad_norm": 3.8409316539764404,
+      "learning_rate": 9.834162496904827e-06,
+      "loss": 2.8176,
+      "step": 245
+    },
+    {
+      "epoch": 0.3380662609871535,
+      "grad_norm": 3.99033522605896,
+      "learning_rate": 9.824716709705584e-06,
+      "loss": 2.7968,
+      "step": 250
+    },
+    {
+      "epoch": 0.3448275862068966,
+      "grad_norm": 3.110179901123047,
+      "learning_rate": 9.815014106796824e-06,
+      "loss": 2.8209,
+      "step": 255
+    },
+    {
+      "epoch": 0.3515889114266396,
+      "grad_norm": 3.054056167602539,
+      "learning_rate": 9.805055204640138e-06,
+      "loss": 2.8437,
+      "step": 260
+    },
+    {
+      "epoch": 0.3583502366463827,
+      "grad_norm": 3.178157091140747,
+      "learning_rate": 9.79484053333971e-06,
+      "loss": 2.808,
+      "step": 265
+    },
+    {
+      "epoch": 0.36511156186612576,
+      "grad_norm": 3.8207147121429443,
+      "learning_rate": 9.784370636614107e-06,
+      "loss": 2.7898,
+      "step": 270
+    },
+    {
+      "epoch": 0.37187288708586885,
+      "grad_norm": 6.010393142700195,
+      "learning_rate": 9.773646071767338e-06,
+      "loss": 2.7979,
+      "step": 275
+    },
+    {
+      "epoch": 0.3786342123056119,
+      "grad_norm": 3.476896047592163,
+      "learning_rate": 9.762667409659183e-06,
+      "loss": 2.7783,
+      "step": 280
+    },
+    {
+      "epoch": 0.385395537525355,
+      "grad_norm": 3.391853094100952,
+      "learning_rate": 9.751435234674813e-06,
+      "loss": 2.7708,
+      "step": 285
+    },
+    {
+      "epoch": 0.39215686274509803,
+      "grad_norm": 3.346642255783081,
+      "learning_rate": 9.739950144693681e-06,
+      "loss": 2.7825,
+      "step": 290
+    },
+    {
+      "epoch": 0.39891818796484113,
+      "grad_norm": 3.348356008529663,
+      "learning_rate": 9.728212751057701e-06,
+      "loss": 2.744,
+      "step": 295
+    },
+    {
+      "epoch": 0.4056795131845842,
+      "grad_norm": 3.324674129486084,
+      "learning_rate": 9.716223678538699e-06,
+      "loss": 2.7587,
+      "step": 300
+    },
+    {
+      "epoch": 0.41244083840432727,
+      "grad_norm": 3.0210790634155273,
+      "learning_rate": 9.70398356530516e-06,
+      "loss": 2.7832,
+      "step": 305
+    },
+    {
+      "epoch": 0.4192021636240703,
+      "grad_norm": 3.3017354011535645,
+      "learning_rate": 9.691493062888264e-06,
+      "loss": 2.7574,
+      "step": 310
+    },
+    {
+      "epoch": 0.4259634888438134,
+      "grad_norm": 2.9814178943634033,
+      "learning_rate": 9.6787528361472e-06,
+      "loss": 2.7481,
+      "step": 315
+    },
+    {
+      "epoch": 0.43272481406355645,
+      "grad_norm": 2.686392307281494,
+      "learning_rate": 9.66576356323378e-06,
+      "loss": 2.7586,
+      "step": 320
+    },
+    {
+      "epoch": 0.43948613928329955,
+      "grad_norm": 3.3055434226989746,
+      "learning_rate": 9.652525935556335e-06,
+      "loss": 2.735,
+      "step": 325
+    },
+    {
+      "epoch": 0.4462474645030426,
+      "grad_norm": 3.388873338699341,
+      "learning_rate": 9.63904065774292e-06,
+      "loss": 2.758,
+      "step": 330
+    },
+    {
+      "epoch": 0.4530087897227857,
+      "grad_norm": 3.3133444786071777,
+      "learning_rate": 9.625308447603801e-06,
+      "loss": 2.734,
+      "step": 335
+    },
+    {
+      "epoch": 0.45977011494252873,
+      "grad_norm": 2.6662232875823975,
+      "learning_rate": 9.611330036093252e-06,
+      "loss": 2.755,
+      "step": 340
+    },
+    {
+      "epoch": 0.4665314401622718,
+      "grad_norm": 3.3504395484924316,
+      "learning_rate": 9.59710616727064e-06,
+      "loss": 2.7532,
+      "step": 345
+    },
+    {
+      "epoch": 0.47329276538201487,
+      "grad_norm": 3.4668731689453125,
+      "learning_rate": 9.582637598260824e-06,
+      "loss": 2.7325,
+      "step": 350
+    },
+    {
+      "epoch": 0.48005409060175797,
+      "grad_norm": 3.009427785873413,
+      "learning_rate": 9.567925099213857e-06,
+      "loss": 2.7179,
+      "step": 355
+    },
+    {
+      "epoch": 0.486815415821501,
+      "grad_norm": 3.1236302852630615,
+      "learning_rate": 9.552969453263983e-06,
+      "loss": 2.7613,
+      "step": 360
+    },
+    {
+      "epoch": 0.4935767410412441,
+      "grad_norm": 3.0972957611083984,
+      "learning_rate": 9.537771456487959e-06,
+      "loss": 2.7275,
+      "step": 365
+    },
+    {
+      "epoch": 0.5003380662609872,
+      "grad_norm": 2.5772786140441895,
+      "learning_rate": 9.522331917862674e-06,
+      "loss": 2.7529,
+      "step": 370
+    },
+    {
+      "epoch": 0.5070993914807302,
+      "grad_norm": 3.032407760620117,
+      "learning_rate": 9.506651659222096e-06,
+      "loss": 2.7065,
+      "step": 375
+    },
+    {
+      "epoch": 0.5138607167004733,
+      "grad_norm": 2.7317681312561035,
+      "learning_rate": 9.49073151521352e-06,
+      "loss": 2.7036,
+      "step": 380
+    },
+    {
+      "epoch": 0.5206220419202163,
+      "grad_norm": 3.070013999938965,
+      "learning_rate": 9.47457233325314e-06,
+      "loss": 2.705,
+      "step": 385
+    },
+    {
+      "epoch": 0.5273833671399595,
+      "grad_norm": 3.1261889934539795,
+      "learning_rate": 9.458174973480945e-06,
+      "loss": 2.7006,
+      "step": 390
+    },
+    {
+      "epoch": 0.5341446923597025,
+      "grad_norm": 2.8663241863250732,
+      "learning_rate": 9.441540308714934e-06,
+      "loss": 2.7197,
+      "step": 395
+    },
+    {
+      "epoch": 0.5409060175794456,
+      "grad_norm": 3.94987416267395,
+      "learning_rate": 9.424669224404659e-06,
+      "loss": 2.7162,
+      "step": 400
+    },
+    {
+      "epoch": 0.5476673427991886,
+      "grad_norm": 2.7462100982666016,
+      "learning_rate": 9.407562618584083e-06,
+      "loss": 2.662,
+      "step": 405
+    },
+    {
+      "epoch": 0.5544286680189318,
+      "grad_norm": 2.9217984676361084,
+      "learning_rate": 9.390221401823793e-06,
+      "loss": 2.6988,
+      "step": 410
+    },
+    {
+      "epoch": 0.5611899932386748,
+      "grad_norm": 3.6465213298797607,
+      "learning_rate": 9.372646497182519e-06,
+      "loss": 2.6772,
+      "step": 415
+    },
+    {
+      "epoch": 0.5679513184584178,
+      "grad_norm": 3.873567581176758,
+      "learning_rate": 9.354838840158006e-06,
+      "loss": 2.6699,
+      "step": 420
+    },
+    {
+      "epoch": 0.5747126436781609,
+      "grad_norm": 3.826775550842285,
+      "learning_rate": 9.33679937863722e-06,
+      "loss": 2.6944,
+      "step": 425
+    },
+    {
+      "epoch": 0.581473968897904,
+      "grad_norm": 3.7498714923858643,
+      "learning_rate": 9.318529072845885e-06,
+      "loss": 2.6897,
+      "step": 430
+    },
+    {
+      "epoch": 0.5882352941176471,
+      "grad_norm": 3.4026408195495605,
+      "learning_rate": 9.300028895297382e-06,
+      "loss": 2.6942,
+      "step": 435
+    },
+    {
+      "epoch": 0.5949966193373901,
+      "grad_norm": 2.911181688308716,
+      "learning_rate": 9.281299830740973e-06,
+      "loss": 2.6993,
+      "step": 440
+    },
+    {
+      "epoch": 0.6017579445571332,
+      "grad_norm": 3.125908613204956,
+      "learning_rate": 9.262342876109387e-06,
+      "loss": 2.6688,
+      "step": 445
+    },
+    {
+      "epoch": 0.6085192697768763,
+      "grad_norm": 2.82594895362854,
+      "learning_rate": 9.243159040465757e-06,
+      "loss": 2.6831,
+      "step": 450
+    },
+    {
+      "epoch": 0.6152805949966194,
+      "grad_norm": 3.2787418365478516,
+      "learning_rate": 9.223749344949904e-06,
+      "loss": 2.66,
+      "step": 455
+    },
+    {
+      "epoch": 0.6220419202163624,
+      "grad_norm": 4.373532772064209,
+      "learning_rate": 9.204114822723986e-06,
+      "loss": 2.6689,
+      "step": 460
+    },
+    {
+      "epoch": 0.6288032454361054,
+      "grad_norm": 2.8605294227600098,
+      "learning_rate": 9.184256518917504e-06,
+      "loss": 2.6509,
+      "step": 465
+    },
+    {
+      "epoch": 0.6355645706558486,
+      "grad_norm": 4.4151387214660645,
+      "learning_rate": 9.164175490571663e-06,
+      "loss": 2.6912,
+      "step": 470
+    },
+    {
+      "epoch": 0.6423258958755916,
+      "grad_norm": 2.953627824783325,
+      "learning_rate": 9.143872806583118e-06,
+      "loss": 2.6415,
+      "step": 475
+    },
+    {
+      "epoch": 0.6490872210953347,
+      "grad_norm": 2.935363531112671,
+      "learning_rate": 9.12334954764707e-06,
+      "loss": 2.6751,
+      "step": 480
+    },
+    {
+      "epoch": 0.6558485463150777,
+      "grad_norm": 2.8612430095672607,
+      "learning_rate": 9.102606806199743e-06,
+      "loss": 2.6564,
+      "step": 485
+    },
+    {
+      "epoch": 0.6626098715348209,
+      "grad_norm": 2.89137864112854,
+      "learning_rate": 9.081645686360234e-06,
+      "loss": 2.662,
+      "step": 490
+    },
+    {
+      "epoch": 0.6693711967545639,
+      "grad_norm": 3.100198745727539,
+      "learning_rate": 9.060467303871746e-06,
+      "loss": 2.6656,
+      "step": 495
+    },
+    {
+      "epoch": 0.676132521974307,
+      "grad_norm": 2.7693257331848145,
+      "learning_rate": 9.039072786042187e-06,
+      "loss": 2.6409,
+      "step": 500
+    },
+    {
+      "epoch": 0.68289384719405,
+      "grad_norm": 3.0061259269714355,
+      "learning_rate": 9.017463271684183e-06,
+      "loss": 2.6472,
+      "step": 505
+    },
+    {
+      "epoch": 0.6896551724137931,
+      "grad_norm": 3.591599941253662,
+      "learning_rate": 8.995639911054437e-06,
+      "loss": 2.6488,
+      "step": 510
+    },
+    {
+      "epoch": 0.6964164976335362,
+      "grad_norm": 3.120270252227783,
+      "learning_rate": 8.973603865792525e-06,
+      "loss": 2.6648,
+      "step": 515
+    },
+    {
+      "epoch": 0.7031778228532792,
+      "grad_norm": 3.2259232997894287,
+      "learning_rate": 8.951356308859041e-06,
+      "loss": 2.6392,
+      "step": 520
+    },
+    {
+      "epoch": 0.7099391480730223,
+      "grad_norm": 2.992349624633789,
+      "learning_rate": 8.928898424473178e-06,
+      "loss": 2.6525,
+      "step": 525
+    },
+    {
+      "epoch": 0.7167004732927654,
+      "grad_norm": 2.7220916748046875,
+      "learning_rate": 8.906231408049682e-06,
+      "loss": 2.6146,
+      "step": 530
+    },
+    {
+      "epoch": 0.7234617985125085,
+      "grad_norm": 2.8423595428466797,
+      "learning_rate": 8.883356466135231e-06,
+      "loss": 2.6213,
+      "step": 535
+    },
+    {
+      "epoch": 0.7302231237322515,
+      "grad_norm": 2.929335594177246,
+      "learning_rate": 8.860274816344203e-06,
+      "loss": 2.6185,
+      "step": 540
+    },
+    {
+      "epoch": 0.7369844489519946,
+      "grad_norm": 2.9922103881835938,
+      "learning_rate": 8.836987687293868e-06,
+      "loss": 2.6042,
+      "step": 545
+    },
+    {
+      "epoch": 0.7437457741717377,
+      "grad_norm": 3.488055467605591,
+      "learning_rate": 8.813496318538985e-06,
+      "loss": 2.6443,
+      "step": 550
+    },
+    {
+      "epoch": 0.7505070993914807,
+      "grad_norm": 7.2074785232543945,
+      "learning_rate": 8.789801960505826e-06,
+      "loss": 2.6226,
+      "step": 555
+    },
+    {
+      "epoch": 0.7572684246112238,
+      "grad_norm": 3.040048360824585,
+      "learning_rate": 8.765905874425622e-06,
+      "loss": 2.6449,
+      "step": 560
+    },
+    {
+      "epoch": 0.7640297498309668,
+      "grad_norm": 2.8889758586883545,
+      "learning_rate": 8.741809332267413e-06,
+      "loss": 2.6391,
+      "step": 565
+    },
+    {
+      "epoch": 0.77079107505071,
+      "grad_norm": 3.0062615871429443,
+      "learning_rate": 8.717513616670357e-06,
+      "loss": 2.6383,
+      "step": 570
+    },
+    {
+      "epoch": 0.777552400270453,
+      "grad_norm": 2.8223934173583984,
+      "learning_rate": 8.693020020875448e-06,
+      "loss": 2.611,
+      "step": 575
+    },
+    {
+      "epoch": 0.7843137254901961,
+      "grad_norm": 2.990478754043579,
+      "learning_rate": 8.668329848656683e-06,
+      "loss": 2.6141,
+      "step": 580
+    },
+    {
+      "epoch": 0.7910750507099391,
+      "grad_norm": 2.6774513721466064,
+      "learning_rate": 8.643444414251659e-06,
+      "loss": 2.6101,
+      "step": 585
+    },
+    {
+      "epoch": 0.7978363759296823,
+      "grad_norm": 2.859231472015381,
+      "learning_rate": 8.61836504229162e-06,
+      "loss": 2.612,
+      "step": 590
+    },
+    {
+      "epoch": 0.8045977011494253,
+      "grad_norm": 4.325037479400635,
+      "learning_rate": 8.593093067730941e-06,
+      "loss": 2.6328,
+      "step": 595
+    },
+    {
+      "epoch": 0.8113590263691683,
+      "grad_norm": 2.8452208042144775,
+      "learning_rate": 8.56762983577609e-06,
+      "loss": 2.6006,
+      "step": 600
+    },
+    {
+      "epoch": 0.8181203515889114,
+      "grad_norm": 3.683469772338867,
+      "learning_rate": 8.54197670181399e-06,
+      "loss": 2.6527,
+      "step": 605
+    },
+    {
+      "epoch": 0.8248816768086545,
+      "grad_norm": 3.5392167568206787,
+      "learning_rate": 8.516135031339908e-06,
+      "loss": 2.6207,
+      "step": 610
+    },
+    {
+      "epoch": 0.8316430020283976,
+      "grad_norm": 2.7016098499298096,
+      "learning_rate": 8.490106199884742e-06,
+      "loss": 2.6212,
+      "step": 615
+    },
+    {
+      "epoch": 0.8384043272481406,
+      "grad_norm": 3.4844460487365723,
+      "learning_rate": 8.463891592941827e-06,
+      "loss": 2.6115,
+      "step": 620
+    },
+    {
+      "epoch": 0.8451656524678837,
+      "grad_norm": 2.874047040939331,
+      "learning_rate": 8.437492605893167e-06,
+      "loss": 2.5725,
+      "step": 625
+    },
+    {
+      "epoch": 0.8519269776876268,
+      "grad_norm": 3.063899040222168,
+      "learning_rate": 8.41091064393517e-06,
+      "loss": 2.6075,
+      "step": 630
+    },
+    {
+      "epoch": 0.8586883029073699,
+      "grad_norm": 2.8391199111938477,
+      "learning_rate": 8.38414712200385e-06,
+      "loss": 2.6123,
+      "step": 635
+    },
+    {
+      "epoch": 0.8654496281271129,
+      "grad_norm": 3.388949394226074,
+      "learning_rate": 8.357203464699502e-06,
+      "loss": 2.5755,
+      "step": 640
+    },
+    {
+      "epoch": 0.8722109533468559,
+      "grad_norm": 3.038212776184082,
+      "learning_rate": 8.330081106210889e-06,
+      "loss": 2.6086,
+      "step": 645
+    },
+    {
+      "epoch": 0.8789722785665991,
+      "grad_norm": 4.426064968109131,
+      "learning_rate": 8.302781490238888e-06,
+      "loss": 2.5918,
+      "step": 650
+    },
+    {
+      "epoch": 0.8857336037863421,
+      "grad_norm": 2.6291635036468506,
+      "learning_rate": 8.275306069919639e-06,
+      "loss": 2.5939,
+      "step": 655
+    },
+    {
+      "epoch": 0.8924949290060852,
+      "grad_norm": 2.598675489425659,
+      "learning_rate": 8.247656307747214e-06,
+      "loss": 2.5961,
+      "step": 660
+    },
+    {
+      "epoch": 0.8992562542258282,
+      "grad_norm": 2.9545516967773438,
+      "learning_rate": 8.219833675495754e-06,
+      "loss": 2.5854,
+      "step": 665
+    },
+    {
+      "epoch": 0.9060175794455714,
+      "grad_norm": 3.168442726135254,
+      "learning_rate": 8.19183965414113e-06,
+      "loss": 2.5691,
+      "step": 670
+    },
+    {
+      "epoch": 0.9127789046653144,
+      "grad_norm": 2.742135763168335,
+      "learning_rate": 8.163675733782114e-06,
+      "loss": 2.5574,
+      "step": 675
+    },
+    {
+      "epoch": 0.9195402298850575,
+      "grad_norm": 2.5857961177825928,
+      "learning_rate": 8.135343413561073e-06,
+      "loss": 2.591,
+      "step": 680
+    },
+    {
+      "epoch": 0.9263015551048005,
+      "grad_norm": 2.525702714920044,
+      "learning_rate": 8.106844201584145e-06,
+      "loss": 2.5952,
+      "step": 685
+    },
+    {
+      "epoch": 0.9330628803245437,
+      "grad_norm": 4.6076860427856445,
+      "learning_rate": 8.07817961484099e-06,
+      "loss": 2.5904,
+      "step": 690
+    },
+    {
+      "epoch": 0.9398242055442867,
+      "grad_norm": 4.223016738891602,
+      "learning_rate": 8.049351179124028e-06,
+      "loss": 2.5729,
+      "step": 695
+    },
+    {
+      "epoch": 0.9465855307640297,
+      "grad_norm": 3.6000704765319824,
+      "learning_rate": 8.020360428947223e-06,
+      "loss": 2.577,
+      "step": 700
+    },
+    {
+      "epoch": 0.9533468559837728,
+      "grad_norm": 2.7744557857513428,
+      "learning_rate": 7.991208907464405e-06,
+      "loss": 2.6169,
+      "step": 705
+    },
+    {
+      "epoch": 0.9601081812035159,
+      "grad_norm": 2.9555327892303467,
+      "learning_rate": 7.961898166387136e-06,
+      "loss": 2.5857,
+      "step": 710
+    },
+    {
+      "epoch": 0.966869506423259,
+      "grad_norm": 3.189509153366089,
+      "learning_rate": 7.932429765902095e-06,
+      "loss": 2.5515,
+      "step": 715
+    },
+    {
+      "epoch": 0.973630831643002,
+      "grad_norm": 3.0023581981658936,
+      "learning_rate": 7.902805274588048e-06,
+      "loss": 2.5478,
+      "step": 720
+    },
+    {
+      "epoch": 0.9803921568627451,
+      "grad_norm": 2.8606345653533936,
+      "learning_rate": 7.873026269332349e-06,
+      "loss": 2.5681,
+      "step": 725
+    },
+    {
+      "epoch": 0.9871534820824882,
+      "grad_norm": 3.11521053314209,
+      "learning_rate": 7.843094335246997e-06,
+      "loss": 2.5664,
+      "step": 730
+    },
+    {
+      "epoch": 0.9939148073022313,
+      "grad_norm": 2.513296604156494,
+      "learning_rate": 7.813011065584273e-06,
+      "loss": 2.5622,
+      "step": 735
+    },
+    {
+      "epoch": 1.0,
+      "grad_norm": 2.6706948280334473,
+      "learning_rate": 7.782778061651922e-06,
+      "loss": 2.3027,
+      "step": 740
+    },
+    {
+      "epoch": 1.0067613252197432,
+      "grad_norm": 2.6789958477020264,
+      "learning_rate": 7.752396932727926e-06,
+      "loss": 2.5502,
+      "step": 745
+    },
+    {
+      "epoch": 1.013522650439486,
+      "grad_norm": 2.6674933433532715,
+      "learning_rate": 7.721869295974831e-06,
+      "loss": 2.5365,
+      "step": 750
+    },
+    {
+      "epoch": 1.0202839756592292,
+      "grad_norm": 2.80924654006958,
+      "learning_rate": 7.691196776353681e-06,
+      "loss": 2.4935,
+      "step": 755
+    },
+    {
+      "epoch": 1.0270453008789722,
+      "grad_norm": 2.9341821670532227,
+      "learning_rate": 7.660381006537516e-06,
+      "loss": 2.548,
+      "step": 760
+    },
+    {
+      "epoch": 1.0338066260987153,
+      "grad_norm": 2.9331564903259277,
+      "learning_rate": 7.629423626824462e-06,
+      "loss": 2.5119,
+      "step": 765
+    },
+    {
+      "epoch": 1.0405679513184585,
+      "grad_norm": 3.1369872093200684,
+      "learning_rate": 7.598326285050429e-06,
+      "loss": 2.5118,
+      "step": 770
+    },
+    {
+      "epoch": 1.0473292765382014,
+      "grad_norm": 3.999962329864502,
+      "learning_rate": 7.567090636501386e-06,
+      "loss": 2.555,
+      "step": 775
+    },
+    {
+      "epoch": 1.0540906017579446,
+      "grad_norm": 2.822645902633667,
+      "learning_rate": 7.535718343825267e-06,
+      "loss": 2.5165,
+      "step": 780
+    },
+    {
+      "epoch": 1.0608519269776877,
+      "grad_norm": 2.6943705081939697,
+      "learning_rate": 7.504211076943448e-06,
+      "loss": 2.485,
+      "step": 785
+    },
+    {
+      "epoch": 1.0676132521974306,
+      "grad_norm": 2.860985040664673,
+      "learning_rate": 7.4725705129618855e-06,
+      "loss": 2.5193,
+      "step": 790
+    },
+    {
+      "epoch": 1.0743745774171738,
+      "grad_norm": 2.7031378746032715,
+      "learning_rate": 7.440798336081821e-06,
+      "loss": 2.511,
+      "step": 795
+    },
+    {
+      "epoch": 1.081135902636917,
+      "grad_norm": 3.222702741622925,
+      "learning_rate": 7.408896237510146e-06,
+      "loss": 2.5307,
+      "step": 800
+    },
+    {
+      "epoch": 1.0878972278566599,
+      "grad_norm": 3.502304792404175,
+      "learning_rate": 7.376865915369378e-06,
+      "loss": 2.5419,
+      "step": 805
+    },
+    {
+      "epoch": 1.094658553076403,
+      "grad_norm": 2.73856520652771,
+      "learning_rate": 7.3447090746072666e-06,
+      "loss": 2.5162,
+      "step": 810
+    },
+    {
+      "epoch": 1.101419878296146,
+      "grad_norm": 2.9833099842071533,
+      "learning_rate": 7.312427426906048e-06,
+      "loss": 2.5051,
+      "step": 815
+    },
+    {
+      "epoch": 1.1081812035158891,
+      "grad_norm": 3.2921409606933594,
+      "learning_rate": 7.280022690591326e-06,
+      "loss": 2.5362,
+      "step": 820
+    },
+    {
+      "epoch": 1.1149425287356323,
+      "grad_norm": 3.6853201389312744,
+      "learning_rate": 7.247496590540612e-06,
+      "loss": 2.4959,
+      "step": 825
+    },
+    {
+      "epoch": 1.1217038539553752,
+      "grad_norm": 3.182521343231201,
+      "learning_rate": 7.214850858091509e-06,
+      "loss": 2.4962,
+      "step": 830
+    },
+    {
+      "epoch": 1.1284651791751183,
+      "grad_norm": 6.297270774841309,
+      "learning_rate": 7.1820872309495534e-06,
+      "loss": 2.5281,
+      "step": 835
+    },
+    {
+      "epoch": 1.1352265043948613,
+      "grad_norm": 3.743927240371704,
+      "learning_rate": 7.14920745309572e-06,
+      "loss": 2.5215,
+      "step": 840
+    },
+    {
+      "epoch": 1.1419878296146044,
+      "grad_norm": 2.9408071041107178,
+      "learning_rate": 7.116213274693591e-06,
+      "loss": 2.4759,
+      "step": 845
+    },
+    {
+      "epoch": 1.1487491548343476,
+      "grad_norm": 2.461620330810547,
+      "learning_rate": 7.083106451996194e-06,
+      "loss": 2.545,
+      "step": 850
+    },
+    {
+      "epoch": 1.1555104800540905,
+      "grad_norm": 3.0902323722839355,
+      "learning_rate": 7.049888747252524e-06,
+      "loss": 2.4778,
+      "step": 855
+    },
+    {
+      "epoch": 1.1622718052738337,
+      "grad_norm": 2.803609848022461,
+      "learning_rate": 7.016561928613733e-06,
+      "loss": 2.4779,
+      "step": 860
+    },
+    {
+      "epoch": 1.1690331304935768,
+      "grad_norm": 3.166276216506958,
+      "learning_rate": 6.983127770039017e-06,
+      "loss": 2.5279,
+      "step": 865
+    },
+    {
+      "epoch": 1.1757944557133198,
+      "grad_norm": 3.0823771953582764,
+      "learning_rate": 6.949588051201187e-06,
+      "loss": 2.5216,
+      "step": 870
+    },
+    {
+      "epoch": 1.182555780933063,
+      "grad_norm": 3.188164472579956,
+      "learning_rate": 6.915944557391942e-06,
+      "loss": 2.5035,
+      "step": 875
+    },
+    {
+      "epoch": 1.189317106152806,
+      "grad_norm": 2.577242612838745,
+      "learning_rate": 6.882199079426842e-06,
+      "loss": 2.4836,
+      "step": 880
+    },
+    {
+      "epoch": 1.196078431372549,
+      "grad_norm": 2.9939839839935303,
+      "learning_rate": 6.8483534135499665e-06,
+      "loss": 2.5142,
+      "step": 885
+    },
+    {
+      "epoch": 1.2028397565922921,
+      "grad_norm": 2.834369421005249,
+      "learning_rate": 6.814409361338331e-06,
+      "loss": 2.5024,
+      "step": 890
+    },
+    {
+      "epoch": 1.209601081812035,
+      "grad_norm": 2.8373208045959473,
+      "learning_rate": 6.780368729605964e-06,
+      "loss": 2.5255,
+      "step": 895
+    },
+    {
+      "epoch": 1.2163624070317782,
+      "grad_norm": 3.0683062076568604,
+      "learning_rate": 6.746233330307748e-06,
+      "loss": 2.5038,
+      "step": 900
+    },
+    {
+      "epoch": 1.2231237322515214,
+      "grad_norm": 2.5889170169830322,
+      "learning_rate": 6.712004980442964e-06,
+      "loss": 2.5026,
+      "step": 905
+    },
+    {
+      "epoch": 1.2298850574712643,
+      "grad_norm": 3.1007654666900635,
+      "learning_rate": 6.677685501958572e-06,
+      "loss": 2.5123,
+      "step": 910
+    },
+    {
+      "epoch": 1.2366463826910075,
+      "grad_norm": 2.7331902980804443,
+      "learning_rate": 6.64327672165224e-06,
+      "loss": 2.5092,
+      "step": 915
+    },
+    {
+      "epoch": 1.2434077079107504,
+      "grad_norm": 2.818505048751831,
+      "learning_rate": 6.608780471075092e-06,
+      "loss": 2.4713,
+      "step": 920
+    },
+    {
+      "epoch": 1.2501690331304935,
+      "grad_norm": 3.1065335273742676,
+      "learning_rate": 6.574198586434229e-06,
+      "loss": 2.519,
+      "step": 925
+    },
+    {
+      "epoch": 1.2569303583502367,
+      "grad_norm": 3.60119366645813,
+      "learning_rate": 6.5395329084949774e-06,
+      "loss": 2.4981,
+      "step": 930
+    },
+    {
+      "epoch": 1.2636916835699799,
+      "grad_norm": 2.812912702560425,
+      "learning_rate": 6.504785282482917e-06,
+      "loss": 2.5046,
+      "step": 935
+    },
+    {
+      "epoch": 1.2704530087897228,
+      "grad_norm": 2.560412645339966,
+      "learning_rate": 6.46995755798565e-06,
+      "loss": 2.4667,
+      "step": 940
+    },
+    {
+      "epoch": 1.277214334009466,
+      "grad_norm": 2.716543436050415,
+      "learning_rate": 6.43505158885436e-06,
+      "loss": 2.483,
+      "step": 945
+    },
+    {
+      "epoch": 1.2839756592292089,
+      "grad_norm": 2.5889837741851807,
+      "learning_rate": 6.400069233105132e-06,
+      "loss": 2.474,
+      "step": 950
+    },
+    {
+      "epoch": 1.290736984448952,
+      "grad_norm": 2.548285722732544,
+      "learning_rate": 6.3650123528200354e-06,
+      "loss": 2.4902,
+      "step": 955
+    },
+    {
+      "epoch": 1.2974983096686952,
+      "grad_norm": 2.6211283206939697,
+      "learning_rate": 6.3298828140480326e-06,
+      "loss": 2.4542,
+      "step": 960
+    },
+    {
+      "epoch": 1.304259634888438,
+      "grad_norm": 2.9109580516815186,
+      "learning_rate": 6.29468248670563e-06,
+      "loss": 2.4925,
+      "step": 965
+    },
+    {
+      "epoch": 1.3110209601081813,
+      "grad_norm": 2.9635941982269287,
+      "learning_rate": 6.259413244477355e-06,
+      "loss": 2.4944,
+      "step": 970
+    },
+    {
+      "epoch": 1.3177822853279242,
+      "grad_norm": 2.7843010425567627,
+      "learning_rate": 6.224076964716015e-06,
+      "loss": 2.4943,
+      "step": 975
+    },
+    {
+      "epoch": 1.3245436105476673,
+      "grad_norm": 2.9173049926757812,
+      "learning_rate": 6.188675528342772e-06,
+      "loss": 2.4869,
+      "step": 980
+    },
+    {
+      "epoch": 1.3313049357674105,
+      "grad_norm": 2.5200576782226562,
+      "learning_rate": 6.153210819747022e-06,
+      "loss": 2.4867,
+      "step": 985
+    },
+    {
+      "epoch": 1.3380662609871534,
+      "grad_norm": 2.5795679092407227,
+      "learning_rate": 6.117684726686086e-06,
+      "loss": 2.5,
+      "step": 990
+    },
+    {
+      "epoch": 1.3448275862068966,
+      "grad_norm": 2.6581368446350098,
+      "learning_rate": 6.082099140184734e-06,
+      "loss": 2.4833,
+      "step": 995
+    },
+    {
+      "epoch": 1.3515889114266395,
+      "grad_norm": 3.1690869331359863,
+      "learning_rate": 6.046455954434518e-06,
+      "loss": 2.4825,
+      "step": 1000
+    },
+    {
+      "epoch": 1.3583502366463827,
+      "grad_norm": 2.5110387802124023,
+      "learning_rate": 6.010757066692957e-06,
+      "loss": 2.4304,
+      "step": 1005
+    },
+    {
+      "epoch": 1.3651115618661258,
+      "grad_norm": 3.335712194442749,
+      "learning_rate": 5.975004377182536e-06,
+      "loss": 2.477,
+      "step": 1010
+    },
+    {
+      "epoch": 1.371872887085869,
+      "grad_norm": 2.6729347705841064,
+      "learning_rate": 5.939199788989566e-06,
+      "loss": 2.4795,
+      "step": 1015
+    },
+    {
+      "epoch": 1.378634212305612,
+      "grad_norm": 3.0427603721618652,
+      "learning_rate": 5.903345207962881e-06,
+      "loss": 2.4701,
+      "step": 1020
+    },
+    {
+      "epoch": 1.385395537525355,
+      "grad_norm": 3.156398057937622,
+      "learning_rate": 5.867442542612394e-06,
+      "loss": 2.4835,
+      "step": 1025
+    },
+    {
+      "epoch": 1.392156862745098,
+      "grad_norm": 2.833634614944458,
+      "learning_rate": 5.831493704007507e-06,
+      "loss": 2.4716,
+      "step": 1030
+    },
+    {
+      "epoch": 1.3989181879648411,
+      "grad_norm": 2.794393539428711,
+      "learning_rate": 5.795500605675388e-06,
+      "loss": 2.5017,
+      "step": 1035
+    },
+    {
+      "epoch": 1.4056795131845843,
+      "grad_norm": 3.1075077056884766,
+      "learning_rate": 5.7594651634991095e-06,
+      "loss": 2.469,
+      "step": 1040
+    },
+    {
+      "epoch": 1.4124408384043272,
+      "grad_norm": 2.718369483947754,
+      "learning_rate": 5.723389295615676e-06,
+      "loss": 2.4749,
+      "step": 1045
+    },
+    {
+      "epoch": 1.4192021636240704,
+      "grad_norm": 2.5099501609802246,
+      "learning_rate": 5.687274922313915e-06,
+      "loss": 2.4739,
+      "step": 1050
+    },
+    {
+      "epoch": 1.4259634888438133,
+      "grad_norm": 2.688002347946167,
+      "learning_rate": 5.651123965932271e-06,
+      "loss": 2.4275,
+      "step": 1055
+    },
+    {
+      "epoch": 1.4327248140635565,
+      "grad_norm": 2.4511373043060303,
+      "learning_rate": 5.614938350756475e-06,
+      "loss": 2.4714,
+      "step": 1060
+    },
+    {
+      "epoch": 1.4394861392832996,
+      "grad_norm": 2.583116054534912,
+      "learning_rate": 5.5787200029171105e-06,
+      "loss": 2.4539,
+      "step": 1065
+    },
+    {
+      "epoch": 1.4462474645030425,
+      "grad_norm": 2.6514389514923096,
+      "learning_rate": 5.542470850287105e-06,
+      "loss": 2.4457,
+      "step": 1070
+    },
+    {
+      "epoch": 1.4530087897227857,
+      "grad_norm": 2.9797661304473877,
+      "learning_rate": 5.506192822379088e-06,
+      "loss": 2.4713,
+      "step": 1075
+    },
+    {
+      "epoch": 1.4597701149425286,
+      "grad_norm": 2.4024760723114014,
+      "learning_rate": 5.469887850242707e-06,
+      "loss": 2.4573,
+      "step": 1080
+    },
+    {
+      "epoch": 1.4665314401622718,
+      "grad_norm": 5.054728031158447,
+      "learning_rate": 5.433557866361818e-06,
+      "loss": 2.4719,
+      "step": 1085
+    },
+    {
+      "epoch": 1.473292765382015,
+      "grad_norm": 2.6720123291015625,
+      "learning_rate": 5.397204804551642e-06,
+      "loss": 2.4781,
+      "step": 1090
+    },
+    {
+      "epoch": 1.480054090601758,
+      "grad_norm": 2.962134838104248,
+      "learning_rate": 5.360830599855812e-06,
+      "loss": 2.4477,
+      "step": 1095
+    },
+    {
+      "epoch": 1.486815415821501,
+      "grad_norm": 2.5753014087677,
+      "learning_rate": 5.324437188443381e-06,
+      "loss": 2.447,
+      "step": 1100
+    },
+    {
+      "epoch": 1.4935767410412442,
+      "grad_norm": 2.300673007965088,
+      "learning_rate": 5.288026507505765e-06,
+      "loss": 2.4847,
+      "step": 1105
+    },
+    {
+      "epoch": 1.500338066260987,
+      "grad_norm": 2.465369701385498,
+      "learning_rate": 5.2516004951536124e-06,
+      "loss": 2.4708,
+      "step": 1110
+    },
+    {
+      "epoch": 1.5070993914807302,
+      "grad_norm": 2.3910348415374756,
+      "learning_rate": 5.21516109031366e-06,
+      "loss": 2.4802,
+      "step": 1115
+    },
+    {
+      "epoch": 1.5138607167004734,
+      "grad_norm": 3.5717501640319824,
+      "learning_rate": 5.178710232625511e-06,
+      "loss": 2.4703,
+      "step": 1120
+    },
+    {
+      "epoch": 1.5206220419202163,
+      "grad_norm": 2.6131703853607178,
+      "learning_rate": 5.142249862338395e-06,
+      "loss": 2.4493,
+      "step": 1125
+    },
+    {
+      "epoch": 1.5273833671399595,
+      "grad_norm": 2.7011117935180664,
+      "learning_rate": 5.10578192020789e-06,
+      "loss": 2.4647,
+      "step": 1130
+    },
+    {
+      "epoch": 1.5341446923597024,
+      "grad_norm": 2.7816450595855713,
+      "learning_rate": 5.069308347392614e-06,
+      "loss": 2.4587,
+      "step": 1135
+    },
+    {
+      "epoch": 1.5409060175794456,
+      "grad_norm": 3.070551633834839,
+      "learning_rate": 5.0328310853509074e-06,
+      "loss": 2.4518,
+      "step": 1140
+    },
+    {
+      "epoch": 1.5476673427991887,
+      "grad_norm": 2.793508291244507,
+      "learning_rate": 4.996352075737481e-06,
+      "loss": 2.4641,
+      "step": 1145
+    },
+    {
+      "epoch": 1.5544286680189319,
+      "grad_norm": 2.4779341220855713,
+      "learning_rate": 4.95987326030007e-06,
+      "loss": 2.4686,
+      "step": 1150
+    },
+    {
+      "epoch": 1.5611899932386748,
+      "grad_norm": 3.4994513988494873,
+      "learning_rate": 4.923396580776068e-06,
+      "loss": 2.4357,
+      "step": 1155
+    },
+    {
+      "epoch": 1.5679513184584177,
+      "grad_norm": 5.211581230163574,
+      "learning_rate": 4.886923978789183e-06,
+      "loss": 2.4893,
+      "step": 1160
+    },
+    {
+      "epoch": 1.5747126436781609,
+      "grad_norm": 2.783217430114746,
+      "learning_rate": 4.850457395746074e-06,
+      "loss": 2.4364,
+      "step": 1165
+    },
+    {
+      "epoch": 1.581473968897904,
+      "grad_norm": 2.6504037380218506,
+      "learning_rate": 4.813998772733022e-06,
+      "loss": 2.4823,
+      "step": 1170
+    },
+    {
+      "epoch": 1.5882352941176472,
+      "grad_norm": 2.455275058746338,
+      "learning_rate": 4.7775500504125945e-06,
+      "loss": 2.4616,
+      "step": 1175
+    },
+    {
+      "epoch": 1.5949966193373901,
+      "grad_norm": 2.9078075885772705,
+      "learning_rate": 4.741113168920359e-06,
+      "loss": 2.4495,
+      "step": 1180
+    },
+    {
+      "epoch": 1.601757944557133,
+      "grad_norm": 3.3229355812072754,
+      "learning_rate": 4.704690067761605e-06,
+      "loss": 2.4375,
+      "step": 1185
+    },
+    {
+      "epoch": 1.6085192697768762,
+      "grad_norm": 2.7873892784118652,
+      "learning_rate": 4.668282685708102e-06,
+      "loss": 2.4541,
+      "step": 1190
+    },
+    {
+      "epoch": 1.6152805949966194,
+      "grad_norm": 2.7972123622894287,
+      "learning_rate": 4.631892960694911e-06,
+      "loss": 2.4567,
+      "step": 1195
+    },
+    {
+      "epoch": 1.6220419202163625,
+      "grad_norm": 2.4656670093536377,
+      "learning_rate": 4.595522829717217e-06,
+      "loss": 2.4463,
+      "step": 1200
+    },
+    {
+      "epoch": 1.6288032454361054,
+      "grad_norm": 3.0135202407836914,
+      "learning_rate": 4.559174228727234e-06,
+      "loss": 2.4437,
+      "step": 1205
+    },
+    {
+      "epoch": 1.6355645706558486,
+      "grad_norm": 2.723539352416992,
+      "learning_rate": 4.5228490925311506e-06,
+      "loss": 2.4397,
+      "step": 1210
+    },
+    {
+      "epoch": 1.6423258958755915,
+      "grad_norm": 2.662691354751587,
+      "learning_rate": 4.486549354686148e-06,
+      "loss": 2.4373,
+      "step": 1215
+    },
+    {
+      "epoch": 1.6490872210953347,
+      "grad_norm": 2.6444590091705322,
+      "learning_rate": 4.450276947397466e-06,
+      "loss": 2.4563,
+      "step": 1220
+    },
+    {
+      "epoch": 1.6558485463150778,
+      "grad_norm": 2.985874652862549,
+      "learning_rate": 4.414033801415569e-06,
+      "loss": 2.459,
+      "step": 1225
+    },
+    {
+      "epoch": 1.662609871534821,
+      "grad_norm": 2.606755495071411,
+      "learning_rate": 4.377821845933364e-06,
+      "loss": 2.4403,
+      "step": 1230
+    },
+    {
+      "epoch": 1.669371196754564,
+      "grad_norm": 2.5894505977630615,
+      "learning_rate": 4.3416430084835095e-06,
+      "loss": 2.4437,
+      "step": 1235
+    },
+    {
+      "epoch": 1.6761325219743068,
+      "grad_norm": 2.9056777954101562,
+      "learning_rate": 4.305499214835825e-06,
+      "loss": 2.4335,
+      "step": 1240
+    },
+    {
+      "epoch": 1.68289384719405,
+      "grad_norm": 2.6746933460235596,
+      "learning_rate": 4.269392388894768e-06,
+      "loss": 2.4317,
+      "step": 1245
+    },
+    {
+      "epoch": 1.6896551724137931,
+      "grad_norm": 2.4791347980499268,
+      "learning_rate": 4.233324452597046e-06,
+      "loss": 2.4187,
+      "step": 1250
+    },
+    {
+      "epoch": 1.6964164976335363,
+      "grad_norm": 2.4208455085754395,
+      "learning_rate": 4.197297325809295e-06,
+      "loss": 2.4328,
+      "step": 1255
+    },
+    {
+      "epoch": 1.7031778228532792,
+      "grad_norm": 2.7689950466156006,
+      "learning_rate": 4.161312926225898e-06,
+      "loss": 2.4606,
+      "step": 1260
+    },
+    {
+      "epoch": 1.7099391480730222,
+      "grad_norm": 2.2801952362060547,
+      "learning_rate": 4.125373169266901e-06,
+      "loss": 2.4256,
+      "step": 1265
+    },
+    {
+      "epoch": 1.7167004732927653,
+      "grad_norm": 2.5390591621398926,
+      "learning_rate": 4.089479967976063e-06,
+      "loss": 2.407,
+      "step": 1270
+    },
+    {
+      "epoch": 1.7234617985125085,
+      "grad_norm": 3.087167501449585,
+      "learning_rate": 4.05363523291902e-06,
+      "loss": 2.438,
+      "step": 1275
+    },
+    {
+      "epoch": 1.7302231237322516,
+      "grad_norm": 2.693443775177002,
+      "learning_rate": 4.017840872081593e-06,
+      "loss": 2.4081,
+      "step": 1280
+    },
+    {
+      "epoch": 1.7369844489519946,
+      "grad_norm": 2.6463053226470947,
+      "learning_rate": 3.982098790768224e-06,
+      "loss": 2.424,
+      "step": 1285
+    },
+    {
+      "epoch": 1.7437457741717377,
+      "grad_norm": 2.368029832839966,
+      "learning_rate": 3.946410891500556e-06,
+      "loss": 2.4678,
+      "step": 1290
+    },
+    {
+      "epoch": 1.7505070993914806,
+      "grad_norm": 2.3758246898651123,
+      "learning_rate": 3.91077907391617e-06,
+      "loss": 2.4215,
+      "step": 1295
+    },
+    {
+      "epoch": 1.7572684246112238,
+      "grad_norm": 2.502300262451172,
+      "learning_rate": 3.875205234667463e-06,
+      "loss": 2.4178,
+      "step": 1300
+    },
+    {
+      "epoch": 1.764029749830967,
+      "grad_norm": 3.306715250015259,
+      "learning_rate": 3.839691267320693e-06,
+      "loss": 2.429,
+      "step": 1305
+    },
+    {
+      "epoch": 1.77079107505071,
+      "grad_norm": 2.5375657081604004,
+      "learning_rate": 3.8042390622551856e-06,
+      "loss": 2.4502,
+      "step": 1310
+    },
+    {
+      "epoch": 1.777552400270453,
+      "grad_norm": 2.51631760597229,
+      "learning_rate": 3.7688505065627114e-06,
+      "loss": 2.4389,
+      "step": 1315
+    },
+    {
+      "epoch": 1.784313725490196,
+      "grad_norm": 4.152524948120117,
+      "learning_rate": 3.7335274839470394e-06,
+      "loss": 2.425,
+      "step": 1320
+    },
+    {
+      "epoch": 1.791075050709939,
+      "grad_norm": 2.4869284629821777,
+      "learning_rate": 3.6982718746236603e-06,
+      "loss": 2.4557,
+      "step": 1325
+    },
+    {
+      "epoch": 1.7978363759296823,
+      "grad_norm": 2.5345115661621094,
+      "learning_rate": 3.6630855552197177e-06,
+      "loss": 2.4146,
+      "step": 1330
+    },
+    {
+      "epoch": 1.8045977011494254,
+      "grad_norm": 2.511878490447998,
+      "learning_rate": 3.627970398674104e-06,
+      "loss": 2.4366,
+      "step": 1335
+    },
+    {
+      "epoch": 1.8113590263691683,
+      "grad_norm": 2.6323208808898926,
+      "learning_rate": 3.5929282741377773e-06,
+      "loss": 2.4518,
+      "step": 1340
+    },
+    {
+      "epoch": 1.8181203515889113,
+      "grad_norm": 2.635624885559082,
+      "learning_rate": 3.557961046874255e-06,
+      "loss": 2.4258,
+      "step": 1345
+    },
+    {
+      "epoch": 1.8248816768086544,
+      "grad_norm": 2.3183505535125732,
+      "learning_rate": 3.523070578160339e-06,
+      "loss": 2.4306,
+      "step": 1350
+    },
+    {
+      "epoch": 1.8316430020283976,
+      "grad_norm": 2.523566961288452,
+      "learning_rate": 3.4882587251870336e-06,
+      "loss": 2.4553,
+      "step": 1355
+    },
+    {
+      "epoch": 1.8384043272481407,
+      "grad_norm": 2.98215651512146,
+      "learning_rate": 3.4535273409606935e-06,
+      "loss": 2.4179,
+      "step": 1360
+    },
+    {
+      "epoch": 1.8451656524678837,
+      "grad_norm": 2.7971649169921875,
+      "learning_rate": 3.4188782742043892e-06,
+      "loss": 2.4225,
+      "step": 1365
+    },
+    {
+      "epoch": 1.8519269776876268,
+      "grad_norm": 2.7544167041778564,
+      "learning_rate": 3.3843133692594966e-06,
+      "loss": 2.4076,
+      "step": 1370
+    },
+    {
+      "epoch": 1.8586883029073697,
+      "grad_norm": 2.5221734046936035,
+      "learning_rate": 3.3498344659875304e-06,
+      "loss": 2.4144,
+      "step": 1375
+    },
+    {
+      "epoch": 1.865449628127113,
+      "grad_norm": 2.2641069889068604,
+      "learning_rate": 3.3154433996722037e-06,
+      "loss": 2.4038,
+      "step": 1380
+    },
+    {
+      "epoch": 1.872210953346856,
+      "grad_norm": 2.3761255741119385,
+      "learning_rate": 3.281142000921745e-06,
+      "loss": 2.4296,
+      "step": 1385
+    },
+    {
+      "epoch": 1.8789722785665992,
+      "grad_norm": 2.4303650856018066,
+      "learning_rate": 3.2469320955714472e-06,
+      "loss": 2.3978,
+      "step": 1390
+    },
+    {
+      "epoch": 1.8857336037863421,
+      "grad_norm": 3.042694091796875,
+      "learning_rate": 3.2128155045864893e-06,
+      "loss": 2.4117,
+      "step": 1395
+    },
+    {
+      "epoch": 1.892494929006085,
+      "grad_norm": 2.3523356914520264,
+      "learning_rate": 3.178794043964998e-06,
+      "loss": 2.4145,
+      "step": 1400
+    },
+    {
+      "epoch": 1.8992562542258282,
+      "grad_norm": 2.467633008956909,
+      "learning_rate": 3.144869524641393e-06,
+      "loss": 2.397,
+      "step": 1405
+    },
+    {
+      "epoch": 1.9060175794455714,
+      "grad_norm": 2.5464890003204346,
+      "learning_rate": 3.1110437523899874e-06,
+      "loss": 2.4064,
+      "step": 1410
+    },
+    {
+      "epoch": 1.9127789046653145,
+      "grad_norm": 2.362809658050537,
+      "learning_rate": 3.077318527728867e-06,
+      "loss": 2.4229,
+      "step": 1415
+    },
+    {
+      "epoch": 1.9195402298850575,
+      "grad_norm": 2.8589625358581543,
+      "learning_rate": 3.043695645824055e-06,
+      "loss": 2.4406,
+      "step": 1420
+    },
+    {
+      "epoch": 1.9263015551048004,
+      "grad_norm": 2.349301338195801,
+      "learning_rate": 3.010176896393949e-06,
+      "loss": 2.4243,
+      "step": 1425
+    },
+    {
+      "epoch": 1.9330628803245435,
+      "grad_norm": 2.634345769882202,
+      "learning_rate": 2.976764063614067e-06,
+      "loss": 2.4362,
+      "step": 1430
+    },
+    {
+      "epoch": 1.9398242055442867,
+      "grad_norm": 2.4867677688598633,
+      "learning_rate": 2.9434589260220637e-06,
+      "loss": 2.427,
+      "step": 1435
+    },
+    {
+      "epoch": 1.9465855307640298,
+      "grad_norm": 5.755255699157715,
+      "learning_rate": 2.9102632564230726e-06,
+      "loss": 2.3987,
+      "step": 1440
+    },
+    {
+      "epoch": 1.9533468559837728,
+      "grad_norm": 2.603665590286255,
+      "learning_rate": 2.8771788217953333e-06,
+      "loss": 2.4174,
+      "step": 1445
+    },
+    {
+      "epoch": 1.960108181203516,
+      "grad_norm": 2.6414690017700195,
+      "learning_rate": 2.844207383196137e-06,
+      "loss": 2.4183,
+      "step": 1450
+    },
+    {
+      "epoch": 1.9668695064232589,
+      "grad_norm": 2.567423105239868,
+      "learning_rate": 2.811350695668097e-06,
+      "loss": 2.4177,
+      "step": 1455
+    },
+    {
+      "epoch": 1.973630831643002,
+      "grad_norm": 2.662747859954834,
+      "learning_rate": 2.7786105081457138e-06,
+      "loss": 2.4356,
+      "step": 1460
+    },
+    {
+      "epoch": 1.9803921568627452,
+      "grad_norm": 2.4450554847717285,
+      "learning_rate": 2.745988563362287e-06,
+      "loss": 2.3956,
+      "step": 1465
+    },
+    {
+      "epoch": 1.9871534820824883,
+      "grad_norm": 2.6526248455047607,
+      "learning_rate": 2.7134865977571622e-06,
+      "loss": 2.428,
+      "step": 1470
+    },
+    {
+      "epoch": 1.9939148073022313,
+      "grad_norm": 2.8979547023773193,
+      "learning_rate": 2.681106341383283e-06,
+      "loss": 2.4408,
+      "step": 1475
+    },
+    {
+      "epoch": 2.0,
+      "grad_norm": 2.4628653526306152,
+      "learning_rate": 2.6488495178151107e-06,
+      "loss": 2.2029,
+      "step": 1480
+    },
+    {
+      "epoch": 2.006761325219743,
+      "grad_norm": 2.8820254802703857,
+      "learning_rate": 2.616717844056882e-06,
+      "loss": 2.394,
+      "step": 1485
+    },
+    {
+      "epoch": 2.0135226504394863,
+      "grad_norm": 4.178656101226807,
+      "learning_rate": 2.5847130304512158e-06,
+      "loss": 2.4002,
+      "step": 1490
+    },
+    {
+      "epoch": 2.020283975659229,
+      "grad_norm": 2.5161166191101074,
+      "learning_rate": 2.5528367805880627e-06,
+      "loss": 2.3765,
+      "step": 1495
+    },
+    {
+      "epoch": 2.027045300878972,
+      "grad_norm": 2.4519433975219727,
+      "learning_rate": 2.5210907912140326e-06,
+      "loss": 2.4097,
+      "step": 1500
+    },
+    {
+      "epoch": 2.0338066260987153,
+      "grad_norm": 2.246267795562744,
+      "learning_rate": 2.4894767521420805e-06,
+      "loss": 2.3988,
+      "step": 1505
+    },
+    {
+      "epoch": 2.0405679513184585,
+      "grad_norm": 2.803452968597412,
+      "learning_rate": 2.4579963461615496e-06,
+      "loss": 2.4005,
+      "step": 1510
+    },
+    {
+      "epoch": 2.0473292765382016,
+      "grad_norm": 3.4639711380004883,
+      "learning_rate": 2.426651248948607e-06,
+      "loss": 2.4229,
+      "step": 1515
+    },
+    {
+      "epoch": 2.0540906017579443,
+      "grad_norm": 2.396606206893921,
+      "learning_rate": 2.395443128977042e-06,
+      "loss": 2.3826,
+      "step": 1520
+    },
+    {
+      "epoch": 2.0608519269776875,
+      "grad_norm": 3.1604995727539062,
+      "learning_rate": 2.3643736474294626e-06,
+      "loss": 2.38,
+      "step": 1525
+    },
+    {
+      "epoch": 2.0676132521974306,
+      "grad_norm": 5.628172397613525,
+      "learning_rate": 2.3334444581088633e-06,
+      "loss": 2.4044,
+      "step": 1530
+    },
+    {
+      "epoch": 2.074374577417174,
+      "grad_norm": 2.5617904663085938,
+      "learning_rate": 2.302657207350599e-06,
+      "loss": 2.3835,
+      "step": 1535
+    },
+    {
+      "epoch": 2.081135902636917,
+      "grad_norm": 2.4702658653259277,
+      "learning_rate": 2.272013533934752e-06,
+      "loss": 2.3919,
+      "step": 1540
+    },
+    {
+      "epoch": 2.08789722785666,
+      "grad_norm": 2.327834129333496,
+      "learning_rate": 2.241515068998903e-06,
+      "loss": 2.383,
+      "step": 1545
+    },
+    {
+      "epoch": 2.094658553076403,
+      "grad_norm": 2.5790200233459473,
+      "learning_rate": 2.2111634359513007e-06,
+      "loss": 2.3888,
+      "step": 1550
+    }
+  ],
+  "logging_steps": 5,
+  "max_steps": 2220,
+  "num_input_tokens_seen": 0,
+  "num_train_epochs": 3,
+  "save_steps": 25,
+  "stateful_callbacks": {
+    "TrainerControl": {
+      "args": {
+        "should_epoch_stop": false,
+        "should_evaluate": false,
+        "should_log": false,
+        "should_save": true,
+        "should_training_stop": false
+      },
+      "attributes": {}
+    }
+  },
+  "total_flos": 1.3720287815503394e+20,
+  "train_batch_size": 32,
+  "trial_name": null,
+  "trial_params": null
+}

training_args.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:dc1f054fb9418a96145fe6dc451e5844758ecb41d8154f513a240f7460914e81
+size 7761

zero_to_fp32.py

@@ -0,0 +1,760 @@
+#!/usr/bin/env python
+
+# Copyright (c) Microsoft Corporation.
+# SPDX-License-Identifier: Apache-2.0
+
+# DeepSpeed Team
+
+# This script extracts fp32 consolidated weights from a zero 1, 2 and 3 DeepSpeed checkpoints. It gets
+# copied into the top level checkpoint dir, so the user can easily do the conversion at any point in
+# the future. Once extracted, the weights don't require DeepSpeed and can be used in any
+# application.
+#
+# example:
+#   python zero_to_fp32.py . output_dir/
+#   or
+#   python zero_to_fp32.py . output_dir/ --safe_serialization
+
+import argparse
+import torch
+import glob
+import math
+import os
+import re
+import gc
+import json
+import numpy as np
+from tqdm import tqdm
+from collections import OrderedDict
+from dataclasses import dataclass
+
+# while this script doesn't use deepspeed to recover data, since the checkpoints are pickled with
+# DeepSpeed data structures it has to be available in the current python environment.
+from deepspeed.utils import logger
+from deepspeed.checkpoint.constants import (DS_VERSION, OPTIMIZER_STATE_DICT, SINGLE_PARTITION_OF_FP32_GROUPS,
+                                            FP32_FLAT_GROUPS, ZERO_STAGE, PARTITION_COUNT, PARAM_SHAPES, BUFFER_NAMES,
+                                            FROZEN_PARAM_SHAPES, FROZEN_PARAM_FRAGMENTS)
+
+
+@dataclass
+class zero_model_state:
+    buffers: dict()
+    param_shapes: dict()
+    shared_params: list
+    ds_version: int
+    frozen_param_shapes: dict()
+    frozen_param_fragments: dict()
+
+
+debug = 0
+
+# load to cpu
+device = torch.device('cpu')
+
+
+def atoi(text):
+    return int(text) if text.isdigit() else text
+
+
+def natural_keys(text):
+    '''
+    alist.sort(key=natural_keys) sorts in human order
+    http://nedbatchelder.com/blog/200712/human_sorting.html
+    (See Toothy's implementation in the comments)
+    '''
+    return [atoi(c) for c in re.split(r'(\d+)', text)]
+
+
+def get_model_state_file(checkpoint_dir, zero_stage):
+    if not os.path.isdir(checkpoint_dir):
+        raise FileNotFoundError(f"Directory '{checkpoint_dir}' doesn't exist")
+
+    # there should be only one file
+    if zero_stage <= 2:
+        file = os.path.join(checkpoint_dir, "mp_rank_00_model_states.pt")
+    elif zero_stage == 3:
+        file = os.path.join(checkpoint_dir, "zero_pp_rank_0_mp_rank_00_model_states.pt")
+
+    if not os.path.exists(file):
+        raise FileNotFoundError(f"can't find model states file at '{file}'")
+
+    return file
+
+
+def get_checkpoint_files(checkpoint_dir, glob_pattern):
+    # XXX: need to test that this simple glob rule works for multi-node setup too
+    ckpt_files = sorted(glob.glob(os.path.join(checkpoint_dir, glob_pattern)), key=natural_keys)
+
+    if len(ckpt_files) == 0:
+        raise FileNotFoundError(f"can't find {glob_pattern} files in directory '{checkpoint_dir}'")
+
+    return ckpt_files
+
+
+def get_optim_files(checkpoint_dir):
+    return get_checkpoint_files(checkpoint_dir, "*_optim_states.pt")
+
+
+def get_model_state_files(checkpoint_dir):
+    return get_checkpoint_files(checkpoint_dir, "*_model_states.pt")
+
+
+def parse_model_states(files):
+    zero_model_states = []
+    for file in files:
+        state_dict = torch.load(file, map_location=device, weights_only=False)
+
+        if BUFFER_NAMES not in state_dict:
+            raise ValueError(f"{file} is not a model state checkpoint")
+        buffer_names = state_dict[BUFFER_NAMES]
+        if debug:
+            print("Found buffers:", buffer_names)
+
+        # recover just the buffers while restoring them to fp32 if they were saved in fp16
+        buffers = {k: v.float() for k, v in state_dict["module"].items() if k in buffer_names}
+        param_shapes = state_dict[PARAM_SHAPES]
+
+        # collect parameters that are included in param_shapes
+        param_names = []
+        for s in param_shapes:
+            for name in s.keys():
+                param_names.append(name)
+
+        # update with frozen parameters
+        frozen_param_shapes = state_dict.get(FROZEN_PARAM_SHAPES, None)
+        if frozen_param_shapes is not None:
+            if debug:
+                print(f"Found frozen_param_shapes: {frozen_param_shapes}")
+            param_names += list(frozen_param_shapes.keys())
+
+        # handle shared params
+        shared_params = [[k, v] for k, v in state_dict["shared_params"].items()]
+
+        ds_version = state_dict.get(DS_VERSION, None)
+
+        frozen_param_fragments = state_dict.get(FROZEN_PARAM_FRAGMENTS, None)
+
+        z_model_state = zero_model_state(buffers=buffers,
+                                         param_shapes=param_shapes,
+                                         shared_params=shared_params,
+                                         ds_version=ds_version,
+                                         frozen_param_shapes=frozen_param_shapes,
+                                         frozen_param_fragments=frozen_param_fragments)
+        zero_model_states.append(z_model_state)
+
+    return zero_model_states
+
+
+def parse_optim_states(files, ds_checkpoint_dir):
+    total_files = len(files)
+    state_dicts = []
+    for f in tqdm(files, desc='Loading checkpoint shards'):
+        state_dict = torch.load(f, map_location=device, mmap=True, weights_only=False)
+        # immediately discard the potentially huge 2 optimizer states as we only care for fp32 master weights
+        # and also handle the case where it was already removed by another helper script
+        state_dict["optimizer_state_dict"].pop("optimizer_state_dict", None)
+        state_dicts.append(state_dict)
+
+    if not ZERO_STAGE in state_dicts[0][OPTIMIZER_STATE_DICT]:
+        raise ValueError(f"{files[0]} is not a zero checkpoint")
+    zero_stage = state_dicts[0][OPTIMIZER_STATE_DICT][ZERO_STAGE]
+    world_size = state_dicts[0][OPTIMIZER_STATE_DICT][PARTITION_COUNT]
+
+    # For ZeRO-2 each param group can have different partition_count as data parallelism for expert
+    # parameters can be different from data parallelism for non-expert parameters. So we can just
+    # use the max of the partition_count to get the dp world_size.
+
+    if type(world_size) is list:
+        world_size = max(world_size)
+
+    if world_size != total_files:
+        raise ValueError(
+            f"Expected {world_size} of '*_optim_states.pt' under '{ds_checkpoint_dir}' but found {total_files} files. "
+            "Possibly due to an overwrite of an old checkpoint, or a checkpoint didn't get saved by one or more processes."
+        )
+
+    # the groups are named differently in each stage
+    if zero_stage <= 2:
+        fp32_groups_key = SINGLE_PARTITION_OF_FP32_GROUPS
+    elif zero_stage == 3:
+        fp32_groups_key = FP32_FLAT_GROUPS
+    else:
+        raise ValueError(f"unknown zero stage {zero_stage}")
+
+    fp32_flat_groups = [state_dicts[i][OPTIMIZER_STATE_DICT][fp32_groups_key] for i in range(len(state_dicts))]
+    return zero_stage, world_size, fp32_flat_groups
+
+
+def _get_fp32_state_dict_from_zero_checkpoint(ds_checkpoint_dir, exclude_frozen_parameters):
+    """
+    Returns fp32 state_dict reconstructed from ds checkpoint
+
+    Args:
+        - ``ds_checkpoint_dir``: path to the deepspeed checkpoint folder (where the optimizer files are)
+
+    """
+    print(f"Processing zero checkpoint '{ds_checkpoint_dir}'")
+
+    optim_files = get_optim_files(ds_checkpoint_dir)
+    zero_stage, world_size, fp32_flat_groups = parse_optim_states(optim_files, ds_checkpoint_dir)
+    print(f"Detected checkpoint of type zero stage {zero_stage}, world_size: {world_size}")
+
+    model_files = get_model_state_files(ds_checkpoint_dir)
+
+    zero_model_states = parse_model_states(model_files)
+    print(f'Parsing checkpoint created by deepspeed=={zero_model_states[0].ds_version}')
+
+    if zero_stage <= 2:
+        return _get_fp32_state_dict_from_zero2_checkpoint(world_size, fp32_flat_groups, zero_model_states,
+                                                          exclude_frozen_parameters)
+    elif zero_stage == 3:
+        return _get_fp32_state_dict_from_zero3_checkpoint(world_size, fp32_flat_groups, zero_model_states,
+                                                          exclude_frozen_parameters)
+
+
+def _zero2_merge_frozen_params(state_dict, zero_model_states):
+    if zero_model_states[0].frozen_param_shapes is None or len(zero_model_states[0].frozen_param_shapes) == 0:
+        return
+
+    frozen_param_shapes = zero_model_states[0].frozen_param_shapes
+    frozen_param_fragments = zero_model_states[0].frozen_param_fragments
+
+    if debug:
+        num_elem = sum(s.numel() for s in frozen_param_shapes.values())
+        print(f'rank 0: {FROZEN_PARAM_SHAPES}.numel = {num_elem}')
+
+        wanted_params = len(frozen_param_shapes)
+        wanted_numel = sum(s.numel() for s in frozen_param_shapes.values())
+        avail_numel = sum([p.numel() for p in frozen_param_fragments.values()])
+        print(f'Frozen params: Have {avail_numel} numels to process.')
+        print(f'Frozen params: Need {wanted_numel} numels in {wanted_params} params')
+
+    total_params = 0
+    total_numel = 0
+    for name, shape in frozen_param_shapes.items():
+        total_params += 1
+        unpartitioned_numel = shape.numel()
+        total_numel += unpartitioned_numel
+
+        state_dict[name] = frozen_param_fragments[name]
+
+        if debug:
+            print(f"{name} full shape: {shape} unpartitioned numel {unpartitioned_numel} ")
+
+    print(f"Reconstructed Frozen fp32 state dict with {total_params} params {total_numel} elements")
+
+
+def _has_callable(obj, fn):
+    attr = getattr(obj, fn, None)
+    return callable(attr)
+
+
+def _zero2_merge_trainable_params(state_dict, world_size, fp32_flat_groups, zero_model_states):
+    param_shapes = zero_model_states[0].param_shapes
+
+    # Reconstruction protocol:
+    #
+    # XXX: document this
+
+    if debug:
+        for i in range(world_size):
+            for j in range(len(fp32_flat_groups[0])):
+                print(f"{FP32_FLAT_GROUPS}[{i}][{j}].shape={fp32_flat_groups[i][j].shape}")
+
+    # XXX: memory usage doubles here (zero2)
+    num_param_groups = len(fp32_flat_groups[0])
+    merged_single_partition_of_fp32_groups = []
+    for i in range(num_param_groups):
+        merged_partitions = [sd[i] for sd in fp32_flat_groups]
+        full_single_fp32_vector = torch.cat(merged_partitions, 0)
+        merged_single_partition_of_fp32_groups.append(full_single_fp32_vector)
+    avail_numel = sum(
+        [full_single_fp32_vector.numel() for full_single_fp32_vector in merged_single_partition_of_fp32_groups])
+
+    if debug:
+        wanted_params = sum([len(shapes) for shapes in param_shapes])
+        wanted_numel = sum([sum(shape.numel() for shape in shapes.values()) for shapes in param_shapes])
+        # not asserting if there is a mismatch due to possible padding
+        print(f"Have {avail_numel} numels to process.")
+        print(f"Need {wanted_numel} numels in {wanted_params} params.")
+
+    # params
+    # XXX: for huge models that can't fit into the host's RAM we will have to recode this to support
+    # out-of-core computing solution
+    total_numel = 0
+    total_params = 0
+    for shapes, full_single_fp32_vector in zip(param_shapes, merged_single_partition_of_fp32_groups):
+        offset = 0
+        avail_numel = full_single_fp32_vector.numel()
+        for name, shape in shapes.items():
+
+            unpartitioned_numel = shape.numel() if _has_callable(shape, 'numel') else math.prod(shape)
+            total_numel += unpartitioned_numel
+            total_params += 1
+
+            if debug:
+                print(f"{name} full shape: {shape} unpartitioned numel {unpartitioned_numel} ")
+            state_dict[name] = full_single_fp32_vector.narrow(0, offset, unpartitioned_numel).view(shape)
+            offset += unpartitioned_numel
+
+        # Z2 started to align to 2*world_size to improve nccl performance. Therefore both offset and
+        # avail_numel can differ by anywhere between 0..2*world_size. Due to two unrelated complex
+        # paddings performed in the code it's almost impossible to predict the exact numbers w/o the
+        # live optimizer object, so we are checking that the numbers are within the right range
+        align_to = 2 * world_size
+
+        def zero2_align(x):
+            return align_to * math.ceil(x / align_to)
+
+        if debug:
+            print(f"original offset={offset}, avail_numel={avail_numel}")
+
+        offset = zero2_align(offset)
+        avail_numel = zero2_align(avail_numel)
+
+        if debug:
+            print(f"aligned  offset={offset}, avail_numel={avail_numel}")
+
+        # Sanity check
+        if offset != avail_numel:
+            raise ValueError(f"consumed {offset} numels out of {avail_numel} - something is wrong")
+
+    print(f"Reconstructed fp32 state dict with {total_params} params {total_numel} elements")
+
+
+def _get_fp32_state_dict_from_zero2_checkpoint(world_size, fp32_flat_groups, zero_model_states,
+                                               exclude_frozen_parameters):
+    state_dict = OrderedDict()
+
+    # buffers
+    buffers = zero_model_states[0].buffers
+    state_dict.update(buffers)
+    if debug:
+        print(f"added {len(buffers)} buffers")
+
+    if not exclude_frozen_parameters:
+        _zero2_merge_frozen_params(state_dict, zero_model_states)
+
+    _zero2_merge_trainable_params(state_dict, world_size, fp32_flat_groups, zero_model_states)
+
+    # recover shared parameters
+    for pair in zero_model_states[0].shared_params:
+        if pair[1] in state_dict:
+            state_dict[pair[0]] = state_dict[pair[1]]
+
+    return state_dict
+
+
+def zero3_partitioned_param_info(unpartitioned_numel, world_size):
+    remainder = unpartitioned_numel % world_size
+    padding_numel = (world_size - remainder) if remainder else 0
+    partitioned_numel = math.ceil(unpartitioned_numel / world_size)
+    return partitioned_numel, padding_numel
+
+
+def _zero3_merge_frozen_params(state_dict, world_size, zero_model_states):
+    if zero_model_states[0].frozen_param_shapes is None or len(zero_model_states[0].frozen_param_shapes) == 0:
+        return
+
+    if debug:
+        for i in range(world_size):
+            num_elem = sum(s.numel() for s in zero_model_states[i].frozen_param_fragments.values())
+            print(f'rank {i}: {FROZEN_PARAM_SHAPES}.numel = {num_elem}')
+
+        frozen_param_shapes = zero_model_states[0].frozen_param_shapes
+        wanted_params = len(frozen_param_shapes)
+        wanted_numel = sum(s.numel() for s in frozen_param_shapes.values())
+        avail_numel = sum([p.numel() for p in zero_model_states[0].frozen_param_fragments.values()]) * world_size
+        print(f'Frozen params: Have {avail_numel} numels to process.')
+        print(f'Frozen params: Need {wanted_numel} numels in {wanted_params} params')
+
+    total_params = 0
+    total_numel = 0
+    for name, shape in zero_model_states[0].frozen_param_shapes.items():
+        total_params += 1
+        unpartitioned_numel = shape.numel()
+        total_numel += unpartitioned_numel
+
+        param_frags = tuple(model_state.frozen_param_fragments[name] for model_state in zero_model_states)
+        state_dict[name] = torch.cat(param_frags, 0).narrow(0, 0, unpartitioned_numel).view(shape)
+
+        partitioned_numel, partitioned_padding_numel = zero3_partitioned_param_info(unpartitioned_numel, world_size)
+
+        if debug:
+            print(
+                f"Frozen params: {total_params} {name} full shape: {shape} partition0 numel={partitioned_numel} partitioned_padding_numel={partitioned_padding_numel}"
+            )
+
+    print(f"Reconstructed Frozen fp32 state dict with {total_params} params {total_numel} elements")
+
+
+class GatheredTensor:
+    """
+    A pseudo tensor that collects partitioned weights.
+    It is more memory efficient when there are multiple groups.
+    """
+
+    def __init__(self, flat_groups, flat_groups_offset, offset, partitioned_numel, shape):
+        self.flat_groups = flat_groups
+        self.flat_groups_offset = flat_groups_offset
+        self.offset = offset
+        self.partitioned_numel = partitioned_numel
+        self.shape = shape
+        self.dtype = self.flat_groups[0][0].dtype
+
+    def contiguous(self):
+        """
+        Merge partitioned weights from flat_groups into a single tensor.
+        """
+        end_idx = self.offset + self.partitioned_numel
+        world_size = len(self.flat_groups)
+        pad_flat_param_chunks = []
+
+        for rank_i in range(world_size):
+            # for each rank, we need to collect weights from related group/groups
+            flat_groups_at_rank_i = self.flat_groups[rank_i]
+            start_group_id = None
+            end_group_id = None
+            for group_id in range(len(self.flat_groups_offset)):
+                if self.flat_groups_offset[group_id] <= self.offset < self.flat_groups_offset[group_id + 1]:
+                    start_group_id = group_id
+                if self.flat_groups_offset[group_id] < end_idx <= self.flat_groups_offset[group_id + 1]:
+                    end_group_id = group_id
+                    break
+            # collect weights from related group/groups
+            for group_id in range(start_group_id, end_group_id + 1):
+                flat_tensor = flat_groups_at_rank_i[group_id]
+                start_offset = self.offset - self.flat_groups_offset[group_id]
+                end_offset = min(end_idx, self.flat_groups_offset[group_id + 1]) - self.flat_groups_offset[group_id]
+                pad_flat_param_chunks.append(flat_tensor[start_offset:end_offset])
+
+        # collect weights from all ranks
+        pad_flat_param = torch.cat(pad_flat_param_chunks, dim=0)
+        param = pad_flat_param[:self.shape.numel()].view(self.shape).contiguous()
+        return param
+
+
+def _zero3_merge_trainable_params(state_dict, world_size, fp32_flat_groups, zero_model_states):
+    param_shapes = zero_model_states[0].param_shapes
+    avail_numel = sum([flat_group.numel() for flat_group in fp32_flat_groups[0]]) * world_size
+
+    # Reconstruction protocol: For zero3 we need to zip the partitions together at boundary of each
+    # param, re-consolidating each param, while dealing with padding if any
+
+    # merge list of dicts, preserving order
+    param_shapes = {k: v for d in param_shapes for k, v in d.items()}
+
+    if debug:
+        for i in range(world_size):
+            print(f"{FP32_FLAT_GROUPS}[{i}].shape={fp32_flat_groups[i].shape}")
+
+        wanted_params = len(param_shapes)
+        wanted_numel = sum(shape.numel() for shape in param_shapes.values())
+        # not asserting if there is a mismatch due to possible padding
+        avail_numel = fp32_flat_groups[0].numel() * world_size
+        print(f"Trainable params: Have {avail_numel} numels to process.")
+        print(f"Trainable params: Need {wanted_numel} numels in {wanted_params} params.")
+
+    # params
+    # XXX: for huge models that can't fit into the host's RAM we will have to recode this to support
+    # out-of-core computing solution
+    offset = 0
+    total_numel = 0
+    total_params = 0
+    flat_groups_offset = [0] + list(np.cumsum([flat_tensor.numel() for flat_tensor in fp32_flat_groups[0]]))
+    for name, shape in tqdm(param_shapes.items(), desc='Gathering sharded weights'):
+        unpartitioned_numel = shape.numel()
+        total_numel += unpartitioned_numel
+        total_params += 1
+        partitioned_numel, partitioned_padding_numel = zero3_partitioned_param_info(unpartitioned_numel, world_size)
+
+        if debug:
+            print(
+                f"Trainable params: {total_params} {name} full shape: {shape} partition0 numel={partitioned_numel} partitioned_padding_numel={partitioned_padding_numel}"
+            )
+
+        # memory efficient tensor
+        tensor = GatheredTensor(fp32_flat_groups, flat_groups_offset, offset, partitioned_numel, shape)
+        state_dict[name] = tensor
+        offset += partitioned_numel
+
+    offset *= world_size
+
+    # Sanity check
+    if offset != avail_numel:
+        raise ValueError(f"consumed {offset} numels out of {avail_numel} - something is wrong")
+
+    print(f"Reconstructed Trainable fp32 state dict with {total_params} params {total_numel} elements")
+
+
+def _get_fp32_state_dict_from_zero3_checkpoint(world_size, fp32_flat_groups, zero_model_states,
+                                               exclude_frozen_parameters):
+    state_dict = OrderedDict()
+
+    # buffers
+    buffers = zero_model_states[0].buffers
+    state_dict.update(buffers)
+    if debug:
+        print(f"added {len(buffers)} buffers")
+
+    if not exclude_frozen_parameters:
+        _zero3_merge_frozen_params(state_dict, world_size, zero_model_states)
+
+    _zero3_merge_trainable_params(state_dict, world_size, fp32_flat_groups, zero_model_states)
+
+    # recover shared parameters
+    for pair in zero_model_states[0].shared_params:
+        if pair[1] in state_dict:
+            state_dict[pair[0]] = state_dict[pair[1]]
+
+    return state_dict
+
+
+def to_torch_tensor(state_dict, return_empty_tensor=False):
+    """
+    Convert state_dict of GatheredTensor to torch tensor
+    """
+    torch_state_dict = {}
+    converted_tensors = {}
+    for name, tensor in state_dict.items():
+        tensor_id = id(tensor)
+        if tensor_id in converted_tensors:  # shared tensors
+            shared_tensor = torch_state_dict[converted_tensors[tensor_id]]
+            torch_state_dict[name] = shared_tensor
+        else:
+            converted_tensors[tensor_id] = name
+            if return_empty_tensor:
+                torch_state_dict[name] = torch.empty(tensor.shape, dtype=tensor.dtype)
+            else:
+                torch_state_dict[name] = tensor.contiguous()
+    return torch_state_dict
+
+
+def get_fp32_state_dict_from_zero_checkpoint(checkpoint_dir,
+                                             tag=None,
+                                             exclude_frozen_parameters=False,
+                                             lazy_mode=False):
+    """
+    Convert ZeRO 2 or 3 checkpoint into a single fp32 consolidated state_dict that can be loaded with
+    ``load_state_dict()`` and used for training without DeepSpeed or shared with others, for example
+    via a model hub.
+
+    Args:
+        - ``checkpoint_dir``: path to the desired checkpoint folder
+        - ``tag``: checkpoint tag used as a unique identifier for checkpoint. If not provided will attempt to load tag in 'latest' file. e.g., ``global_step14``
+        - ``exclude_frozen_parameters``: exclude frozen parameters
+        - ``lazy_mode``: get state_dict in lazy mode. It returns a dict of pesduo tensor instead of torch tensor, which is more memory efficient.
+          Convert the pesduo tensor to torch tensor by ``.contiguous()``
+
+    Returns:
+        - pytorch ``state_dict``
+
+    A typical usage might be ::
+
+        from deepspeed.utils.zero_to_fp32 import get_fp32_state_dict_from_zero_checkpoint
+        # do the training and checkpoint saving
+        state_dict = get_fp32_state_dict_from_zero_checkpoint(checkpoint_dir) # already on cpu
+        model = model.cpu() # move to cpu
+        model.load_state_dict(state_dict)
+        # submit to model hub or save the model to share with others
+
+    In this example the ``model`` will no longer be usable in the deepspeed context of the same
+    application. i.e. you will need to re-initialize the deepspeed engine, since
+    ``model.load_state_dict(state_dict)`` will remove all the deepspeed magic from it.
+
+    If you want it all done for you, use ``load_state_dict_from_zero_checkpoint`` instead.
+
+    Note: the above usage may not work if your application doesn't have sufficient free CPU memory.
+    You may need to use the offline approach using the ``zero_to_fp32.py`` script that is saved with
+    the checkpoint. Or you can load state_dict in lazy mode ::
+
+        from deepspeed.utils.zero_to_fp32 import get_fp32_state_dict_from_zero_checkpoint
+        state_dict = get_fp32_state_dict_from_zero_checkpoint(checkpoint_dir, lazy_mode=True) # not on cpu
+        for name, lazy_tensor in state_dict.item():
+            tensor = lazy_tensor.contiguous()  # to cpu
+            print(name, tensor)
+            # del tensor to release memory if it no longer in use
+    """
+    if tag is None:
+        latest_path = os.path.join(checkpoint_dir, 'latest')
+        if os.path.isfile(latest_path):
+            with open(latest_path, 'r') as fd:
+                tag = fd.read().strip()
+        else:
+            raise ValueError(f"Unable to find 'latest' file at {latest_path}")
+
+    ds_checkpoint_dir = os.path.join(checkpoint_dir, tag)
+
+    if not os.path.isdir(ds_checkpoint_dir):
+        raise FileNotFoundError(f"Directory '{ds_checkpoint_dir}' doesn't exist")
+
+    state_dict = _get_fp32_state_dict_from_zero_checkpoint(ds_checkpoint_dir, exclude_frozen_parameters)
+    if lazy_mode:
+        return state_dict
+    else:
+        return to_torch_tensor(state_dict)
+
+
+def convert_zero_checkpoint_to_fp32_state_dict(checkpoint_dir,
+                                               output_dir,
+                                               max_shard_size="5GB",
+                                               safe_serialization=False,
+                                               tag=None,
+                                               exclude_frozen_parameters=False):
+    """
+    Convert ZeRO 2 or 3 checkpoint into a single fp32 consolidated ``state_dict`` file that can be
+    loaded with ``torch.load(file)`` + ``load_state_dict()`` and used for training without DeepSpeed.
+
+    Args:
+        - ``checkpoint_dir``: path to the desired checkpoint folder. (one that contains the tag-folder, like ``global_step14``)
+        - ``output_dir``: directory to the pytorch fp32 state_dict output files
+        - ``max_shard_size``: the maximum size for a checkpoint before being sharded, default value is 5GB
+        - ``safe_serialization``:  whether to save the model using `safetensors` or the traditional PyTorch way (that uses `pickle`).
+        - ``tag``: checkpoint tag used as a unique identifier for checkpoint. If not provided will attempt to load tag in the file named ``latest`` in the checkpoint folder, e.g., ``global_step14``
+        - ``exclude_frozen_parameters``: exclude frozen parameters
+    """
+
+    # Dependency pre-check
+    if safe_serialization:
+        try:
+            from safetensors.torch import save_file
+        except ImportError:
+            print('If you want to use `safe_serialization`, please `pip install safetensors`')
+            raise
+    if max_shard_size is not None:
+        try:
+            from huggingface_hub import split_torch_state_dict_into_shards
+        except ImportError:
+            print('If you want to use `max_shard_size`, please `pip install huggingface_hub`')
+            raise
+
+    # Convert zero checkpoint to state_dict
+    state_dict = get_fp32_state_dict_from_zero_checkpoint(checkpoint_dir,
+                                                          tag,
+                                                          exclude_frozen_parameters,
+                                                          lazy_mode=True)
+
+    # Shard the model if it is too big.
+    weights_name = "model.safetensors" if safe_serialization else "pytorch_model.bin"
+    if max_shard_size is not None:
+        filename_pattern = weights_name.replace(".bin", "{suffix}.bin").replace(".safetensors", "{suffix}.safetensors")
+        # an memory-efficient approach for sharding
+        empty_state_dict = to_torch_tensor(state_dict, return_empty_tensor=True)
+        state_dict_split = split_torch_state_dict_into_shards(empty_state_dict,
+                                                              filename_pattern=filename_pattern,
+                                                              max_shard_size=max_shard_size)
+    else:
+        from collections import namedtuple
+        StateDictSplit = namedtuple("StateDictSplit", ["is_sharded", "filename_to_tensors"])
+        state_dict_split = StateDictSplit(is_sharded=False,
+                                          filename_to_tensors={weights_name: list(state_dict.keys())})
+
+    # Save the model by shard
+    os.makedirs(output_dir, exist_ok=True)
+    filename_to_tensors = state_dict_split.filename_to_tensors.items()
+    for shard_file, tensors in tqdm(filename_to_tensors, desc="Saving checkpoint shards"):
+        shard_state_dict = {tensor_name: state_dict[tensor_name] for tensor_name in tensors}
+        shard_state_dict = to_torch_tensor(shard_state_dict)
+        output_path = os.path.join(output_dir, shard_file)
+        if safe_serialization:
+            save_file(shard_state_dict, output_path, metadata={"format": "pt"})
+        else:
+            torch.save(shard_state_dict, output_path)
+        # release the memory of current shard
+        for tensor_name in list(shard_state_dict.keys()):
+            del state_dict[tensor_name]
+            del shard_state_dict[tensor_name]
+        del shard_state_dict
+        gc.collect()
+
+    # Save index if sharded
+    if state_dict_split.is_sharded:
+        index = {
+            "metadata": state_dict_split.metadata,
+            "weight_map": state_dict_split.tensor_to_filename,
+        }
+        save_index_file = "model.safetensors.index.json" if safe_serialization else "pytorch_model.bin.index.json"
+        save_index_file = os.path.join(output_dir, save_index_file)
+        with open(save_index_file, "w", encoding="utf-8") as f:
+            content = json.dumps(index, indent=2, sort_keys=True) + "\n"
+            f.write(content)
+
+
+def load_state_dict_from_zero_checkpoint(model, checkpoint_dir, tag=None):
+    """
+    1. Put the provided model to cpu
+    2. Convert ZeRO 2 or 3 checkpoint into a single fp32 consolidated ``state_dict``
+    3. Load it into the provided model
+
+    Args:
+        - ``model``: the model object to update
+        - ``checkpoint_dir``: path to the desired checkpoint folder. (one that contains the tag-folder, like ``global_step14``)
+        - ``tag``: checkpoint tag used as a unique identifier for checkpoint. If not provided will attempt to load tag in the file named ``latest`` in the checkpoint folder, e.g., ``global_step14``
+
+    Returns:
+        - ``model`: modified model
+
+    Make sure you have plenty of CPU memory available before you call this function. If you don't
+    have enough use the ``zero_to_fp32.py`` utility to do the conversion. You will find it
+    conveniently placed for you in the checkpoint folder.
+
+    A typical usage might be ::
+
+        from deepspeed.utils.zero_to_fp32 import load_state_dict_from_zero_checkpoint
+        model = load_state_dict_from_zero_checkpoint(trainer.model, checkpoint_dir)
+        # submit to model hub or save the model to share with others
+
+    Note, that once this was run, the ``model`` will no longer be usable in the deepspeed context
+    of the same application. i.e. you will need to re-initialize the deepspeed engine, since
+    ``model.load_state_dict(state_dict)`` will remove all the deepspeed magic from it.
+
+    """
+    logger.info(f"Extracting fp32 weights")
+    state_dict = get_fp32_state_dict_from_zero_checkpoint(checkpoint_dir, tag)
+
+    logger.info(f"Overwriting model with fp32 weights")
+    model = model.cpu()
+    model.load_state_dict(state_dict, strict=False)
+
+    return model
+
+
+if __name__ == "__main__":
+    parser = argparse.ArgumentParser()
+    parser.add_argument("checkpoint_dir",
+                        type=str,
+                        help="path to the desired checkpoint folder, e.g., path/checkpoint-12")
+    parser.add_argument("output_dir",
+                        type=str,
+                        help="directory to the pytorch fp32 state_dict output files"
+                        "(e.g. path/checkpoint-12-output/)")
+    parser.add_argument(
+        "--max_shard_size",
+        type=str,
+        default="5GB",
+        help="The maximum size for a checkpoint before being sharded. Checkpoints shard will then be each of size"
+        "lower than this size. If expressed as a string, needs to be digits followed by a unit (like `5MB`"
+        "We default it to 5GB in order for models to be able to run easily on free-tier google colab instances"
+        "without CPU OOM issues.")
+    parser.add_argument(
+        "--safe_serialization",
+        default=False,
+        action='store_true',
+        help="Whether to save the model using `safetensors` or the traditional PyTorch way (that uses `pickle`).")
+    parser.add_argument("-t",
+                        "--tag",
+                        type=str,
+                        default=None,
+                        help="checkpoint tag used as a unique identifier for checkpoint. e.g., global_step1")
+    parser.add_argument("--exclude_frozen_parameters", action='store_true', help="exclude frozen parameters")
+    parser.add_argument("-d", "--debug", action='store_true', help="enable debug")
+    args = parser.parse_args()
+
+    debug = args.debug
+
+    convert_zero_checkpoint_to_fp32_state_dict(args.checkpoint_dir,
+                                               args.output_dir,
+                                               max_shard_size=args.max_shard_size,
+                                               safe_serialization=args.safe_serialization,
+                                               tag=args.tag,
+                                               exclude_frozen_parameters=args.exclude_frozen_parameters)