Upload 3 files

configuration_rwkv_hybrid.py

@@ -0,0 +1,252 @@
+# coding=utf-8
+# Copyright 2025 RWKV team. All rights reserved.
+# Copyright 2024 The Qwen team, Alibaba Group and the HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""RwkvHybrid model configuration"""
+
+from transformers.configuration_utils import PretrainedConfig
+from transformers.modeling_rope_utils import rope_config_validation
+from transformers.utils import logging
+from typing import Optional, Union, List
+
+
+logger = logging.get_logger(__name__)
+
+
+class RwkvHybridConfig(PretrainedConfig):
+    r"""
+    This is the configuration class to store the configuration of a [`RwkvHybridModel`]. It is used to instantiate a
+    RwkvHybrid model according to the specified arguments, defining the model architecture. Instantiating a configuration
+    with the defaults will yield a similar configuration to that of
+    RwkvHybrid-7B-beta.
+
+    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 RwkvHybrid model. Defines the number of different tokens that can be represented by the
+            `inputs_ids` passed when calling [`RwkvHybridModel`]
+        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
+        use_sliding_window (`bool`, *optional*, defaults to `False`):
+            Whether to use sliding window attention.
+        sliding_window (`int`, *optional*, defaults to 4096):
+            Sliding window attention (SWA) window size. If not specified, will default to `4096`.
+        max_window_layers (`int`, *optional*, defaults to 28):
+            The number of layers that use SWA (Sliding Window Attention). The bottom layers use SWA while the top use full attention.
+        attention_dropout (`float`, *optional*, defaults to 0.0):
+            The dropout ratio for the attention probabilities.
+        head_size (`int`, *optional*, defaults to 64):
+            Dimensionality of each RWKV attention head. Defines the hidden dimension size for RWKV attention mechanisms.
+        head_size_divisor (`int`, *optional*, defaults to 8):
+            Constraint for head_size initialization, typically set to the square root of head_size. Ensures divisibility
+            between hidden_size and head_size.
+        wkv_version (`int`, *optional*, defaults to 7):
+            Version of RWKV attention implementation. Currently supports:
+            - 6: Original implementation requiring `wkv_has_gate=True` and `wkv_use_vfirst=False`
+            - 7: Improved version requiring `wkv_use_vfirst=True`
+        wkv_has_gate (`bool`, *optional*, defaults to False):
+            Whether to include gating mechanism in RWKV attention. Required for version 6.
+        wkv_has_group_norm (`bool`, *optional*, defaults to True):
+            Whether to apply group normalization in RWKV attention layers.
+        wkv_use_vfirst (`bool`, *optional*, defaults to True):
+            Whether to prioritize value projection in RWKV attention computation. Required for version 7.
+        wkv_layers (`Union[str, List[int]]`, *optional*, defaults to None):
+            Specifies which layers use RWKV attention:
+            - `"full"` or `None`: All layers use RWKV
+            - List of integers: Only specified layers (e.g., `[0,1,2]`) use RWKV attention
+
+    ```python
+    >>> from transformers import RwkvHybridModel, RwkvHybridConfig
+
+    >>> # Initializing a RwkvHybrid style configuration
+    >>> configuration = RwkvHybridConfig()
+
+    >>> # Initializing a model from the RwkvHybrid-7B style configuration
+    >>> model = RwkvHybridModel(configuration)
+
+    >>> # Accessing the model configuration
+    >>> configuration = model.config
+    ```"""
+
+    model_type = "rwkv_hybrid"
+    keys_to_ignore_at_inference = ["past_key_values"]
+
+    # Default tensor parallel plan for base model `RwkvHybrid`
+    base_model_tp_plan = {
+        "layers.*.self_attn.q_proj": "colwise",
+        "layers.*.self_attn.k_proj": "colwise",
+        "layers.*.self_attn.v_proj": "colwise",
+        "layers.*.self_attn.o_proj": "rowwise",
+        "layers.*.mlp.gate_proj": "colwise",
+        "layers.*.mlp.up_proj": "colwise",
+        "layers.*.mlp.down_proj": "rowwise",
+    }
+
+    def __init__(
+        self,
+        vocab_size: int = 151936,
+        hidden_size: int = 4096,
+        intermediate_size: int = 22016,
+        num_hidden_layers: int = 32,
+        num_attention_heads: int = 32,
+        num_key_value_heads: int = 32,
+        head_size: int = 64,
+        head_size_divisor: int = 8,
+        hidden_act: str = "silu",
+        max_position_embeddings: int = 32768,
+        initializer_range: float = 0.02,
+        rms_norm_eps: float = 1e-6,
+        use_cache: bool = True,
+        tie_word_embeddings: bool = False,
+        rope_theta: float = 10000.0,
+        rope_scaling: Optional[dict] = None,
+        use_sliding_window: bool = False,
+        sliding_window: int = 4096,
+        max_window_layers: int = 28,
+        attention_dropout: float = 0.0,
+        wkv_version: int = 7,
+        wkv_has_gate: bool = False,
+        wkv_has_group_norm: bool = True,
+        wkv_use_vfirst: bool = True,
+        wkv_layers: Optional[Union[str, List[int]]] = 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_wkv_heads = hidden_size // head_size
+        assert hidden_size % head_size == 0, "hidden_size must be divisible by head_size"
+        self.num_attention_heads = num_attention_heads
+        self.use_sliding_window = use_sliding_window
+        self.sliding_window = sliding_window if use_sliding_window else None
+        self.max_window_layers = max_window_layers
+        self.head_size = head_size
+        self.head_size_divisor = head_size_divisor
+        self.wkv_version = wkv_version
+
+        self.wkv_has_gate = wkv_has_gate
+        self.wkv_has_group_norm = wkv_has_group_norm
+        self.wkv_use_vfirst = wkv_use_vfirst
+
+        if self.wkv_version == 7:
+            assert self.wkv_use_vfirst, "wkv_use_vfirst must be True for wkv_version 7"
+        elif self.wkv_version == 6:
+            assert self.wkv_has_gate, "wkv_has_gate must be True for wkv_version 6"
+            assert not self.wkv_use_vfirst, "wkv_use_vfirst must be False for wkv_version 6"
+        else:
+            raise NotImplementedError(f"Unsupported wkv_version: {self.wkv_version}, \
+                                        wkv_version must be 6 or 7")
+
+        if wkv_layers == "full" or wkv_layers == None:
+            self.wkv_layers = list(range(num_hidden_layers))
+        elif isinstance(wkv_layers, list):
+            if all(isinstance(layer, int) for layer in wkv_layers):
+                self.wkv_layers = wkv_layers
+            else:
+                raise ValueError("All elements in wkv_layers must be integers.")
+        else:
+            raise TypeError("wkv_layers must be either 'full', None, or a list of integers.")
+
+        # 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
+        # Validate the correctness of rotary position embeddings parameters
+        # BC: if there is a 'type' field, move it to 'rope_type'.
+        if self.rope_scaling is not None and "type" in self.rope_scaling:
+            self.rope_scaling["rope_type"] = self.rope_scaling["type"]
+        rope_config_validation(self)
+
+        super().__init__(
+            tie_word_embeddings=tie_word_embeddings,
+            **kwargs,
+        )

hybrid_cache.py

@@ -0,0 +1,154 @@
+import torch
+from typing import Any, Dict, Optional, Union
+from transformers.cache_utils import DynamicCache
+
+
+class TimeMixState:
+    def __init__(self, shift_state: torch.Tensor, wkv_state: torch.Tensor):
+        self.shift_state = shift_state
+        self.wkv_state = wkv_state
+
+
+class ChannelMixState:
+    def __init__(self, shift_state: torch.Tensor):
+        self.shift_state = shift_state
+
+
+class BlockState:
+    def __init__(self, time_mix_state: TimeMixState,
+                 channel_mix_state: ChannelMixState):
+        self.time_mix_state = time_mix_state
+        self.channel_mix_state = channel_mix_state
+
+
+class BlockStateList:
+    def __init__(self, shift_states, wkv_states):
+        self.wkv_states = wkv_states
+        self.shift_states = shift_states
+
+    @staticmethod
+    def create(N, B, C, H, device, dtype):
+        result = BlockStateList.empty(N, B, C, H, device, dtype)
+        result.wkv_states[:] = 0
+        result.wkv_states[:] = 0
+        result.shift_states[:] = 0
+        return result
+
+    @staticmethod
+    def empty(N, B, C, H, device, dtype):
+        wkv_states = torch.empty((N, B, H, C//H, C//H),
+                                 device=device,
+                                 dtype=torch.bfloat16)
+        shift_states = torch.empty((N, 2, B, C), device=device, dtype=dtype)
+        return BlockStateList(shift_states, wkv_states)
+
+    def __getitem__(self, layer: int):
+        return BlockState(
+            TimeMixState(self.shift_states[layer, 0], self.wkv_states[layer]),
+            ChannelMixState(self.shift_states[layer, 1]))
+
+    def __setitem__(self, layer: int, state: BlockState):
+        self.shift_states[layer, 0] = state.time_mix_state.shift_state
+        self.wkv_states[layer] = state.time_mix_state.wkv_state
+        self.shift_states[layer, 1] = state.channel_mix_state.shift_state
+
+
+class HybridCache(DynamicCache):
+    def __init__(self) -> None:
+        super().__init__()
+        self.rwkv_layers = set()
+
+    def __repr__(self) -> str:
+        rwkv_layers = f"HybridCache(rwkv_layers={self.rwkv_layers})"
+        # count the number of key_cache and value_cache
+        key_cache_count = sum(len(cache) for cache in self.key_cache)
+        value_cache_count = sum(len(cache) for cache in self.value_cache)
+        count_info = rwkv_layers + \
+            f", key_cache_count={key_cache_count}, value_cache_count={value_cache_count}"
+        memories = 0
+        seq_length = self.get_seq_length()
+        for cache in self.value_cache:
+            for data in cache:
+                if not isinstance(data, torch.Tensor):
+                    memories += data.time_mix_state.wkv_state.numel()
+                else:
+                    memories += data.numel()
+        count_info += f", memories={memories / 1024/1024}MB, seq_length={seq_length}"
+        return count_info
+
+    def update(self,
+               key_states: Union[int, torch.Tensor],
+               value_states: Union[torch.Tensor, BlockState],
+               layer_idx: int,
+               cache_kwargs: Optional[Dict[str, Any]] = None):
+        if isinstance(key_states, int) and not isinstance(value_states, torch.Tensor):
+            self.rwkv_layers.add(layer_idx)
+            if layer_idx >= len(self.key_cache):
+                self.key_cache.append([])
+                self.value_cache.append([])
+
+            if len(self.key_cache[layer_idx]) == 0:
+                self.key_cache[layer_idx].append(key_states)
+                self.value_cache[layer_idx].append(value_states)
+            else:
+                self.key_cache[layer_idx][0] = self.key_cache[layer_idx][0]+key_states
+                self.value_cache[layer_idx][0] = value_states
+
+            return key_states, value_states
+
+        return super().update(key_states, value_states, layer_idx, cache_kwargs)
+
+    def get_seq_length(self, layer_idx: Optional[int] = 0):
+        if layer_idx in self.rwkv_layers:
+            return self.key_cache[layer_idx][0]
+        return super().get_seq_length(layer_idx)
+
+    def get_max_length(self):
+        return super().get_max_length()
+
+    def reorder_cache(self, beam_idx):
+        return super().reorder_cache(beam_idx)
+
+    def __getitem__(self, item):
+        if item in self.rwkv_layers:
+            return self.value_cache[item]
+        return super().__getitem__(item)
+
+    def offload_to_cpu(self):
+        for cache in self.value_cache:
+            for data in cache:
+                if isinstance(data, torch.Tensor):
+                    data.cpu()
+                else:
+                    data.time_mix_state.wkv_state.cpu()
+                    data.time_mix_state.shift_state.cpu()
+
+    def offload_to_cuda(self, device: str):
+        for cache in self.value_cache:
+            for data in cache:
+                if isinstance(data, torch.Tensor):
+                    data.cuda(device)
+                else:
+                    data.time_mix_state.wkv_state.cuda(device)
+                    data.time_mix_state.shift_state.cuda(device)
+
+    def offload_to_device(self, device_type: str, device_id: int = 0):
+        for cache in self.value_cache:
+            for data in cache:
+                if isinstance(data, torch.Tensor):
+                    method = getattr(data, device_type)
+                    if device_type == 'cpu':
+                        method()
+                    else:
+                        method(device_id)
+                else:
+                    wkv_state_method = getattr(
+                        data.time_mix_state.wkv_state, device_type)
+                    shift_state_method = getattr(
+                        data.time_mix_state.shift_state, device_type)
+                    if device_type == 'cpu':
+                        wkv_state_method()
+                        shift_state_method()
+                    else:
+                        wkv_state_method(device_id)
+                        shift_state_method(device_id)

wkv.py

@@ -0,0 +1,604 @@
+import torch
+from einops import rearrange
+
+import math
+import torch.nn as nn
+from torch.nn import functional as F
+from .configuration_rwkv_hybrid import RwkvHybridConfig
+from typing import Optional
+from .hybrid_cache import HybridCache, AttnState, BlockState
+
+try:
+    import triton  # pylint: disable=F401
+    from rwkvfla.ops.rwkv7 import (
+        fused_recurrent_rwkv7,
+        chunk_rwkv7,
+        native_recurrent_rwkv7,
+        fused_addcmul_rwkv7,
+    )  # pylint: disable=C0411
+    from rwkvfla.ops.rwkv6 import (
+        fused_recurrent_rwkv6,
+        chunk_rwkv6,
+        native_recurrent_rwkv6,
+    )
+except ImportError:
+    from rwkvfla.ops.rwkv7 import native_recurrent_rwkv7  # pylint: disable=C0411
+    from rwkvfla.ops.rwkv6 import native_recurrent_rwkv6
+    from rwkvfla.ops.rwkv7 import torch_addcmul_rwkv7
+
+    fused_recurrent_rwkv7 = native_recurrent_rwkv7
+    chunk_rwkv7 = native_recurrent_rwkv7
+    chunk_rwkv6 = native_recurrent_rwkv6
+    fused_recurrent_rwkv6 = native_recurrent_rwkv6
+    fused_addcmul_rwkv7 = torch_addcmul_rwkv7
+
+from rwkvfla.utils import check_pytorch_version
+
+if check_pytorch_version("2.6"):
+    compile_decorator = torch.compile
+    torch._dynamo.config.cache_size_limit = 512
+else:
+    def compile_decorator(func):
+        return func
+
+
+class Rwkv_Tmix_x070(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id, **kwargs):
+        super().__init__()
+        self.args = args
+        self.layer_id = layer_id
+        self.hidden_size = args.hidden_size
+
+        self.head_size = args.head_size
+        self.n_head = args.num_wkv_heads
+        assert args.hidden_size % self.n_head == 0
+        H = self.n_head
+        N = self.head_size
+
+        self.x_r = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.x_w = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.x_k = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.x_v = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.x_a = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+
+        D_DECAY_LORA = 64
+        D_AAA_LORA = 64
+        D_MV_LORA = 32
+        D_GATE_LORA = 128
+
+        self.w1 = nn.Parameter(torch.Tensor(args.hidden_size, D_DECAY_LORA))
+        self.w2 = nn.Parameter(torch.Tensor(D_DECAY_LORA, args.hidden_size))
+        self.w0 = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+
+        self.a1 = nn.Parameter(torch.Tensor(args.hidden_size, D_AAA_LORA))
+        self.a2 = nn.Parameter(torch.Tensor(D_AAA_LORA, args.hidden_size))
+        self.a0 = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+
+        self.v1 = nn.Parameter(torch.Tensor(args.hidden_size, D_MV_LORA))
+        self.v2 = nn.Parameter(torch.Tensor(D_MV_LORA, args.hidden_size))
+        self.v0 = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+
+        if self.args.wkv_has_gate:
+            self.x_g = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+            self.g1 = nn.Parameter(torch.Tensor(args.hidden_size, D_GATE_LORA))
+            self.g2 = nn.Parameter(torch.Tensor(D_GATE_LORA, args.hidden_size))
+
+        self.k_k = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.k_a = nn.Parameter(torch.Tensor(1, 1, args.hidden_size))
+        self.r_k = nn.Parameter(torch.Tensor(H, N))
+
+        self.time_shift = nn.ZeroPad2d((0, 0, 1, -1))
+        self.receptance = nn.Linear(
+            args.hidden_size, args.hidden_size, bias=False)
+        self.key = nn.Linear(args.hidden_size, args.hidden_size, bias=False)
+        self.value = nn.Linear(args.hidden_size, args.hidden_size, bias=False)
+        self.output = nn.Linear(args.hidden_size, args.hidden_size, bias=False)
+
+        if self.args.wkv_has_group_norm:
+            self.ln_x = nn.GroupNorm(
+                H, args.hidden_size, eps=(1e-5) * (args.head_size_divisor**2)
+            )
+
+    def post_init(self):
+        with torch.no_grad():
+            ratio_0_to_1 = self.layer_id / \
+                (self.args.num_hidden_layers - 1)  # 0 to 1
+            ratio_1_to_almost0 = 1.0 - (
+                self.layer_id / self.args.num_hidden_layers
+            )  # 1 to ~0
+
+            ddd = torch.ones(1, 1, self.args.hidden_size)
+            for i in range(self.args.hidden_size):
+                ddd[0, 0, i] = i / self.args.hidden_size
+
+            nn.init.constant_(
+                self.x_r, 1.0 - torch.pow(ddd, 0.2 * ratio_1_to_almost0))
+            nn.init.constant_(
+                self.x_w, 1.0 - torch.pow(ddd, 0.9 * ratio_1_to_almost0))
+            nn.init.constant_(
+                self.x_k,
+                1.0 - (torch.pow(ddd, 0.9 * ratio_1_to_almost0) +
+                       0.4 * ratio_0_to_1),
+            )
+            nn.init.constant_(
+                self.x_v,
+                1.0 - (torch.pow(ddd, 0.4 * ratio_1_to_almost0) +
+                       0.6 * ratio_0_to_1),
+            )
+            nn.init.constant_(
+                self.x_a, 1.0 - torch.pow(ddd, 0.9 * ratio_1_to_almost0))
+
+            def ortho_init(x, scale):
+                shape = x.shape
+                original_dtype = x.dtype
+                x_fp32 = x.float()
+                if len(shape) == 2:
+                    gain = math.sqrt(shape[0] / shape[1]
+                                     ) if shape[0] > shape[1] else 1
+                    nn.init.orthogonal_(x_fp32, gain=gain * scale)
+                elif len(shape) == 3:
+                    gain = math.sqrt(shape[1] / shape[2]
+                                     ) if shape[1] > shape[2] else 1
+                    for i in range(shape[0]):
+                        nn.init.orthogonal_(x_fp32[i], gain=gain * scale)
+                else:
+                    raise ValueError(
+                        "ortho_init only supports 2D or 3D tensors")
+                x.data.copy_(x_fp32.to(original_dtype))
+                return x
+
+            D_DECAY_LORA = 64
+            nn.init.zeros_(self.w1)
+            self.w2 = nn.Parameter(
+                ortho_init(torch.zeros(
+                    D_DECAY_LORA, self.args.hidden_size), 0.1)
+            )
+
+            decay_speed = torch.ones(self.args.hidden_size)
+            for n in range(self.args.hidden_size):
+                decay_speed[n] = -7 + 5 * (n / (self.args.hidden_size - 1)) ** (
+                    0.85 + 1.0 * ratio_0_to_1**0.5
+                )
+            nn.init.constant_(
+                self.w0, decay_speed.reshape(1, 1, self.args.hidden_size) + 0.5
+            )
+
+            D_AAA_LORA = 64
+            nn.init.zeros_(self.a1)
+            self.a2 = nn.Parameter(
+                ortho_init(torch.zeros(D_AAA_LORA, self.args.hidden_size), 0.1)
+            )
+            nn.init.zeros_(self.a0)
+
+            D_MV_LORA = 32
+            nn.init.zeros_(self.v1)
+            self.v2 = nn.Parameter(
+                ortho_init(torch.zeros(D_MV_LORA, self.args.hidden_size), 0.1)
+            )
+            nn.init.constant_(self.v0, 1.0)
+
+            D_GATE_LORA = 128
+            if self.args.wkv_has_gate:
+                nn.init.zeros_(self.g1)
+                self.g2 = nn.Parameter(
+                    ortho_init(torch.zeros(
+                        D_GATE_LORA, self.args.hidden_size), 0.1)
+                )
+                nn.init.constant_(
+                    self.x_g, 1.0 - torch.pow(ddd, 0.2 * ratio_1_to_almost0))
+
+            nn.init.constant_(self.k_k, 0.85)
+            nn.init.constant_(self.k_a, 1.0)
+            nn.init.zeros_(self.r_k)
+
+            nn.init.zeros_(self.receptance.weight)
+            nn.init.zeros_(self.key.weight)
+            nn.init.zeros_(self.value.weight)
+            nn.init.zeros_(self.output.weight)
+
+            if self.args.wkv_has_group_norm:
+                nn.init.ones_(self.ln_x.weight)
+                nn.init.zeros_(self.ln_x.bias)
+
+    def apply_wkv7_state(
+        self, r, k, v, w, a, b, s,
+        output_final_state,
+        cu_seqlens
+    ):
+        if r.device.type == "cpu":
+            r, w, k, v, a, b = map(lambda x: rearrange(
+                x, 'b l (h d) -> b h l d', h=self.n_head), (r, w, k, v, a, b))
+            o, state = native_recurrent_rwkv7(
+                r=r, k=k, v=v, w=w,
+                a=a, b=b,
+                scale=1.0,
+                initial_state=s.transpose(-1, -2),
+                output_final_state=True,
+                head_first=True,
+            )
+            state = state.transpose(-1, -2)
+            x = rearrange(o, "b h l d -> b l (h d)")
+        else:
+            r, w, k, v, a, b = map(lambda x: rearrange(
+                x, 'b l (h d) -> b l h d', h=self.n_head), (r, w, k, v, a, b))
+            wkv7_func = chunk_rwkv7 if r.shape[1] != 1 else fused_recurrent_rwkv7
+            o, state = wkv7_func(
+                r=r, k=k, v=v, w=w,
+                a=a, b=b,
+                scale=1.0,
+                initial_state=s,
+                output_final_state=output_final_state,
+                cu_seqlens=cu_seqlens,
+                head_first=False,
+            )
+            x = rearrange(o, "b l h d -> b l (h d)")
+        return x, state
+
+    @compile_decorator
+    def forward(
+            self,
+            hidden_states,
+            last_state: AttnState,
+            use_cache: Optional[bool] = False,
+            cu_seqlens: Optional[torch.Tensor] = None,
+            v_first: Optional[torch.Tensor] = None,
+            attention_mask: Optional[torch.Tensor] = None,
+            **kwargs
+    ):
+        shift_state = last_state.shift_state
+        B, T, C = hidden_states.size()
+
+        xx = torch.concat((shift_state.unsqueeze(
+            1), hidden_states[:, :-1]), dim=1) - hidden_states
+
+        lx = hidden_states[:, -1]
+
+        if self.args.wkv_has_gate:
+            xr, xw, xk, xv, xa, xg = fused_addcmul_rwkv7(
+                hidden_states, xx, self.x_r, self.x_w, self.x_k, self.x_v, self.x_a, self.x_g)
+        else:
+            xr, xw, xk, xv, xa, _ = fused_addcmul_rwkv7(
+                hidden_states, xx, self.x_r, self.x_w, self.x_k, self.x_v, self.x_a)
+
+        r = self.receptance(xr)
+        w = (
+            -F.softplus(-(self.w0 + torch.tanh(xw @ self.w1) @ self.w2)) - 0.5
+        )  # soft-clamp to (-inf, -0.5)
+        k = self.key(xk)
+        v = self.value(xv)
+        if self.layer_id == 0:
+            v_first = v
+        else:
+            v = torch.lerp(v, v_first, torch.sigmoid(
+                self.v0 + (xv @ self.v1) @ self.v2
+            ))  # add value residual
+
+        if attention_mask is not None:
+            v = v.mul(attention_mask[:, -v.shape[-2]:, None])
+        a = torch.sigmoid(
+            self.a0 + (xa @ self.a1) @ self.a2
+        )  # a is "in-context learning rate"
+        if self.args.wkv_has_gate:
+            g_delta = torch.sigmoid(xg @ self.g1) @ self.g2
+            g = 1.0 + g_delta
+        kk = k * self.k_k
+        kk = F.normalize(kk.view(B, T, self.n_head, -1),
+                         p=2.0, dim=-1, eps=1e-4 if kk.dtype == torch.float16 else 1e-12).view(B, T, C)
+        k = torch.lerp(k, k * a, self.k_a)
+
+        wkv_state = last_state.wkv_state
+        hidden_states, wkv_state = self.apply_wkv7_state(
+            r,
+            k,
+            v,
+            w,
+            -kk,
+            (kk * a),
+            s=wkv_state,
+            output_final_state=use_cache,
+            cu_seqlens=cu_seqlens
+        )
+        if self.args.wkv_has_group_norm:
+            hidden_states = self.ln_x(
+                hidden_states.view(B * T, C)).view(B, T, C)
+
+        # original code:
+        # weighted_sum_rk = (r.view(B, T, self.n_head, -1) * k.view(B, T, self.n_head, -1) * self.r_k).sum(
+        #         dim=-1, keepdim=True
+        #     )
+        weighted_sum_rk = torch.einsum('btij,btij,ij->btij', r.view(B, T, self.n_head, -1),
+                                       k.view(B, T, self.n_head, -1), self.r_k).sum(dim=-1, keepdim=True)
+        hidden_states = hidden_states + \
+            (weighted_sum_rk * v.view(B, T, self.n_head, -1)).view(B, T, C)
+        hidden_states = self.output(
+            hidden_states * g) if self.args.wkv_has_gate else self.output(hidden_states)
+        return hidden_states, AttnState(lx, wkv_state), v_first
+
+
+class Rwkv7Attention(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id):
+        super().__init__()
+        self.args = args
+        self.layer_idx = layer_id
+        self.time_mixer = Rwkv_Tmix_x070(args, layer_id)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.Tensor] = None,
+        past_key_value: Optional[HybridCache] = None,
+        output_attentions: Optional[bool] = False,
+        use_cache: Optional[bool] = False,
+        cache_position: Optional[torch.Tensor] = None,
+        position_embeddings: Optional[torch.Tensor] = None,
+        cu_seqlens: Optional[torch.Tensor] = None,
+        v_first: Optional[torch.Tensor] = None,
+        **kwargs
+    ):
+
+        batch_size, token_length, _ = hidden_states.shape
+
+        if use_cache and len(past_key_value) > self.layer_idx:
+            last_state = past_key_value[self.layer_idx][0]
+        else:
+            last_state = self.init_state(
+                batch_size, hidden_states.device, hidden_states.dtype
+            )
+
+        attn_output, states, v_first = self.time_mixer(hidden_states=hidden_states,
+                                                       last_state=last_state.attn_state,
+                                                       use_cache=use_cache,
+                                                       cu_seqlens=cu_seqlens,
+                                                       v_first=v_first,
+                                                       **kwargs)
+
+        if use_cache:
+            last_state.attn_state = states
+            past_key_value.update(token_length, last_state, self.layer_idx)
+
+        return attn_output, None, v_first
+
+    def init_state(self, batch_size, device, dtype) -> BlockState:
+        wkv_states = torch.zeros(
+            (
+                batch_size,
+                self.args.num_wkv_heads,
+                self.args.head_size,
+                self.args.head_size,
+            ),
+            device=device,
+            dtype=torch.float32,
+        )
+        shift_states = torch.zeros(
+            (batch_size, self.args.hidden_size), device=device, dtype=dtype
+        )
+        return BlockState(AttnState(shift_states, wkv_states), None)
+
+
+class Rwkv_Tmix_x060(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id, **kwargs):
+        super().__init__()
+        self.args = args
+        self.layer_id = layer_id
+        self.hidden_size = args.hidden_size
+
+        self.head_size = args.head_size
+        self.n_head = args.num_wkv_heads
+        assert args.hidden_size % self.n_head == 0
+
+        with torch.no_grad():
+            ratio_0_to_1 = layer_id / (args.n_layer - 1)  # 0 to 1
+            ratio_1_to_almost0 = 1.0 - (layer_id / args.n_layer)  # 1 to ~0
+            ddd = torch.ones(1, 1, args.hidden_size)
+            for i in range(args.hidden_size):
+                ddd[0, 0, i] = i / args.hidden_size
+
+            # fancy time_mix
+            self.time_maa_x = nn.Parameter(
+                1.0 - torch.pow(ddd, ratio_1_to_almost0))
+            self.time_maa_w = nn.Parameter(
+                1.0 - torch.pow(ddd, ratio_1_to_almost0))
+            self.time_maa_k = nn.Parameter(
+                1.0 - torch.pow(ddd, ratio_1_to_almost0))
+            self.time_maa_v = nn.Parameter(
+                1.0 - (torch.pow(ddd, ratio_1_to_almost0) + 0.3 * ratio_0_to_1)
+            )
+            self.time_maa_r = nn.Parameter(
+                1.0 - torch.pow(ddd, 0.5 * ratio_1_to_almost0)
+            )
+            self.time_maa_g = nn.Parameter(
+                1.0 - torch.pow(ddd, 0.5 * ratio_1_to_almost0)
+            )
+
+            D_MIX_LORA = 32  # generate TIME_MIX for w,k,v,r,g
+            if args.hidden_size == 4096:
+                D_MIX_LORA = D_MIX_LORA * 2
+            self.time_maa_w1 = nn.Parameter(
+                torch.zeros(args.hidden_size, D_MIX_LORA * 5)
+            )
+            self.time_maa_w2 = nn.Parameter(
+                torch.zeros(5, D_MIX_LORA,
+                            args.hidden_size).uniform_(-0.01, 0.01)
+            )
+
+            # fancy time_decay
+            decay_speed = torch.ones(args.head_size)
+            for n in range(args.head_size):
+                decay_speed[n] = -6 + 5 * (n / (args.head_size - 1)) ** (
+                    0.7 + 1.3 * ratio_0_to_1
+                )
+            self.time_decay = nn.Parameter(
+                decay_speed.reshape(1, 1, args.head_size))
+
+            D_DECAY_LORA = 64
+            if args.hidden_size == 4096:
+                D_DECAY_LORA = D_DECAY_LORA * 2
+            self.time_decay_w1 = nn.Parameter(
+                torch.zeros(args.hidden_size, D_DECAY_LORA)
+            )
+            self.time_decay_w2 = nn.Parameter(
+                torch.zeros(D_DECAY_LORA, args.head_size).uniform_(-0.01, 0.01)
+            )
+
+            tmp = torch.zeros(args.head_size)
+            for n in range(args.head_size):
+                zigzag = ((n + 1) % 3 - 1) * 0.1
+                tmp[n] = ratio_0_to_1 * \
+                    (1 - (n / (args.head_size - 1))) + zigzag
+
+            self.time_faaaa = nn.Parameter(
+                tmp.reshape(self.n_head, self.head_size))
+
+        self.time_shift = nn.ZeroPad2d((0, 0, 1, -1))
+        self.receptance = nn.Linear(
+            args.hidden_size, args.head_size, bias=False)
+        self.key = nn.Linear(args.hidden_size, args.head_size, bias=False)
+
+        self.value = nn.Linear(args.hidden_size, args.head_size, bias=False)
+        self.output = nn.Linear(args.head_size, args.hidden_size, bias=False)
+        self.gate = nn.Linear(args.hidden_size, args.head_size, bias=False)
+
+        if self.args.wkv_has_group_norm:
+            self.ln_x = nn.GroupNorm(
+                self.n_head, args.head_size, eps=(
+                    1e-5) * (args.head_size_divisor**2)
+            )
+
+    def post_init(self):
+        pass
+
+    @compile_decorator
+    def forward(
+        self,
+        hidden_states,
+        last_state: AttnState,
+        use_cache: Optional[bool] = False,
+        cu_seqlens: Optional[torch.Tensor] = None,
+        v_first: Optional[torch.Tensor] = None,
+        **kwargs
+    ):
+        shift_state = last_state.shift_state
+        B, T, C = hidden_states.size()
+        H = self.n_head
+
+        xx = torch.concat((shift_state.unsqueeze(
+            1), hidden_states[:, :-1]), dim=1) - hidden_states
+
+        lx = hidden_states[:, -1]
+
+        xxx = hidden_states + xx * self.time_maa_x
+        xxx = torch.tanh(xxx @ self.time_maa_w1).view(B *
+                                                      T, 5, -1).transpose(0, 1)
+        xxx = torch.bmm(xxx, self.time_maa_w2).view(5, B, T, -1)
+        mw, mk, mv, mr, mg = xxx.unbind(dim=0)
+
+        xw = hidden_states + xx * (self.time_maa_w + mw)
+        xk = hidden_states + xx * (self.time_maa_k + mk)
+        xv = hidden_states + xx * (self.time_maa_v + mv)
+        xr = hidden_states + xx * (self.time_maa_r + mr)
+        xg = hidden_states + xx * (self.time_maa_g + mg)
+
+        r = self.receptance(xr)
+        k = self.key(xk)
+        v = self.value(xv)
+        g = F.silu(self.gate(xg))
+
+        ww = torch.tanh(xw @ self.time_decay_w1) @ self.time_decay_w2
+        w = self.time_decay + ww
+
+        wkv_state = last_state.wkv_state
+        hidden_states, wkv_state = self.apply_wkv6_state(
+            B, T, C, H, r, k, v, w, u=self.time_faaaa, s=wkv_state
+        )
+        if self.args.wkv_has_group_norm:
+            hidden_states = self.ln_x(
+                hidden_states.view(B * T, C)).view(B, T, C)
+        hidden_states = self.output(hidden_states * g)
+        return hidden_states, AttnState(lx, wkv_state), None
+
+    def apply_wkv6_state(self, B, T, C, H, r, k, v, w, u, s):
+        r, w, k, v = map(lambda x: rearrange(
+            x, 'b l (h d) -> b h l d', h=self.n_head), (r, w, k, v))
+
+        if r.device.type == "cpu":
+            wkv6_func = native_recurrent_rwkv6
+        elif self.training:
+            wkv6_func = chunk_rwkv6
+        else:
+            wkv6_func = fused_recurrent_rwkv6
+
+        o, state = wkv6_func(
+            r,
+            k,
+            v,
+            -torch.exp(w),
+            u=u,
+            scale=1.0,
+            initial_state=s,
+            output_final_state=True,
+        )
+        x = rearrange(o, "b h l d -> b l (h d)")
+        return x, state
+
+
+class Rwkv6Attention(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id, **kwargs):
+        super().__init__()
+        self.args = args
+        self.layer_idx = layer_id
+        self.time_mixer = Rwkv_Tmix_x060(args, layer_id, **kwargs)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.Tensor] = None,
+        past_key_value: Optional[HybridCache] = None,
+        output_attentions: Optional[bool] = False,
+        use_cache: Optional[bool] = False,
+        cache_position: Optional[torch.Tensor] = None,
+        position_embeddings: Optional[torch.Tensor] = None,
+        cu_seqlens: Optional[torch.Tensor] = None,
+        v_first: Optional[torch.Tensor] = None,
+        **kwargs
+    ):
+        attn_output = hidden_states
+
+        batch_size, token_length, _ = hidden_states.shape
+
+        if use_cache and len(past_key_value) > self.layer_idx:
+            last_state = past_key_value[self.layer_idx][0]
+        else:
+            last_state = self.init_state(
+                batch_size, hidden_states.device, hidden_states.dtype
+            )
+        
+        attn_output, states, v_first = self.time_mixer(hidden_states=hidden_states,
+                                                       last_state=last_state.attn_state,
+                                                       use_cache=use_cache,
+                                                       cu_seqlens=cu_seqlens,
+                                                       v_first=v_first,
+                                                       **kwargs)
+
+        if use_cache:
+            last_state.attn_state = states
+            past_key_value.update(token_length, last_state, self.layer_idx)
+
+        return attn_output, None, v_first
+    
+    def init_state(self, batch_size, device, dtype) -> BlockState:
+        wkv_states = torch.zeros(
+            (
+                batch_size,
+                self.args.num_wkv_heads,
+                self.args.head_size,
+                self.args.head_size,
+            ),
+            device=device,
+            dtype=torch.float32,
+        )
+        shift_states = torch.zeros(
+            (batch_size, self.args.hidden_size), device=device, dtype=dtype
+        )
+        return BlockState(AttnState(shift_states, wkv_states), None)