Upload folder using huggingface_hub

.gitattributes

@@ -33,3 +33,4 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zip filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
+tokenizer.json filter=lfs diff=lfs merge=lfs -text

added_tokens.json

@@ -0,0 +1,24 @@
+{
+  "</tool_call>": 151658,
+  "<tool_call>": 151657,
+  "<|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,54 @@
+{%- if tools %}
+    {{- '<|im_start|>system\n' }}
+    {%- if messages[0]['role'] == 'system' %}
+        {{- messages[0]['content'] }}
+    {%- else %}
+        {{- 'You are a helpful assistant.' }}
+    {%- endif %}
+    {{- "\n\n# 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' }}
+    {%- else %}
+        {{- '<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n' }}
+    {%- endif %}
+{%- endif %}
+{%- for message in messages %}
+    {%- if (message.role == "user") or (message.role == "system" and not loop.first) or (message.role == "assistant" and not message.tool_calls) %}
+        {{- '<|im_start|>' + message.role + '\n' + message.content + '<|im_end|>' + '\n' }}
+    {%- elif message.role == "assistant" %}
+        {{- '<|im_start|>' + message.role }}
+        {%- if message.content %}
+            {{- '\n' + message.content }}
+        {%- endif %}
+        {%- for tool_call in message.tool_calls %}
+            {%- if tool_call.function is defined %}
+                {%- set tool_call = tool_call.function %}
+            {%- endif %}
+            {{- '\n<tool_call>\n{"name": "' }}
+            {{- tool_call.name }}
+            {{- '", "arguments": ' }}
+            {{- tool_call.arguments | tojson }}
+            {{- '}\n</tool_call>' }}
+        {%- endfor %}
+        {{- '<|im_end|>\n' }}
+    {%- elif message.role == "tool" %}
+        {%- if (loop.index0 == 0) 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' }}
+{%- endif %}

config.json

@@ -0,0 +1,84 @@
+{
+  "architectures": [
+    "JetNemotronForCausalLM"
+  ],
+  "attention_dropout": 0.0,
+  "auto_map": {
+    "AutoConfig": "configuration_jet_nemotron.JetNemotronConfig",
+    "AutoModelForCausalLM": "modeling_jet_nemotron.JetNemotronForCausalLM"
+  },
+  "bos_token_id": 151643,
+  "efficient_attention_config": {
+    "jet": {
+      "conv_size": 4,
+      "dconv_generator_reduction": 8,
+      "dconv_implementation": "triton",
+      "expand_v": 2,
+      "head_dim": 128,
+      "mode": "chunk",
+      "norm_eps": "1e-5",
+      "num_heads": 16
+    },
+    "swa": {
+      "window_size": 2048
+    }
+  },
+  "eos_token_id": 151643,
+  "hidden_act": "silu",
+  "hidden_size": 2048,
+  "initializer_range": 0.02,
+  "intermediate_size": 11008,
+  "layer_types": [
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "swa",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "swa",
+    "attn",
+    "jet",
+    "swa",
+    "attn",
+    "swa",
+    "swa",
+    "jet",
+    "jet",
+    "swa",
+    "jet",
+    "swa",
+    "jet",
+    "jet",
+    "jet",
+    "jet",
+    "attn",
+    "jet",
+    "jet",
+    "jet"
+  ],
+  "max_position_embeddings": 32768,
+  "max_window_layers": 36,
+  "model_type": "jet_nemotron",
+  "num_attention_heads": 16,
+  "num_hidden_layers": 36,
+  "num_key_value_heads": 2,
+  "rms_norm_eps": 1e-06,
+  "rope_scaling": null,
+  "rope_theta": 1000000.0,
+  "tie_word_embeddings": true,
+  "torch_dtype": "bfloat16",
+  "transformers_version": "4.51.3",
+  "use_cache": true,
+  "use_mrope": false,
+  "vocab_size": 151936
+}

configuration_jet_nemotron.py

@@ -0,0 +1,185 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+# This file is modified from https://github.com/huggingface/transformers/blob/main/src/transformers/models/qwen2/configuration_qwen2.py
+
+"""Jet-Nemotron 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 JetNemotronConfig(PretrainedConfig):
+    r"""
+    This is the configuration class to store the configuration of a [`JetModel`]. It is used to instantiate a
+    Jet model according to the specified arguments, defining the model architecture. 
+
+    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 Jet-Nemotron model. Defines the number of different tokens that can be represented by the
+            `inputs_ids` passed when calling [`JetModel`]
+        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`.
+        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_dropout (`float`, *optional*, defaults to 0.0):
+            The dropout ratio for the attention probabilities.
+
+    ```python
+    >>> from transformers import JetModel, JetConfig
+
+    >>> # Initializing a Jet-Nemotron style configuration
+    >>> configuration = JetConfig()
+
+    >>> # Initializing a model from the Jet-Nemotron-2B style configuration
+    >>> model = Jet-Nemotron-Model(configuration)
+
+    >>> # Accessing the model configuration
+    >>> configuration = model.config
+    ```"""
+
+    model_type = "jet_nemotron"
+    keys_to_ignore_at_inference = ["past_key_values"]
+
+    def __init__(
+        self,
+        vocab_size=151936,
+        hidden_size=1536,
+        intermediate_size=8960,
+        num_hidden_layers=28,
+        num_attention_heads=12,
+        num_key_value_heads=2,
+        hidden_act="silu",
+        max_position_embeddings=32768,
+        initializer_range=0.02,
+        rms_norm_eps=1e-6,
+        use_cache=True,
+        tie_word_embeddings=True,
+        rope_theta=1000000.0,
+        rope_scaling=None,
+        attention_dropout=0.0,
+        layer_types=None,
+        efficient_attention_config=None,        
+        **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
+
+        # 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.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_dropout = attention_dropout
+        if layer_types is None:
+            self.layer_types = ["attn"] * num_hidden_layers
+        elif isinstance(layer_types, str):
+            self.layer_types = eval(layer_types)
+        else:
+            self.layer_types = layer_types
+        # 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)
+
+        self.efficient_attention_config = efficient_attention_config
+        
+        super().__init__(
+            tie_word_embeddings=tie_word_embeddings,
+            **kwargs,
+        )

dconv_fwd_cache.py

@@ -0,0 +1,330 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import torch
+import triton
+import triton.language as tl
+from einops import rearrange
+import torch.nn.functional as F
+from torch.autograd import Function
+
+# Helper function to ensure tensors are contiguous for Triton
+def ensure_contiguous(t: torch.Tensor) -> torch.Tensor:
+    return t if t.is_contiguous() else t.contiguous()
+
+# --- Forward Kernel (modified for optional cache) ---
+@triton.jit
+def _dynamic_conv_fwd_kernel(
+    X_ptr, K_ptr, Out_ptr,
+    Cache_ptr, # New: Pointer to cache tensor
+    B, T, D, T_CACHE: tl.constexpr, # New: T is shape of x, T_CACHE is shape of cache
+    X_stride_b, X_stride_t, X_stride_d,
+    K_stride_b, K_stride_t, K_stride_d, K_stride_w,
+    Out_stride_b, Out_stride_t, Out_stride_d,
+    Cache_stride_b, Cache_stride_t, Cache_stride_d, # New: Strides for cache tensor
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    pid_batch_time = tl.program_id(0) # Covers B * T_out
+    pid_d_block = tl.program_id(1)
+
+    # T here is the time dimension of x and Out
+    batch_idx = tl.cast(pid_batch_time // T, tl.int64)
+    time_idx = pid_batch_time % T # Current output time step for x (0 to T-1)
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    accumulator = tl.zeros((BLOCK_SIZE_D,), dtype=tl.float32)
+    offs_w = tl.arange(0, W) # Kernel window offsets [0, 1, ..., W-1]
+
+    # Load Kernels (kernels are aligned with x's T dimension)
+    # K_ptr is indexed by time_idx which is the output time relative to x's start
+    k_ptrs = K_ptr + (batch_idx * K_stride_b + time_idx * K_stride_t +
+                      offs_d[:, None] * K_stride_d + offs_w[None, :] * K_stride_w)
+    k_vals = tl.load(k_ptrs, mask=d_mask[:, None], other=0.0) # Shape: [BLOCK_SIZE_D, W]
+
+    # --- Load Input from conceptual [Cache, X] tensor ---
+    # `time_idx` is the current output time step (0 to T-1, where T is x.shape[1])
+    # `offs_w` is [0, ..., W-1]
+    # Convolution input time indices relative to the *start of x*:
+    # e.g., for W=3, offs_w - W + 1 gives [-2, -1, 0]
+    # so input_time_indices_rel_to_x_start are [time_idx-2, time_idx-1, time_idx]
+    input_time_indices_rel_to_x_start = time_idx + offs_w - W + 1 # Shape: [W]
+
+    # Effective input time indices in the conceptual [Cache, X] sequence:
+    # These indices range from 0 (start of cache) to T_CACHE + T - 1 (end of x)
+    eff_t_indices = input_time_indices_rel_to_x_start + T_CACHE # Shape: [W]
+
+    # Overall mask for valid time indices within the conceptual [Cache, X] tensor
+    # Total effective length is T_CACHE (for cache) + T (for x)
+    eff_t_valid_mask = (eff_t_indices >= 0) & (eff_t_indices < (T_CACHE + T)) # Shape: [W]
+
+    # --- Load from Cache ---
+    # Condition for loading from cache: index is valid AND index < T_CACHE
+    # (eff_t_indices are 0-indexed from the start of the cache)
+    cache_load_time_mask = eff_t_valid_mask & (eff_t_indices < T_CACHE) # Shape: [W]
+    cache_ptr_indices = eff_t_indices # Use directly if in cache range
+
+    cache_ptrs = Cache_ptr + (batch_idx * Cache_stride_b +
+                              cache_ptr_indices[None, :] * Cache_stride_t +
+                              offs_d[:, None] * Cache_stride_d)
+    cache_final_load_mask = d_mask[:, None] & cache_load_time_mask[None, :] # Shape: [BLOCK_SIZE_D, W]
+    vals_from_cache = tl.load(cache_ptrs, mask=cache_final_load_mask, other=0.0) # Shape: [BLOCK_SIZE_D, W]
+
+    # --- Load from X ---
+    # Condition for loading from X: index is valid AND index >= T_CACHE
+    x_load_time_mask = eff_t_valid_mask & (eff_t_indices >= T_CACHE) # Shape: [W]
+    # Adjust indices for X_ptr: X_ptr expects indices from 0 to T-1 (relative to start of x)
+    x_ptr_indices = eff_t_indices - T_CACHE # Shape: [W]
+
+    x_ptrs = X_ptr + (batch_idx * X_stride_b +
+                      x_ptr_indices[None, :] * X_stride_t +
+                      offs_d[:, None] * X_stride_d)
+    x_final_load_mask = d_mask[:, None] & x_load_time_mask[None, :] # Shape: [BLOCK_SIZE_D, W]
+    vals_from_x = tl.load(x_ptrs, mask=x_final_load_mask, other=0.0) # Shape: [BLOCK_SIZE_D, W]
+
+    # Combine values. Masks ensure only one source contributes non-zero per element.
+    # If T_CACHE == 0, cache_load_time_mask is all False, so vals_from_cache is 0.0.
+    x_input_vals = vals_from_cache + vals_from_x # Shape: [BLOCK_SIZE_D, W]
+
+    # Compute and Accumulate
+    product = k_vals * x_input_vals # Element-wise product
+    accumulator += tl.sum(product, axis=1) # Sum over W dimension
+
+    # Store Result
+    out_ptrs = Out_ptr + (batch_idx * Out_stride_b + time_idx * Out_stride_t +
+                          offs_d * Out_stride_d)
+    tl.store(out_ptrs, accumulator, mask=d_mask)
+
+# --- Backward Kernel for Input Gradient (dX) ---
+@triton.jit
+def _dynamic_conv_bwd_dx_kernel(
+    GradOut_ptr, K_ptr, GradX_ptr, # Note: GradX is accumulated into
+    B, T, D,
+    GradOut_stride_b, GradOut_stride_t, GradOut_stride_d,
+    K_stride_b, K_stride_t, K_stride_d, K_stride_w,
+    GradX_stride_b, GradX_stride_t, GradX_stride_d,
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    """
+    Computes gradient w.r.t. input X.
+    Grid: (B * T, cdiv(D, BLOCK_SIZE_D)) - covering GradX output
+    GradX[b, t_x, d] = sum_{w=0}^{W-1} GradOut[b, t, d] * K[b, t, d, w]
+                       where t = t_x + W - 1 - w
+    """
+    pid_batch_time_x = tl.program_id(0) # Covers B * T for output GradX
+    pid_d_block = tl.program_id(1)
+
+    batch_idx = pid_batch_time_x // T
+    time_idx_x = pid_batch_time_x % T # This is t_x
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    # Accumulator for GradX elements
+    accumulator = tl.zeros((BLOCK_SIZE_D,), dtype=tl.float32)
+    offs_w = tl.arange(0, W) # [W]
+
+    # Loop over W to accumulate contributions
+    # Calculate the 't' index needed for GradOut and K based on t_x and w
+    # t = t_x + W - 1 - w
+    t_k_gradout_offs = time_idx_x + W - 1 - offs_w # Shape [W]
+
+    # Mask for valid 't' indices [0, T)
+    t_k_gradout_mask = (t_k_gradout_offs >= 0) & (t_k_gradout_offs < T) # Shape [W]
+
+    # --- Load GradOut ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    gradout_ptrs = GradOut_ptr + (batch_idx * GradOut_stride_b +
+                                  t_k_gradout_offs[None, :] * GradOut_stride_t +
+                                  offs_d[:, None] * GradOut_stride_d)
+    # Combined mask for loading GradOut (valid D and valid t)
+    gradout_load_mask = d_mask[:, None] & t_k_gradout_mask[None, :]
+    # Shape: [BLOCK_SIZE_D, W]
+    gradout_vals = tl.load(gradout_ptrs, mask=gradout_load_mask, other=0.0)
+
+    # --- Load Kernels ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    k_ptrs = K_ptr + (batch_idx * K_stride_b +
+                      t_k_gradout_offs[None, :] * K_stride_t +
+                      offs_d[:, None] * K_stride_d +
+                      offs_w[None, :] * K_stride_w) # Index K with 't' and 'w'
+    # Combined mask for loading K (valid D and valid t)
+    k_load_mask = d_mask[:, None] & t_k_gradout_mask[None, :]
+    # Shape: [BLOCK_SIZE_D, W]
+    k_vals = tl.load(k_ptrs, mask=k_load_mask, other=0.0)
+
+    # --- Compute product and accumulate ---
+    # Shape: [BLOCK_SIZE_D, W]
+    product = gradout_vals * k_vals
+    # Sum contributions over the W dimension
+    accumulator += tl.sum(product, axis=1) # Shape: [BLOCK_SIZE_D]
+
+    # --- Store accumulated gradients ---
+    gradx_ptrs = GradX_ptr + (batch_idx * GradX_stride_b +
+                              time_idx_x * GradX_stride_t +
+                              offs_d * GradX_stride_d)
+    tl.store(gradx_ptrs, accumulator, mask=d_mask)
+
+
+# --- Backward Kernel for Kernel Gradient (dK) ---
+@triton.jit
+def _dynamic_conv_bwd_dk_kernel(
+    GradOut_ptr, X_ptr, GradK_ptr, # Note: GradK is written directly
+    B, T, D,
+    GradOut_stride_b, GradOut_stride_t, GradOut_stride_d,
+    X_stride_b, X_stride_t, X_stride_d,
+    GradK_stride_b, GradK_stride_t, GradK_stride_d, GradK_stride_w,
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    """
+    Computes gradient w.r.t. kernels K.
+    Grid: (B * T, cdiv(D, BLOCK_SIZE_D)) - covering GradK output dims B, T, D
+    GradK[b, t, d, w] = GradOut[b, t, d] * X[b, t + w - W + 1, d]
+    """
+    pid_batch_time = tl.program_id(0) # Covers B * T for output GradK
+    pid_d_block = tl.program_id(1)
+
+    batch_idx = pid_batch_time // T
+    time_idx = pid_batch_time % T # This is 't' for GradK and GradOut
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    offs_w = tl.arange(0, W) # [W]
+
+    # --- Load GradOut ---
+    # Pointers shape: [BLOCK_SIZE_D] (only depends on b, t, d)
+    gradout_ptrs = GradOut_ptr + (batch_idx * GradOut_stride_b +
+                                  time_idx * GradOut_stride_t +
+                                  offs_d * GradOut_stride_d)
+    # Shape: [BLOCK_SIZE_D]
+    gradout_vals = tl.load(gradout_ptrs, mask=d_mask, other=0.0)
+
+    # --- Load Input X with implicit padding ---
+    # Calculate X's time index: t_x = t + w - W + 1
+    t_in_offs = time_idx + offs_w - W + 1 # Shape [W]
+    # Mask for valid t_x index [0, T)
+    t_in_mask = (t_in_offs >= 0) & (t_in_offs < T) # Shape [W]
+
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    x_ptrs = X_ptr + (batch_idx * X_stride_b +
+                      t_in_offs[None, :] * X_stride_t +
+                      offs_d[:, None] * X_stride_d)
+    # Combined mask for loading X (valid D and valid t_x)
+    x_load_mask = d_mask[:, None] & t_in_mask[None, :] # Shape [BLOCK_SIZE_D, W]
+    # Shape: [BLOCK_SIZE_D, W]
+    x_vals = tl.load(x_ptrs, mask=x_load_mask, other=0.0)
+
+    # --- Compute GradK = GradOut * X ---
+    # Broadcast gradout_vals: [BLOCK_SIZE_D, 1] * [BLOCK_SIZE_D, W] -> [BLOCK_SIZE_D, W]
+    gradk_vals = gradout_vals[:, None] * x_vals # Shape [BLOCK_SIZE_D, W]
+
+    # --- Store gradients for Kernels ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    gradk_ptrs = GradK_ptr + (batch_idx * GradK_stride_b +
+                              time_idx * GradK_stride_t +
+                              offs_d[:, None] * GradK_stride_d +
+                              offs_w[None, :] * GradK_stride_w)
+    # Mask only needed for D dimension (W is fully computed)
+    # Store computed gradient values.
+    tl.store(gradk_ptrs, gradk_vals, mask=d_mask[:, None])
+
+
+# --- Autograd Function ---
+class DynamicConvTritonFunc(Function):
+
+    @staticmethod
+    def forward(ctx, x, kernels, cache=None): # Added cache argument
+        """
+        Args:
+            x: Input tensor [B, T, D]
+            kernels: Kernels tensor [B, T, D, W]
+            cache: Optional past context tensor [B, T_cache, D]
+        """
+        x = ensure_contiguous(x)
+        kernels = ensure_contiguous(kernels)
+
+        B, T, D = x.shape # T is the time dimension of the current input x
+        _B_k, _T_k, _D_k, W = kernels.shape # Kernels are [B, T_x, D, W]
+        assert B == _B_k and T == _T_k and D == _D_k, \
+            f"Shape mismatch between x ({x.shape}) and kernels ({kernels.shape}) on B, T, or D dims"
+        assert W <= 4, "Kernel W > 4 not expected for this version"
+
+        out = torch.empty_like(x) # Output shape [B, T, D], corresponds to x
+
+        T_cache_val = 0
+        # Use x's data pointer and zero strides as placeholders if cache is None.
+        # These won't be used by the kernel if T_CACHE_VAL is 0 due to masking.
+        cache_ptr_val = x 
+        cache_s_b, cache_s_t, cache_s_d = 0, 0, 0
+
+        if cache is not None:
+            cache = ensure_contiguous(cache)
+            B_c, T_c, D_c = cache.shape
+            assert B_c == B, f"Batch size mismatch: x ({B}) vs cache ({B_c})"
+            assert D_c == D, f"Dimension mismatch: x ({D}) vs cache ({D_c})"
+            T_cache_val = T_c
+            cache_ptr_val = cache
+            cache_s_b, cache_s_t, cache_s_d = cache.stride(0), cache.stride(1), cache.stride(2)
+
+        grid = lambda meta: (B * T, triton.cdiv(D, meta['BLOCK_SIZE_D']))
+        BLOCK_SIZE_D = 128 # Consider tuning
+
+        _dynamic_conv_fwd_kernel[grid](
+            x, kernels, out,                                                           # X, K, Out pointers
+            cache_ptr_val,                                                              # Cache pointer
+            B, T, D, T_cache_val,                                                       # Shapes: B, T_x, D, T_cache
+            x.stride(0), x.stride(1), x.stride(2),                                      # X strides
+            kernels.stride(0), kernels.stride(1), kernels.stride(2), kernels.stride(3), # K strides
+            out.stride(0), out.stride(1), out.stride(2),                                # Out strides
+            cache_s_b, cache_s_t, cache_s_d,                                            # Cache strides
+            W=W,
+            BLOCK_SIZE_D=BLOCK_SIZE_D,
+        )
+
+        # Save tensors needed for backward (cache is not needed for current backward)
+        ctx.save_for_backward(x, kernels)
+        ctx.W = W
+        ctx.BLOCK_SIZE_D = BLOCK_SIZE_D
+        # ctx.T_cache = T_cache_val # Not needed for current backward
+
+        return out
+
+    @staticmethod
+    def backward(ctx, grad_out):
+        raise NotImplementedError("Backward of cached fwdbwd is not implemented")
+
+
+# --- User-facing function ---
+def dynamic_conv_triton_cache(x: torch.Tensor, kernels: torch.Tensor, cache: torch.Tensor = None) -> torch.Tensor:
+    """
+    Fused dynamic convolution with autograd support using Triton kernels.
+    Assumes W <= 4.
+
+    Args:
+        x: Input tensor of shape [B, T, D].
+        kernels: Dynamic kernels of shape [B, T, D, W].
+        cache: Optional past context tensor of shape [B, T_cache, D].
+               If provided, treated as concatenated before x for convolution input.
+
+    Returns:
+        Output tensor of shape [B, T, D].
+    """
+    return DynamicConvTritonFunc.apply(x, kernels, cache)

dconv_fwdbwd.py

@@ -0,0 +1,318 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import torch
+import triton
+import triton.language as tl
+from einops import rearrange
+import torch.nn.functional as F
+from torch.autograd import Function
+
+# Helper function to ensure tensors are contiguous for Triton
+def ensure_contiguous(t: torch.Tensor) -> torch.Tensor:
+    return t if t.is_contiguous() else t.contiguous()
+
+# --- Forward Kernel (from previous step, slight modifications for autograd context) ---
+@triton.jit
+def _dynamic_conv_fwd_kernel(
+    X_ptr, K_ptr, Out_ptr,
+    B, T, D,
+    X_stride_b, X_stride_t, X_stride_d,
+    K_stride_b, K_stride_t, K_stride_d, K_stride_w,
+    Out_stride_b, Out_stride_t, Out_stride_d,
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    pid_batch_time = tl.program_id(0)
+    pid_d_block = tl.program_id(1)
+
+    batch_idx = tl.cast(pid_batch_time // T, tl.int64)
+    time_idx = pid_batch_time % T
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    accumulator = tl.zeros((BLOCK_SIZE_D,), dtype=tl.float32)
+    offs_w = tl.arange(0, W)
+
+    # Load Kernels
+    k_ptrs = K_ptr + (batch_idx * K_stride_b + time_idx * K_stride_t +
+                      offs_d[:, None] * K_stride_d + offs_w[None, :] * K_stride_w)
+    k_vals = tl.load(k_ptrs, mask=d_mask[:, None], other=0.0)
+
+    # Load Input X with implicit padding
+    t_in_offs = time_idx + offs_w - W + 1
+    t_in_mask = (t_in_offs >= 0) & (t_in_offs < T)
+    x_ptrs = X_ptr + (batch_idx * X_stride_b + t_in_offs[None, :] * X_stride_t +
+                      offs_d[:, None] * X_stride_d)
+    x_load_mask = d_mask[:, None] & t_in_mask[None, :]
+    x_vals = tl.load(x_ptrs, mask=x_load_mask, other=0.0)
+
+    # Compute and Accumulate
+    product = k_vals * x_vals
+    accumulator += tl.sum(product, axis=1)
+
+    # Store Result
+    out_ptrs = Out_ptr + (batch_idx * Out_stride_b + time_idx * Out_stride_t +
+                          offs_d * Out_stride_d)
+    tl.store(out_ptrs, accumulator, mask=d_mask)
+
+# --- Backward Kernel for Input Gradient (dX) ---
+@triton.jit
+def _dynamic_conv_bwd_dx_kernel(
+    GradOut_ptr, K_ptr, GradX_ptr, # Note: GradX is accumulated into
+    B, T, D,
+    GradOut_stride_b, GradOut_stride_t, GradOut_stride_d,
+    K_stride_b, K_stride_t, K_stride_d, K_stride_w,
+    GradX_stride_b, GradX_stride_t, GradX_stride_d,
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    """
+    Computes gradient w.r.t. input X.
+    Grid: (B * T, cdiv(D, BLOCK_SIZE_D)) - covering GradX output
+    GradX[b, t_x, d] = sum_{w=0}^{W-1} GradOut[b, t, d] * K[b, t, d, w]
+                       where t = t_x + W - 1 - w
+    """
+    pid_batch_time_x = tl.program_id(0) # Covers B * T for output GradX
+    pid_d_block = tl.program_id(1)
+
+    batch_idx = pid_batch_time_x // T
+    time_idx_x = pid_batch_time_x % T # This is t_x
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    # Accumulator for GradX elements
+    accumulator = tl.zeros((BLOCK_SIZE_D,), dtype=tl.float32)
+    offs_w = tl.arange(0, W) # [W]
+
+    # Loop over W to accumulate contributions
+    # Calculate the 't' index needed for GradOut and K based on t_x and w
+    # t = t_x + W - 1 - w
+    t_k_gradout_offs = time_idx_x + W - 1 - offs_w # Shape [W]
+
+    # Mask for valid 't' indices [0, T)
+    t_k_gradout_mask = (t_k_gradout_offs >= 0) & (t_k_gradout_offs < T) # Shape [W]
+
+    # --- Load GradOut ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    gradout_ptrs = GradOut_ptr + (batch_idx * GradOut_stride_b +
+                                  t_k_gradout_offs[None, :] * GradOut_stride_t +
+                                  offs_d[:, None] * GradOut_stride_d)
+    # Combined mask for loading GradOut (valid D and valid t)
+    gradout_load_mask = d_mask[:, None] & t_k_gradout_mask[None, :]
+    # Shape: [BLOCK_SIZE_D, W]
+    gradout_vals = tl.load(gradout_ptrs, mask=gradout_load_mask, other=0.0)
+
+    # --- Load Kernels ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    k_ptrs = K_ptr + (batch_idx * K_stride_b +
+                      t_k_gradout_offs[None, :] * K_stride_t +
+                      offs_d[:, None] * K_stride_d +
+                      offs_w[None, :] * K_stride_w) # Index K with 't' and 'w'
+    # Combined mask for loading K (valid D and valid t)
+    k_load_mask = d_mask[:, None] & t_k_gradout_mask[None, :]
+    # Shape: [BLOCK_SIZE_D, W]
+    k_vals = tl.load(k_ptrs, mask=k_load_mask, other=0.0)
+
+    # --- Compute product and accumulate ---
+    # Shape: [BLOCK_SIZE_D, W]
+    product = gradout_vals * k_vals
+    # Sum contributions over the W dimension
+    accumulator += tl.sum(product, axis=1) # Shape: [BLOCK_SIZE_D]
+
+    # --- Store accumulated gradients ---
+    # Note: This kernel computes the *entire* gradient value for GradX[b, t_x, d_block].
+    # If this kernel could potentially be called multiple times for the same GradX elements
+    # (e.g., in complex graphs), atomic adds would be needed. Here, it seems direct store is fine.
+    gradx_ptrs = GradX_ptr + (batch_idx * GradX_stride_b +
+                              time_idx_x * GradX_stride_t +
+                              offs_d * GradX_stride_d)
+    tl.store(gradx_ptrs, accumulator, mask=d_mask)
+
+
+# --- Backward Kernel for Kernel Gradient (dK) ---
+@triton.jit
+def _dynamic_conv_bwd_dk_kernel(
+    GradOut_ptr, X_ptr, GradK_ptr, # Note: GradK is written directly
+    B, T, D,
+    GradOut_stride_b, GradOut_stride_t, GradOut_stride_d,
+    X_stride_b, X_stride_t, X_stride_d,
+    GradK_stride_b, GradK_stride_t, GradK_stride_d, GradK_stride_w,
+    W: tl.constexpr,
+    BLOCK_SIZE_D: tl.constexpr,
+):
+    """
+    Computes gradient w.r.t. kernels K.
+    Grid: (B * T, cdiv(D, BLOCK_SIZE_D)) - covering GradK output dims B, T, D
+    GradK[b, t, d, w] = GradOut[b, t, d] * X[b, t + w - W + 1, d]
+    """
+    pid_batch_time = tl.program_id(0) # Covers B * T for output GradK
+    pid_d_block = tl.program_id(1)
+
+    batch_idx = pid_batch_time // T
+    time_idx = pid_batch_time % T # This is 't' for GradK and GradOut
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D
+
+    offs_w = tl.arange(0, W) # [W]
+
+    # --- Load GradOut ---
+    # Pointers shape: [BLOCK_SIZE_D] (only depends on b, t, d)
+    gradout_ptrs = GradOut_ptr + (batch_idx * GradOut_stride_b +
+                                  time_idx * GradOut_stride_t +
+                                  offs_d * GradOut_stride_d)
+    # Shape: [BLOCK_SIZE_D]
+    gradout_vals = tl.load(gradout_ptrs, mask=d_mask, other=0.0)
+
+    # --- Load Input X with implicit padding ---
+    # Calculate X's time index: t_x = t + w - W + 1
+    t_in_offs = time_idx + offs_w - W + 1 # Shape [W]
+    # Mask for valid t_x index [0, T)
+    t_in_mask = (t_in_offs >= 0) & (t_in_offs < T) # Shape [W]
+
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    x_ptrs = X_ptr + (batch_idx * X_stride_b +
+                      t_in_offs[None, :] * X_stride_t +
+                      offs_d[:, None] * X_stride_d)
+    # Combined mask for loading X (valid D and valid t_x)
+    x_load_mask = d_mask[:, None] & t_in_mask[None, :] # Shape [BLOCK_SIZE_D, W]
+    # Shape: [BLOCK_SIZE_D, W]
+    x_vals = tl.load(x_ptrs, mask=x_load_mask, other=0.0)
+
+    # --- Compute GradK = GradOut * X ---
+    # Broadcast gradout_vals: [BLOCK_SIZE_D, 1] * [BLOCK_SIZE_D, W] -> [BLOCK_SIZE_D, W]
+    gradk_vals = gradout_vals[:, None] * x_vals # Shape [BLOCK_SIZE_D, W]
+
+    # --- Store gradients for Kernels ---
+    # Pointers shape: [BLOCK_SIZE_D, W]
+    gradk_ptrs = GradK_ptr + (batch_idx * GradK_stride_b +
+                              time_idx * GradK_stride_t +
+                              offs_d[:, None] * GradK_stride_d +
+                              offs_w[None, :] * GradK_stride_w)
+    # Mask only needed for D dimension (W is fully computed)
+    # Store computed gradient values.
+    tl.store(gradk_ptrs, gradk_vals, mask=d_mask[:, None])
+
+
+# --- Autograd Function ---
+class DynamicConvTritonFunc(Function):
+
+    @staticmethod
+    def forward(ctx, x, kernels):
+        """
+        Args:
+            x: Input tensor [B, T, D]
+            kernels: Kernels tensor [B, T, D, W]
+        """
+        x = ensure_contiguous(x)
+        kernels = ensure_contiguous(kernels)
+
+        B, T, D = x.shape
+        W = kernels.shape[3]
+        assert W <= 4, "Kernel W > 4 not expected for this version"
+
+        out = torch.empty_like(x) # Output shape [B, T, D]
+
+        grid = lambda meta: (B * T, triton.cdiv(D, meta['BLOCK_SIZE_D']))
+        BLOCK_SIZE_D = 128 # Consider tuning
+
+        _dynamic_conv_fwd_kernel[grid](
+            x, kernels, out,
+            B, T, D,
+            x.stride(0), x.stride(1), x.stride(2),
+            kernels.stride(0), kernels.stride(1), kernels.stride(2), kernels.stride(3),
+            out.stride(0), out.stride(1), out.stride(2),
+            W=W,
+            BLOCK_SIZE_D=BLOCK_SIZE_D,
+        )
+
+        # Save tensors needed for backward
+        # Need x for dK, need kernels for dX
+        ctx.save_for_backward(x, kernels)
+        # Store W and BLOCK_SIZE_D needed for backward kernel calls
+        ctx.W = W
+        ctx.BLOCK_SIZE_D = BLOCK_SIZE_D
+
+        return out
+
+    @staticmethod
+    def backward(ctx, grad_out):
+        """
+        Args:
+            grad_out: Gradient w.r.t. the output tensor [B, T, D]
+        Returns:
+            grad_x: Gradient w.r.t. input x [B, T, D]
+            grad_kernels: Gradient w.r.t. kernels [B, T, D, W]
+        """
+        grad_out = ensure_contiguous(grad_out)
+        x, kernels = ctx.saved_tensors
+        W = ctx.W
+        BLOCK_SIZE_D = ctx.BLOCK_SIZE_D
+
+        B, T, D = x.shape
+
+        # Initialize gradients
+        # grad_x needs accumulation, start with zeros.
+        grad_x = torch.zeros_like(x)
+        # grad_kernels is computed directly, can use empty_like if kernel handles all writes.
+        # Using empty and relying on kernel writing zeros via masking/other=0.0.
+        grad_kernels = torch.empty_like(kernels)
+
+        # Define grid (can often be the same as forward or similar)
+        grid = lambda meta: (B * T, triton.cdiv(D, meta['BLOCK_SIZE_D']))
+
+        # Kernel call for grad_x
+        _dynamic_conv_bwd_dx_kernel[grid](
+            grad_out, kernels, grad_x,
+            B, T, D,
+            grad_out.stride(0), grad_out.stride(1), grad_out.stride(2),
+            kernels.stride(0), kernels.stride(1), kernels.stride(2), kernels.stride(3),
+            grad_x.stride(0), grad_x.stride(1), grad_x.stride(2),
+            W=W,
+            BLOCK_SIZE_D=BLOCK_SIZE_D,
+        )
+
+        # Kernel call for grad_kernels
+        _dynamic_conv_bwd_dk_kernel[grid](
+            grad_out, x, grad_kernels,
+            B, T, D,
+            grad_out.stride(0), grad_out.stride(1), grad_out.stride(2),
+            x.stride(0), x.stride(1), x.stride(2),
+            grad_kernels.stride(0), grad_kernels.stride(1), grad_kernels.stride(2), grad_kernels.stride(3),
+            W=W,
+            BLOCK_SIZE_D=BLOCK_SIZE_D,
+        )
+
+        # Return gradients in the order inputs were received by forward
+        return grad_x, grad_kernels
+
+# --- User-facing function ---
+def dynamic_conv_triton_autograd(x: torch.Tensor, kernels: torch.Tensor) -> torch.Tensor:
+    """
+    Fused dynamic convolution with autograd support using Triton kernels.
+    Assumes W <= 4.
+
+    Args:
+        x: Input tensor of shape [B, T, D].
+        kernels: Dynamic kernels of shape [B, T, D, W].
+
+    Returns:
+        Output tensor of shape [B, T, D].
+    """
+    return DynamicConvTritonFunc.apply(x, kernels)

dconv_step.py

@@ -0,0 +1,225 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import torch
+import triton
+import triton.language as tl
+from einops import rearrange
+import torch.nn.functional as F
+from typing import Tuple, Optional
+
+# Helper function to ensure tensors are contiguous for Triton
+def ensure_contiguous(t: torch.Tensor) -> torch.Tensor:
+    # Ensure tensor is contiguous in memory.
+    return t if t.is_contiguous() else t.contiguous()
+
+# Removed _apply_activation helper function
+
+@triton.jit
+def _causal_conv_step_kernel(
+    # --- Input/Output Pointers ---
+    X_ptr,         # Pointer to current input x [B, D] (after squeeze)
+    Cache_ptr,     # Pointer to cache [B, D, W], updated IN-PLACE
+    Kernels_ptr,   # Pointer to generated kernels [B, D, W]
+    Out_ptr,       # Pointer to output tensor [B, D]
+
+    # --- Tensor Dimensions ---
+    B, D,          # Batch size, Feature dimension
+
+    # --- Tensor Strides ---
+    X_stride_b, X_stride_d,
+    Cache_stride_b, Cache_stride_d, Cache_stride_w,
+    Kernels_stride_b, Kernels_stride_d, Kernels_stride_w,
+    Out_stride_b, Out_stride_d,
+
+    # --- Kernel Meta-Parameters ---
+    W: tl.constexpr,               # Kernel width (Cache size), passed as compile-time constant (1 < W <= 4)
+    BLOCK_SIZE_D: tl.constexpr,    # Block size for D dimension (tuning parameter)
+    # Removed ACTIVATION: tl.constexpr
+):
+    """
+    Triton kernel for a single step (T=1) of causal dynamic convolution.
+    Updates the cache in-place and computes the output (without activation).
+    Optimized for small W (1 < W <= 4) by manually unrolling the W dimension.
+    Does NOT handle separate static bias.
+
+    Grid: (B, cdiv(D, BLOCK_SIZE_D))
+    Updates Cache[b, d, :] and computes Out[b, d].
+    """
+    # 1. --- Get Program IDs and Calculate Indices ---
+    pid_b = tl.program_id(0)       # Program ID for batch dimension
+    pid_d_block = tl.program_id(1) # Program ID for dimension block
+
+    offs_d = pid_d_block * BLOCK_SIZE_D + tl.arange(0, BLOCK_SIZE_D)
+    d_mask = offs_d < D # Shape: [BLOCK_SIZE_D]
+
+    # 2. --- Load Current Input X ---
+    x_ptrs = X_ptr + pid_b * X_stride_b + offs_d * X_stride_d
+    x_curr = tl.load(x_ptrs, mask=d_mask, other=0.0) # Shape: [BLOCK_SIZE_D]
+
+    # --- Initialize Accumulator ---
+    accumulator = tl.zeros((BLOCK_SIZE_D,), dtype=x_curr.dtype) # Use input dtype
+
+    # --- Manually Unroll Operations for W ---
+    # We will load kernel values and cache values step-by-step
+    # and perform the calculation and cache update.
+
+    # --- Step w = 0 ---
+    # Compute: cache_val_1 * k_val_0 (part 1)
+    # Cache Update: store cache_val_1 at index 0
+    if tl.constexpr(W > 1):
+        # Load k_val_0
+        k_ptr_0 = Kernels_ptr + pid_b * Kernels_stride_b + offs_d * Kernels_stride_d + 0 * Kernels_stride_w
+        k_val_0 = tl.load(k_ptr_0, mask=d_mask, other=0.0)
+
+        # Load cache_val_1 (needed for computation and storing at index 0)
+        cache_ptr_1 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 1 * Cache_stride_w
+        cache_val_1 = tl.load(cache_ptr_1, mask=d_mask, other=0.0)
+
+        # Accumulate Part 1
+        accumulator += cache_val_1 * k_val_0
+
+        # Cache Update: Store cache_val_1 -> cache_ptr_0
+        cache_ptr_0 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 0 * Cache_stride_w
+        tl.store(cache_ptr_0, cache_val_1, mask=d_mask)
+
+    # --- Step w = 1 ---
+    # Compute: cache_val_2 * k_val_1 (part 1)
+    # Cache Update: store cache_val_2 at index 1
+    if tl.constexpr(W > 2):
+        # Load k_val_1
+        k_ptr_1 = Kernels_ptr + pid_b * Kernels_stride_b + offs_d * Kernels_stride_d + 1 * Kernels_stride_w
+        k_val_1 = tl.load(k_ptr_1, mask=d_mask, other=0.0)
+
+        # Load cache_val_2
+        cache_ptr_2 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 2 * Cache_stride_w
+        cache_val_2 = tl.load(cache_ptr_2, mask=d_mask, other=0.0)
+
+        # Accumulate Part 1
+        accumulator += cache_val_2 * k_val_1
+
+        # Cache Update: Store cache_val_2 -> cache_ptr_1
+        cache_ptr_1 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 1 * Cache_stride_w
+        tl.store(cache_ptr_1, cache_val_2, mask=d_mask)
+
+    # --- Step w = 2 ---
+    # Compute: cache_val_3 * k_val_2 (part 1)
+    # Cache Update: store cache_val_3 at index 2
+    if tl.constexpr(W > 3):
+        # Load k_val_2
+        k_ptr_2 = Kernels_ptr + pid_b * Kernels_stride_b + offs_d * Kernels_stride_d + 2 * Kernels_stride_w
+        k_val_2 = tl.load(k_ptr_2, mask=d_mask, other=0.0)
+
+        # Load cache_val_3
+        cache_ptr_3 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 3 * Cache_stride_w
+        cache_val_3 = tl.load(cache_ptr_3, mask=d_mask, other=0.0)
+
+        # Accumulate Part 1
+        accumulator += cache_val_3 * k_val_2
+
+        # Cache Update: Store cache_val_3 -> cache_ptr_2
+        cache_ptr_2 = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + 2 * Cache_stride_w
+        tl.store(cache_ptr_2, cache_val_3, mask=d_mask)
+
+    # --- Final Step (Part 2 and Final Cache Update) ---
+    # Compute: x_curr * k_val_{W-1} (part 2)
+    # Cache Update: store x_curr at index W-1
+
+    # Load k_val_{W-1}
+    k_ptr_last = Kernels_ptr + pid_b * Kernels_stride_b + offs_d * Kernels_stride_d + (W - 1) * Kernels_stride_w
+    k_val_last = tl.load(k_ptr_last, mask=d_mask, other=0.0)
+
+    # Accumulate Part 2
+    accumulator += x_curr * k_val_last
+
+    # Final Cache Update: Store x_curr -> cache_ptr_{W-1}
+    cache_ptr_last = Cache_ptr + pid_b * Cache_stride_b + offs_d * Cache_stride_d + (W - 1) * Cache_stride_w
+    tl.store(cache_ptr_last, x_curr, mask=d_mask)
+
+    # Removed activation application: accumulator = _apply_activation(accumulator, ACTIVATION)
+
+    # 6. --- Store Output ---
+    out_ptrs = Out_ptr + pid_b * Out_stride_b + offs_d * Out_stride_d
+    tl.store(out_ptrs, accumulator, mask=d_mask) # Store result without activation
+
+    # Cache update is now fully handled within the unrolled steps.
+
+
+# --- Python Wrapper Function ---
+def causal_conv_step_triton(
+    x: torch.Tensor,           # Input tensor [B, 1, D]
+    cache: torch.Tensor,       # Cache tensor [B, D, W], modified in-place
+    kernels: torch.Tensor,     # Kernels tensor [B, D, W]
+    # Removed activation parameter
+) -> torch.Tensor:             # Returns output tensor [B, D] (before activation)
+    """
+    Performs one step of causal dynamic convolution using Triton.
+    Updates the cache in-place. Does NOT fuse activation. Assumes 1 < W <= 4.
+    Uses manually unrolled kernel for W dimension.
+
+    Args:
+        x: Current input token tensor of shape [B, 1, D].
+        cache: Cache tensor of shape [B, D, W]. Will be updated in-place.
+        kernels: Dynamically generated kernels tensor of shape [B, D, W].
+
+    Returns:
+        Output tensor of shape [B, D] for the current step (before activation).
+    """
+    # --- Input Validation and Preparation ---
+    assert x.dim() == 3 and x.shape[1] == 1, "Input x must have shape [B, 1, D]"
+    assert cache.dim() == 3, "Cache must have shape [B, D, W]"
+    assert kernels.dim() == 3, "Kernels must have shape [B, D, W]"
+    B, _, D = x.shape
+    W = cache.shape[2]
+    # Updated assertion: W must be > 1 and <= 4
+    assert 1 < W <= 4, f"Kernel W={W}, this optimized version assumes 1 < W <= 4"
+    assert cache.shape[0] == B and cache.shape[1] == D, f"Cache shape mismatch: {cache.shape}"
+    assert kernels.shape == cache.shape, f"Kernels shape mismatch: {kernels.shape}"
+    assert x.is_cuda and cache.is_cuda and kernels.is_cuda, "Inputs must be CUDA tensors"
+    # Allow different input dtypes, but ensure they are compatible or handled
+    # assert x.dtype == cache.dtype == kernels.dtype, "Input dtypes must match"
+
+    # Squeeze the time dimension from input x
+    x_squeezed = x.squeeze(1) # Shape [B, D]
+
+    # Ensure tensors are contiguous for correct stride calculations in Triton
+    x_squeezed = ensure_contiguous(x_squeezed)
+    # Cache MUST be contiguous for in-place updates and loads/stores to work reliably
+    cache = ensure_contiguous(cache)
+    kernels = ensure_contiguous(kernels)
+
+    # Create output tensor with the same dtype as input x
+    out = torch.empty_like(x_squeezed) # Shape [B, D]
+
+    # --- Triton Kernel Launch ---
+    grid = lambda meta: (B, triton.cdiv(D, meta['BLOCK_SIZE_D']))
+    BLOCK_SIZE_D = 64 # Example, tune this value
+
+    # Launch the kernel
+    _causal_conv_step_kernel[grid](
+        x_squeezed, cache, kernels, out,   # Tensor pointers
+        B, D,                              # Dimensions
+        x_squeezed.stride(0), x_squeezed.stride(1), # x strides
+        cache.stride(0), cache.stride(1), cache.stride(2), # cache strides
+        kernels.stride(0), kernels.stride(1), kernels.stride(2), # kernels strides
+        out.stride(0), out.stride(1),      # out strides
+        # --- Meta-parameters ---
+        W=W,                               # Pass W as constexpr
+        BLOCK_SIZE_D=BLOCK_SIZE_D,         # Pass BLOCK_SIZE_D as constexpr
+        # Removed ACTIVATION=activation
+    )
+
+    return out # Return the computed output [B, D] (before activation)

dynamic_conv.py

@@ -0,0 +1,274 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from __future__ import annotations
+
+from collections import OrderedDict
+from typing import Optional, Tuple, Callable
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+from einops import rearrange
+
+from transformers.activations import ACT2FN
+
+from .dconv_fwdbwd import dynamic_conv_triton_autograd
+from .dconv_fwd_cache import dynamic_conv_triton_cache
+from .dconv_step import causal_conv_step_triton
+
+
+class DynamicShortConvolution(nn.Module):
+    """
+    Simple wrapper around `nn.Conv1d` that accepts dimension last.
+    """
+
+    def __init__(
+        self,
+        hidden_size: int,
+        kernel_size: int,
+        generator_input_size: Optional[int] = None,
+        generator_reduction: Optional[int] = None,
+        generator_activation: str = 'silu',
+        activation: Optional[str] = 'silu',
+        static_conv_init: Callable = None,
+        use_fast_conv1d: bool = True,
+        implementation: str = "naive",
+    ) -> DynamicShortConvolution:
+        super().__init__()
+
+        self.hidden_size = hidden_size
+        self.generator_input_size = hidden_size if generator_input_size is None else generator_input_size
+        self.generator_hidden_size = hidden_size if generator_reduction is None else (hidden_size // generator_reduction)
+        self.kernel_size = kernel_size
+        self.activation = None
+        self.use_fast_conv1d = use_fast_conv1d
+        self.implementation = implementation
+
+        if activation is not None:
+            assert activation in ['silu', 'swish'], f"Activation `{activation}` not supported yet."
+            self.activation = activation
+        
+        self.static_conv_init = static_conv_init
+        
+        self.kernel_generator = nn.Sequential(
+            OrderedDict([
+                ("w1", nn.Linear(self.generator_input_size, self.generator_hidden_size, bias=False)),
+                ("act", ACT2FN[generator_activation]),
+                ("w2", nn.Linear(self.generator_hidden_size, self.hidden_size * self.kernel_size, bias=True)),
+            ])
+        )
+        self._init_kernel_generator()
+
+    def _init_kernel_generator(self):
+        """
+        Initialize the kernel generator.
+        """
+        for layer in self.kernel_generator:
+            if isinstance(layer, nn.Linear):
+                layer.weight.data.zero_()
+                if layer.bias is not None:
+                    layer.bias.data.zero_()
+        
+        if self.static_conv_init is not None:
+            # init for static_bias
+            self.static_conv_init(self.kernel_generator.w2.bias)
+
+    def get_kernel(self, x: torch.Tensor) -> torch.Tensor:
+        flat_kernels = self.kernel_generator(x)
+        if flat_kernels.dim() == 3:
+            kernels = rearrange(flat_kernels, 'b t (d w) -> b t d w', w=self.kernel_size)
+        elif flat_kernels.dim() == 2:
+            kernels = rearrange(flat_kernels, 'b (d w) -> b d w', w=self.kernel_size)
+        else:
+            raise ValueError(f"Invalid kernel shape: {flat_kernels.shape}")
+        return kernels
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        mask: Optional[torch.Tensor] = None,
+        cache: Optional[torch.Tensor] = None,
+        output_final_state: bool = False,
+        cu_seqlens: Optional[torch.LongTensor] = None,
+        generator_input: Optional[torch.Tensor] = None,
+        **kwargs,
+    ) -> Tuple[torch.Tensor, torch.Tensor]:
+        """
+        Args:
+            x (`torch.Tensor`):
+                Tensor of shape `[B, T, D]`.
+                If `seq_idx` is provided, `B` must be 1.
+            mask (`Optional[torch.Tensor]`):
+                Attention mask dealing with padded positions.
+            cache (`Optional[torch.Tensor]`):
+                Previous cache tensor of shape `[N, D, W]`, where `W` is the kernel size.
+                If provided, the cache is updated **inplace**.
+            output_final_state (Optional[bool]):
+                Whether to output the final state of shape `[N, D, W]`. Default: `False`.
+            cu_seqlens (Optional[torch.LongTensor]):
+                Cumulative sequence lengths for each batch. Used for varlen. Default: `None`.
+                Shape: [B+1]
+
+        Returns:
+            Tensor of shape `[B, T, D]`.
+        """
+
+        """
+        x: [B, T, D]
+        return: [B, T, D]
+        """
+        
+        assert cu_seqlens is None, "cu_seqlens not supported yet."
+        
+        B, T, D, W = *x.shape, self.kernel_size
+        N = B
+
+        input_dtype = x.dtype
+
+        if mask is not None:
+            x = x.mul_(mask.unsqueeze(-1))
+
+        implementation = self.implementation
+        if implementation == "triton" and not self.training:
+            implementation = "triton_cache"
+
+        # during the decoding phase, we assume the batch is composed of sequences of length 1
+        if cache is not None and B * T == N:
+            assert T == 1
+            if implementation in ["naive", "triton_training"]:
+                x, cache = self._step_naive(x, cache, cu_seqlens, generator_input=generator_input)
+            elif implementation in ["triton", "triton_cache", "triton_decoding"]:
+                x, cache = self._step_triton(x, cache, cu_seqlens, generator_input=generator_input)
+            else:
+                raise ValueError(f"Unknown implementation: {implementation}")
+            return x, cache
+
+        if output_final_state:
+            new_cache = rearrange(x[..., -min(W, T):, :], 'n w d -> n d w')
+        else:
+            new_cache = None
+        
+        if implementation in ["naive", "triton_decoding"]:
+            x = self._forward_naive(x, generator_input=generator_input)  # [B, T, D]
+        elif implementation in ["triton", "triton_training"]:
+            assert cache is None, "Cache not supported in pure triton mode. Please set model.eval() or use triton_cache mode."
+            x = self._forward_triton(x, generator_input=generator_input)
+        elif implementation == "triton_cache":
+            x = self._forward_triton_cache(x, generator_input=generator_input, cache=cache)
+        else:
+            raise ValueError(f"Unknown implementation: {implementation}")
+
+        if self.activation is not None:
+            x = ACT2FN[self.activation](x)
+        
+        x = x.to(input_dtype)
+        if output_final_state:
+            if cache is None:
+                cache = x.new_zeros(N, D, W)
+            cache[:, :, -min(W, T):].copy_(new_cache)
+
+        return x, cache
+
+    def _forward_naive(self, x: torch.Tensor, generator_input: Optional[torch.Tensor] = None) -> torch.Tensor:
+        W = self.kernel_size
+        generator_input = x if generator_input is None else generator_input
+        kernels = self.get_kernel(generator_input)
+        x = F.pad(x.transpose(1, 2), (W - 1, 0))  # [B, D, T+W-1]
+        x = x.unfold(dimension=2, size=W, step=1)  # [B, D, T, W]
+        x = x.permute(0, 2, 1, 3)  # [B, T, D, W]
+        x = (x * kernels).sum(dim=-1)  # [B, T, D]
+        return x
+
+    def _forward_triton(self, x: torch.Tensor, generator_input: Optional[torch.Tensor] = None) -> torch.Tensor:
+        generator_input = x if generator_input is None else generator_input
+        kernels = self.get_kernel(generator_input)
+        output_triton = dynamic_conv_triton_autograd(x, kernels)
+        return output_triton
+
+    @torch.no_grad
+    def _forward_triton_cache(self, x: torch.Tensor, generator_input: Optional[torch.Tensor] = None, cache: Optional[torch.Tensor] = None) -> torch.Tensor:
+        generator_input = x if generator_input is None else generator_input
+        assert not self.training, "Triton implementation is only available in eval mode."
+        # cache: [B, D, T(W)]
+        CHUNK_SIZE = 2048
+        n_chunk = (x.shape[1] + CHUNK_SIZE - 1) // CHUNK_SIZE
+        output_triton = torch.zeros_like(x)
+        if cache is not None:
+            cache = rearrange(cache, "b d t -> b t d")  # [B, T(W), D]
+        for i in range(n_chunk):
+            start = i * CHUNK_SIZE
+            end = min((i + 1) * CHUNK_SIZE, x.shape[1])
+            kernels = self.get_kernel(generator_input[:, start:end])
+            out = dynamic_conv_triton_cache(x[:, start:end], kernels, cache=cache)
+            output_triton[:, i*CHUNK_SIZE:end, :] = out
+            cache = x[:, end-self.kernel_size:end, :]
+        return output_triton
+
+    def _step_naive(
+        self,
+        x: torch.Tensor,
+        cache: torch.Tensor,
+        cu_seqlens: Optional[torch.LongTensor] = None,
+        generator_input: Optional[torch.Tensor] = None
+    ) -> Tuple[torch.Tensor, torch.Tensor]:
+        assert x.shape[1] == 1, "x must be of shape [B, 1, D]"
+        shape = x.shape
+        generator_input = x if generator_input is None else generator_input
+        x = x.squeeze(1)
+        generator_input = generator_input.squeeze(1) # Shape [B, D]
+        B, D, W = *x.shape, self.kernel_size
+
+        # we follow the fast mode that updates the cache in-place
+        cache.copy_(cache.roll(shifts=-1, dims=-1))
+        cache[:, :, -1] = x # [B, D, T(W)]
+        
+        kernels = self.get_kernel(generator_input) # [B, D, W]
+        x = torch.sum(cache * kernels, dim=-1)
+        
+        if self.activation is not None:
+            x = ACT2FN[self.activation](x)
+        
+        return x.view(shape), cache
+    
+    def _step_triton(
+        self,
+        x: torch.Tensor,
+        cache: torch.Tensor,
+        cu_seqlens: Optional[torch.LongTensor] = None,
+        generator_input: Optional[torch.Tensor] = None
+    ) -> Tuple[torch.Tensor, torch.Tensor]:
+        # --- Triton Implementation ---
+        assert x.shape[1] == 1, "x must be of shape [B, 1, D]"
+        shape = x.shape # Keep original shape [B, 1, D] for return
+        generator_input = x if generator_input is None else generator_input
+
+        # 1. Generate kernels
+        kernels_triton = self.get_kernel(generator_input.squeeze(1)) # [B, D, W]
+
+        # 2. Call Triton kernel without activation
+        x_out_triton = causal_conv_step_triton(
+            x,
+            cache,
+            kernels_triton,
+        )
+        
+        # Apply activation (if any) after kernel execution
+        if self.activation is not None:
+            x_out_triton = ACT2FN[self.activation](x_out_triton)
+
+        # 3. Return reshaped output and the *same cache tensor* (it was updated in-place)
+        return x_out_triton.view(shape), cache

generation_config.json

@@ -0,0 +1,6 @@
+{
+  "_from_model_config": true,
+  "bos_token_id": 151643,
+  "eos_token_id": 151643,
+  "transformers_version": "4.52.0.dev0"
+}

jet_block.py

@@ -0,0 +1,258 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+# This file is modified from https://github.com/fla-org/flash-linear-attention/blob/main/fla/layers/gated_deltanet.py
+
+from __future__ import annotations
+
+import math
+from dataclasses import dataclass
+from typing import Optional, Tuple
+
+import torch
+import torch.nn as nn
+from torch.nn import functional as F
+from einops import rearrange
+
+from fla.layers.utils import get_unpad_data, index_first_axis, pad_input
+from fla.modules import FusedRMSNormGated
+from fla.ops.gated_delta_rule import (chunk_gated_delta_rule,
+                                      fused_recurrent_gated_delta_rule)
+
+
+from .dynamic_conv import DynamicShortConvolution
+from .configuration_jet_nemotron import JetNemotronConfig
+from .kv_cache import JetNemotronCache
+
+
+@dataclass
+class JetBlockConfig():
+    mode: str = 'chunk'
+    expand_v: int = 2.0
+    num_heads: int = 6
+    head_dim: int = 256
+    norm_eps: float = 1e-5
+    conv_size: int = 4
+    dconv_generator_reduction: int = None
+    dconv_implementation: str = 'triton'
+
+
+def init_linear_conv1d(weight: torch.Tensor, std: float, bias: Optional[torch.Tensor] = None) -> None:
+    weight.data.normal_(mean=0.0, std=std)
+    if bias is not None:
+        if not getattr(bias, "_no_reinit", False):
+            nn.init.zeros_(bias)
+
+
+class JetBlock(nn.Module):
+    def __init__(
+        self,
+        config: Optional[JetNemotronConfig] = None,
+        layer_type: str = 'jet',
+        layer_idx: Optional[int] = None,
+        hidden_size: Optional[int] = None,
+        initializer_range: Optional[float] = None,
+        jet_block_config: Optional[JetBlockConfig] = None
+    ) -> JetBlock:
+        super().__init__()
+
+        if jet_block_config is None:
+            assert config.efficient_attention_config is not None, "Efficient attention config must be provided in JetConfig."
+            assert layer_type in config.efficient_attention_config, \
+                f"{layer_type} configuration must be provided in efficient_attention_config."
+            jet_block_config = JetBlockConfig(**config.efficient_attention_config[layer_type])
+
+        hidden_size = hidden_size or config.hidden_size
+        initializer_range = initializer_range or config.initializer_range
+
+        self.mode = jet_block_config.mode
+
+        self.hidden_size = hidden_size
+        self.expand_v = jet_block_config.expand_v
+
+        self.conv_size = jet_block_config.conv_size
+
+        self.head_dim = jet_block_config.head_dim
+        self.num_heads = jet_block_config.num_heads
+
+        self.key_dim = int(self.num_heads * self.head_dim)
+        self.value_dim = int(self.key_dim * self.expand_v)
+        self.head_k_dim = jet_block_config.head_dim
+        self.head_v_dim = int(jet_block_config.head_dim * self.expand_v)
+        self.layer_idx = layer_idx
+
+        self.autotune_interval = 32 * 16 * 1024 # 32 batch size * 16 num head * 1024 sequence length
+
+        # Consistency check: Ensure expand_v produces integer values
+        if not math.isclose(self.key_dim * self.expand_v, self.value_dim, rel_tol=1e-5):
+            raise ValueError(
+                f"expand_v={self.expand_v} does not produce an integer value when multiplied by key_dim={self.key_dim}. "
+                f"Resulting value_dim would be {self.key_dim * self.expand_v}, which is invalid for nn.Linear."
+            )
+        if not math.isclose(self.head_dim * self.expand_v, self.head_v_dim, rel_tol=1e-5):
+            raise ValueError(
+                f"expand_v={self.expand_v} does not produce an integer value when multiplied by head_dim={self.head_dim}. "
+                f"Resulting head_v_dim would be {self.head_dim * self.expand_v}, which is invalid for FusedRMSNormGated."
+            )
+        assert self.mode in ['chunk', 'fused_recurrent'], f"Not suppoerted mode `{jet_block_config.mode}`."
+
+        self.q_proj = nn.Linear(hidden_size, self.key_dim, bias=False)
+        self.k_proj = nn.Linear(hidden_size, self.key_dim, bias=False)
+        self.v_proj = nn.Linear(hidden_size, self.value_dim, bias=False)
+        self.b_proj = nn.Linear(hidden_size, self.num_heads, bias=False)
+        self.a_proj = nn.Linear(hidden_size, self.num_heads, bias=False)
+
+        A = torch.empty(self.num_heads, dtype=torch.float32).uniform_(0, 16)
+        self.A_log = nn.Parameter(torch.log(A))
+        self.A_log._no_weight_decay = True
+        # hard coded for now
+        dt_min = 0.001
+        dt_max = 0.1
+        dt_init_floor = 1e-4
+        dt = torch.exp(
+            torch.rand(self.num_heads) * (math.log(dt_max) - math.log(dt_min))
+            + math.log(dt_min)
+        )
+        dt = torch.clamp(dt, min=dt_init_floor)
+        # Inverse of softplus: https://github.com/pytorch/pytorch/issues/72759
+        inv_dt = dt + torch.log(-torch.expm1(-dt))
+        self.dt_bias = nn.Parameter(inv_dt)
+        # Just to be explicit. Without this we already don't put wd on dt_bias because of the check
+        # name.endswith("bias") in param_grouping.py
+        self.dt_bias._no_weight_decay = True
+
+        self.dynamic_conv1d = DynamicShortConvolution(
+            hidden_size=self.value_dim,
+            kernel_size=self.conv_size,
+            generator_input_size=self.hidden_size,
+            generator_reduction=jet_block_config.dconv_generator_reduction,
+            static_conv_init=lambda x: init_linear_conv1d(x, std=initializer_range),
+            implementation=jet_block_config.dconv_implementation,
+        )
+
+        self.g_proj = nn.Linear(hidden_size, self.value_dim, bias=False)
+        self.o_norm = FusedRMSNormGated(self.head_v_dim, eps=float(jet_block_config.norm_eps), autotune_interval=self.autotune_interval)
+        self.o_proj = nn.Linear(self.value_dim, hidden_size, bias=False)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        past_key_value: Optional[JetNemotronCache] = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        use_cache: Optional[bool] = False,
+        **kwargs
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[JetNemotronCache]]:
+        if attention_mask is not None:
+            if len(attention_mask.shape) > 2:
+                attention_mask = attention_mask.squeeze(1)
+                attention_mask = torch.where(attention_mask[:, -1] > -1, 1, 0)
+
+            assert len(attention_mask.shape) == 2, (
+                "Expected attention_mask as a 0-1 matrix with shape [batch_size, seq_len] "
+                "for padding purposes (0 indicating padding). "
+                "Arbitrary attention masks of shape [batch_size, seq_len, seq_len] are not allowed."
+            )
+
+        batch_size, q_len, _ = hidden_states.shape
+        # change to inference mode.
+        mode = 'fused_recurrent' if q_len <= 64 else self.mode
+        if self.training:
+            assert mode == 'chunk', "Only chunk mode is supported in training."
+
+        last_state = None
+        if past_key_value is not None and len(past_key_value) > self.layer_idx:
+            last_state = past_key_value[self.layer_idx]
+
+        cu_seqlens = kwargs.get('cu_seqlens', None)
+        if attention_mask is not None and q_len > 1:
+            indices, cu_seqlens, _ = get_unpad_data(attention_mask[:, -q_len:])
+
+        conv_state = None
+
+        conv_mask = attention_mask[:, -hidden_states.shape[1]:] if attention_mask is not None else None
+
+        q = F.silu(self.q_proj(hidden_states))
+        k = F.silu(self.k_proj(hidden_states))
+
+        conv_state_v = None
+        if last_state is not None:
+            conv_state_v = last_state['conv_state'][-1]
+        v, conv_state_v = self.dynamic_conv1d(
+            x=self.v_proj(hidden_states),
+            generator_input=hidden_states,
+            mask=conv_mask,
+            cache=conv_state_v,
+            output_final_state=use_cache,
+        )
+        conv_state = conv_state + (conv_state_v,) if conv_state is not None else (conv_state_v,)
+
+        if attention_mask is not None and q_len > 1:
+            q = index_first_axis(rearrange(q, "b s ... -> (b s) ..."), indices).unsqueeze(0)
+            k = index_first_axis(rearrange(k, "b s ... -> (b s) ..."), indices).unsqueeze(0)
+            v = index_first_axis(rearrange(v, "b s ... -> (b s) ..."), indices).unsqueeze(0)
+            hidden_states = index_first_axis(rearrange(hidden_states, "b s ... -> (b s) ..."), indices).unsqueeze(0)
+        
+        q, k = map(lambda x: rearrange(x, '... (h d) -> ... h d', d=self.head_k_dim), (q, k))
+        v = rearrange(v, '... (h d) -> ... h d', d=self.head_v_dim)
+        beta = self.b_proj(hidden_states).sigmoid()
+
+        g = -self.A_log.float().exp() * F.softplus(self.a_proj(hidden_states).float() + self.dt_bias)
+
+        recurrent_state = last_state['recurrent_state'] if last_state is not None else None
+        if mode == 'chunk':
+            o, recurrent_state = chunk_gated_delta_rule(
+                q=q,
+                k=k,
+                v=v,
+                g=g,
+                beta=beta,
+                initial_state=recurrent_state,
+                output_final_state=use_cache,
+                cu_seqlens=cu_seqlens,
+                use_qk_l2norm_in_kernel=True,
+                autotune_interval=self.autotune_interval
+            )
+        elif mode == 'fused_recurrent':
+            o, recurrent_state = fused_recurrent_gated_delta_rule(
+                q=q,
+                k=k,
+                v=v,
+                g=g,
+                beta=beta,
+                initial_state=recurrent_state,
+                output_final_state=use_cache,
+                cu_seqlens=cu_seqlens,
+                use_qk_l2norm_in_kernel=True
+            )
+        else:
+            raise NotImplementedError(f"Not supported mode `{mode}`.")
+
+        if past_key_value is not None:
+            past_key_value.update(
+                recurrent_state=recurrent_state,
+                conv_state=conv_state,
+                layer_idx=self.layer_idx,
+                offset=q_len
+            )
+
+        g = rearrange(self.g_proj(hidden_states), '... (h d) -> ... h d', d=self.head_v_dim)
+        o = self.o_norm(o, g)
+        o = rearrange(o, 'b t h d -> b t (h d)')
+        o = self.o_proj(o)
+        if attention_mask is not None and q_len > 1:
+            o = pad_input(o.squeeze(0), indices, batch_size, q_len)
+
+        return o, past_key_value

kv_cache.py

@@ -0,0 +1,212 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+from __future__ import annotations
+
+from typing import Any, Optional
+
+import torch
+import transformers
+
+__all__ = ["JetCache"]
+
+
+class JetNemotronCache(transformers.cache_utils.Cache):
+
+    def __init__(
+        self,
+        seen_tokens: int = 0
+    ) -> JetNemotronCache:
+
+        self.states: list[dict[str, Any]] = []
+        self.layer_wise_states: dict[str, Any] = {}
+
+        self._base_seen_tokens = seen_tokens 
+        self._seen_tokens = []  # Used in `generate` to keep tally of how many tokens the cache has seen
+
+    def __getitem__(self, layer_idx: int) -> dict[str, Any]:
+        if layer_idx < len(self):
+            return self.states[layer_idx]
+        else:
+            raise KeyError(f"Cache only has {len(self)} layers, attempted to access layer with index {layer_idx}")
+
+    def __iter__(self):
+        for state in self.states:
+            yield state
+
+    def __len__(self):
+        return len(self.states)
+
+    def update(
+        self,
+        recurrent_state: torch.Tensor = None,
+        attn_state: tuple[torch.Tensor, torch.Tensor] = None,
+        conv_state: tuple[torch.Tensor] = None,
+        ffn_state: torch.Tensor = None,
+        layer_idx: int = 0,
+        offset: Optional[int] = 1,
+        increase_seen_tokens: bool = True,
+        cache_kwargs: dict[str, Any] = {},
+    ) -> dict[str, Any]:
+        """
+        Updates the cache with the new `recurrent_state`/`attn_state`/`conv_state` for the layer `layer_idx`.
+
+        Args:
+            recurrent_state (`torch.Tensor`, `optional`):
+                The new recurrent state to cache.
+            attn_state (`Tuple[torch.Tensor, torch.Tensor]`, `optional`):
+                The new attention key/value states to cache.
+            conv_state (`Tuple[torch.Tensor]`, `optional`):
+                The new convolution state to cache.
+            layer_idx (`int`, defaults to 0):
+                The index of the layer to cache the states for.
+            offset (`int`, `optional`, defaults to 1):
+                The number of new tokens being processed.
+            cache_kwargs (`Dict[str, Any]`, `optional`):
+                Additional arguments for the cache subclass.
+
+        Return:
+            Dictionary of the updated state.
+        """
+        if len(self._seen_tokens) <= layer_idx:
+            self._seen_tokens.append(self._base_seen_tokens)
+
+        # Update the number of seen tokens
+        if increase_seen_tokens:
+            self.increase_seen_tokens(layer_idx, offset)
+            
+        if attn_state is not None:
+            input_size = attn_state[0].shape[-2]
+            window_size = cache_kwargs.get('window_size', None)
+            if not isinstance(attn_state, tuple) or len(attn_state) != 2:
+                raise ValueError("`attn_state` must be a tuple of two tensors for key/value states")
+        if len(self.states) <= layer_idx:
+            # in prefilling stage
+            state = dict(
+                recurrent_state=recurrent_state,
+                attn_state=attn_state,
+                conv_state=conv_state,
+                ffn_state=ffn_state
+            )
+            if attn_state is not None and window_size is not None:
+                # in prefilling stage, the cached and returned key/value states are different
+                # original key/value states are returned, but the cached states are the last `window_size` tokens
+                _key_state = attn_state[0][..., -window_size:, :]
+                _value_state = attn_state[1][..., -window_size:, :]
+
+                _attn_state = (_key_state, _value_state)
+                _state = dict(
+                    recurrent_state=recurrent_state,
+                    attn_state=_attn_state,
+                    conv_state=conv_state,
+                    ffn_state=ffn_state
+                )
+                self.states.append(_state)
+            else:
+                self.states.append(state)
+        else:
+            state = self.states[layer_idx]
+            if recurrent_state is not None:
+                state['recurrent_state'] = recurrent_state
+            if attn_state is not None:
+                key_state, value_state = state['attn_state']
+                assert window_size is None or key_state.shape[-2] <= window_size
+                if window_size is not None and key_state.shape[-2] == window_size and input_size == 1:
+                    # DO NOT allocate new memory if the cache is full
+                    # only works in decoding stage
+                    # roll the key/value states to the left by `input_size`
+                                        
+                    key_state = key_state.roll(-input_size, -2)
+                    value_state = value_state.roll(-input_size, -2)
+                                        
+                    # replace the last `input_size` tokens with the new key/value states
+                    key_state[..., -input_size:, :] = attn_state[0]
+                    value_state[..., -input_size:, :] = attn_state[1]
+                    
+                    attn_state = (key_state, value_state)
+                else:
+                    # <= window_size or not sliding window or chunk-prefilling (input_size > 1)
+                    attn_state = (torch.cat([key_state, attn_state[0]], -2),
+                                  torch.cat([value_state, attn_state[1]], -2),)
+                state['attn_state'] = attn_state
+            if conv_state is not None:
+                state['conv_state'] = conv_state
+            if ffn_state is not None:
+                state['ffn_state'] = ffn_state
+
+        assert len(self.states) == len(self._seen_tokens)
+
+        return state
+
+    def trim_attn_state(self, layer_idx: int, window_size: int) -> None:
+        # handle the case when the input length of SWA > 1 and has a cache, especially the chunk-prefilling case
+        # this function is called after attention is donw
+        assert layer_idx < len(self.states), f"Layer index {layer_idx} out of range for states with length {len(self.states)}"
+        state = self.states[layer_idx]
+        assert state["attn_state"] is not None, f"Layer {layer_idx} does not have an attention state"
+        key_state, value_state = state["attn_state"]
+        if key_state.shape[-2] > window_size:
+            state["attn_state"] = (
+                key_state[..., -window_size:, :],
+                value_state[..., -window_size:, :],
+            )
+
+    def increase_seen_tokens(self, layer_idx: int, offset: int = 1) -> None:
+        """Increases the number of seen tokens for the layer `layer_idx` by `offset`."""
+        self._seen_tokens[layer_idx] += offset
+
+    def get_seq_length(self, layer_idx: Optional[int] = 0) -> int:
+        """Returns the sequence length of the cached states. A layer index can be optionally passed."""
+        if len(self._seen_tokens) <= layer_idx:
+            return self._base_seen_tokens
+        return self._seen_tokens[layer_idx]
+
+    def get_max_length(self) -> Optional[int]:
+        """Returns the maximum sequence length of the cached states. Cache does not have a maximum length."""
+        return None
+
+    def to_legacy_cache(self) -> tuple:
+        return tuple(self.states)
+
+    def print_kv_sizes(self) -> None:
+        """Returns the size of the cached key/value states."""
+        for layer_idx, state in enumerate(self.states):
+            if state.get("attn_state", None) is not None:
+                key_state, value_state = state["attn_state"]
+                # compute state size in MB
+                key_size = key_state.element_size() * key_state.nelement() / (1024**2)
+                value_size = value_state.element_size() * value_state.nelement() / (1024**2)
+                print(key_state.shape, value_state.shape)
+                print(f"Layer {layer_idx}: Attention. cache size: {key_size + value_size:.2f} MB")
+            if state.get("conv_state", None) is not None:
+                conv_state = state["conv_state"]
+                # compute state size in MB
+                conv_sizes = []
+                for conv in conv_state:
+                    conv_size = conv.element_size() * conv.nelement() / (1024**2)
+                    conv_sizes.append(conv_size)
+                conv_size = sum(conv_sizes)
+                print(f"Layer {layer_idx}: Convolution. cache size: {conv_size:.2f} MB")
+            if state.get("ffn_state", None) is not None:
+                ffn_state = state["ffn_state"]
+                # compute state size in MB
+                ffn_size = ffn_state.element_size() * ffn_state.nelement() / (1024**2)
+                print(f"Layer {layer_idx}: FFN. cache size: {ffn_size:.2f} MB")
+            if state.get("recurrent_state", None) is not None:
+                recurrent_state = state["recurrent_state"]
+                # compute state size in MB
+                recurrent_size = recurrent_state.element_size() * recurrent_state.nelement() / (1024**2)
+                print(f"Layer {layer_idx}: Recurrent. cache size: {recurrent_size:.2f} MB")

model-00001-of-00002.safetensors

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

model-00002-of-00002.safetensors

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

model.safetensors.index.json

@@ -0,0 +1,597 @@
+{
+  "metadata": {
+    "total_size": 7920850368
+  },
+  "weight_map": {
+    "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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.0.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.1.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.10.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.11.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.12.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.13.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.14.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.15.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.18.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.2.self_attn.g_proj.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_norm.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_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-00001-of-00002.safetensors",
+    "model.layers.20.mlp.down_proj.weight": "model-00001-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-00001-of-00002.safetensors",
+    "model.layers.20.self_attn.k_proj.bias": "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_proj.bias": "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.bias": "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-00001-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_proj.bias": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "model.layers.21.self_attn.v_proj.weight": "model-00001-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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.23.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.24.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.26.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.28.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.29.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.3.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.30.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.31.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.33.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.34.self_attn.g_proj.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_norm.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_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.A_log": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.a_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.b_proj.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.dt_bias": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00002-of-00002.safetensors",
+    "model.layers.35.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.4.self_attn.g_proj.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_norm.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_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_proj.bias": "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_proj.bias": "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.bias": "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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.6.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.7.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.8.self_attn.g_proj.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_norm.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_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.A_log": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.a_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.b_proj.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.dt_bias": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.dynamic_conv1d.kernel_generator.w1.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.dynamic_conv1d.kernel_generator.w2.bias": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.dynamic_conv1d.kernel_generator.w2.weight": "model-00001-of-00002.safetensors",
+    "model.layers.9.self_attn.g_proj.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_norm.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_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_jet_nemotron.py

@@ -0,0 +1,967 @@
+# Copyright 2025 NVIDIA CORPORATION & AFFILIATES
+#
+# 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.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+# This file is modified from https://github.com/huggingface/transformers/blob/main/src/transformers/models/qwen2/modeling_qwen2.py
+
+from functools import partial
+from typing import Callable, Optional, Tuple, Union
+
+import torch
+from torch import nn
+
+from transformers.activations import ACT2FN
+from transformers.generation import GenerationMixin, GenerationConfig
+from transformers.modeling_attn_mask_utils import AttentionMaskConverter
+from transformers.modeling_flash_attention_utils import FlashAttentionKwargs
+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, ALL_ATTENTION_FUNCTIONS
+from transformers.processing_utils import Unpack
+from transformers.utils import (
+    LossKwargs,
+    add_start_docstrings,
+    add_start_docstrings_to_model_forward,
+    can_return_tuple,
+    logging,
+    replace_return_docstrings,
+)
+from transformers.utils.deprecation import deprecate_kwarg
+
+
+from .configuration_jet_nemotron import JetNemotronConfig
+from .jet_block import JetBlock
+from .kv_cache import JetNemotronCache
+
+try:
+    from .dynamic_conv import DynamicShortConvolution
+    from .dconv_fwdbwd import dynamic_conv_triton_autograd
+    from .dconv_fwd_cache import dynamic_conv_triton_cache
+    from .dconv_step import causal_conv_step_triton
+except ImportError:
+    raise ImportError(
+        "Dynamic convolution is not available. Please install the required dependencies to use this feature."
+    )
+
+logger = logging.get_logger(__name__)
+
+_CHECKPOINT_FOR_DOC = "jet-ai/Jet-Nemotron-2B"
+_CONFIG_FOR_DOC = "JetNemotronConfig"
+
+
+class JetNemotronMLP(nn.Module):
+    def __init__(self, config: JetNemotronConfig):
+        super().__init__()
+        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, hidden_state):
+        return self.down_proj(self.act_fn(self.gate_proj(hidden_state)) * self.up_proj(hidden_state))
+
+
+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 JetNemotronAttention(nn.Module):
+    """Multi-headed attention from 'Attention Is All You Need' paper"""
+
+    def __init__(self, config: JetNemotronConfig, layer_idx: Optional[int] = None, sliding_window: Optional[int] = None):
+        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.sliding_window = sliding_window
+        self.q_proj = nn.Linear(config.hidden_size, config.num_attention_heads * self.head_dim, bias=True)
+        self.k_proj = nn.Linear(config.hidden_size, config.num_key_value_heads * self.head_dim, bias=True)
+        self.v_proj = nn.Linear(config.hidden_size, config.num_key_value_heads * self.head_dim, bias=True)
+        self.o_proj = nn.Linear(config.num_attention_heads * self.head_dim, config.hidden_size, bias=False)
+
+    def _get_target_length(
+        self,
+        sequence_length: int,
+        past_key_values: JetNemotronCache,
+    ):
+        past_seen_tokens = past_key_values.get_seq_length(self.layer_idx) if past_key_values is not None else 0
+        target_length = sequence_length + min(past_seen_tokens, self.sliding_window - 1)
+        return target_length
+
+    def _update_causal_mask_for_sliding_window(
+        self,
+        attention_mask: torch.Tensor,
+        input_tensor: torch.Tensor,
+        past_key_values: JetNemotronCache,
+    ) -> torch.Tensor:
+
+        dtype, device = input_tensor.dtype, input_tensor.device
+        min_dtype = torch.finfo(dtype).min
+        sequence_length = input_tensor.shape[1]
+
+        target_length = self._get_target_length(sequence_length, past_key_values)
+
+        past_seen_tokens = past_key_values.get_seq_length(self.layer_idx) if past_key_values is not None else 0
+        cache_position = torch.arange(
+            past_seen_tokens, past_seen_tokens + sequence_length, device=input_tensor.device
+        )
+        
+        if attention_mask is not None:
+            # left padding
+            assert attention_mask.dim() == 4, "Attention mask must be 4D"
+            diagonal_attend_mask = attention_mask < -1
+            diagonal_attend_mask = diagonal_attend_mask[:, :, :, -target_length:]
+        else:
+            diagonal_attend_mask = torch.arange(target_length, device=device) > cache_position.reshape(-1, 1)
+            diagonal_attend_mask = diagonal_attend_mask[None, None, :, :]
+                
+        if past_key_values is None or target_length > self.sliding_window:
+            # training mode or prefill mode when dealing with long prefix)
+            sliding_attend_mask = torch.arange(past_seen_tokens + sequence_length, device=device)[-target_length:] <= (
+                cache_position.reshape(-1, 1) - self.sliding_window
+            ) # bs, sequence_length, target_length
+            sliding_attend_mask = sliding_attend_mask[None, None, :, :]            
+            
+            diagonal_attend_mask = diagonal_attend_mask | sliding_attend_mask
+
+        # training
+        causal_mask = torch.full(
+            (1, 1, sequence_length, target_length), fill_value=min_dtype, dtype=dtype, device=device
+        )
+        causal_mask = causal_mask * diagonal_attend_mask
+        causal_mask = AttentionMaskConverter._unmask_unattended(causal_mask, min_dtype)
+
+        return causal_mask
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        position_embeddings: Tuple[torch.Tensor, torch.Tensor],
+        attention_mask: Optional[torch.Tensor],
+        past_key_value: Optional[JetNemotronCache] = 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]
+        hidden_shape = (*input_shape, -1, self.head_dim)
+
+        if self.sliding_window is not None and self.config._attn_implementation != "flash_attention_2":            
+            attention_mask = self._update_causal_mask_for_sliding_window(attention_mask, hidden_states, past_key_value)
+
+        query_states = self.q_proj(hidden_states).view(hidden_shape).transpose(1, 2)
+        key_states = 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:
+            # sin and cos are specific to RoPE models; cache_position needed for the static cache
+            state = past_key_value.update(
+                attn_state=(key_states, value_states), layer_idx=self.layer_idx, 
+                offset = hidden_states.shape[1],
+                cache_kwargs={"window_size": self.sliding_window})
+            key_states, value_states = state["attn_state"]
+
+        fa2_sliding_window = None
+        if self.sliding_window is not None:
+            fa2_sliding_window = self.sliding_window - 1
+
+        attention_interface: Callable = eager_attention_forward
+        if self.config._attn_implementation != "eager":
+            if self.config._attn_implementation == "sdpa":                
+                past_seen_tokens = past_key_value.get_seq_length() if past_key_value is not None else 0
+                if self.sliding_window is None:
+                    # 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.
+                    if AttentionMaskConverter._ignore_causal_mask_sdpa(
+                        attention_mask,
+                        inputs_embeds=hidden_states,
+                        past_key_values_length=past_seen_tokens,
+                        sliding_window=self.sliding_window,
+                        is_training=self.training,
+                    ):
+                        attention_mask = None
+
+            elif self.config._attn_implementation == "flash_attention_2":
+                if attention_mask is not None:
+                    assert len(attention_mask.shape) == 2, "Attention mask must be 2D"
+                    attention_mask = attention_mask[:, -key_states.shape[2]:]
+            else:
+                raise ValueError(
+                    f"Unsupported attention implementation: {self.config._attn_implementation}. "
+                    "Supported implementations are: eager, sdpa, flash_attention_2."
+                )
+                
+            if self.config._attn_implementation == "sdpa" and kwargs.get("output_attentions", False):
+                # When output attentions is True, sdpa implementation's forward method calls the eager implementation's forward
+                logger.warning_once(
+                    "`torch.nn.functional.scaled_dot_product_attention` does not support `output_attentions=True`. Falling back to "
+                    'eager attention. This warning can be removed using the argument `attn_implementation="eager"` when loading the model.'
+                )
+            else:
+                attention_interface = ALL_ATTENTION_FUNCTIONS[self.config._attn_implementation]
+
+        attn_output, attn_weights = attention_interface(
+            self,
+            query_states,
+            key_states,
+            value_states,
+            attention_mask,
+            dropout=0.0 if not self.training else self.attention_dropout,
+            scaling=self.scaling,
+            sliding_window=fa2_sliding_window,  # main diff with Llama
+            **kwargs,
+        )
+
+        if self.sliding_window is not None and past_key_value is not None:
+            past_key_value.trim_attn_state(self.layer_idx, self.sliding_window)
+
+        attn_output = attn_output.reshape(*input_shape, -1).contiguous()
+        attn_output = self.o_proj(attn_output)
+        
+        return attn_output, attn_weights
+
+
+class JetNemotronRMSNorm(nn.Module):
+    def __init__(self, hidden_size, eps=1e-6):
+        """
+        JetNemotronRMSNorm is equivalent to T5LayerNorm
+        """
+        super().__init__()
+        self.weight = nn.Parameter(torch.ones(hidden_size))
+        self.variance_epsilon = eps
+
+    def forward(self, hidden_states):
+        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}"
+
+
+EFFICIENT_ATTENTION_CLASSES = {
+    "jet": JetBlock,
+}
+
+
+class JetNemotronDecoderLayer(nn.Module):
+    def __init__(self, config: JetNemotronConfig, layer_idx: int):
+        super().__init__()
+        self.hidden_size = config.hidden_size
+        
+        if config.layer_types[layer_idx] == "attn":
+            self.self_attn = JetNemotronAttention(config, layer_idx)
+        elif config.layer_types[layer_idx] == "swa":
+            assert config.efficient_attention_config is not None, "Efficient attention config must be provided in JetNemotronConfig."
+            assert "swa" in config.efficient_attention_config, (
+                "Sliding Window Attention is enabled but no `swa` configuration found in `efficient_attention_config`."
+            )
+            self.self_attn = JetNemotronAttention(config, layer_idx, sliding_window=config.efficient_attention_config["swa"]["window_size"])
+        else:
+            assert config.layer_types[layer_idx] in EFFICIENT_ATTENTION_CLASSES, (
+                f"Layer type {config.layer_types[layer_idx]} not supported. Supported types are: "
+                f"{['attn', 'swa'] + list(EFFICIENT_ATTENTION_CLASSES.keys())}"
+            )
+            self.self_attn = EFFICIENT_ATTENTION_CLASSES[config.layer_types[layer_idx]](config, config.layer_types[layer_idx], layer_idx)
+
+        self.mlp = JetNemotronMLP(config)
+        self.input_layernorm = JetNemotronRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+        self.post_attention_layernorm = JetNemotronRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Tuple[torch.Tensor]] = None,
+        output_attentions: Optional[bool] = False,
+        use_cache: Optional[bool] = False,
+        cache_position: Optional[torch.LongTensor] = None,
+        position_embeddings: Optional[Tuple[torch.Tensor, torch.Tensor]] = None,  # will become mandatory in v4.46
+        **kwargs,
+    ) -> 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,
+            cache_position=cache_position,
+            position_embeddings=position_embeddings,
+        )
+                
+        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
+
+
+class JetNemotronRotaryEmbedding(nn.Module):
+    def __init__(self, config: JetNemotronConfig, 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()
+    @dynamic_rope_update  # power user: used with advanced RoPE types (e.g. dynamic rope)
+    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)
+
+
+JET_START_DOCSTRING = r"""
+    This model inherits from [`PreTrainedModel`]. Check the superclass documentation for the generic methods the
+    library implements for all its model (such as downloading or saving, resizing the input embeddings, pruning heads
+    etc.)
+
+    This model is also a PyTorch [torch.nn.Module](https://pytorch.org/docs/stable/nn.html#torch.nn.Module) subclass.
+    Use it as a regular PyTorch Module and refer to the PyTorch documentation for all matter related to general usage
+    and behavior.
+
+    Parameters:
+        config ([`JetNemotronConfig`]):
+            Model configuration class with all the parameters of the model. Initializing with a config file does not
+            load the weights associated with the model, only the configuration. Check out the
+            [`~PreTrainedModel.from_pretrained`] method to load the model weights.
+"""
+
+
+@add_start_docstrings(
+    "The bare Jet-Nemotron Model outputting raw hidden-states without any specific head on top.",
+    JET_START_DOCSTRING,
+)
+class JetNemotronPreTrainedModel(PreTrainedModel):
+    config_class = JetNemotronConfig
+    base_model_prefix = "model"
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["JetNemotronDecoderLayer"]
+    _skip_keys_device_placement = "past_key_values"
+    _supports_flash_attn_2 = True
+    _supports_sdpa = True
+    _supports_flex_attn = False
+    _supports_cache_class = True
+    _supports_quantized_cache = False
+    _supports_static_cache = False
+    _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_()
+
+
+JET_INPUTS_DOCSTRING = r"""
+    Args:
+        input_ids (`torch.LongTensor` of shape `(batch_size, sequence_length)`):
+            Indices of input sequence tokens in the vocabulary. Padding will be ignored by default should you provide
+            it.
+
+            Indices can be obtained using [`AutoTokenizer`]. See [`PreTrainedTokenizer.encode`] and
+            [`PreTrainedTokenizer.__call__`] for details.
+
+            [What are input IDs?](../glossary#input-ids)
+        attention_mask (`torch.Tensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Mask to avoid performing attention on padding token indices. Mask values selected in `[0, 1]`:
+
+            - 1 for tokens that are **not masked**,
+            - 0 for tokens that are **masked**.
+
+            [What are attention masks?](../glossary#attention-mask)
+
+            Indices can be obtained using [`AutoTokenizer`]. See [`PreTrainedTokenizer.encode`] and
+            [`PreTrainedTokenizer.__call__`] for details.
+
+            If `past_key_values` is used, optionally only the last `input_ids` have to be input (see
+            `past_key_values`).
+
+            If you want to change padding behavior, you should read [`modeling_opt._prepare_decoder_attention_mask`]
+            and modify to your needs. See diagram 1 in [the paper](https://arxiv.org/abs/1910.13461) for more
+            information on the default strategy.
+
+            - 1 indicates the head is **not masked**,
+            - 0 indicates the head is **masked**.
+        position_ids (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Indices of positions of each input sequence tokens in the position embeddings. Selected in the range `[0,
+            config.n_positions - 1]`.
+
+            [What are position IDs?](../glossary#position-ids)
+        past_key_values (`Cache`, *optional*):
+            Pre-computed hidden-states (key and values in the self-attention blocks and in the cross-attention
+            blocks) that can be used to speed up sequential decoding. This typically consists in the `past_key_values`
+            returned by the model at a previous stage of decoding, when `use_cache=True` or `config.use_cache=True`.
+
+            It is a [`~cache_utils.Cache`] instance. For more details, see our [kv cache guide](https://huggingface.co/docs/transformers/en/kv_cache).
+
+            If `past_key_values` are used, the user can optionally input only the last `input_ids` (those that don't
+            have their past key value states given to this model) of shape `(batch_size, 1)` instead of all `input_ids`
+            of shape `(batch_size, sequence_length)`.
+        inputs_embeds (`torch.FloatTensor` of shape `(batch_size, sequence_length, hidden_size)`, *optional*):
+            Optionally, instead of passing `input_ids` you can choose to directly pass an embedded representation. This
+            is useful if you want more control over how to convert `input_ids` indices into associated vectors than the
+            model's internal embedding lookup matrix.
+        use_cache (`bool`, *optional*):
+            If set to `True`, `past_key_values` key value states are returned and can be used to speed up decoding (see
+            `past_key_values`).
+        output_attentions (`bool`, *optional*):
+            Whether or not to return the attentions tensors of all attention layers. See `attentions` under returned
+            tensors for more detail.
+        output_hidden_states (`bool`, *optional*):
+            Whether or not to return the hidden states of all layers. See `hidden_states` under returned tensors for
+            more detail.
+        return_dict (`bool`, *optional*):
+            Whether or not to return a [`~utils.ModelOutput`] instead of a plain tuple.
+        cache_position (`torch.LongTensor` of shape `(sequence_length)`, *optional*):
+            Indices depicting the position of the input sequence tokens in the sequence. Contrarily to `position_ids`,
+            this tensor is not affected by padding. It is used to update the cache in the correct position and to infer
+            the complete sequence length.
+"""
+
+
+@add_start_docstrings(
+    "The bare Jet Nemotron Model outputting raw hidden-states without any specific head on top.",
+    JET_START_DOCSTRING,
+)
+class JetNemotronModel(JetNemotronPreTrainedModel):
+    """
+    Transformer decoder consisting of *config.num_hidden_layers* layers. Each layer is a [`JetNemotronDecoderLayer`]
+
+    Args:
+        config: JetNemotronConfig
+    """
+
+    def __init__(self, config: JetNemotronConfig):
+        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(
+            [JetNemotronDecoderLayer(config, layer_idx) for layer_idx in range(config.num_hidden_layers)]
+        )
+        self._attn_implementation = config._attn_implementation
+        assert self._attn_implementation in ["sdpa", "flash_attention_2"]
+        self.norm = JetNemotronRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+        self.rotary_emb = JetNemotronRotaryEmbedding(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
+    @add_start_docstrings_to_model_forward(JET_INPUTS_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[JetNemotronCache] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        use_cache: 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
+
+        if use_cache and past_key_values is None:
+            past_key_values = JetNemotronCache()
+
+        if inputs_embeds is None:
+            inputs_embeds = self.embed_tokens(input_ids)
+
+        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,)
+
+            if self.gradient_checkpointing and self.training:
+                layer_outputs = self._gradient_checkpointing_func(
+                    partial(decoder_layer.__call__, **flash_attn_kwargs),
+                    hidden_states,
+                    causal_mask,
+                    position_ids,
+                    past_key_values,
+                    output_attentions,
+                    use_cache,
+                    cache_position,
+                    position_embeddings,
+                )
+            else:
+                layer_outputs = decoder_layer(
+                    hidden_states,
+                    attention_mask=causal_mask,
+                    position_ids=position_ids,
+                    past_key_value=past_key_values,
+                    output_attentions=output_attentions,
+                    use_cache=use_cache,
+                    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: torch.Tensor,
+        input_tensor: torch.Tensor,
+        cache_position: torch.Tensor,
+        past_key_values: JetNemotronCache,
+        output_attentions: bool,
+    ):
+        if self.config._attn_implementation == "flash_attention_2":
+            if attention_mask is not None and past_key_values is not None:
+                is_empty = attention_mask.sum(dim=-1).long() == 0
+                last_is_1 = (attention_mask[:, -1].long() == 1) | is_empty
+                is_padding_right = last_is_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 Jet-Nemotron. 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":
+            raise NotImplementedError(
+                "Flex attention is not supported yet. Please use `flash_attention_2`, `eager`, or `sdpa` instead."
+            )
+
+        past_seen_tokens = past_key_values.get_seq_length() if past_key_values is not None else 0
+        dtype, device = input_tensor.dtype, input_tensor.device
+        min_dtype = torch.finfo(dtype).min
+        sequence_length = input_tensor.shape[1]
+
+        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 == "cuda"
+            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: JetNemotronConfig,
+        past_key_values: JetNemotronCache,
+    ):
+        """
+        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 (`JetNemotronConfig`):
+                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
+            )
+            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): ...
+
+
+class JetNemotronForCausalLM(JetNemotronPreTrainedModel, GenerationMixin):
+    _tied_weights_keys = ["lm_head.weight"]
+    _tp_plan = {"lm_head": "colwise_rep"}
+    _pp_plan = {"lm_head": (["hidden_states"], ["logits"])}
+
+    def __init__(self, config: JetNemotronConfig):
+        super().__init__(config)
+        self.model = JetNemotronModel(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
+
+    @can_return_tuple
+    @deprecate_kwarg("num_logits_to_keep", version="4.50", new_name="logits_to_keep")
+    @add_start_docstrings_to_model_forward(JET_INPUTS_DOCSTRING)
+    @replace_return_docstrings(output_type=CausalLMOutputWithPast, config_class=_CONFIG_FOR_DOC)
+    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[JetNemotronCache] = 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]`.
+
+            logits_to_keep (`int` or `torch.Tensor`, *optional*):
+                If an `int`, compute logits for the last `logits_to_keep` tokens. If `0`, calculate logits for all
+                `input_ids` (special case). Only last token logits are needed for generation, and calculating them only for that
+                token can save memory, which becomes pretty significant for long sequences or large vocabulary size.
+                If a `torch.Tensor`, must be 1D corresponding to the indices to keep in the sequence length dimension.
+                This is useful when using packed tensor format (single dimension for batch and sequence length).
+
+        Returns:
+
+        Example:
+
+        ```python
+        >>> from transformers import AutoTokenizer, AutoModelForCausalLM
+
+        >>> model = AutoModelForCausalLM.from_pretrained("jet-ai/Jet-Nemotron-2B")
+        >>> tokenizer = AutoTokenizer.from_pretrained("jet-ai/Jet-Nemotron-2B")
+
+        >>> 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
+        )
+
+        # 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
+        logits = self.lm_head(hidden_states[:, slice_indices, :])
+
+        loss = None
+        if labels is not None:
+            loss = self.loss_function(logits=logits, labels=labels, vocab_size=self.config.vocab_size, **kwargs)
+
+        return CausalLMOutputWithPast(
+            loss=loss,
+            logits=logits,
+            past_key_values=outputs.past_key_values,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )
+
+    def _prepare_cache_for_generation(
+        self,
+        generation_config: GenerationConfig,
+        model_kwargs: dict,
+        assistant_model: "PreTrainedModel",
+        batch_size: int,
+        max_cache_length: int,
+        device: torch.device,
+    ) -> bool:
+        assert not generation_config.return_legacy_cache, "Legacy cache is not supported for generation."
+        if generation_config.use_cache is False:
+            return
+        model_kwargs["past_key_values"] = JetNemotronCache()
+    
+    def _beam_search(self, *args, **kwargs):
+        raise NotImplementedError("Beam search is not supported for Jet-Nemotron models.")
+
+    def _contrastive_search(self, *args, **kwargs):
+        raise NotImplementedError("Contrastive search is not supported for Jet-Nemotron models.")
+    
+    def _group_beam_search(self, *args, **kwargs):
+        raise NotImplementedError("Group beam search is not supported for Jet-Nemotron models.")
+    
+    def _constrained_beam_search(self, *args, **kwargs):
+        raise NotImplementedError("Constrained beam search is not supported for Jet-Nemotron models.")

special_tokens_map.json

@@ -0,0 +1,31 @@
+{
+  "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|>"
+  ],
+  "eos_token": {
+    "content": "<|endoftext|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  },
+  "pad_token": {
+    "content": "<|endoftext|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  }
+}

tokenizer.json

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

tokenizer_config.json

@@ -0,0 +1,207 @@
+{
+  "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
+    }
+  },
+  "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|>"
+  ],
+  "bos_token": null,
+  "clean_up_tokenization_spaces": false,
+  "eos_token": "<|endoftext|>",
+  "errors": "replace",
+  "extra_special_tokens": {},
+  "model_max_length": 131072,
+  "pad_token": "<|endoftext|>",
+  "split_special_tokens": false,
+  "tokenizer_class": "Qwen2Tokenizer",
+  "unk_token": null
+}