Upload 3 files

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)

modeling_rwkv_hybrid.py

@@ -0,0 +1,632 @@
+from typing import Callable, List, Optional, Tuple, Union
+
+import torch
+import torch.nn as nn
+from transformers.cache_utils import Cache
+
+from transformers.activations import ACT2FN
+from transformers.cache_utils import Cache, StaticCache
+from .hybrid_cache import HybridCache
+from transformers.generation import GenerationMixin
+from transformers.modeling_attn_mask_utils import AttentionMaskConverter
+from transformers.modeling_flash_attention_utils import FlashAttentionKwargs
+from transformers.modeling_utils import ALL_ATTENTION_FUNCTIONS, PreTrainedModel
+
+from transformers.modeling_outputs import (
+    BaseModelOutputWithPast,
+    CausalLMOutputWithPast,
+)
+from transformers.processing_utils import Unpack
+from transformers.utils import (
+    LossKwargs,
+    add_start_docstrings,
+    add_start_docstrings_to_model_forward,
+    logging,
+)
+
+import threading
+from .wkv import Rwkv7Attention, Rwkv6Attention
+from .configuration_rwkv_hybrid import RwkvHybridConfig
+
+from transformers.models.qwen2.modeling_qwen2 import (Qwen2MLP, 
+                                    Qwen2RMSNorm, 
+                                    Qwen2RotaryEmbedding, 
+                                    Qwen2Attention)
+
+logger = logging.get_logger(__name__)
+
+_CONFIG_FOR_DOC = "RwkvHybridConfig"
+
+class RwkvHybridDecoderLayer(nn.Module):
+    def __init__(self, config: RwkvHybridConfig, layer_idx: int, update_v_first, get_v_first):
+        super().__init__()
+        self.hidden_size = config.hidden_size
+
+        self.is_rwkv = True if layer_idx in config.wkv_layers else False
+        if self.is_rwkv:
+            if config.wkv_version == 7:
+                self.self_attn = Rwkv7Attention(args=config, layer_id=layer_idx,
+                                                update_v_first=update_v_first,
+                                                get_v_first=get_v_first)
+            elif config.wkv_version == 6:
+                self.self_attn = Rwkv6Attention(args=config, layer_id=layer_idx,
+                                                update_v_first=update_v_first,
+                                                get_v_first=get_v_first)
+            else:
+                raise NotImplementedError
+        elif not self.is_rwkv:
+            self.self_attn = Qwen2Attention(config=config, layer_idx=layer_idx)
+        else:
+            self.self_attn = None
+            raise NotImplementedError
+
+        self.mlp = Qwen2MLP(config)
+        self.input_layernorm = Qwen2RMSNorm(
+            config.hidden_size, eps=config.rms_norm_eps)
+        self.post_attention_layernorm = Qwen2RMSNorm(
+            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[Cache] = 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,  # necessary, but kept here for BC
+        **kwargs,
+    ) -> Tuple[torch.FloatTensor, Optional[Tuple[torch.FloatTensor, torch.FloatTensor]]]:
+        residual = hidden_states
+
+        hidden_states = self.input_layernorm(hidden_states)
+
+        # RWKV 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
+
+RWKV_HYBRID_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 ([`RwkvHybridConfig`]):
+            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 RWKV Hybrid Model outputting raw hidden-states without any specific head on top.",
+    RWKV_HYBRID_START_DOCSTRING,
+)
+class RwkvHybridPreTrainedModel(PreTrainedModel):
+    config_class = RwkvHybridConfig
+    base_model_prefix = "rwkv_hybrid"
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["RwkvHybridDecoderLayer"]
+    _skip_keys_device_placement = ["past_key_values"]
+
+    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_()
+
+RWKV_HYBRID_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` or `tuple(tuple(torch.FloatTensor))`, *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`.
+
+            Two formats are allowed:
+            - a [`~cache_utils.Cache`] instance, see our
+            [kv cache guide](https://huggingface.co/docs/transformers/en/kv_cache);
+            - Tuple of `tuple(torch.FloatTensor)` of length `config.n_layers`, with each tuple having 2 tensors of
+            shape `(batch_size, num_heads, sequence_length, embed_size_per_head)`). This is also known as the legacy
+            cache format.
+
+            The model will output the same cache format that is fed as input. If no `past_key_values` are passed, the
+            legacy cache format will be returned.
+
+            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 RWKV Hybrid Model outputting raw hidden-states without any specific head on top.",
+    RWKV_HYBRID_START_DOCSTRING,
+)
+class RwkvHybridModel(RwkvHybridPreTrainedModel):
+    """
+    RWKV and Transformer hybrid decoder consisting of *config.num_hidden_layers* layers. Each layer is a [`RwkvHybridDecoderLayer`]
+
+    Args:
+        config: RwkvHybridConfig
+    """
+
+    def __init__(self, config: RwkvHybridConfig):
+        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.thread_local = threading.local()
+        self.thread_local.v_first = None
+        self.layers = nn.ModuleList(
+            [RwkvHybridDecoderLayer(config, layer_idx, self.update_v_first, self.get_v_first) for layer_idx in range(config.num_hidden_layers)]
+        )
+        self.norm = Qwen2RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+        self.rotary_emb = Qwen2RotaryEmbedding(config=config)
+        self.gradient_checkpointing = False
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    def post_init(self):
+        """
+        A method executed at the end of each Transformer model initialization, to execute code that needs the model's
+        modules properly initialized (such as weight initialization).
+        """
+        self.init_weights()
+        self._backward_compatibility_gradient_checkpointing()
+        # If current model is a base model, attach `base_model_tp_plan` from config
+        if self.base_model is self:
+            self._tp_plan = self.config.base_model_tp_plan
+        from transformers.modeling_utils import _init_weights
+        if _init_weights:
+            for layer in self.layers:
+                layer.self_attn.time_mixer.post_init()
+
+    def update_v_first(self, new_v_first):
+        """Callback function to update v_first in HybridModel."""
+        self.thread_local.v_first = new_v_first
+
+    def get_v_first(self):
+        return self.thread_local.v_first
+
+    def get_input_embeddings(self):
+        return self.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.embed_tokens = value
+
+    @add_start_docstrings_to_model_forward(RWKV_HYBRID_INPUTS_DOCSTRING)
+    def forward(
+        self,
+        input_ids: torch.LongTensor = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Cache] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        use_cache: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        return_dict: Optional[bool] = None,
+        cache_position: Optional[torch.LongTensor] = None,
+        **flash_attn_kwargs: Unpack[FlashAttentionKwargs],
+    ) -> Union[Tuple, 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
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        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 inputs_embeds is None:
+            inputs_embeds = self.embed_tokens(input_ids)
+
+        if use_cache and past_key_values is None:
+            past_key_values = HybridCache()
+
+        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(
+                    decoder_layer.__call__,
+                    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,)
+
+        output = 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,
+        )
+        return output if return_dict else output.to_tuple()
+
+    def _update_causal_mask(
+        self,
+        attention_mask: torch.Tensor,
+        input_tensor: torch.Tensor,
+        cache_position: torch.Tensor,
+        past_key_values: Cache,
+        output_attentions: bool,
+    ):
+        if self.config._attn_implementation == "flash_attention_2":
+            if attention_mask is not None and (attention_mask == 0.0).any():
+                return attention_mask
+            return 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.
+        past_seen_tokens = past_key_values.get_seq_length() if past_key_values is not None else 0
+        using_static_cache = isinstance(past_key_values, StaticCache)
+
+        # When output attentions is True, sdpa implementation's forward method calls the eager implementation's forward
+        if self.config._attn_implementation == "sdpa" and not using_static_cache and not output_attentions:
+            if AttentionMaskConverter._ignore_causal_mask_sdpa(
+                attention_mask,
+                inputs_embeds=input_tensor,
+                past_key_values_length=past_seen_tokens,
+                is_training=self.training,
+            ):
+                return None
+
+        dtype, device = input_tensor.dtype, input_tensor.device
+        sequence_length = input_tensor.shape[1]
+        if using_static_cache:
+            target_length = past_key_values.get_max_cache_shape()
+        else:
+            target_length = (
+                attention_mask.shape[-1]
+                if isinstance(attention_mask, torch.Tensor)
+                else past_seen_tokens + sequence_length + 1
+            )
+
+        # In case the provided `attention` mask is 2D, we generate a causal mask here (4D).
+        causal_mask = self._prepare_4d_causal_attention_mask_with_cache_position(
+            attention_mask,
+            sequence_length=sequence_length,
+            target_length=target_length,
+            dtype=dtype,
+            device=device,
+            cache_position=cache_position,
+            batch_size=input_tensor.shape[0],
+        )
+
+        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
+            min_dtype = torch.finfo(dtype).min
+            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,
+        device: torch.device,
+        cache_position: torch.Tensor,
+        batch_size: int,
+        **kwargs,
+    ):
+        """
+        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.
+            device (`torch.device`):
+                The device to plcae the 4D attention mask on.
+            cache_position (`torch.Tensor`):
+                Indices depicting the position of the input sequence tokens in the sequence.
+            batch_size (`torch.Tensor`):
+                Batch size.
+        """
+        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=device
+            )
+            if sequence_length != 1:
+                causal_mask = torch.triu(causal_mask, diagonal=1)
+            causal_mask *= torch.arange(target_length, device=device) > cache_position.reshape(-1, 1)
+            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
+                mask_length = attention_mask.shape[-1]
+                padding_mask = causal_mask[:, :, :, :mask_length] + attention_mask[:, None, None, :]
+                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 RwkvHybridForCausalLM(RwkvHybridPreTrainedModel, GenerationMixin):
+    _tied_weights_keys = ["lm_head.weight"]
+    _tp_plan = {"lm_head": "colwise_rep"}
+
+    def __init__(self, config):
+        super().__init__(config)
+        self.model = RwkvHybridModel(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
+
+    # @add_start_docstrings_to_model_forward(QWEN2_INPUTS_DOCSTRING)
+    # @replace_return_docstrings(output_type=CausalLMOutputWithPast, config_class=_CONFIG_FOR_DOC)
+    def forward(
+        self,
+        input_ids: torch.LongTensor = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Union[Cache, List[torch.FloatTensor]]] = 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,
+        return_dict: Optional[bool] = None,
+        cache_position: Optional[torch.LongTensor] = None,
+        num_logits_to_keep: int = 0,
+        **kwargs: Unpack[KwargsForCausalLM],
+    ) -> Union[Tuple, CausalLMOutputWithPast]:
+        r"""
+        Args:
+            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]`.
+
+            num_logits_to_keep (`int`, *optional*):
+                Calculate logits for the last `num_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.
+
+        Returns:
+
+        Example:
+
+        ```python
+        >>> from transformers import AutoTokenizer, RwkvHybridForCausalLM
+
+        >>> model = Qwen2ForCausalLM.from_pretrained("RWKV-Red-Team/ARWKV-7B-Preview-0.1")
+        >>> tokenizer = AutoTokenizer.from_pretrained("RWKV-Red-Team/ARWKV-7B-Preview-0.1")
+
+        >>> 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
+        )
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        # decoder outputs consists of (dec_features, layer_state, dec_hidden, dec_attn)
+        outputs = 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,
+            return_dict=return_dict,
+            cache_position=cache_position,
+            **kwargs,
+        )
+
+        hidden_states = outputs[0]
+        # Only compute necessary logits, and do not upcast them to float if we are not computing the loss
+        logits = self.lm_head(hidden_states[:, -num_logits_to_keep:, :])
+
+        loss = None
+        if labels is not None:
+            loss = self.loss_function(logits=logits, labels=labels, vocab_size=self.config.vocab_size, **kwargs)
+
+        if not return_dict:
+            output = (logits,) + outputs[1:]
+            return (loss,) + output if loss is not None else output
+
+        return CausalLMOutputWithPast(
+            loss=loss,
+            logits=logits,
+            past_key_values=outputs.past_key_values,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )
+

wkv.py

@@ -0,0 +1,565 @@
+import torch
+from einops import rearrange
+
+from .hybrid_cache import TimeMixState, BlockState
+import math
+import torch.nn as nn
+from torch.nn import functional as F
+from .configuration_rwkv_hybrid import RwkvHybridConfig
+from typing import TYPE_CHECKING, Optional
+from transformers.cache_utils import Cache
+
+try:
+    import triton
+    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
+
+
+class Rwkv_Tmix_x070(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id, update_v_first, get_v_first):
+        super().__init__()
+        self.args = args
+        self.layer_id = layer_id
+        self.hidden_size = args.hidden_size
+
+        self.update_v_first = update_v_first
+        self.get_v_first = get_v_first
+
+        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,
+                         head_first
+                         ):
+
+        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 self.training 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=head_first,
+            )
+            x = rearrange(o, "b l h d -> b l (h d)")
+        return x, state
+
+    def forward(
+            self,
+            hidden_states,
+            last_state: TimeMixState,
+            sequence_mask: Optional[torch.Tensor] = None,
+            use_cache: Optional[bool] = False,
+            cu_seqlens: Optional[torch.Tensor] = None,
+            **kwargs
+    ):
+        if sequence_mask is not None:
+            hidden_states = hidden_states.mul(
+                sequence_mask[:, -hidden_states.shape[-2]:, None])
+
+        shift_state = last_state.shift_state
+        B, T, C = hidden_states.size()
+
+        if shift_state is not None:
+            xx = torch.concat((shift_state.unsqueeze(
+                1), hidden_states[:, :-1]), dim=1) - hidden_states
+        else:
+            xx = self.time_shift(hidden_states) - 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:
+            self.update_v_first(v)
+        else:
+            # Original implementation
+            v = v + (self.get_v_first().to(v.device) - v) * torch.sigmoid(
+                self.v0 + (xv @ self.v1) @ self.v2
+            )  # add value residual
+
+        a = torch.sigmoid(
+            self.a0 + (xa @ self.a1) @ self.a2
+        )  # a is "in-context learning rate"
+        if self.args.wkv_has_gate:
+            g = torch.sigmoid(xg @ self.g1) @ self.g2
+        kk = k * self.k_k
+        kk = F.normalize(kk.view(B, T, self.n_head, -1), dim=-1, p=2.0).view(B, T, C)
+        k = k * (1 + (a - 1) * 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,
+            head_first=False
+        )
+        if self.args.wkv_has_group_norm:
+            hidden_states = self.ln_x(
+                hidden_states.view(B * T, C)).view(B, T, C)
+        hidden_states = hidden_states + (
+            (r.view(B, T, self.n_head, -1) * k.view(B, T, self.n_head, -1) * self.r_k).sum(
+                dim=-1, keepdim=True
+            )
+            * 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, TimeMixState(lx, wkv_state)
+
+
+class Rwkv7Attention(nn.Module):
+    def __init__(self, args: RwkvHybridConfig, layer_id, update_v_first, get_v_first):
+        super().__init__()
+        self.args = args
+        self.layer_idx = layer_id
+        self.time_mixer = Rwkv_Tmix_x070(
+            args, layer_id, update_v_first, get_v_first)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        sequence_mask: Optional[torch.Tensor] = None,
+        past_key_value: Optional[Cache] = None,
+        use_cache: Optional[bool] = False,
+        output_attentions: Optional[bool] = False,
+        **kwargs
+    ):
+        if sequence_mask is not None:
+            assert len(sequence_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, token_length, _ = hidden_states.shape
+
+        if past_key_value is not None 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 = self.time_mixer(hidden_states=hidden_states,
+                                              last_state=last_state.time_mix_state,
+                                              sequence_mask=sequence_mask,
+                                              use_cache=use_cache,
+                                              **kwargs)
+        last_state.time_mix_state = states
+
+        if past_key_value is not None:
+            past_key_value.update(token_length, last_state, self.layer_idx)
+
+        return attn_output, None
+
+    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,
+        )
+        token_shift = torch.zeros(
+            (batch_size, self.args.hidden_size), device=device, dtype=dtype
+        )
+        return BlockState(TimeMixState(token_shift, 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
+        H = self.n_head
+        N = self.head_size
+
+        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_state = nn.Parameter(torch.zeros(self.n_head, self.head_size, 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
+
+    def forward(self, x, last_state: TimeMixState):
+        shift_state = last_state.shift_state
+        B, T, C = x.size()
+        H = self.n_head
+        if shift_state is not None:
+            xx = torch.concat((shift_state.unsqueeze(1), x[:, :-1]), dim=1) - x
+        else:
+            xx = self.time_shift(x) - x
+        lx = x[:, -1]
+
+        xxx = x + 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 = x + xx * (self.time_maa_w + mw)
+        xk = x + xx * (self.time_maa_k + mk)
+        xv = x + xx * (self.time_maa_v + mv)
+        xr = x + xx * (self.time_maa_r + mr)
+        xg = x + 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
+        x, 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:
+            x = self.ln_x(x.view(B * T, C)).view(B, T, C)
+        x = self.output(x * g)
+        return x, TimeMixState(lx, wkv_state)
+
+    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, past_key_value, **kwargs):
+        attn_output = hidden_states
+        B, T, C = attn_output.size()
+        if past_key_value is not None:
+            if len(past_key_value) <= self.layer_idx:
+                last_state = None
+            else:
+                last_state = past_key_value[self.layer_idx][0]
+        if last_state is None:
+            wkv_states = torch.zeros(
+                (B, self.args.num_wkv_heads,
+                 self.args.head_size, self.args.head_size),
+                device=attn_output.device,
+                dtype=torch.float32,
+            )
+            token_shift = torch.zeros(
+                (B, C), device=attn_output.device, dtype=attn_output.dtype
+            )
+            time_state = TimeMixState(token_shift, wkv_states)
+            channel_state = None
+            last_state = BlockState(time_state, channel_state)
+        attn_output, states = self.time_mixer(
+            attn_output, last_state.time_mix_state)
+        last_state.time_mix_state = states
+
+        if past_key_value is not None:
+            past_key_value.update(T, last_state, self.layer_idx)
+        return attn_output, None