modeling_rize.py

""" PyTorch Rize model."""
from contextlib import nullcontext
import os
try:
    # works with DeepSpeed ZeRO-3 (parameter partitioning)
    from deepspeed.runtime.zero.partition_parameters import (
        GatheredParameters,
        is_zero_param,
    )
    _HAS_DEEPSPEED = True
except Exception:
    _HAS_DEEPSPEED = False
    GatheredParameters = None
    def is_zero_param(p): return False

def _maybe_zero_gather(params):
    """Return a context that gathers ZeRO-sharded params if (and only if) needed."""
    if _HAS_DEEPSPEED and any(is_zero_param(p) for p in params):
        return GatheredParameters(list(params))
    return nullcontext()

import math
import warnings
from typing import List, Optional, Tuple, Union

import torch
import torch.nn.functional as F
import torch.utils.checkpoint
from torch import nn
from torch.nn import BCEWithLogitsLoss, CrossEntropyLoss, MSELoss
from cut_cross_entropy import linear_cross_entropy

from transformers.activations import ACT2FN
from transformers.cache_utils import Cache, DynamicCache
from transformers.modeling_attn_mask_utils import (
    AttentionMaskConverter,
    _prepare_4d_attention_mask,
    _prepare_4d_causal_attention_mask,
)
from transformers.modeling_outputs import (
    BaseModelOutputWithPast,
    CausalLMOutputWithPast,
    SequenceClassifierOutputWithPast,
)
from transformers.modeling_utils import PreTrainedModel
from transformers.pytorch_utils import (
    ALL_LAYERNORM_LAYERS,
    is_torch_greater_or_equal_than_1_13,
)
from transformers.utils import (
    add_start_docstrings,
    add_start_docstrings_to_model_forward,
    is_flash_attn_2_available,
    is_flash_attn_greater_or_equal_2_10,
    logging,
    replace_return_docstrings,
)
from transformers.utils.import_utils import is_torch_fx_available
from transformers.generation.utils import GenerationMixin
from .configuration_rize import RizeConfig
import torch.distributed as dist
import numpy as np

# -----------------------------------------------------------------------------
# Triton grouped GEMM (MoE training acceleration)
# -----------------------------------------------------------------------------
# This replaces the "Python loop over experts + many small GEMMs" pattern with
# grouped GEMMs that run in a small number of kernels.
#
# Requirements:
#   - CUDA device
#   - triton installed
#   - bf16 or fp16 activations/weights
#
# Notes:
#   - ZeRO-3 parameter partitioning is NOT supported by pointer-based grouped GEMM
#     (weights may be gathered/freed dynamically). We automatically fall back.
# -----------------------------------------------------------------------------
_TRITON_MOE_AVAILABLE = False
try:
    import triton
    import triton.language as tl
    _TRITON_MOE_AVAILABLE = True
except Exception:
    triton = None
    tl = None
    _TRITON_MOE_AVAILABLE = False


def _moe_triton_num_sms(device: Optional[Union[int, torch.device]] = None) -> int:
    if not torch.cuda.is_available():
        return 1
    if device is None:
        dev_idx = torch.cuda.current_device()
    elif isinstance(device, int):
        dev_idx = device
    else:
        dev_idx = device.index
        if dev_idx is None:
            dev_idx = torch.cuda.current_device()
    return int(torch.cuda.get_device_properties(dev_idx).multi_processor_count)


if _TRITON_MOE_AVAILABLE:

    # ---------------------------
    # Forward: Y = X @ W^T
    #   X: [M, K]
    #   W: [N, K]  (nn.Linear weight layout)
    #   Y: [M, N]
    # ---------------------------
    @triton.jit
    def _moe_grouped_linear_fwd_bf16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
    ):
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)
            num_m_tiles = tl.cdiv(gm, BLOCK_M)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_tiles = num_m_tiles * num_n_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_m = tile_in_g // num_n_tiles
                pid_n = tile_in_g - pid_m * num_n_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.bfloat16))
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.bfloat16))
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.bfloat16))

                offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)

                acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

                for k_start in range(0, K, BLOCK_K):
                    offs_k = k_start + tl.arange(0, BLOCK_K)

                    a_tile_ptrs = a_ptr + (offs_m[:, None] * lda + offs_k[None, :])
                    b_tile_ptrs = b_ptr + (offs_n[:, None] * ldb + offs_k[None, :])

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_n[:, None] < N) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    acc += tl.dot(a, b.T)

                c_tile_ptrs = c_ptr + (offs_m[:, None] * ldc + offs_n[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_m[:, None] < gm) & (offs_n[None, :] < N),
                )

                tile_idx += num_pids

            last_end += num_tiles

    @triton.jit
    def _moe_grouped_linear_fwd_fp16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
    ):
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)
            num_m_tiles = tl.cdiv(gm, BLOCK_M)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_tiles = num_m_tiles * num_n_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_m = tile_in_g // num_n_tiles
                pid_n = tile_in_g - pid_m * num_n_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.float16))
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.float16))
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.float16))

                offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)

                acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

                for k_start in range(0, K, BLOCK_K):
                    offs_k = k_start + tl.arange(0, BLOCK_K)

                    a_tile_ptrs = a_ptr + (offs_m[:, None] * lda + offs_k[None, :])
                    b_tile_ptrs = b_ptr + (offs_n[:, None] * ldb + offs_k[None, :])

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_n[:, None] < N) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    acc += tl.dot(a, b.T)

                c_tile_ptrs = c_ptr + (offs_m[:, None] * ldc + offs_n[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_m[:, None] < gm) & (offs_n[None, :] < N),
                )

                tile_idx += num_pids

            last_end += num_tiles

    # -----------------------------------------
    # Backward (input): dX = dY @ W
    #   dY: [M, N]
    #   W : [N, K]
    #   dX: [M, K]
    # -----------------------------------------
    @triton.jit
    def _moe_grouped_linear_bwd_x_bf16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
    ):
        # Here:
        #   A = dY   shape [M, K] where K == N_out (shared dim)
        #   B = W    shape [K, N] where N == K_in  (output cols)
        #   C = dX   shape [M, N]
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)
            num_m_tiles = tl.cdiv(gm, BLOCK_M)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_tiles = num_m_tiles * num_n_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_m = tile_in_g // num_n_tiles
                pid_n = tile_in_g - pid_m * num_n_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.bfloat16))
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.bfloat16))
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.bfloat16))

                offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)

                acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

                for k_start in range(0, K, BLOCK_K):
                    offs_k = k_start + tl.arange(0, BLOCK_K)

                    a_tile_ptrs = a_ptr + (offs_m[:, None] * lda + offs_k[None, :])
                    b_tile_ptrs = b_ptr + (offs_k[:, None] * ldb + offs_n[None, :])

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_k[:, None] < K) & (offs_n[None, :] < N),
                        other=0.0,
                    )
                    acc += tl.dot(a, b)

                c_tile_ptrs = c_ptr + (offs_m[:, None] * ldc + offs_n[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_m[:, None] < gm) & (offs_n[None, :] < N),
                )

                tile_idx += num_pids

            last_end += num_tiles

    @triton.jit
    def _moe_grouped_linear_bwd_x_fp16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
    ):
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)
            num_m_tiles = tl.cdiv(gm, BLOCK_M)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_tiles = num_m_tiles * num_n_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_m = tile_in_g // num_n_tiles
                pid_n = tile_in_g - pid_m * num_n_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.float16))
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.float16))
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.float16))

                offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)

                acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)

                for k_start in range(0, K, BLOCK_K):
                    offs_k = k_start + tl.arange(0, BLOCK_K)

                    a_tile_ptrs = a_ptr + (offs_m[:, None] * lda + offs_k[None, :])
                    b_tile_ptrs = b_ptr + (offs_k[:, None] * ldb + offs_n[None, :])

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_k[:, None] < K) & (offs_n[None, :] < N),
                        other=0.0,
                    )
                    acc += tl.dot(a, b)

                c_tile_ptrs = c_ptr + (offs_m[:, None] * ldc + offs_n[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_m[:, None] < gm) & (offs_n[None, :] < N),
                )

                tile_idx += num_pids

            last_end += num_tiles

    # -----------------------------------------
    # Backward (weight): dW = dY^T @ X
    #   dY: [M, N]
    #   X : [M, K]
    #   dW: [N, K]
    # -----------------------------------------
    @triton.jit
    def _moe_grouped_linear_bwd_w_bf16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, BLOCK_M: tl.constexpr,
    ):
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)  # shared dim (M)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_k_tiles = tl.cdiv(K, BLOCK_K)
            num_tiles = num_n_tiles * num_k_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_n = tile_in_g // num_k_tiles
                pid_k = tile_in_g - pid_n * num_k_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.bfloat16))  # dY [M, N]
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.bfloat16))  # X  [M, K]
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.bfloat16))  # dW [N, K]

                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
                offs_k = pid_k * BLOCK_K + tl.arange(0, BLOCK_K)

                acc = tl.zeros((BLOCK_N, BLOCK_K), dtype=tl.float32)

                m_start = 0
                while m_start < gm:
                    # NOTE: We cannot loop to gm directly as tl.constexpr; use mask and break via tl.multiple_of not possible.
                    # Instead, we guard with mask on m_offsets and rely on masks to zero out beyond gm.
                    offs_m = m_start + tl.arange(0, BLOCK_M)
                    # If m_start >= gm, all masks are false and dot does nothing; but we'd still waste time.
                    # Triton currently requires static range() bounds; keep this kernel simple and rely on bwd_w
                    # being relatively small vs GEMM time when gm is non-trivial.

                    a_tile_ptrs = a_ptr + (offs_m[None, :] * lda + offs_n[:, None])  # dY[m, n] -> [N, M]
                    b_tile_ptrs = b_ptr + (offs_m[:, None] * ldb + offs_k[None, :])  # X[m, k]  -> [M, K]

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_n[:, None] < N) & (offs_m[None, :] < gm),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    acc += tl.dot(a, b)

                    m_start += BLOCK_M

                c_tile_ptrs = c_ptr + (offs_n[:, None] * ldc + offs_k[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_n[:, None] < N) & (offs_k[None, :] < K),
                )

                tile_idx += num_pids

            last_end += num_tiles

    @triton.jit
    def _moe_grouped_linear_bwd_w_fp16(
        a_ptrs, b_ptrs, c_ptrs,
        m_sizes_ptr,
        lda: tl.constexpr, ldb: tl.constexpr, ldc: tl.constexpr,
        GROUP_SIZE,
        N: tl.constexpr, K: tl.constexpr,
        BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, BLOCK_M: tl.constexpr,
    ):
        pid = tl.program_id(0)
        num_pids = tl.num_programs(0)
        tile_idx = pid
        last_end = 0

        for g in range(GROUP_SIZE):
            gm = tl.load(m_sizes_ptr + g)
            num_n_tiles = tl.cdiv(N, BLOCK_N)
            num_k_tiles = tl.cdiv(K, BLOCK_K)
            num_tiles = num_n_tiles * num_k_tiles

            while (tile_idx >= last_end) & (tile_idx < last_end + num_tiles):
                tile_in_g = tile_idx - last_end
                pid_n = tile_in_g // num_k_tiles
                pid_k = tile_in_g - pid_n * num_k_tiles

                a_ptr = tl.load(a_ptrs + g).to(tl.pointer_type(tl.float16))
                b_ptr = tl.load(b_ptrs + g).to(tl.pointer_type(tl.float16))
                c_ptr = tl.load(c_ptrs + g).to(tl.pointer_type(tl.float16))

                offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
                offs_k = pid_k * BLOCK_K + tl.arange(0, BLOCK_K)

                acc = tl.zeros((BLOCK_N, BLOCK_K), dtype=tl.float32)

                m_start = 0
                while m_start < gm:
                    offs_m = m_start + tl.arange(0, BLOCK_M)

                    a_tile_ptrs = a_ptr + (offs_m[None, :] * lda + offs_n[:, None])
                    b_tile_ptrs = b_ptr + (offs_m[:, None] * ldb + offs_k[None, :])

                    a = tl.load(
                        a_tile_ptrs,
                        mask=(offs_n[:, None] < N) & (offs_m[None, :] < gm),
                        other=0.0,
                    )
                    b = tl.load(
                        b_tile_ptrs,
                        mask=(offs_m[:, None] < gm) & (offs_k[None, :] < K),
                        other=0.0,
                    )
                    acc += tl.dot(a, b)

                    if m_start + BLOCK_M >= gm:
                        break

                c_tile_ptrs = c_ptr + (offs_n[:, None] * ldc + offs_k[None, :])
                tl.store(
                    c_tile_ptrs,
                    acc,
                    mask=(offs_n[:, None] < N) & (offs_k[None, :] < K),
                )

                tile_idx += num_pids

            last_end += num_tiles


    class _TritonMoEGroupedLinear(torch.autograd.Function):
        @staticmethod
        def forward(ctx, x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64, *weights):
            if x_sorted.dtype not in (torch.float16, torch.bfloat16):
                raise TypeError(f"Triton grouped GEMM expects bf16/fp16, got {x_sorted.dtype}")
            if not x_sorted.is_cuda:
                raise RuntimeError("Triton grouped GEMM requires CUDA tensors")

            group_size = len(weights)
            # Expect weight layout [N, K] contiguous.
            N = int(weights[0].shape[0])
            K = int(weights[0].shape[1])
            for w in weights:
                if w.dtype != x_sorted.dtype:
                    raise TypeError(f"Weight dtype {w.dtype} must match x dtype {x_sorted.dtype} for Triton grouped GEMM")
                if w.shape != (N, K):
                    raise ValueError("All expert weights must share the same shape [N, K] for grouped GEMM")
                if w.stride(1) != 1:
                    raise ValueError("Expert weights must be contiguous in the last dimension (stride(1)==1)")

            # Ensure x_sorted contiguous [M, K]
            if x_sorted.stride(1) != 1:
                x_sorted = x_sorted.contiguous()
            if x_sorted.shape[1] != K:
                raise ValueError(f"x_sorted has wrong K: {x_sorted.shape[1]} vs weight K {K}")

            y_sorted = torch.empty((x_sorted.shape[0], N), device=x_sorted.device, dtype=x_sorted.dtype)

            # Build per-group A/C pointers on GPU: base + offsets * stride0 * element_size
            elem = x_sorted.element_size()
            offs_i64 = offsets_i32[:-1].to(torch.int64)
            a_ptrs = offs_i64 * (x_sorted.stride(0) * elem) + int(x_sorted.data_ptr())
            c_ptrs = offs_i64 * (y_sorted.stride(0) * elem) + int(y_sorted.data_ptr())

            # Launch
            num_sms = _moe_triton_num_sms(x_sorted.device)
            grid = (num_sms,)

            if x_sorted.dtype == torch.bfloat16:
                _moe_grouped_linear_fwd_bf16[grid](
                    a_ptrs, w_ptrs_i64, c_ptrs,
                    m_sizes_i32,
                    lda=x_sorted.stride(0), ldb=K, ldc=y_sorted.stride(0),
                    GROUP_SIZE=group_size,
                    N=N, K=K,
                    BLOCK_M=64, BLOCK_N=128, BLOCK_K=32,
                    num_warps=4,
                )
            else:
                _moe_grouped_linear_fwd_fp16[grid](
                    a_ptrs, w_ptrs_i64, c_ptrs,
                    m_sizes_i32,
                    lda=x_sorted.stride(0), ldb=K, ldc=y_sorted.stride(0),
                    GROUP_SIZE=group_size,
                    N=N, K=K,
                    BLOCK_M=64, BLOCK_N=128, BLOCK_K=32,
                    num_warps=4,
                )

            # Save for backward
            ctx.save_for_backward(x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64)
            ctx.N = N
            ctx.K = K
            ctx.group_size = group_size
            ctx.dtype = x_sorted.dtype

            return y_sorted

        @staticmethod
        def backward(ctx, grad_y_sorted):
            x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64 = ctx.saved_tensors
            N = int(ctx.N)
            K = int(ctx.K)
            group_size = int(ctx.group_size)
            dtype = ctx.dtype

            # grad wrt x_sorted: dX = dY @ W
            grad_x = None
            if ctx.needs_input_grad[0]:
                grad_x = torch.empty((x_sorted.shape[0], K), device=x_sorted.device, dtype=dtype)

                elem = grad_y_sorted.element_size()
                offs_i64 = offsets_i32[:-1].to(torch.int64)
                a_ptrs = offs_i64 * (grad_y_sorted.stride(0) * elem) + int(grad_y_sorted.data_ptr())
                c_ptrs = offs_i64 * (grad_x.stride(0) * elem) + int(grad_x.data_ptr())

                num_sms = _moe_triton_num_sms(x_sorted.device)
                grid = (num_sms,)

                # For bwd_x kernel: A = dY [M, N] => K(shared)=N, output cols = K_in
                if dtype == torch.bfloat16:
                    _moe_grouped_linear_bwd_x_bf16[grid](
                        a_ptrs, w_ptrs_i64, c_ptrs,
                        m_sizes_i32,
                        lda=grad_y_sorted.stride(0), ldb=K, ldc=grad_x.stride(0),
                        GROUP_SIZE=group_size,
                        N=K, K=N,
                        BLOCK_M=64, BLOCK_N=128, BLOCK_K=32,
                        num_warps=4,
                    )
                else:
                    _moe_grouped_linear_bwd_x_fp16[grid](
                        a_ptrs, w_ptrs_i64, c_ptrs,
                        m_sizes_i32,
                        lda=grad_y_sorted.stride(0), ldb=K, ldc=grad_x.stride(0),
                        GROUP_SIZE=group_size,
                        N=K, K=N,
                        BLOCK_M=64, BLOCK_N=128, BLOCK_K=32,
                        num_warps=4,
                    )

            # grad wrt weights: dW = dY^T @ X
            grad_weights = []
            if any(ctx.needs_input_grad[4:]):
                # Allocate grads for each expert weight [N, K]
                grads = [torch.empty((N, K), device=x_sorted.device, dtype=dtype) for _ in range(group_size)]
                dw_ptrs = torch.tensor([int(g.data_ptr()) for g in grads], device=x_sorted.device, dtype=torch.int64)

                elem = x_sorted.element_size()
                offs_i64 = offsets_i32[:-1].to(torch.int64)
                dy_ptrs = offs_i64 * (grad_y_sorted.stride(0) * elem) + int(grad_y_sorted.data_ptr())
                x_ptrs = offs_i64 * (x_sorted.stride(0) * elem) + int(x_sorted.data_ptr())

                num_sms = _moe_triton_num_sms(x_sorted.device)
                grid = (num_sms,)

                if dtype == torch.bfloat16:
                    _moe_grouped_linear_bwd_w_bf16[grid](
                        dy_ptrs, x_ptrs, dw_ptrs,
                        m_sizes_i32,
                        lda=grad_y_sorted.stride(0), ldb=x_sorted.stride(0), ldc=K,
                        GROUP_SIZE=group_size,
                        N=N, K=K,
                        BLOCK_N=128, BLOCK_K=64, BLOCK_M=32,
                        num_warps=4,
                    )
                else:
                    _moe_grouped_linear_bwd_w_fp16[grid](
                        dy_ptrs, x_ptrs, dw_ptrs,
                        m_sizes_i32,
                        lda=grad_y_sorted.stride(0), ldb=x_sorted.stride(0), ldc=K,
                        GROUP_SIZE=group_size,
                        N=N, K=K,
                        BLOCK_N=128, BLOCK_K=64, BLOCK_M=32,
                        num_warps=4,
                    )

                # Map grads to inputs: return grads for weight inputs, None for non-tensor args
                for i in range(group_size):
                    grad_weights.append(grads[i])
            else:
                grad_weights = [None] * group_size

            # Inputs to forward: (x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64, *weights)
            return (grad_x, None, None, None, *grad_weights)


    def triton_moe_grouped_linear(x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64, weights: List[torch.Tensor]):
        # Wrapper to keep call sites clean.
        return _TritonMoEGroupedLinear.apply(x_sorted, m_sizes_i32, offsets_i32, w_ptrs_i64, *weights)





def _dist_is_initialized() -> bool:
    """Return True if torch.distributed is available and initialized."""
    return dist.is_available() and dist.is_initialized()

if is_flash_attn_2_available():
    from flash_attn import flash_attn_func, flash_attn_varlen_func
    from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input  # noqa


# -----------------------------------------------------------------------------
# Fallback padding helpers (used by Kimi-Linear style Delta Attention / KDA)
# -----------------------------------------------------------------------------
# flash-attn provides fast implementations of these helpers. If flash-attn isn't
# available, we provide simple PyTorch fallbacks.
if "index_first_axis" not in globals():
    def index_first_axis(x: torch.Tensor, indices: torch.Tensor) -> torch.Tensor:
        """Gather `x` on the first axis using 1D `indices`."""
        return x.index_select(0, indices)

    def pad_input(
        hidden_states: torch.Tensor,
        indices: torch.Tensor,
        batch_size: int,
        seq_len: int,
    ) -> torch.Tensor:
        """Inverse of `index_first_axis` for (batch, seq) token selection."""
        first_axis_dim = batch_size * seq_len
        out = torch.zeros(
            first_axis_dim,
            *hidden_states.shape[1:],
            device=hidden_states.device,
            dtype=hidden_states.dtype,
        )
        out.index_copy_(0, indices, hidden_states)
        return out.view(batch_size, seq_len, *hidden_states.shape[1:])


# -----------------------------------------------------------------------------
# Optional dependency: fla-core for Kimi Delta Attention (KDA)
# -----------------------------------------------------------------------------
_FLA_AVAILABLE = False
try:
    from fla.modules import FusedRMSNormGated, ShortConvolution
    from fla.ops.kda import chunk_kda, fused_recurrent_kda
    from fla.ops.kda.gate import fused_kda_gate
    _FLA_AVAILABLE = True
except Exception:
    FusedRMSNormGated = None
    ShortConvolution = None
    chunk_kda = None
    fused_recurrent_kda = None
    fused_kda_gate = None
    _FLA_AVAILABLE = False


# This makes `_prepare_4d_causal_attention_mask` a leaf function in the FX graph.
# It means that the function will not be traced through and simply appear as a node in the graph.
if is_torch_fx_available():
    if not is_torch_greater_or_equal_than_1_13:
        import torch.fx

    _prepare_4d_causal_attention_mask = torch.fx.wrap(_prepare_4d_causal_attention_mask)


logger = logging.get_logger(__name__)

_CONFIG_FOR_DOC = "RizeConfig"


def _get_unpad_data(attention_mask):
    seqlens_in_batch = attention_mask.sum(dim=-1, dtype=torch.int32)
    indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
    max_seqlen_in_batch = seqlens_in_batch.max().item()
    cu_seqlens = F.pad(
        torch.cumsum(seqlens_in_batch, dim=0, dtype=torch.torch.int32), (1, 0)
    )
    cu_seqlens = cu_seqlens.to(dtype=torch.int32, device=attention_mask.device).contiguous()
    return (
        indices,
        cu_seqlens,
        max_seqlen_in_batch,
    )


def _get_unpad_data_from_sequence_ids(
    sequence_ids: torch.Tensor,
    attention_mask: Optional[torch.Tensor] = None,
):
    """Build FlashAttention varlen metadata from packed-sample ids.

    `sequence_ids` is expected to be shape [B, S] and constant on each contiguous
    packed sample span. Tokens from different packed samples become different
    variable-length sequences, which yields an exact block-diagonal causal mask.
    """
    if sequence_ids is None:
        raise ValueError("sequence_ids must not be None")

    if attention_mask is None:
        active_mask = torch.ones_like(sequence_ids, dtype=torch.bool)
    else:
        active_mask = attention_mask.to(dtype=torch.bool)

    flat_active = active_mask.reshape(-1)
    indices = torch.nonzero(flat_active, as_tuple=False).flatten()
    if indices.numel() == 0:
        cu = torch.zeros((1,), device=sequence_ids.device, dtype=torch.int32)
        return indices, cu, 0

    seq_i64 = sequence_ids.to(dtype=torch.int64)
    safe_seq = torch.where(active_mask, seq_i64, torch.zeros_like(seq_i64))
    max_seq_id = int(safe_seq.max().item()) if safe_seq.numel() > 0 else 0
    stride = max(1, max_seq_id + 1)

    batch_offsets = (
        torch.arange(sequence_ids.shape[0], device=sequence_ids.device, dtype=torch.int64).unsqueeze(1)
        * int(stride)
    )
    global_seq = torch.where(
        active_mask,
        seq_i64 + batch_offsets,
        torch.full_like(seq_i64, -1),
    )

    active_seq = global_seq.reshape(-1).index_select(0, indices)
    starts = torch.ones((active_seq.shape[0],), device=active_seq.device, dtype=torch.bool)
    if active_seq.numel() > 1:
        starts[1:] = active_seq[1:] != active_seq[:-1]

    start_pos = torch.nonzero(starts, as_tuple=False).flatten()
    bounds = torch.cat(
        [start_pos, torch.tensor([active_seq.numel()], device=active_seq.device, dtype=start_pos.dtype)]
    )
    seqlens = bounds[1:] - bounds[:-1]
    cu_seqlens = F.pad(torch.cumsum(seqlens.to(torch.int32), dim=0), (1, 0))
    cu_seqlens = cu_seqlens.to(dtype=torch.int32, device=sequence_ids.device).contiguous()
    max_seqlen = int(seqlens.max().item()) if seqlens.numel() > 0 else 0
    return indices, cu_seqlens, max_seqlen


def _build_reset_position_ids_from_sequence_ids(
    sequence_ids: torch.Tensor,
    attention_mask: Optional[torch.Tensor] = None,
) -> torch.LongTensor:
    """Return per-token position ids reset to zero at each packed-sample boundary."""
    if sequence_ids is None:
        raise ValueError("sequence_ids must not be None")

    if attention_mask is None:
        active_mask = torch.ones_like(sequence_ids, dtype=torch.bool)
    else:
        active_mask = attention_mask.to(dtype=torch.bool)

    batch_size, seq_len = sequence_ids.shape
    token_pos = torch.arange(seq_len, device=sequence_ids.device, dtype=torch.long).unsqueeze(0).expand(batch_size, -1)

    starts = active_mask.clone()
    if seq_len > 1:
        starts[:, 1:] = active_mask[:, 1:] & (
            (~active_mask[:, :-1]) | (sequence_ids[:, 1:] != sequence_ids[:, :-1])
        )

    last_start = torch.where(starts, token_pos, torch.zeros_like(token_pos))
    last_start = torch.cummax(last_start, dim=-1).values
    position_ids = token_pos - last_start
    position_ids = torch.where(active_mask, position_ids, torch.zeros_like(position_ids))
    return position_ids


def _prepare_4d_block_diagonal_causal_attention_mask(
    attention_mask: Optional[torch.Tensor],
    sequence_ids: torch.Tensor,
    input_shape: Tuple[int, int],
    inputs_embeds: torch.Tensor,
    past_key_values_length: int = 0,
):
    """Create a 4D additive mask that blocks attention across packed-sample boundaries."""
    if past_key_values_length != 0:
        raise NotImplementedError(
            "block-diagonal packed attention does not support cached decoding/past_key_values yet"
        )

    bsz, tgt_len = input_shape
    device = inputs_embeds.device
    dtype = inputs_embeds.dtype

    if attention_mask is None:
        active_mask = torch.ones((bsz, tgt_len), device=device, dtype=torch.bool)
    else:
        active_mask = attention_mask.to(device=device, dtype=torch.bool)

    seq = sequence_ids.to(device=device)
    causal = torch.tril(torch.ones((tgt_len, tgt_len), device=device, dtype=torch.bool)).unsqueeze(0)
    same_seq = seq[:, :, None] == seq[:, None, :]
    allowed = same_seq & causal & active_mask[:, :, None] & active_mask[:, None, :]

    min_value = torch.finfo(dtype).min
    mask = torch.full((bsz, 1, tgt_len, tgt_len), min_value, device=device, dtype=dtype)
    mask = mask.masked_fill(allowed.unsqueeze(1), 0)

    # Avoid NaNs on padded query rows by zeroing the whole row (their outputs are ignored anyway).
    if not torch.all(active_mask):
        mask = torch.where(
            active_mask[:, None, :, None],
            mask,
            torch.zeros_like(mask),
        )

    return mask


class RizeRMSNorm(nn.Module):
    def __init__(self, hidden_size, eps=1e-6):
        """
        RizeRMSNorm is equivalent to T5LayerNorm

        Fast path:
            - Prefer `torch.nn.functional.rms_norm` when available. It dispatches to a
            fused kernel on recent PyTorch/CUDA stacks and avoids materializing a full
            fp32 copy of the activation tensor on every call.
        
        Fallback:
            - Keep the original fp32 implementation for older runtimes or when the fast
            path is explicitly disabled via `USE_FAST_RMSNORM=0`
        """
        super().__init__()
        self.weight = nn.Parameter(torch.ones(hidden_size))
        self.variance_epsilon = eps
        self.use_fast_rmsnorm = os.getenv("USE_FAST_RMSNORM", "1").lower() not in ("0", "false", "no")

    def forward(self, hidden_states):
        if self.use_fast_rmsnorm and hasattr(F, "rms_norm"):
            try:
                return F.rms_norm(
                    hidden_states,
                    (self.hidden_size,),
                    self.weight.to(dtype=hidden_states.dtype),
                    self.variance_epsilon,
                )
            except Exception:
                # Fall back to the numerically conservative implementation below.
                pass

        input_dtype = hidden_states.dtype
        hidden_states = hidden_states.to(torch.float32)
        variance = hidden_states.pow(2).mean(-1, keepdim=True)
        hidden_states = hidden_states * torch.rsqrt(variance + self.variance_epsilon)
        return self.weight.to(dtype=input_dtype) * hidden_states.to(input_dtype)


ALL_LAYERNORM_LAYERS.append(RizeRMSNorm)


class RizeRotaryEmbedding(nn.Module):
    def __init__(self, dim, max_position_embeddings=2048, base=10000, device=None):
        super().__init__()

        self.dim = dim
        self.max_position_embeddings = max_position_embeddings
        self.base = base
        inv_freq = 1.0 / (
            self.base ** (torch.arange(0, self.dim, 2).float().to(device) / self.dim)
        )
        self.register_buffer("inv_freq", inv_freq, persistent=False)

        # Build here to make `torch.jit.trace` work.
        self._set_cos_sin_cache(
            seq_len=max_position_embeddings,
            device=self.inv_freq.device,
            dtype=torch.get_default_dtype(),
        )
        self.max_seq_len_cached = None

    def _set_cos_sin_cache(self, seq_len, device, dtype):
        self.max_seq_len_cached = seq_len
        t = torch.arange(
            self.max_seq_len_cached, device=device, dtype=self.inv_freq.dtype
        )

        freqs = torch.outer(t, self.inv_freq.to(t.device))
        # Different from paper, but it uses a different permutation in order to obtain the same calculation
        emb = torch.cat((freqs, freqs), dim=-1)
        self.register_buffer("cos_cached", emb.cos().to(dtype), persistent=False)
        self.register_buffer("sin_cached", emb.sin().to(dtype), persistent=False)

    def forward(self, x, seq_len=None):
        # x: [bs, num_attention_heads, seq_len, head_size]
        if self.max_seq_len_cached is None or seq_len > self.max_seq_len_cached:
            self._set_cos_sin_cache(seq_len=seq_len, device=x.device, dtype=x.dtype)

        return (
            self.cos_cached[:seq_len].to(dtype=x.dtype),
            self.sin_cached[:seq_len].to(dtype=x.dtype),
        )


# Copied from transformers.models.llama.modeling_llama.LlamaLinearScalingRotaryEmbedding with Llama->Rize
class RizeLinearScalingRotaryEmbedding(RizeRotaryEmbedding):
    """RizeRotaryEmbedding extended with linear scaling. Credits to the Reddit user /u/kaiokendev"""

    def __init__(
        self,
        dim,
        max_position_embeddings=2048,
        base=10000,
        device=None,
        scaling_factor=1.0,
    ):
        self.scaling_factor = scaling_factor
        super().__init__(dim, max_position_embeddings, base, device)

    def _set_cos_sin_cache(self, seq_len, device, dtype):
        self.max_seq_len_cached = seq_len
        t = torch.arange(
            self.max_seq_len_cached, device=device, dtype=self.inv_freq.dtype
        )
        t = t / self.scaling_factor

        freqs = torch.outer(t, self.inv_freq)
        # Different from paper, but it uses a different permutation in order to obtain the same calculation
        emb = torch.cat((freqs, freqs), dim=-1)
        self.register_buffer("cos_cached", emb.cos().to(dtype), persistent=False)
        self.register_buffer("sin_cached", emb.sin().to(dtype), persistent=False)


# Copied from transformers.models.llama.modeling_llama.LlamaDynamicNTKScalingRotaryEmbedding with Llama->Rize
class RizeDynamicNTKScalingRotaryEmbedding(RizeRotaryEmbedding):
    """RizeRotaryEmbedding extended with Dynamic NTK scaling. Credits to the Reddit users /u/bloc97 and /u/emozilla"""

    def __init__(
        self,
        dim,
        max_position_embeddings=2048,
        base=10000,
        device=None,
        scaling_factor=1.0,
    ):
        self.scaling_factor = scaling_factor
        super().__init__(dim, max_position_embeddings, base, device)

    def _set_cos_sin_cache(self, seq_len, device, dtype):
        self.max_seq_len_cached = seq_len

        if seq_len > self.max_position_embeddings:
            base = self.base * (
                (self.scaling_factor * seq_len / self.max_position_embeddings)
                - (self.scaling_factor - 1)
            ) ** (self.dim / (self.dim - 2))
            inv_freq = 1.0 / (
                base ** (torch.arange(0, self.dim, 2).float().to(device) / self.dim)
            )
            self.register_buffer("inv_freq", inv_freq, persistent=False)

        t = torch.arange(
            self.max_seq_len_cached, device=device, dtype=self.inv_freq.dtype
        )

        freqs = torch.outer(t, self.inv_freq)
        # Different from paper, but it uses a different permutation in order to obtain the same calculation
        emb = torch.cat((freqs, freqs), dim=-1)
        self.register_buffer("cos_cached", emb.cos().to(dtype), persistent=False)
        self.register_buffer("sin_cached", emb.sin().to(dtype), persistent=False)


# Inverse dim formula to find dim based on number of rotations
def yarn_find_correction_dim(
    num_rotations, dim, base=10000, max_position_embeddings=2048
):
    return (dim * math.log(max_position_embeddings / (num_rotations * 2 * math.pi))) / (
        2 * math.log(base)
    )


# Find dim range bounds based on rotations
def yarn_find_correction_range(
    low_rot, high_rot, dim, base=10000, max_position_embeddings=2048
):
    low = math.floor(
        yarn_find_correction_dim(low_rot, dim, base, max_position_embeddings)
    )
    high = math.ceil(
        yarn_find_correction_dim(high_rot, dim, base, max_position_embeddings)
    )
    return max(low, 0), min(high, dim - 1)  # Clamp values just in case


def yarn_get_mscale(scale=1, mscale=1):
    if scale <= 1:
        return 1.0
    return 0.1 * mscale * math.log(scale) + 1.0


def yarn_linear_ramp_mask(min, max, dim):
    if min == max:
        max += 0.001  # Prevent singularity

    linear_func = (torch.arange(dim, dtype=torch.float32) - min) / (max - min)
    ramp_func = torch.clamp(linear_func, 0, 1)
    return ramp_func


class RizeYarnRotaryEmbedding(RizeRotaryEmbedding):

    def __init__(
        self,
        dim,
        max_position_embeddings=2048,
        base=10000,
        device=None,
        scaling_factor=1.0,
        original_max_position_embeddings=4096,
        beta_fast=32,
        beta_slow=1,
        mscale=1,
        mscale_all_dim=0,
    ):
        self.scaling_factor = scaling_factor
        self.original_max_position_embeddings = original_max_position_embeddings
        self.beta_fast = beta_fast
        self.beta_slow = beta_slow
        self.mscale = mscale
        self.mscale_all_dim = mscale_all_dim
        super().__init__(dim, max_position_embeddings, base, device)

    def _set_cos_sin_cache(self, seq_len, device, dtype):
        self.max_seq_len_cached = seq_len
        dim = self.dim

        freq_extra = 1.0 / (
            self.base
            ** (torch.arange(0, dim, 2, dtype=torch.float32, device=device) / dim)
        )
        freq_inter = 1.0 / (
            self.scaling_factor
            * self.base
            ** (torch.arange(0, dim, 2, dtype=torch.float32, device=device) / dim)
        )

        low, high = yarn_find_correction_range(
            self.beta_fast,
            self.beta_slow,
            dim,
            self.base,
            self.original_max_position_embeddings,
        )
        inv_freq_mask = 1.0 - yarn_linear_ramp_mask(low, high, dim // 2).to(
            device=device, dtype=torch.float32
        )
        inv_freq = freq_inter * (1 - inv_freq_mask) + freq_extra * inv_freq_mask
        self.register_buffer("inv_freq", inv_freq, persistent=False)

        t = torch.arange(seq_len, device=device, dtype=torch.float32)

        freqs = torch.outer(t, inv_freq)

        _mscale = float(
            yarn_get_mscale(self.scaling_factor, self.mscale)
            / yarn_get_mscale(self.scaling_factor, self.mscale_all_dim)
        )

        emb = torch.cat((freqs, freqs), dim=-1)
        self.register_buffer(
            "cos_cached", (emb.cos() * _mscale).to(dtype), persistent=False
        )
        self.register_buffer(
            "sin_cached", (emb.sin() * _mscale).to(dtype), persistent=False
        )


# Copied from transformers.models.llama.modeling_llama.rotate_half
def rotate_half(x):
    """Rotates half the hidden dims of the input."""
    x1 = x[..., : x.shape[-1] // 2]
    x2 = x[..., x.shape[-1] // 2 :]
    return torch.cat((-x2, x1), dim=-1)


# Copied from transformers.models.llama.modeling_llama.apply_rotary_pos_emb
def apply_rotary_pos_emb(q, k, cos, sin, position_ids, unsqueeze_dim=1):
    """Applies Rotary Position Embedding to the query and key tensors.

    Args:
        q (`torch.Tensor`): The query tensor.
        k (`torch.Tensor`): The key tensor.
        cos (`torch.Tensor`): The cosine part of the rotary embedding.
        sin (`torch.Tensor`): The sine part of the rotary embedding.
        position_ids (`torch.Tensor`):
            The position indices of the tokens corresponding to the query and key tensors. For example, this can be
            used to pass offsetted position ids when working with a KV-cache.
        unsqueeze_dim (`int`, *optional*, defaults to 1):
            The 'unsqueeze_dim' argument specifies the dimension along which to unsqueeze cos[position_ids] and
            sin[position_ids] so that they can be properly broadcasted to the dimensions of q and k. For example, note
            that cos[position_ids] and sin[position_ids] have the shape [batch_size, seq_len, head_dim]. Then, if q and
            k have the shape [batch_size, heads, seq_len, head_dim], then setting unsqueeze_dim=1 makes
            cos[position_ids] and sin[position_ids] broadcastable to the shapes of q and k. Similarly, if q and k have
            the shape [batch_size, seq_len, heads, head_dim], then set unsqueeze_dim=2.
    Returns:
        `tuple(torch.Tensor)` comprising of the query and key tensors rotated using the Rotary Position Embedding.
    """
    cos = cos[position_ids].unsqueeze(unsqueeze_dim)
    sin = sin[position_ids].unsqueeze(unsqueeze_dim)

    b, h, s, d = q.shape
    q = q.view(b, h, s, d // 2, 2).transpose(4, 3).reshape(b, h, s, d)

    b, h, s, d = k.shape
    k = k.view(b, h, s, d // 2, 2).transpose(4, 3).reshape(b, h, s, d)

    q_embed = (q * cos) + (rotate_half(q) * sin)
    k_embed = (k * cos) + (rotate_half(k) * sin)
    return q_embed, k_embed


class RizeMLP(nn.Module):
    def __init__(self, config, hidden_size=None, intermediate_size=None):
        super().__init__()
        self.config = config
        self.hidden_size = config.hidden_size if hidden_size is None else hidden_size
        self.intermediate_size = (
            config.intermediate_size if intermediate_size is None else intermediate_size
        )

        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
        self.act_fn = ACT2FN[config.hidden_act]

    def forward(self, x):
        # Ensure dtype matches expert weights (important for bf16 kernels / Triton grouped GEMM).
        x = x.to(self.gate_proj.weight.dtype)
        down_proj = self.down_proj(self.act_fn(self.gate_proj(x)) * self.up_proj(x))
        return down_proj


class MoEGate(nn.Module):
    """Rize/Moonlight-style router (gate) with paper-aligned fixes.

    Paper-aligned points implemented here:
      - Gate scaling factor (Appendix C): routed_scaling_factor is set (or auto-estimated)
        so that MoE output RMS matches dense models.
      - Aux-free bias update (Appendix C): bias is updated by
            b_i <- b_i + u * (sign(e_i) - mean(sign(e)))
        where e_i measures *underuse* vs the desired uniform load.
      - Bias/EMA buffers are kept in FP32 even when the model runs in bf16.
      - Bias update is synchronized across ranks (all_reduce of counts), so every rank
        applies the same update (critical because this bias has no gradients).
    """

    _GATE_SCALE_CACHE = {}

    @staticmethod
    @torch.no_grad()
    def _estimate_gate_scaling_factor(
        num_experts: int,
        top_k: int,
        *,
        num_samples: int = 20000,
        seed: int = 0,
    ) -> float:
        """Monte-Carlo estimate used in the Moonlight paper (Figure 6).

        Moonlight computes the *per-sample* scaling factor as:

            factor = 1 / sqrt(sum_i p_i^2)

        where p is the top-k routing probability vector after sigmoid and renormalization,
        then returns mean(factor) across samples. (This differs from 1/E[||p||].)

        We follow that implementation exactly.

        Notes:
          - For (num_experts=64, top_k=6) this yields ~2.446.
          - This is deterministic given (seed, num_samples).
        """
        key = (int(num_experts), int(top_k), int(num_samples), int(seed))
        if key in MoEGate._GATE_SCALE_CACHE:
            return MoEGate._GATE_SCALE_CACHE[key]

        g = torch.Generator(device='cpu')
        g.manual_seed(int(seed))
        logits = torch.randn(int(num_samples), int(num_experts), generator=g, dtype=torch.float32, device='cpu')
        scores = torch.sigmoid(logits)
        topk_scores, _ = torch.topk(scores, k=int(top_k), dim=-1)
        topk_scores = topk_scores / topk_scores.sum(dim=-1, keepdim=True).clamp_min(1e-12)

        # Paper's Figure-6 logic: mean( 1 / sqrt(sum(p^2)) ).
        factors = 1.0 / torch.sqrt(torch.sum(topk_scores * topk_scores, dim=-1).clamp_min(1e-20))
        scale = float(torch.mean(factors).item())
        MoEGate._GATE_SCALE_CACHE[key] = scale
        return scale

    def __init__(self, config):
        super().__init__()
        self.config = config
        self.top_k = int(config.num_experts_per_tok)
        self.n_routed_experts = int(config.n_routed_experts)

        # --- Gate scaling factor (Appendix C) ---
        # If config.routed_scaling_factor is missing / <=0 / "auto", estimate it once.
        rsf = getattr(config, "routed_scaling_factor", None)
        if rsf is None:
            self.routed_scaling_factor = self._estimate_gate_scaling_factor(self.n_routed_experts, self.top_k)
        elif isinstance(rsf, str) and rsf.lower() == "auto":
            self.routed_scaling_factor = self._estimate_gate_scaling_factor(self.n_routed_experts, self.top_k)
        else:
            rsf_f = float(rsf)
            self.routed_scaling_factor = (
                rsf_f if rsf_f > 0.0 else self._estimate_gate_scaling_factor(self.n_routed_experts, self.top_k)
            )

        self.scoring_func = config.scoring_func
        self.n_group = int(getattr(config, "n_group", 1))
        self.topk_group = int(getattr(config, "topk_group", 1))
        self.topk_method = getattr(config, "topk_method", "noaux_tc")
        self.norm_topk_prob = bool(getattr(config, "norm_topk_prob", True))
        self.gating_dim = int(config.hidden_size)
        if self.n_group > 1 and (self.n_routed_experts % self.n_group) != 0:
            raise ValueError(
                "n_routed_experts must be divisible by n_group for grouped routing; "
                f"got n_routed_experts={self.n_routed_experts}, n_group={self.n_group}."
            )

        # router weight
        self.weight = nn.Parameter(torch.empty((self.n_routed_experts, self.gating_dim)))
        nn.init.kaiming_uniform_(self.weight, a=math.sqrt(5))

        # --- Aux-loss-free bias + EMA (Appendix C) ---
        # IMPORTANT: keep these in FP32 (bias updates are ~1e-3 and would be quantized away in bf16).
        if self.topk_method == "noaux_tc":
            self.register_buffer(
                "e_score_correction_bias",
                torch.zeros(self.n_routed_experts, dtype=torch.float32),
                persistent=True,
            )

        # Accumulator for bias update under gradient accumulation.
        # We sum per-microbatch routing counts and apply the aux-free bias update once per
        # `auxfree_update_interval` (typically == gradient_accumulation_steps).
        self.register_buffer(
            "auxfree_count_accum",
            torch.zeros(self.n_routed_experts, dtype=torch.float32),
            persistent=False,
        )

        # Hyperparams (default to paper-friendly values if missing)
        self.auxfree_bias_lr = float(getattr(config, "auxfree_bias_lr", 0.0))          # u in Appendix C (e.g., 1e-3)
        self.auxfree_tau = float(getattr(config, "auxfree_tau", 1.0))                 # temperature for sigmoid
        self.auxfree_momentum = float(getattr(config, "auxfree_momentum", 0.0))       # EMA smoothing (optional)
        self.auxfree_update_interval = int(getattr(config, "auxfree_update_interval", 1))
        self.auxfree_bias_clip = float(getattr(config, "auxfree_bias_clip", 5.0))

        self.register_buffer("ema_count", torch.zeros(self.n_routed_experts, dtype=torch.float32), persistent=False)
        self.register_buffer("_auxfree_step", torch.zeros((), dtype=torch.long), persistent=False)

    def _ensure_fp32_aux_buffers(self):
        # Buffers may get cast to bf16 by model.to(bfloat16) or DeepSpeed bf16.
        # We force them back to fp32 in-place once they appear.
        if hasattr(self, "e_score_correction_bias") and self.e_score_correction_bias.dtype != torch.float32:
            self.e_score_correction_bias.data = self.e_score_correction_bias.data.float()
        if self.ema_count.dtype != torch.float32:
            self.ema_count.data = self.ema_count.data.float()
        if self.auxfree_count_accum.dtype != torch.float32:
            self.auxfree_count_accum.data = self.auxfree_count_accum.data.float()


    @torch.no_grad()
    def _maybe_update_auxfree_bias(
        self,
        topk_idx: torch.Tensor,
        active_mask: Optional[torch.Tensor] = None,
    ) -> None:
        """Apply the aux-free bias update (Moonlight Appendix C) in-place.

        When `active_mask` is given, only non-pad / active tokens contribute to the
        aux-free bias update. This is important for packed SFT where padding tokens
        should not perturb router load statistics.
        """
        self._auxfree_step.add_(1)

        if active_mask is not None:
            valid = active_mask.reshape(-1).to(dtype=torch.bool, device=topk_idx.device)
            topk_idx = topk_idx[valid]

        if topk_idx.numel() == 0:
            batch_count = torch.zeros(
                (self.n_routed_experts,),
                dtype=torch.float32,
                device=self.auxfree_count_accum.device,
            )
        else:
            flat = topk_idx.reshape(-1)
            batch_count = torch.bincount(flat, minlength=self.n_routed_experts).to(torch.float32)

        self.auxfree_count_accum.add_(batch_count)

        # Apply the update once per auxfree_update_interval (typically per optimizer step).
        if (int(self._auxfree_step.item()) % max(1, self.auxfree_update_interval)) != 0:
            return

        count_total = self.auxfree_count_accum

        # Synchronize across ranks so the bias update is identical everywhere.
        if _dist_is_initialized():
            dist.all_reduce(count_total, op=dist.ReduceOp.SUM)

        m = float(self.auxfree_momentum)
        if m > 0.0:
            self.ema_count.mul_(m).add_(count_total, alpha=1.0 - m)
            count_for_update = self.ema_count
        else:
            count_for_update = count_total

        desired = count_for_update.mean()
        # Underuse is positive: e_i = desired - count_i
        e = desired - count_for_update
        sgn = torch.sign(e)
        sgn = sgn - sgn.mean()
        self.e_score_correction_bias.add_(sgn * float(self.auxfree_bias_lr))

        if self.auxfree_bias_clip > 0.0:
            self.e_score_correction_bias.clamp_(-self.auxfree_bias_clip, self.auxfree_bias_clip)

        # Reset accumulator for the next interval.
        self.auxfree_count_accum.zero_()

    def forward(
        self,
        hidden_states,
        active_mask: Optional[torch.Tensor] = None,
    ):
        # hidden_states: [B, S, H]
        B, S, H = hidden_states.shape
        N = B * S

        # Router computation uses weight dtype for matmul, then casts to fp32 for stable top-k and bias update.
        g_dtype = self.weight.dtype
        logits = F.linear(hidden_states.view(-1, H).to(g_dtype), self.weight, None)

        if self.scoring_func == "sigmoid":
            tau = max(self.auxfree_tau, 1e-6)
            scores = torch.sigmoid(logits.float() / tau)  # fp32
        else:
            raise NotImplementedError(f"Unsupported gating scoring_func: {self.scoring_func}")

        if self.topk_method == "noaux_tc":
            # FP32 aux buffers (critical)
            if self.training and self.auxfree_bias_lr > 0.0:
                self._ensure_fp32_aux_buffers()

            scores_for_choice = scores.view(N, -1)  # [N, E], fp32
            if hasattr(self, "e_score_correction_bias"):
                # Bias is only for *selection*, not for combining weights.
                scores_for_choice = scores_for_choice + self.e_score_correction_bias.unsqueeze(0)

            if self.n_group > 1:
                # group selection then top-k within selected groups
                group_size = self.n_routed_experts // self.n_group
                k_in_group = 2 if group_size >= 2 else 1
                group_scores = scores_for_choice.view(N, self.n_group, group_size).topk(k_in_group, dim=-1)[0].sum(dim=-1)
                group_idx = torch.topk(group_scores, k=self.topk_group, dim=-1, sorted=False)[1]
                group_mask = torch.zeros_like(group_scores)
                group_mask.scatter_(1, group_idx, 1)
                score_mask = group_mask.unsqueeze(-1).expand(N, self.n_group, group_size).reshape(N, -1)
                masked_scores = scores_for_choice.masked_fill(~score_mask.bool(), 0.0)
                _, topk_idx = torch.topk(masked_scores, k=self.top_k, dim=-1, sorted=False)
            else:
                _, topk_idx = torch.topk(scores_for_choice, k=self.top_k, dim=-1, sorted=False)

            # Combine weights come from the original scores (so gradients flow to router weights normally)
            topk_weight = scores.gather(1, topk_idx)

            # --- Aux-free bias update (Appendix C) ---
            # b_i <- b_i + u * (sign(e_i) - mean(sign(e)))
            # where we define e_i as *underuse* relative to uniform desired load.
            if self.training and self.auxfree_bias_lr > 0.0 and hasattr(self, "e_score_correction_bias"):
                # Apply once per optimizer step (typically == gradient_accumulation_steps).
                self._maybe_update_auxfree_bias(topk_idx, active_mask=active_mask)

        else:
            # Fallback: standard top-k (no aux-free bias)
            topk_weight, topk_idx = torch.topk(scores, k=self.top_k, dim=-1, sorted=False)

        # Normalize & gate scaling factor
        if self.top_k > 1 and self.norm_topk_prob:
            topk_weight = topk_weight / topk_weight.sum(dim=-1, keepdim=True).clamp_min(1e-20)
        topk_weight = topk_weight * float(self.routed_scaling_factor)

        return topk_idx, topk_weight


class RizeMoE(nn.Module):
    """
    A mixed expert module containing shared experts.
    """

    def __init__(self, config):
        super().__init__()
        self.config = config
        self.num_experts_per_tok = config.num_experts_per_tok

        if hasattr(config, "ep_size") and config.ep_size > 1:
            # 簡易実装: 学習時 EP 未対応（推論は既存実装）
            self.ep_size = config.ep_size
            assert self.ep_size == dist.get_world_size()
            self.experts_per_rank = config.n_routed_experts // config.ep_size
            self.ep_rank = dist.get_rank()
            self.experts = nn.ModuleList(
                [
                    (
                        RizeMLP(config, intermediate_size=config.moe_intermediate_size)
                        if i >= self.ep_rank * self.experts_per_rank
                        and i < (self.ep_rank + 1) * self.experts_per_rank
                        else None
                    )
                    for i in range(config.n_routed_experts)
                ]
            )
        else:
            self.ep_size = 1
            self.experts_per_rank = config.n_routed_experts
            self.ep_rank = 0
            self.experts = nn.ModuleList(
                [RizeMLP(config, intermediate_size=config.moe_intermediate_size) for _ in range(config.n_routed_experts)]
            )

        # Cached constants (avoid relying on config attribute names at call sites).
        # NOTE: `n_routed_experts` is the global expert count (including non-local experts under EP).
        self.n_routed_experts = int(getattr(config, "n_routed_experts", len(self.experts)))

        self.gate = MoEGate(config)

        # Triton grouped-GEMM acceleration (training)
        self.use_triton_moe = bool(getattr(config, "use_triton_moe", True))
        self._triton_failed = False  # sticky fallback if Triton compilation/execution fails
        self._triton_gate_ptrs = None
        self._triton_up_ptrs = None
        self._triton_down_ptrs = None
        self._triton_ptrs_device = None
        if config.n_shared_experts is not None:
            intermediate_size = config.moe_intermediate_size * config.n_shared_experts
            self.shared_experts = RizeMLP(config=config, intermediate_size=intermediate_size)

        # Global-Batch LBL (Demons in the Detail; ACL 2025):
        # keep a per-layer buffer of synchronized expert counts within the current
        # optimizer step so gradient accumulation approximates a larger balance batch.
        self.global_lbl_enabled = bool(getattr(config, "global_lbl_enabled", False))
        self.global_lbl_sync_across_ranks = bool(getattr(config, "global_lbl_sync_across_ranks", False))
        self.global_lbl_buffer_across_ga = bool(getattr(config, "global_lbl_buffer_across_ga", False))
        if self.global_lbl_buffer_across_ga:
            self.register_buffer(
                "global_lbl_tokens_per_expert_buffer",
                torch.zeros(self.n_routed_experts, dtype=torch.float32),
                persistent=False,
            )

    def _ensure_fp32_global_lbl_buffers(self):
        if hasattr(self, "global_lbl_tokens_per_expert_buffer") and self.global_lbl_tokens_per_expert_buffer.dtype != torch.float32:
            self.global_lbl_tokens_per_expert_buffer.data = self.global_lbl_tokens_per_expert_buffer.data.float()

    @torch.no_grad()
    def reset_global_lbl_buffer(self):
        if hasattr(self, "global_lbl_tokens_per_expert_buffer"):
            self._ensure_fp32_global_lbl_buffers()
            self.global_lbl_tokens_per_expert_buffer.zero_()

    @torch.no_grad()
    def _global_lbl_load_from_counts(
        self,
        idx_flat: torch.Tensor,
        *,
        n_valid: int,
    ) -> tuple[torch.Tensor, torch.Tensor]:
        """Return (load, rank_scale) for Global-Batch LBL.

        - `load` is the expert frequency vector f_i built from synchronized counts.
        - `rank_scale` compensates for unequal numbers of active tokens across ranks so
          that DDP's equal-rank gradient averaging matches a token-weighted global P_i.

        When buffering across gradient accumulation is enabled, synchronized counts are
        accumulated into a per-layer buffer that is reset once per optimizer step by a
        Trainer callback.
        """
        E = int(getattr(self.config, "n_routed_experts"))
        K = int(self.num_experts_per_tok)

        counts = torch.bincount(idx_flat, minlength=E).to(device=idx_flat.device, dtype=torch.float32)

        if self.global_lbl_sync_across_ranks and _dist_is_initialized():
            dist.all_reduce(counts, op=dist.ReduceOp.SUM)

        counts_for_loss = counts
        if self.global_lbl_buffer_across_ga and hasattr(self, "global_lbl_tokens_per_expert_buffer"):
            self._ensure_fp32_global_lbl_buffers()
            self.global_lbl_tokens_per_expert_buffer.add_(counts.to(self.global_lbl_tokens_per_expert_buffer.device))
            counts_for_loss = self.global_lbl_tokens_per_expert_buffer.to(device=counts.device)

        denom = counts_for_loss.sum().clamp_min(1e-12)
        load = counts_for_loss / denom

        rank_scale = counts.new_tensor(1.0)
        if self.global_lbl_sync_across_ranks and _dist_is_initialized():
            global_assignments = float(counts.sum().item())
            if global_assignments > 0.0:
                world = float(dist.get_world_size())
                local_assignments = float(max(0, int(n_valid)) * K)
                rank_scale = counts.new_tensor(world * local_assignments / global_assignments)

        return load, rank_scale

    def _router_aux_loss(
        self,
        topk_idx: torch.Tensor,
        topk_weight: torch.Tensor,
        B: int,
        S: int,
        active_mask: Optional[torch.Tensor] = None,
    ) -> torch.Tensor:
        """
        Compute a standard MoE load-balancing auxiliary loss (importance * load).
        Returns a scalar tensor (fp32).

        Definitions (per batch or per sequence):
          - importance_i = sum of routing probs to expert i / (#tokens)
          - load_i       = count of assignments to expert i / (#tokens * K)
          - aux          = E * sum_i importance_i * load_i

        When `active_mask` is provided, only non-pad / active tokens contribute.
        """
        E = int(getattr(self.config, "n_routed_experts"))
        K = int(self.num_experts_per_tok)
        device = topk_weight.device

        # Convert scaled combine-weights back to probabilities.
        rsf = float(getattr(self.gate, "routed_scaling_factor", 1.0))
        rsf = max(rsf, 1e-12)
        topk_prob = (topk_weight / rsf).to(torch.float32)  # [N, K]
        topk_prob = topk_prob / topk_prob.sum(dim=-1, keepdim=True).clamp_min(1e-12)

        if active_mask is not None:
            active_mask = active_mask.to(device=topk_idx.device, dtype=torch.bool)

        if active_mask is None:
            valid = torch.ones((B * S,), device=topk_idx.device, dtype=torch.bool)
        else:
            valid = active_mask.reshape(-1)

        global_lbl_enabled = bool(getattr(self.config, "global_lbl_enabled", False))
        seq_aux = bool(getattr(self.config, "seq_aux", True)) and (not global_lbl_enabled)

        if not seq_aux:
            n_valid = int(valid.sum().item())
            if n_valid <= 0:
                return torch.zeros((), device=device, dtype=torch.float32)

            idx_flat = topk_idx[valid].reshape(-1)
            prob_flat = topk_prob[valid].reshape(-1)

            if global_lbl_enabled:
                load, rank_scale = self._global_lbl_load_from_counts(idx_flat, n_valid=n_valid)
            else:
                counts = torch.bincount(idx_flat, minlength=E).to(device=device, dtype=torch.float32)
                load = counts / counts.sum().clamp_min(1e-12)
                rank_scale = counts.new_tensor(1.0)

            imp = torch.zeros(E, device=device, dtype=torch.float32)
            imp.index_add_(0, idx_flat, prob_flat)  # sum probs per expert
            imp = imp / float(n_valid)  # sums to ~1

            return (rank_scale * float(E) * (imp * load.to(device=device)).sum()).to(torch.float32)

        idx = topk_idx.view(B, S, K)
        prob = topk_prob.view(B, S, K)
        mask = active_mask.view(B, S) if active_mask is not None else None
        aux_list = []
        for b in range(B):
            if mask is None:
                valid_b = torch.ones((S,), device=idx.device, dtype=torch.bool)
            else:
                valid_b = mask[b]
            n_valid = int(valid_b.sum().item())
            if n_valid <= 0:
                continue

            idx_flat = idx[b][valid_b].reshape(-1)
            prob_flat = prob[b][valid_b].reshape(-1)

            counts = torch.bincount(idx_flat, minlength=E).to(torch.float32)
            load = counts / float(n_valid * K)

            imp = torch.zeros(E, device=device, dtype=torch.float32)
            imp.index_add_(0, idx_flat, prob_flat)
            imp = imp / float(n_valid)

            aux_list.append(float(E) * (imp * load).sum())

        if not aux_list:
            return torch.zeros((), device=device, dtype=torch.float32)

        return torch.stack(aux_list, dim=0).mean().to(torch.float32)


    @torch.no_grad()
    def _compute_moe_stats(
        self,
        topk_idx: torch.Tensor,
        topk_weight: torch.Tensor,
        *,
        B: int,
        S: int,
        active_mask: Optional[torch.Tensor] = None,
    ) -> dict:
        """Compute and return MoE router statistics for logging (best-effort).

        When `active_mask` is provided, only non-pad / active tokens contribute.
        """

        E = int(self.config.n_routed_experts)
        K = int(self.num_experts_per_tok)

        if active_mask is None:
            valid = torch.ones((B * S,), device=topk_idx.device, dtype=torch.bool)
        else:
            valid = active_mask.reshape(-1).to(device=topk_idx.device, dtype=torch.bool)

        valid_topk_idx = topk_idx[valid]
        valid_topk_weight = topk_weight[valid]

        if valid_topk_idx.numel() == 0:
            counts = torch.zeros((E,), device=topk_idx.device, dtype=torch.float32)
            total = counts.sum()
            mean = counts.mean()
            std = counts.std(unbiased=False)
            w_prob = torch.zeros((0, K), device=topk_weight.device, dtype=topk_weight.dtype)
            top1_counts = torch.zeros((E,), device=topk_idx.device, dtype=torch.float32)
            topk_entropy_mean = 0.0
            topk_entropy_std = 0.0
            topk_maxprob_mean = 0.0
            topk_margin_mean = 0.0
            top1_cv_tokens = 0.0
            top1_entropy = 0.0
            avg_topk_weight = 0.0
        else:
            # ---- expert load (token counts) ----
            flat = valid_topk_idx.reshape(-1)
            counts = torch.bincount(flat, minlength=E).to(torch.float32)
            total = counts.sum()
            mean = counts.mean()
            std = counts.std(unbiased=False)

            # ---- top-k mixing entropy per token (how uniform within chosen top-k) ----
            w_prob = valid_topk_weight / valid_topk_weight.sum(dim=-1, keepdim=True).clamp_min(1e-12)
            ent_tok = -(w_prob * w_prob.clamp_min(1e-12).log()).sum(dim=-1)
            if K > 1:
                topk_entropy_mean = float((ent_tok.mean() / math.log(K)).item())
                topk_entropy_std = float((ent_tok.std(unbiased=False) / math.log(K)).item())
            else:
                topk_entropy_mean = 0.0
                topk_entropy_std = 0.0

            topk_maxprob_mean = float(w_prob.max(dim=-1).values.mean().item())
            if K >= 2:
                top2 = torch.topk(w_prob, k=2, dim=-1).values
                topk_margin_mean = float((top2[:, 0] - top2[:, 1]).mean().item())
            else:
                topk_margin_mean = 0.0

            # ---- top-1 expert distribution (based on max prob inside top-k) ----
            top1_pos = w_prob.argmax(dim=-1)
            top1_expert = valid_topk_idx.gather(1, top1_pos.unsqueeze(1)).squeeze(1).reshape(-1)
            top1_counts = torch.bincount(top1_expert, minlength=E).to(torch.float32)
            top1_mean = top1_counts.mean()
            top1_std = top1_counts.std(unbiased=False)
            top1_cv_tokens = float((top1_std / (top1_mean + 1e-6)).item()) if float(top1_mean.item()) > 0.0 else 0.0
            if E > 1 and float(top1_counts.sum().item()) > 0.0:
                p1 = top1_counts / top1_counts.sum().clamp_min(1e-12)
                top1_entropy = float((-(p1 * p1.clamp_min(1e-12).log()).sum() / math.log(E)).item())
            else:
                top1_entropy = 0.0

            avg_topk_weight = float(valid_topk_weight.mean().item())

        if _dist_is_initialized():
            dist.all_reduce(counts, op=dist.ReduceOp.SUM)
            dist.all_reduce(top1_counts, op=dist.ReduceOp.SUM)
            total = counts.sum()
            mean = counts.mean()
            std = counts.std(unbiased=False)
            if valid_topk_idx.numel() > 0:
                top1_mean = top1_counts.mean()
                top1_std = top1_counts.std(unbiased=False)
                top1_cv_tokens = float((top1_std / (top1_mean + 1e-6)).item()) if float(top1_mean.item()) > 0.0 else 0.0
                if E > 1 and float(top1_counts.sum().item()) > 0.0:
                    p1 = top1_counts / top1_counts.sum().clamp_min(1e-12)
                    top1_entropy = float((-(p1 * p1.clamp_min(1e-12).log()).sum() / math.log(E)).item())
                else:
                    top1_entropy = 0.0

        cv_tokens = float((std / (mean + 1e-6)).item()) if float(mean.item()) > 0.0 else 0.0
        n_used = int((counts > 0).sum().item())
        usage_ratio = float(n_used / E) if E > 0 else 0.0

        max_tokens = float(counts.max().item()) if E > 0 else 0.0
        min_tokens = float(counts.min().item()) if E > 0 else 0.0
        min_tokens_nz = float(counts[counts > 0].min().item()) if n_used > 0 else 0.0

        mean_f = float(mean.item()) if E > 0 else 0.0
        tokens_max_to_mean = float(max_tokens / (mean_f + 1e-6)) if mean_f > 0.0 else 0.0
        tokens_min_nonzero_to_mean = float(min_tokens_nz / (mean_f + 1e-6)) if mean_f > 0.0 else 0.0
        tokens_max_to_min_nonzero = float(max_tokens / (min_tokens_nz + 1e-6)) if min_tokens_nz > 0.0 else 0.0

        # ---- load entropy / gini (how uniform across experts) ----
        if E > 1 and float(total.item()) > 0.0:
            p = counts / total.clamp_min(1e-12)
            load_entropy = float((-(p * p.clamp_min(1e-12).log()).sum() / math.log(E)).item())
            effective_experts = float(math.exp(load_entropy * math.log(E)))
            sp, _ = torch.sort(p)
            idx = torch.arange(1, E + 1, device=sp.device, dtype=sp.dtype)
            load_gini = float(((2.0 * (idx * sp).sum() / E) - (E + 1.0) / E).item())
        else:
            load_entropy = 0.0
            effective_experts = float(E)
            load_gini = 0.0

        # ---- aux-free bias distribution (per expert) ----
        bias_per_expert = None
        bias_mean = bias_std = bias_min = bias_max = bias_maxabs = bias_l2 = 0.0
        bias_clip = float(getattr(self.gate, "auxfree_bias_clip", 0.0))
        bias_clip_frac = 0.0
        auxfree_bias_lr = float(getattr(self.gate, "auxfree_bias_lr", 0.0))
        auxfree_tau = float(getattr(self.gate, "auxfree_tau", 1.0))
        routed_scaling_factor = float(getattr(self.gate, "routed_scaling_factor", 1.0))

        if hasattr(self.gate, "e_score_correction_bias"):
            b = self.gate.e_score_correction_bias.detach().float()
            bias_per_expert = b.cpu().tolist()
            bias_mean = float(b.mean().item())
            bias_std = float(b.std(unbiased=False).item())
            bias_min = float(b.min().item())
            bias_max = float(b.max().item())
            bias_maxabs = float(b.abs().max().item())
            bias_l2 = float(b.norm().item())
            if bias_clip > 0.0:
                bias_clip_frac = float((b.abs() >= (bias_clip - 1e-6)).float().mean().item())

        stats = {
            # backwards-compatible keys
            "cv_tokens": cv_tokens,
            "n_used": n_used,
            "usage_ratio": usage_ratio,
            "avg_topk_weight": avg_topk_weight,
            "num_experts": E,
            "k_per_token": K,

            # expert load summaries
            "tokens_total": float(total.item()),
            "tokens_max": max_tokens,
            "tokens_min": min_tokens,
            "tokens_min_nonzero": min_tokens_nz,
            "tokens_max_to_mean": tokens_max_to_mean,
            "tokens_min_nonzero_to_mean": tokens_min_nonzero_to_mean,
            "tokens_max_to_min_nonzero": tokens_max_to_min_nonzero,
            "load_entropy": load_entropy,
            "load_gini": load_gini,
            "effective_experts": effective_experts,

            # within-topk mixing
            "topk_entropy_mean": topk_entropy_mean,
            "topk_entropy_std": topk_entropy_std,
            "topk_maxprob_mean": topk_maxprob_mean,
            "topk_margin_mean": topk_margin_mean,

            # top-1 expert distribution
            "top1_cv_tokens": top1_cv_tokens,
            "top1_entropy": top1_entropy,

            # aux-free bias summaries
            "bias_mean": bias_mean,
            "bias_std": bias_std,
            "bias_min": bias_min,
            "bias_max": bias_max,
            "bias_maxabs": bias_maxabs,
            "bias_l2": bias_l2,
            "bias_clip": bias_clip,
            "bias_clip_frac": bias_clip_frac,
            "auxfree_bias_lr": auxfree_bias_lr,
            "auxfree_tau": auxfree_tau,
            "routed_scaling_factor": routed_scaling_factor,

            # arrays for histogram logging (do NOT expand into per-expert scalar metrics)
            "tokens_per_expert": counts.cpu().tolist(),
            "top1_tokens_per_expert": top1_counts.cpu().tolist(),
        }
        if bias_per_expert is not None:
            stats["bias_per_expert"] = bias_per_expert

        return stats

    def forward(
        self,
        hidden_states,
        token_mask: Optional[torch.Tensor] = None,
    ):
        # hidden_states: [B, S, H]
        identity = hidden_states
        B, S, H = hidden_states.shape
        orig_shape = hidden_states.shape

        if token_mask is not None and bool(getattr(self.config, "moe_router_active_only", True)):
            token_mask = token_mask.to(device=hidden_states.device, dtype=torch.bool)
        else:
            token_mask = None

        topk_idx, topk_weight = self.gate(hidden_states, active_mask=token_mask)  # [B*S, K], [B*S, K]
        aux = None
        if self.training and float(getattr(self.config, "aux_loss_alpha", 0.0)) > 0.0:
            aux = self._router_aux_loss(topk_idx, topk_weight, B=B, S=S, active_mask=token_mask)

        # ===== MoE 利用統計の収集（デバッグ用途。必要なときだけ有効化推奨） =====
        # NOTE: forward内all_reduceはレイヤ数に比例して通信回数が増え、学習速度を悪化させ得る。
        self._last_moe_stats = None
        if bool(getattr(self.config, "collect_moe_stats", False)):
            self._last_moe_stats = self._compute_moe_stats(
                topk_idx,
                topk_weight,
                B=B,
                S=S,
                active_mask=token_mask,
            )

        x_flat = hidden_states.view(-1, hidden_states.shape[-1])
        if self.training:
            y = self.moe_train(x_flat, topk_idx, topk_weight).view(*orig_shape)
        else:
            y = self.moe_infer(x_flat, topk_idx, topk_weight).view(*orig_shape)

        if self.config.n_shared_experts is not None:
            # 共有 MLP への入力も重み dtype に合わせる
            y = y + self.shared_experts(identity)

        if aux is not None:
            return (y, aux)
        return y

    def moe_train(self, x, topk_ids, topk_weight):
        # x: [T, H] (flattened tokens)
        # Ensure dtype matches expert weights (important for bf16 kernels / Triton grouped GEMM).
        if len(self.experts) > 0 and self.experts[0] is not None:
            x = x.to(self.experts[0].gate_proj.weight.dtype)

        num_experts_per_tok = topk_ids.shape[1]
        flat_e = topk_ids.reshape(-1)  # [T * K]
        E = int(getattr(self, "n_routed_experts", getattr(self.config, "n_routed_experts", len(self.experts))))
        counts = torch.bincount(flat_e, minlength=E)  # [E]
        order = flat_e.argsort()
        tok_idx = order // num_experts_per_tok
        x_sorted = x.index_select(0, tok_idx)  # [T*K, H], grouped by expert id

        # Offsets into x_sorted for each expert.
        counts_i32 = counts.to(torch.int32)
        offsets_i32 = torch.empty((E + 1,), device=x.device, dtype=torch.int32)
        offsets_i32[0] = 0
        offsets_i32[1:] = torch.cumsum(counts_i32, dim=0)

        use_triton = (
            _TRITON_MOE_AVAILABLE
            and self.use_triton_moe
            and x_sorted.is_cuda
            and self.ep_size == 1
            and (x_sorted.dtype in (torch.float16, torch.bfloat16))
            and triton is not None
        )
        # Additional safety: disable Triton if weights are not contiguous in the last dim.
        if use_triton:
            try:
                w0 = self.experts[0].gate_proj.weight
                if (w0.dtype not in (torch.float16, torch.bfloat16)) or (w0.stride(1) != 1):
                    use_triton = False
                w1 = self.experts[0].up_proj.weight
                w2 = self.experts[0].down_proj.weight
                if (w1.dtype != w0.dtype) or (w2.dtype != w0.dtype) or (w1.stride(1) != 1) or (w2.stride(1) != 1):
                    use_triton = False
            except Exception:
                use_triton = False
        # Pointer-based grouped GEMM does not work with ZeRO-3 partitioned params.
        # We detect ZeRO-3 heuristically via DeepSpeed's `is_zero_param` flag on expert weights.
        if use_triton and _HAS_DEEPSPEED:
            try:
                if any(
                    (expert is not None) and (
                        is_zero_param(expert.gate_proj.weight)
                        or is_zero_param(expert.up_proj.weight)
                        or is_zero_param(expert.down_proj.weight)
                    )
                    for expert in self.experts
                ):
                    use_triton = False
            except Exception:
                # Be conservative if we cannot determine ZeRO-3 state.
                use_triton = False

        out_sorted = None

        if use_triton and (not getattr(self, "_triton_failed", False)):
            # Cache expert weight pointer arrays per device (cheap to rebuild if device changes).
            if (self._triton_ptrs_device != x.device) or (self._triton_gate_ptrs is None):
                self._triton_ptrs_device = x.device
                self._triton_gate_ptrs = torch.tensor(
                    [int(expert.gate_proj.weight.data_ptr()) for expert in self.experts],
                    device=x.device,
                    dtype=torch.int64,
                )
                self._triton_up_ptrs = torch.tensor(
                    [int(expert.up_proj.weight.data_ptr()) for expert in self.experts],
                    device=x.device,
                    dtype=torch.int64,
                )
                self._triton_down_ptrs = torch.tensor(
                    [int(expert.down_proj.weight.data_ptr()) for expert in self.experts],
                    device=x.device,
                    dtype=torch.int64,
                )

            gate_ws = [expert.gate_proj.weight for expert in self.experts]
            up_ws = [expert.up_proj.weight for expert in self.experts]
            down_ws = [expert.down_proj.weight for expert in self.experts]

            try:
                gate_out = triton_moe_grouped_linear(x_sorted, counts_i32, offsets_i32, self._triton_gate_ptrs, gate_ws)
                up_out = triton_moe_grouped_linear(x_sorted, counts_i32, offsets_i32, self._triton_up_ptrs, up_ws)
                inter = self.experts[0].act_fn(gate_out) * up_out
                out_sorted = triton_moe_grouped_linear(inter, counts_i32, offsets_i32, self._triton_down_ptrs, down_ws)
            except Exception as e:
                # Sticky fallback: if Triton compilation/execution fails once on this setup,
                # keep using the reference implementation to avoid breaking training mid-run.
                self._triton_failed = True
                use_triton = False
                try:
                    _rank = dist.get_rank() if (dist.is_available() and dist.is_initialized()) else 0
                except Exception:
                    _rank = 0
                if _rank == 0:
                    warnings.warn(f"[MoE][Triton] disabled due to error: {type(e).__name__}: {e}")

        if out_sorted is None:
            # Fallback (reference): per-expert slicing + sequential expert forward.
            counts_cpu = counts.tolist()
            # Under DeepSpeed ZeRO-3, some expert params may be partitioned; in that case,
            # make `counts_cpu` consistent across ranks so each rank iterates the same expert set.
            if _HAS_DEEPSPEED:
                try:
                    zero3 = any((ex is not None) and is_zero_param(ex.gate_proj.weight) for ex in self.experts)
                except Exception:
                    zero3 = False
                if zero3 and dist.is_available() and dist.is_initialized():
                    counts_gpu = counts.to(x.device)
                    dist.all_reduce(counts_gpu, op=dist.ReduceOp.MAX)
                    counts_cpu = counts_gpu.tolist()

            outs = []
            start = 0
            for e, cnt in enumerate(counts_cpu):
                end = start + cnt
                local_e = e - self.ep_rank * self.experts_per_rank
                if local_e >= 0 and local_e < self.experts_per_rank:
                    expert = self.experts[local_e]
                else:
                    expert = None
                if expert is not None and cnt > 0:
                    out_e = expert(x_sorted[start:end])
                else:
                    out_e = torch.zeros((cnt, self.config.hidden_size), device=x_sorted.device, dtype=x_sorted.dtype)
                outs.append(out_e)
                start = end
            out_sorted = torch.cat(outs, dim=0)

        # Unsort back and apply top-k gate weights.
        y_all = torch.empty_like(out_sorted)
        y_all[order] = out_sorted
        y = y_all.view(x.shape[0], num_experts_per_tok, -1)
        y = (y * topk_weight.to(y.dtype).unsqueeze(-1)).sum(dim=1)
        return y

    @torch.no_grad()
    def moe_infer(self, x, topk_ids, topk_weight):
        # 既存の推論実装（変更なし）
        cnts = topk_ids.new_zeros((topk_ids.shape[0], len(self.experts)))
        cnts.scatter_(1, topk_ids, 1)
        tokens_per_expert = cnts.sum(dim=0)
        idxs = topk_ids.view(-1).argsort()
        sorted_tokens = x[idxs // topk_ids.shape[1]]
        sorted_tokens_shape = sorted_tokens.shape
        if self.ep_size > 1:
            tokens_per_ep_rank = tokens_per_expert.view(self.ep_size, -1).sum(dim=1)
            tokens_per_expert_group = tokens_per_expert.new_empty(tokens_per_expert.shape[0])
            dist.all_to_all_single(tokens_per_expert_group, tokens_per_expert)
            output_splits = (
                tokens_per_expert_group.view(self.ep_size, -1).sum(1).cpu().numpy().tolist()
            )
            gathered_tokens = sorted_tokens.new_empty(tokens_per_expert_group.sum(dim=0).cpu().item(), sorted_tokens.shape[1])
            input_split_sizes = tokens_per_ep_rank.cpu().numpy().tolist()
            dist.all_to_all(list(gathered_tokens.split(output_splits)), list(sorted_tokens.split(input_split_sizes)))
            tokens_per_expert_post_gather = tokens_per_expert_group.view(self.ep_size, self.experts_per_rank).sum(dim=0)
            gatherd_idxs = np.zeros(shape=(gathered_tokens.shape[0],), dtype=np.int32)
            s = 0
            for i, k in enumerate(tokens_per_expert_group.cpu().numpy()):
                gatherd_idxs[s : s + k] = i % self.experts_per_rank
                s += k
            gatherd_idxs = gatherd_idxs.argsort()
            sorted_tokens = gathered_tokens[gatherd_idxs]
            tokens_per_expert = tokens_per_expert_post_gather
        tokens_per_expert = tokens_per_expert.cpu().numpy()

        outputs = []
        start_idx = 0
        for i, num_tokens in enumerate(tokens_per_expert):
            end_idx = start_idx + num_tokens
            if num_tokens == 0:
                continue
            expert = self.experts[i + self.ep_rank * self.experts_per_rank]
            tokens_for_this_expert = sorted_tokens[start_idx:end_idx]
            expert_out = expert(tokens_for_this_expert)
            outputs.append(expert_out)
            start_idx = end_idx

        outs = torch.cat(outputs, dim=0) if len(outputs) else sorted_tokens.new_empty(0)
        if self.ep_size > 1:
            new_x = torch.empty_like(outs)
            new_x[gatherd_idxs] = outs
            gathered_tokens = new_x.new_empty(*sorted_tokens_shape)
            dist.all_to_all(list(gathered_tokens.split(input_split_sizes)), list(new_x.split(output_splits)))
            outs = gathered_tokens

        new_x = torch.empty_like(outs)
        new_x[idxs] = outs
        final_out = (
            new_x.view(*topk_ids.shape, -1)
            .type(topk_weight.dtype)
            .mul_(topk_weight.unsqueeze(dim=-1))
            .sum(dim=1)
            .type(new_x.dtype)
        )
        return final_out


# Copied from transformers.models.llama.modeling_llama.repeat_kv
def repeat_kv(hidden_states: torch.Tensor, n_rep: int) -> torch.Tensor:
    """
    This is the equivalent of torch.repeat_interleave(x, dim=1, repeats=n_rep). The hidden states go from (batch,
    num_key_value_heads, seqlen, head_dim) to (batch, num_attention_heads, seqlen, head_dim)
    """
    batch, num_key_value_heads, slen, head_dim = hidden_states.shape
    if n_rep == 1:
        return hidden_states
    hidden_states = hidden_states[:, :, None, :, :].expand(
        batch, num_key_value_heads, n_rep, slen, head_dim
    )
    return hidden_states.reshape(batch, num_key_value_heads * n_rep, slen, head_dim)


# Copied from transformers.models.llama.modeling_llama.LlamaAttention with Llama->Rize
class RizeAttention(nn.Module):
    """Multi-headed attention from 'Attention Is All You Need' paper"""

    def __init__(self, config: RizeConfig, layer_idx: Optional[int] = None):
        super().__init__()
        self.config = config
        self.layer_idx = layer_idx
        if layer_idx is None:
            logger.warning_once(
                f"Instantiating {self.__class__.__name__} without passing `layer_idx` is not recommended and will "
                "to errors during the forward call, if caching is used. Please make sure to provide a `layer_idx` "
                "when creating this class."
            )

        self.attention_dropout = config.attention_dropout
        self.hidden_size = config.hidden_size
        self.num_heads = config.num_attention_heads

        self.max_position_embeddings = config.max_position_embeddings
        self.rope_theta = config.rope_theta
        self.q_lora_rank = config.q_lora_rank
        self.qk_rope_head_dim = config.qk_rope_head_dim
        self.kv_lora_rank = config.kv_lora_rank
        self.v_head_dim = config.v_head_dim
        self.qk_nope_head_dim = config.qk_nope_head_dim
        self.q_head_dim = config.qk_nope_head_dim + config.qk_rope_head_dim

        self.is_causal = True

        if self.q_lora_rank is None:
            self.q_proj = nn.Linear(
                self.hidden_size, self.num_heads * self.q_head_dim, bias=False
            )
        else:
            self.q_a_proj = nn.Linear(
                self.hidden_size, config.q_lora_rank, bias=config.attention_bias
            )
            self.q_a_layernorm = RizeRMSNorm(config.q_lora_rank)
            self.q_b_proj = nn.Linear(
                config.q_lora_rank, self.num_heads * self.q_head_dim, bias=False
            )

        self.kv_a_proj_with_mqa = nn.Linear(
            self.hidden_size,
            config.kv_lora_rank + config.qk_rope_head_dim,
            bias=config.attention_bias,
        )
        self.kv_a_layernorm = RizeRMSNorm(config.kv_lora_rank)
        self.kv_b_proj = nn.Linear(
            config.kv_lora_rank,
            self.num_heads
            * (self.q_head_dim - self.qk_rope_head_dim + self.v_head_dim),
            bias=False,
        )

        self.o_proj = nn.Linear(
            self.num_heads * self.v_head_dim,
            self.hidden_size,
            bias=config.attention_bias,
        )
        self._init_rope()

        self.softmax_scale = self.q_head_dim ** (-0.5)
        if self.config.rope_scaling is not None:
            mscale_all_dim = self.config.rope_scaling.get("mscale_all_dim", 0)
            scaling_factor = self.config.rope_scaling["factor"]
            if mscale_all_dim:
                mscale = yarn_get_mscale(scaling_factor, mscale_all_dim)
                self.softmax_scale = self.softmax_scale * mscale * mscale

    def _init_rope(self):
        if self.config.rope_scaling is None:
            self.rotary_emb = RizeRotaryEmbedding(
                self.qk_rope_head_dim,
                max_position_embeddings=self.max_position_embeddings,
                base=self.rope_theta,
            )
        else:
            scaling_type = self.config.rope_scaling["type"]
            scaling_factor = self.config.rope_scaling["factor"]
            if scaling_type == "linear":
                self.rotary_emb = RizeLinearScalingRotaryEmbedding(
                    self.qk_rope_head_dim,
                    max_position_embeddings=self.max_position_embeddings,
                    scaling_factor=scaling_factor,
                    base=self.rope_theta,
                )
            elif scaling_type == "dynamic":
                self.rotary_emb = RizeDynamicNTKScalingRotaryEmbedding(
                    self.qk_rope_head_dim,
                    max_position_embeddings=self.max_position_embeddings,
                    scaling_factor=scaling_factor,
                    base=self.rope_theta,
                )
            elif scaling_type == "yarn":
                kwargs = {
                    key: self.config.rope_scaling[key]
                    for key in [
                        "original_max_position_embeddings",
                        "beta_fast",
                        "beta_slow",
                        "mscale",
                        "mscale_all_dim",
                    ]
                    if key in self.config.rope_scaling
                }
                self.rotary_emb = RizeYarnRotaryEmbedding(
                    self.qk_rope_head_dim,
                    max_position_embeddings=self.max_position_embeddings,
                    scaling_factor=scaling_factor,
                    base=self.rope_theta,
                    **kwargs,
                )
            else:
                raise ValueError(f"Unknown RoPE scaling type {scaling_type}")

    def _shape(self, tensor: torch.Tensor, seq_len: int, bsz: int):
        return (
            tensor.view(bsz, seq_len, self.num_heads, self.v_head_dim)
            .transpose(1, 2)
            .contiguous()
        )

    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: bool = False,
        use_cache: bool = False,
        **kwargs,
    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
        if "padding_mask" in kwargs:
            warnings.warn(
                "Passing `padding_mask` is deprecated and will be removed in v4.37. Please make sure use `attention_mask` instead.`"
            )

        bsz, q_len, _ = hidden_states.size()
        # q_proj の重み dtype（BF16 等）に入力を合わせる
        target_dtype = (self.q_proj.weight.dtype if self.q_lora_rank is None
                        else self.q_a_proj.weight.dtype)
        hidden_states = hidden_states.to(target_dtype)

        if self.q_lora_rank is None:
            q = self.q_proj(hidden_states)
        else:
            q = self.q_b_proj(self.q_a_layernorm(self.q_a_proj(hidden_states)))
        q = q.view(bsz, q_len, self.num_heads, self.q_head_dim).transpose(1, 2)
        q_nope, q_pe = torch.split(
            q, [self.qk_nope_head_dim, self.qk_rope_head_dim], dim=-1
        )

        compressed_kv = self.kv_a_proj_with_mqa(hidden_states)
        compressed_kv, k_pe = torch.split(
            compressed_kv, [self.kv_lora_rank, self.qk_rope_head_dim], dim=-1
        )
        k_pe = k_pe.view(bsz, q_len, 1, self.qk_rope_head_dim).transpose(1, 2)
        kv = (
            self.kv_b_proj(self.kv_a_layernorm(compressed_kv))
            .view(bsz, q_len, self.num_heads, self.qk_nope_head_dim + self.v_head_dim)
            .transpose(1, 2)
        )

        k_nope, value_states = torch.split(
            kv, [self.qk_nope_head_dim, self.v_head_dim], dim=-1
        )

        kv_seq_len = value_states.shape[-2]
        if past_key_value is not None:
            if self.layer_idx is None:
                raise ValueError("... initialize with layer_idx ...")
            if hasattr(past_key_value, "get_usable_length"):
                kv_seq_len += past_key_value.get_usable_length(kv_seq_len, self.layer_idx)
            elif hasattr(past_key_value, "get_seq_length"):
                kv_seq_len += past_key_value.get_seq_length()

        cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)

        q_pe, k_pe = apply_rotary_pos_emb(q_pe, k_pe, cos, sin, position_ids)

        query_states = k_pe.new_empty(bsz, self.num_heads, q_len, self.q_head_dim)
        query_states[:, :, :, : self.qk_nope_head_dim] = q_nope
        query_states[:, :, :, self.qk_nope_head_dim :] = q_pe

        key_states = k_pe.new_empty(bsz, self.num_heads, q_len, self.q_head_dim)
        key_states[:, :, :, : self.qk_nope_head_dim] = k_nope
        key_states[:, :, :, self.qk_nope_head_dim :] = k_pe
        if past_key_value is not None:
            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
            key_states, value_states = past_key_value.update(
                key_states, value_states, self.layer_idx, cache_kwargs
            )

        attn_weights = (
            torch.matmul(query_states, key_states.transpose(2, 3)) * self.softmax_scale
        )

        if attn_weights.size() != (bsz, self.num_heads, q_len, kv_seq_len):
            raise ValueError(
                f"Attention weights should be of size {(bsz, self.num_heads, q_len, kv_seq_len)}, but is"
                f" {attn_weights.size()}"
            )
        assert attention_mask is not None
        if attention_mask is not None:
            if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
                raise ValueError(
                    f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, but is {attention_mask.size()}"
                )
            attn_weights = attn_weights + attention_mask

        # upcast attention to fp32
        attn_weights = nn.functional.softmax(
            attn_weights, dim=-1, dtype=torch.float32
        ).to(query_states.dtype)
        attn_weights = nn.functional.dropout(
            attn_weights, p=self.attention_dropout, training=self.training
        )
        attn_output = torch.matmul(attn_weights, value_states)

        if attn_output.size() != (bsz, self.num_heads, q_len, self.v_head_dim):
            raise ValueError(
                f"`attn_output` should be of size {(bsz, self.num_heads, q_len, self.v_head_dim)}, but is"
                f" {attn_output.size()}"
            )

        attn_output = attn_output.transpose(1, 2).contiguous()

        attn_output = attn_output.reshape(bsz, q_len, self.num_heads * self.v_head_dim)

        attn_output = self.o_proj(attn_output)

        if not output_attentions:
            attn_weights = None

        return attn_output, attn_weights, past_key_value


# Copied from transformers.models.llama.modeling_llama.LlamaFlashAttention2 with Llama->Rize
class RizeFlashAttention2(RizeAttention):
    """
    Rize flash attention module. This module inherits from `RizeAttention` as the weights of the module stays
    untouched. The only required change would be on the forward pass where it needs to correctly call the public API of
    flash attention and deal with padding tokens in case the input contains any of them.
    """

    def __init__(self, *args, **kwargs):
        super().__init__(*args, **kwargs)

        # TODO: Should be removed once Flash Attention for RoCm is bumped to 2.1.
        # flash_attn<2.1 generates top-left aligned causal mask, while what is needed here is bottom-right alignement, that was made default for flash_attn>=2.1. This attribute is used to handle this difference. Reference: https://github.com/Dao-AILab/flash-attention/releases/tag/v2.1.0.
        # Beware that with flash_attn<2.1, using q_seqlen != k_seqlen (except for the case q_seqlen == 1) produces a wrong mask (top-left).
        self._flash_attn_uses_top_left_mask = not is_flash_attn_greater_or_equal_2_10()

    def forward(
        self,
        hidden_states: torch.Tensor,
        attention_mask: Optional[torch.LongTensor] = None,
        position_ids: Optional[torch.LongTensor] = None,
        past_key_value: Optional[Cache] = None,
        output_attentions: bool = False,
        use_cache: bool = False,
        **kwargs,
    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
        # RizeFlashAttention2 attention does not support output_attentions
        if "padding_mask" in kwargs:
            warnings.warn(
                "Passing `padding_mask` is deprecated and will be removed in v4.37. Please make sure use `attention_mask` instead.`"
            )

            # overwrite attention_mask with padding_mask
            attention_mask = kwargs.pop("padding_mask")

        sequence_ids = kwargs.pop("sequence_ids", None)
        output_attentions = False

        bsz, q_len, _ = hidden_states.size()

        if self.q_lora_rank is None:
            q = self.q_proj(hidden_states)
        else:
            q = self.q_b_proj(self.q_a_layernorm(self.q_a_proj(hidden_states)))
        q = q.view(bsz, q_len, self.num_heads, self.q_head_dim).transpose(1, 2)
        q_nope, q_pe = torch.split(
            q, [self.qk_nope_head_dim, self.qk_rope_head_dim], dim=-1
        )

        # Flash attention requires the input to have the shape
        # batch_size x seq_length x head_dim x hidden_dim
        # therefore we just need to keep the original shape
        compressed_kv = self.kv_a_proj_with_mqa(hidden_states)
        compressed_kv, k_pe = torch.split(
            compressed_kv, [self.kv_lora_rank, self.qk_rope_head_dim], dim=-1
        )
        k_pe = k_pe.view(bsz, q_len, 1, self.qk_rope_head_dim).transpose(1, 2)
        kv = (
            self.kv_b_proj(self.kv_a_layernorm(compressed_kv))
            .view(bsz, q_len, self.num_heads, self.qk_nope_head_dim + self.v_head_dim)
            .transpose(1, 2)
        )

        k_nope, value_states = torch.split(
            kv, [self.qk_nope_head_dim, self.v_head_dim], dim=-1
        )
        kv_seq_len = value_states.shape[-2]
        if past_key_value is not None:
            kv_seq_len += past_key_value.get_usable_length(kv_seq_len, self.layer_idx)

        cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
        q_pe, k_pe = apply_rotary_pos_emb(q_pe, k_pe, cos, sin, position_ids)

        query_states = k_pe.new_empty(bsz, self.num_heads, q_len, self.q_head_dim)
        query_states[:, :, :, : self.qk_nope_head_dim] = q_nope
        query_states[:, :, :, self.qk_nope_head_dim :] = q_pe

        key_states = k_pe.new_empty(bsz, self.num_heads, q_len, self.q_head_dim)
        key_states[:, :, :, : self.qk_nope_head_dim] = k_nope
        key_states[:, :, :, self.qk_nope_head_dim :] = k_pe

        if self.q_head_dim != self.v_head_dim:
            value_states = F.pad(value_states, [0, self.q_head_dim - self.v_head_dim])

        if past_key_value is not None:
            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
            key_states, value_states = past_key_value.update(
                key_states, value_states, self.layer_idx, cache_kwargs
            )

        # TODO: These transpose are quite inefficient but Flash Attention requires the layout [batch_size, sequence_length, num_heads, head_dim]. We would need to refactor the KV cache
        # to be able to avoid many of these transpose/reshape/view.
        query_states = query_states.transpose(1, 2)
        key_states = key_states.transpose(1, 2)
        value_states = value_states.transpose(1, 2)

        dropout_rate = self.attention_dropout if self.training else 0.0

        # In PEFT, usually we cast the layer norms in float32 for training stability reasons
        # therefore the input hidden states gets silently casted in float32. Hence, we need
        # cast them back in the correct dtype just to be sure everything works as expected.
        # This might slowdown training & inference so it is recommended to not cast the LayerNorms
        # in fp32. (RizeRMSNorm handles it correctly)

        input_dtype = query_states.dtype
        if input_dtype == torch.float32:
            # Handle the case where the model is quantized
            if hasattr(self.config, "_pre_quantization_dtype"):
                target_dtype = self.config._pre_quantization_dtype
            elif torch.is_autocast_enabled():
                target_dtype = torch.get_autocast_gpu_dtype()
            else:
                target_dtype = (
                    self.q_proj.weight.dtype
                    if self.q_lora_rank is None
                    else self.q_a_proj.weight.dtype
                )

            logger.warning_once(
                f"The input hidden states seems to be silently casted in float32, this might be related to"
                f" the fact you have upcasted embedding or layer norm layers in float32. We will cast back the input in"
                f" {target_dtype}."
            )

            query_states = query_states.to(target_dtype)
            key_states = key_states.to(target_dtype)
            value_states = value_states.to(target_dtype)

        attn_output = self._flash_attention_forward(
            query_states,
            key_states,
            value_states,
            attention_mask,
            q_len,
            dropout=dropout_rate,
            softmax_scale=self.softmax_scale,
            sequence_ids=sequence_ids,
        )
        if self.q_head_dim != self.v_head_dim:
            attn_output = attn_output[:, :, :, : self.v_head_dim]

        attn_output = attn_output.reshape(
            bsz, q_len, self.num_heads * self.v_head_dim
        ).contiguous()
        attn_output = self.o_proj(attn_output)

        if not output_attentions:
            attn_weights = None

        return attn_output, attn_weights, past_key_value

    def _flash_attention_forward(
        self,
        query_states,
        key_states,
        value_states,
        attention_mask,
        query_length,
        dropout=0.0,
        softmax_scale=None,
        sequence_ids: Optional[torch.Tensor] = None,
    ):
        """
        Calls the forward method of Flash Attention - if the input hidden states contain at least one padding token
        first unpad the input, then computes the attention scores and pad the final attention scores.

        Args:
            query_states (`torch.Tensor`):
                Input query states to be passed to Flash Attention API
            key_states (`torch.Tensor`):
                Input key states to be passed to Flash Attention API
            value_states (`torch.Tensor`):
                Input value states to be passed to Flash Attention API
            attention_mask (`torch.Tensor`):
                The padding mask - corresponds to a tensor of size `(batch_size, seq_len)` where 0 stands for the
                position of padding tokens and 1 for the position of non-padding tokens.
            dropout (`int`, *optional*):
                Attention dropout
            softmax_scale (`float`, *optional*):
                The scaling of QK^T before applying softmax. Default to 1 / sqrt(head_dim)
        """
        if not self._flash_attn_uses_top_left_mask:
            causal = self.is_causal
        else:
            # TODO: Remove the `query_length != 1` check once Flash Attention for RoCm is bumped to 2.1. For details, please see the comment in RizeFlashAttention2 __init__.
            causal = self.is_causal and query_length != 1

        # Packed sample ids take precedence: they define exact block-diagonal causal attention.
        if sequence_ids is not None:
            batch_size = query_states.shape[0]
            (
                query_states,
                key_states,
                value_states,
                indices_q,
                cu_seq_lens,
                max_seq_lens,
            ) = self._upad_input_with_sequence_ids(
                query_states,
                key_states,
                value_states,
                attention_mask,
                sequence_ids,
                query_length,
            )

            cu_seqlens_q, cu_seqlens_k = cu_seq_lens
            cu_seqlens_q = cu_seqlens_q.to(dtype=torch.int32, device=query_states.device).contiguous()
            cu_seqlens_k = cu_seqlens_k.to(dtype=torch.int32, device=key_states.device).contiguous()
            max_seqlen_in_batch_q, max_seqlen_in_batch_k = max_seq_lens

            attn_output_unpad = flash_attn_varlen_func(
                query_states,
                key_states,
                value_states,
                cu_seqlens_q=cu_seqlens_q,
                cu_seqlens_k=cu_seqlens_k,
                max_seqlen_q=max_seqlen_in_batch_q,
                max_seqlen_k=max_seqlen_in_batch_k,
                dropout_p=dropout,
                softmax_scale=softmax_scale,
                causal=causal,
            )

            attn_output = pad_input(attn_output_unpad, indices_q, batch_size, query_length)
        elif attention_mask is not None:
            batch_size = query_states.shape[0]
            (
                query_states,
                key_states,
                value_states,
                indices_q,
                cu_seq_lens,
                max_seq_lens,
            ) = self._upad_input(
                query_states, key_states, value_states, attention_mask, query_length
            )

            cu_seqlens_q, cu_seqlens_k = cu_seq_lens
            cu_seqlens_q = cu_seqlens_q.to(dtype=torch.int32, device=query_states.device).contiguous()
            cu_seqlens_k = cu_seqlens_k.to(dtype=torch.int32, device=key_states.device).contiguous()
            max_seqlen_in_batch_q, max_seqlen_in_batch_k = max_seq_lens

            attn_output_unpad = flash_attn_varlen_func(
                query_states,
                key_states,
                value_states,
                cu_seqlens_q=cu_seqlens_q,
                cu_seqlens_k=cu_seqlens_k,
                max_seqlen_q=max_seqlen_in_batch_q,
                max_seqlen_k=max_seqlen_in_batch_k,
                dropout_p=dropout,
                softmax_scale=softmax_scale,
                causal=causal,
            )

            attn_output = pad_input(
                attn_output_unpad, indices_q, batch_size, query_length
            )
        else:
            attn_output = flash_attn_func(
                query_states,
                key_states,
                value_states,
                dropout,
                softmax_scale=softmax_scale,
                causal=causal,
            )

        return attn_output

    def _upad_input_with_sequence_ids(
        self,
        query_layer,
        key_layer,
        value_layer,
        attention_mask,
        sequence_ids,
        query_length,
    ):
        batch_size, kv_seq_len, num_key_value_heads, head_dim = key_layer.shape
        if query_length != kv_seq_len:
            raise NotImplementedError(
                "packed sequence_ids with FlashAttention2 currently require q_len == kv_seq_len"
            )

        indices_k, cu_seqlens_k, max_seqlen_in_batch_k = _get_unpad_data_from_sequence_ids(
            sequence_ids,
            attention_mask=attention_mask,
        )

        key_layer = index_first_axis(
            key_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim),
            indices_k,
        )
        value_layer = index_first_axis(
            value_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim),
            indices_k,
        )
        query_layer = index_first_axis(
            query_layer.reshape(batch_size * kv_seq_len, self.num_heads, head_dim),
            indices_k,
        )

        return (
            query_layer,
            key_layer,
            value_layer,
            indices_k,
            (cu_seqlens_k, cu_seqlens_k),
            (max_seqlen_in_batch_k, max_seqlen_in_batch_k),
        )

    def _upad_input(
        self, query_layer, key_layer, value_layer, attention_mask, query_length
    ):
        indices_k, cu_seqlens_k, max_seqlen_in_batch_k = _get_unpad_data(attention_mask)
        batch_size, kv_seq_len, num_key_value_heads, head_dim = key_layer.shape

        key_layer = index_first_axis(
            key_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim),
            indices_k,
        )
        value_layer = index_first_axis(
            value_layer.reshape(batch_size * kv_seq_len, num_key_value_heads, head_dim),
            indices_k,
        )
        if query_length == kv_seq_len:
            query_layer = index_first_axis(
                query_layer.reshape(batch_size * kv_seq_len, self.num_heads, head_dim),
                indices_k,
            )
            cu_seqlens_q = cu_seqlens_k
            max_seqlen_in_batch_q = max_seqlen_in_batch_k
            indices_q = indices_k
        elif query_length == 1:
            max_seqlen_in_batch_q = 1
            cu_seqlens_q = torch.arange(
                batch_size + 1, dtype=torch.int32, device=query_layer.device
            )  # There is a memcpy here, that is very bad.
            indices_q = cu_seqlens_q[:-1]
            query_layer = query_layer.squeeze(1)
        else:
            # The -q_len: slice assumes left padding.
            attention_mask = attention_mask[:, -query_length:]
            query_layer, indices_q, cu_seqlens_q, max_seqlen_in_batch_q = unpad_input(
                query_layer, attention_mask
            )

        return (
            query_layer,
            key_layer,
            value_layer,
            indices_q,
            (cu_seqlens_q, cu_seqlens_k),
            (max_seqlen_in_batch_q, max_seqlen_in_batch_k),
        )


ATTENTION_CLASSES = {
    "eager": RizeAttention,
    "flash_attention_2": RizeFlashAttention2,
}


class RizeHybridCache(DynamicCache):
    """Hybrid cache for Kimi-Linear style attention (KDA + full attention).

    `DynamicCache` only stores key/value tensors. Kimi-Linear style Delta Attention (KDA)
    additionally requires per-layer *short convolution* states and recurrent states.

    This cache:
      - keeps the standard key/value cache for *full-attention* layers via `DynamicCache`
      - adds `conv_states` and `recurrent_states` for *KDA* layers
      - fixes `seen_tokens` tracking when early layers are KDA by tracking tokens from
        the first full-attention layer.
    """

    def __init__(self, config: RizeConfig):
        super().__init__()
        self.config = config
        self.num_layers = int(config.num_hidden_layers)

        # Identify which layers are KDA vs full attention.
        self.layer_types: List[str] = []
        for i in range(self.num_layers):
            self.layer_types.append("linear_attention" if config.is_kda_layer(i) else "full_attention")

        self.transformer_layers: List[int] = [
            i for i in range(self.num_layers) if self.layer_types[i] == "full_attention"
        ]

        self.last_linear_layer: int = -1
        for i in range(self.num_layers - 1, -1, -1):
            if self.layer_types[i] == "linear_attention":
                self.last_linear_layer = i
                break

        # Track tokens from the first full-attention layer (matches transformers' expectation
        # that `seen_tokens` advances as decoding proceeds).
        if len(self.transformer_layers) > 0:
            self._token_tracker_layer = int(self.transformer_layers[0])
        elif self.last_linear_layer != -1:
            self._token_tracker_layer = int(self.last_linear_layer)
        else:
            self._token_tracker_layer = 0

        # Extra states for KDA layers.
        self.conv_states: List[Optional[Tuple[torch.Tensor, torch.Tensor, torch.Tensor]]] = [
            None for _ in range(self.num_layers)
        ]
        self.recurrent_states: List[Optional[torch.Tensor]] = [
            None for _ in range(self.num_layers)
        ]

    @property
    def has_previous_state(self) -> bool:
        """True if at least one KDA layer has an initialized conv state."""
        if self.last_linear_layer == -1:
            return False
        return self.conv_states[self.last_linear_layer] is not None

    def _pick_seq_len_layer(self, layer_idx: Optional[int] = None) -> int:
        # Prefer a valid full-attention layer (KV cache) for seq-length queries.
        if layer_idx is None:
            layer_idx = self._token_tracker_layer
        layer_idx = int(layer_idx)
        if len(getattr(self, "transformer_layers", [])) > 0 and layer_idx not in self.transformer_layers:
            layer_idx = int(self.transformer_layers[0])
        return layer_idx

    def get_seq_length(self, layer_idx: Optional[int] = None) -> int:
        layer_idx = self._pick_seq_len_layer(layer_idx)
        # `DynamicCache` stores KV caches in `key_cache`.
        try:
            if len(self.key_cache) > layer_idx and self.key_cache[layer_idx] is not None:
                return int(self.key_cache[layer_idx].shape[-2])
        except Exception:
            pass

        # Fallback: find the first non-empty KV cache.
        for i in getattr(self, "transformer_layers", range(len(getattr(self, "key_cache", [])))):
            try:
                if len(self.key_cache) > i and self.key_cache[i] is not None:
                    return int(self.key_cache[i].shape[-2])
            except Exception:
                continue
        return 0

    def get_usable_length(self, new_seq_len: int, layer_idx: Optional[int] = None) -> int:
        # For our purposes, usable length == currently cached length.
        _ = new_seq_len
        return self.get_seq_length(layer_idx)

    def update(self, key_states, value_states, layer_idx, cache_kwargs=None):
        # Let DynamicCache handle KV concatenation.
        key_states, value_states = super().update(key_states, value_states, layer_idx, cache_kwargs)

        # Sync `seen_tokens` from the designated token-tracker layer (important when early layers are KDA).
        try:
            if int(layer_idx) == self._token_tracker_layer:
                self.seen_tokens = int(key_states.shape[-2])
        except Exception:
            pass
        return key_states, value_states

    def reorder_cache(self, beam_idx: torch.LongTensor):
        # Reorder KV cache first.
        out = super().reorder_cache(beam_idx)

        # Reorder KDA-specific states.
        for i in range(self.num_layers):
            cs = self.conv_states[i]
            if cs is not None:
                self.conv_states[i] = tuple(
                    None if c is None else c.index_select(0, beam_idx)
                    for c in cs
                )
            rs = self.recurrent_states[i]
            if rs is not None:
                self.recurrent_states[i] = rs.index_select(0, beam_idx)
        return out

    @classmethod
    def from_dynamic_cache(cls, base_cache: DynamicCache, config: RizeConfig):
        """Upgrade an existing `DynamicCache` into a `RizeHybridCache`."""
        new = cls(config)
        # Best-effort state transfer (works for transformers' DynamicCache).
        for attr in ("key_cache", "value_cache", "seen_tokens"):
            if hasattr(base_cache, attr):
                setattr(new, attr, getattr(base_cache, attr))
        return new


class RizeKimiDeltaAttention(nn.Module):
    """Kimi-Linear style Delta Attention (KDA).

    This implementation mirrors Moonshot's Kimi-Linear reference and relies on
    `fla-core` for efficient kernels.
    """

    def __init__(self, config: RizeConfig, layer_idx: int):
        super().__init__()
        if not getattr(config, "is_linear_attn", False) or config.linear_attn_config is None:
            raise ValueError("RizeKimiDeltaAttention requires config.linear_attn_config")
        if not _FLA_AVAILABLE:
            raise ImportError(
                "Kimi-Linear (Delta Attention / KDA) requires 'fla-core'. "
                "Install with: pip install -U fla-core"
            )

        self.config = config
        self.mode = "chunk"
        self.hidden_size = config.hidden_size

        la_cfg = config.linear_attn_config
        self.conv_size = int(la_cfg["short_conv_kernel_size"])
        self.head_dim = int(la_cfg["head_dim"])
        self.num_heads = int(la_cfg["num_heads"])
        self.head_k_dim = self.head_dim
        self.num_k_heads = self.num_heads
        self.layer_idx = int(layer_idx)

        assert self.mode in ["chunk", "fused_recurrent"], f"Not supported mode `{self.mode}`."

        projection_k_size = self.head_k_dim * self.num_k_heads
        projection_size = self.head_dim * self.num_heads

        self.q_proj = nn.Linear(self.hidden_size, projection_k_size, bias=False)
        self.k_proj = nn.Linear(self.hidden_size, projection_k_size, bias=False)
        self.v_proj = nn.Linear(self.hidden_size, projection_size, bias=False)

        self.q_conv1d = ShortConvolution(hidden_size=projection_k_size, kernel_size=self.conv_size, activation="silu")
        self.k_conv1d = ShortConvolution(hidden_size=projection_k_size, kernel_size=self.conv_size, activation="silu")
        self.v_conv1d = ShortConvolution(hidden_size=projection_size, kernel_size=self.conv_size, activation="silu")

        # Gate / dynamics parameters
        self.A_log = nn.Parameter(
            torch.log(torch.empty(self.num_heads, dtype=torch.float32).uniform_(1, 16)).view(1, 1, -1, 1)
        )
        self.f_a_proj = nn.Linear(self.hidden_size, self.head_dim, bias=False)
        self.f_b_proj = nn.Linear(self.head_dim, projection_size, bias=False)
        # NOTE: `torch.empty` would leave this uninitialized when training from scratch.
        # Kimi checkpoints will override it on load, but for stability keep a deterministic init.
        self.dt_bias = nn.Parameter(torch.zeros(projection_size, dtype=torch.float32))
        self.b_proj = nn.Linear(self.hidden_size, self.num_heads, bias=False)

        self.g_a_proj = nn.Linear(self.hidden_size, self.head_dim, bias=False)
        self.g_b_proj = nn.Linear(self.head_dim, projection_size, bias=False)

        self.o_norm = FusedRMSNormGated(self.head_dim, eps=config.rms_norm_eps, activation="sigmoid")
        self.o_proj = nn.Linear(projection_size, self.hidden_size, bias=False)

    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,
        **kwargs,
    ):
        _ = position_ids
        _ = output_attentions

        sequence_ids = kwargs.pop("sequence_ids", None)

        # KDA only supports 2D padding masks (0/1). If a 4D mask is provided, try `padding_mask`.
        if attention_mask is not None and attention_mask.dim() != 2:
            attention_mask = kwargs.get("padding_mask", None)
        if attention_mask is not None and attention_mask.dim() != 2:
            raise ValueError(
                "KDA expects a 2D attention_mask of shape [batch, seq] with 0=pad, 1=token."
            )

        cache_params = past_key_value
        use_cache = bool(use_cache) and cache_params is not None

        batch_size, q_len, _hs = hidden_states.shape
        if sequence_ids is not None:
            sequence_ids = sequence_ids.to(device=hidden_states.device, dtype=torch.long)
            if sequence_ids.shape != (batch_size, q_len):
                raise ValueError(
                    f"sequence_ids should have shape {(batch_size, q_len)}, got {tuple(sequence_ids.shape)}"
                )
        mode = "fused_recurrent" if q_len <= 64 else self.mode
        if self.training:
            assert mode == "chunk", "Only chunk mode is supported in training."

        cu_seqlens = kwargs.get("cu_seqlens", None)
        indices = None

        # Unpad / segment for variable-length batches.
        # When packed-sample `sequence_ids` are provided, they take precedence so that
        # recurrent KDA state is reset exactly at each packed-sample boundary.
        if sequence_ids is not None:
            active_mask = attention_mask[:, -q_len:] if attention_mask is not None else None
            indices, cu_seqlens, _ = _get_unpad_data_from_sequence_ids(
                sequence_ids[:, -q_len:],
                attention_mask=active_mask,
            )
            flat = hidden_states.reshape(batch_size * q_len, -1)
            hidden_states = index_first_axis(flat, indices).unsqueeze(0)
        elif attention_mask is not None:
            indices, cu_seqlens, _ = _get_unpad_data(attention_mask[:, -q_len:])
            flat = hidden_states.reshape(batch_size * q_len, -1)
            hidden_states = index_first_axis(flat, indices).unsqueeze(0)

        conv_state_q = conv_state_k = conv_state_v = None
        recurrent_state = None
        if cache_params is not None:
            if not hasattr(cache_params, "conv_states") or not hasattr(cache_params, "recurrent_states"):
                raise ValueError("past_key_value must be a RizeHybridCache when using KDA")
            if cache_params.conv_states[self.layer_idx] is not None:
                conv_state_q, conv_state_k, conv_state_v = cache_params.conv_states[self.layer_idx]
            recurrent_state = cache_params.recurrent_states[self.layer_idx]

        q, conv_state_q = self.q_conv1d(
            x=self.q_proj(hidden_states),
            cache=conv_state_q,
            output_final_state=use_cache,
            cu_seqlens=cu_seqlens,
        )
        k, conv_state_k = self.k_conv1d(
            x=self.k_proj(hidden_states),
            cache=conv_state_k,
            output_final_state=use_cache,
            cu_seqlens=cu_seqlens,
        )
        v, conv_state_v = self.v_conv1d(
            x=self.v_proj(hidden_states),
            cache=conv_state_v,
            output_final_state=use_cache,
            cu_seqlens=cu_seqlens,
        )

        g = self.f_b_proj(self.f_a_proj(hidden_states))  # [B, T, H*D]
        g = g.view(*g.shape[:-1], self.num_heads, self.head_dim)  # [B, T, H, D]
        g = fused_kda_gate(g, self.A_log, dt_bias=self.dt_bias, output_dtype=g.dtype)

        beta = self.b_proj(hidden_states).float().sigmoid()

        # Reshape to [..., h, d]
        q = q.view(*q.shape[:-1], self.num_heads, self.head_k_dim)
        k = k.view(*k.shape[:-1], self.num_heads, self.head_k_dim)
        v = v.view(*v.shape[:-1], self.num_heads, self.head_dim)

        if mode == "chunk":
            o, recurrent_state = chunk_kda(
                q=q,
                k=k,
                v=v,
                g=g,
                beta=beta,
                initial_state=recurrent_state,
                output_final_state=True,
                use_qk_l2norm_in_kernel=True,
                cu_seqlens=cu_seqlens,
            )
        else:
            o, recurrent_state = fused_recurrent_kda(
                q=q,
                k=k,
                v=v,
                g=g,
                beta=beta,
                initial_state=recurrent_state,
                output_final_state=True,
                use_qk_l2norm_in_kernel=True,
                cu_seqlens=cu_seqlens,
            )

        if cache_params is not None and use_cache:
            cache_params.recurrent_states[self.layer_idx] = recurrent_state
            cache_params.conv_states[self.layer_idx] = (conv_state_q, conv_state_k, conv_state_v)

        g2 = self.g_b_proj(self.g_a_proj(hidden_states))
        g2 = g2.view(*g2.shape[:-1], self.num_heads, self.head_dim)
        o = self.o_norm(o, g2)
        o = o.reshape(o.shape[0], o.shape[1], -1)
        o = self.o_proj(o)

        if indices is not None:
            o = pad_input(o.squeeze(0), indices, batch_size, q_len)

        return o, None, cache_params


class RizeDecoderLayer(nn.Module):
    def __init__(self, config: RizeConfig, layer_idx: int):
        super().__init__()
        self.hidden_size = config.hidden_size

        # Kimi-Linear style: some layers use KDA (linear attention) instead of full attention.
        self.is_linear_attn = bool(getattr(config, "is_linear_attn", False) and config.is_kda_layer(layer_idx))
        if self.is_linear_attn:
            self.self_attn = RizeKimiDeltaAttention(config=config, layer_idx=layer_idx)
        else:
            self.self_attn = ATTENTION_CLASSES[config._attn_implementation](
                config=config, layer_idx=layer_idx
            )
        self.mlp = (
            RizeMoE(config)
            if (
                config.n_routed_experts is not None
                and layer_idx >= config.first_k_dense_replace
                and layer_idx % config.moe_layer_freq == 0
            )
            else RizeMLP(config)
        )
        self.input_layernorm = RizeRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
        self.post_attention_layernorm = RizeRMSNorm(config.hidden_size, eps=config.rms_norm_eps)

    def forward(
        self,
        hidden_states: torch.Tensor,
        attention_mask: Optional[torch.Tensor] = None,
        position_ids: Optional[torch.LongTensor] = None,
        past_key_value: Optional[Tuple[torch.Tensor]] = None,
        output_attentions: Optional[bool] = False,
        use_cache: Optional[bool] = False,
        sequence_ids: Optional[torch.LongTensor] = None,
        router_token_mask: Optional[torch.Tensor] = None,
        **kwargs,
    ):
        if "padding_mask" in kwargs:
            warnings.warn("Passing `padding_mask` is deprecated and will be removed in v4.37. Please make sure use `attention_mask` instead.`")
        residual = hidden_states
        hidden_states = self.input_layernorm(hidden_states)

        hidden_states, self_attn_weights, present_key_value = 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,
            sequence_ids=sequence_ids,
            **kwargs,
        )
        hidden_states = residual + hidden_states

        residual = hidden_states
        hidden_states = self.post_attention_layernorm(hidden_states)
        if isinstance(self.mlp, RizeMoE):
            mlp_out = self.mlp(hidden_states, token_mask=router_token_mask)
        else:
            mlp_out = self.mlp(hidden_states)

        router_aux_loss = None
        if isinstance(mlp_out, tuple):
            hidden_states, router_aux_loss = mlp_out
        else:
            hidden_states = mlp_out
        hidden_states = residual + hidden_states

        outputs = (hidden_states,)
        if output_attentions:
            outputs += (self_attn_weights,)
        if use_cache:
            outputs += (present_key_value,)
        if router_aux_loss is not None:
            outputs += (router_aux_loss,)
        return outputs


Rize_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 ([`RizeConfig`]):
            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 Rize Model outputting raw hidden-states without any specific head on top.",
    Rize_START_DOCSTRING,
)
class RizePreTrainedModel(PreTrainedModel):
    config_class = RizeConfig
    base_model_prefix = "model"
    supports_gradient_checkpointing = True
    _no_split_modules = ["RizeDecoderLayer"]
    _skip_keys_device_placement = "past_key_values"
    _supports_flash_attn_2 = True
    _supports_cache_class = True

    def _init_weights(self, module):
        std = self.config.initializer_range
        if isinstance(module, nn.Linear):
            module.weight.data.normal_(mean=0.0, std=std)
            if module.bias is not None:
                module.bias.data.zero_()
        elif isinstance(module, nn.Embedding):
            module.weight.data.normal_(mean=0.0, std=std)
            if module.padding_idx is not None:
                module.weight.data[module.padding_idx].zero_()


Rize_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;
            - 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.
"""


@add_start_docstrings(
    "The bare Rize Model outputting raw hidden-states without any specific head on top.",
    Rize_START_DOCSTRING,
)
class RizeModel(RizePreTrainedModel):
    """
    Transformer decoder consisting of *config.num_hidden_layers* layers. Each layer is a [`RizeDecoderLayer`]

    Args:
        config: RizeConfig
    """

    def __init__(self, config: RizeConfig):
        super().__init__(config)
        self.padding_idx = config.pad_token_id
        self.vocab_size = config.vocab_size

        self.embed_tokens = nn.Embedding(
            config.vocab_size, config.hidden_size, self.padding_idx
        )
        self.layers = nn.ModuleList(
            [
                RizeDecoderLayer(config, layer_idx)
                for layer_idx in range(config.num_hidden_layers)
            ]
        )
        self._use_flash_attention_2 = config._attn_implementation == "flash_attention_2"
        self.norm = RizeRMSNorm(config.hidden_size, eps=config.rms_norm_eps)
        self.router_aux_loss = None  # 直近 forward の合計 aux loss を保持
        self._last_moe_logging = []  # 直近 step の MoE 統計（層ごと）バッファ


        self.gradient_checkpointing = False
        # Initialize weights and apply final processing
        self.post_init()

    def get_input_embeddings(self):
        return self.embed_tokens

    def set_input_embeddings(self, value):
        self.embed_tokens = value

    @add_start_docstrings_to_model_forward(Rize_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[List[torch.FloatTensor]] = 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,
        sequence_ids: Optional[torch.LongTensor] = None,
    ) -> 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
        )

        # retrieve input_ids and inputs_embeds
        if input_ids is not None and inputs_embeds is not None:
            raise ValueError(
                "You cannot specify both input_ids and inputs_embeds at the same time"
            )
        elif input_ids is not None:
            batch_size, seq_length = input_ids.shape[:2]
            device = input_ids.device
        elif inputs_embeds is not None:
            batch_size, seq_length = inputs_embeds.shape[:2]
            device = inputs_embeds.device
        else:
            raise ValueError("You have to specify either input_ids or inputs_embeds")

        if sequence_ids is not None:
            sequence_ids = sequence_ids.to(device=device, dtype=torch.long)
            if sequence_ids.shape != (batch_size, seq_length):
                raise ValueError(
                    f"sequence_ids should have shape {(batch_size, seq_length)}, got {tuple(sequence_ids.shape)}"
                )

        # Keep the raw 2D padding mask around. Some layers (e.g. KDA) only support 2D.
        padding_mask_2d = attention_mask

        past_key_values_length = 0
        if use_cache:
            # If generation doesn't pre-create a cache object, create one here.
            if past_key_values is None:
                if getattr(self.config, "is_linear_attn", False):
                    past_key_values = RizeHybridCache(self.config)
                else:
                    past_key_values = DynamicCache()

            # Legacy tuple cache -> DynamicCache
            if past_key_values is not None and not isinstance(past_key_values, Cache):
                past_key_values = DynamicCache.from_legacy_cache(past_key_values)

            # Hybrid linear attention mode: upgrade DynamicCache to RizeHybridCache.
            if (
                getattr(self.config, "is_linear_attn", False)
                and past_key_values is not None
                and not isinstance(past_key_values, RizeHybridCache)
            ):
                if isinstance(past_key_values, DynamicCache):
                    past_key_values = RizeHybridCache.from_dynamic_cache(past_key_values, self.config)

            # Determine how many tokens were already cached.
            if isinstance(past_key_values, Cache):
                if hasattr(past_key_values, "get_usable_length"):
                    try:
                        past_key_values_length = past_key_values.get_usable_length(seq_length)
                    except TypeError:
                        # Some cache implementations require a layer index.
                        past_key_values_length = past_key_values.get_usable_length(seq_length, 0)
                elif hasattr(past_key_values, "get_seq_length"):
                    try:
                        past_key_values_length = past_key_values.get_seq_length()
                    except TypeError:
                        past_key_values_length = past_key_values.get_seq_length(0)

        if sequence_ids is not None and past_key_values_length != 0:
            raise NotImplementedError(
                "packed sequence_ids with block-diagonal masking do not support cached decoding/past_key_values yet"
            )

        if position_ids is None:
            if sequence_ids is not None:
                position_ids = _build_reset_position_ids_from_sequence_ids(
                    sequence_ids,
                    attention_mask=padding_mask_2d,
                )
            else:
                position_ids = torch.arange(
                    past_key_values_length,
                    seq_length + past_key_values_length,
                    dtype=torch.long,
                    device=device,
                )
                position_ids = position_ids.unsqueeze(0)

        if inputs_embeds is None:
            inputs_embeds = self.embed_tokens(input_ids)

        # Attention mask preparation
        #  - full attention layers: follow existing behavior (2D for flash-attn2, otherwise 4D causal)
        #  - packed training with sequence_ids: build exact block-diagonal causal masking
        #  - KDA layers: always use a 2D padding mask (or None)
        use_packed_block_mask = sequence_ids is not None
        if self._use_flash_attention_2:
            if use_packed_block_mask:
                full_attention_mask = padding_mask_2d
                if full_attention_mask is None:
                    full_attention_mask = torch.ones(
                        (batch_size, seq_length),
                        device=inputs_embeds.device,
                        dtype=torch.long,
                    )
            else:
                full_attention_mask = (
                    padding_mask_2d
                    if (padding_mask_2d is not None and 0 in padding_mask_2d)
                    else None
                )
        else:
            if use_packed_block_mask:
                full_attention_mask = _prepare_4d_block_diagonal_causal_attention_mask(
                    padding_mask_2d,
                    sequence_ids,
                    (batch_size, seq_length),
                    inputs_embeds,
                    past_key_values_length,
                )
            else:
                full_attention_mask = _prepare_4d_causal_attention_mask(
                    padding_mask_2d,
                    (batch_size, seq_length),
                    inputs_embeds,
                    past_key_values_length,
                )

        linear_attention_mask = padding_mask_2d
        if getattr(self.config, "is_linear_attn", False):
            # For decoding steps (cache already has tokens), KDA doesn't need the padding mask.
            cache_position = torch.arange(
                past_key_values_length,
                past_key_values_length + seq_length,
                device=inputs_embeds.device,
            )
            if cache_position.numel() > 0 and int(cache_position[0].item()) > 0:
                linear_attention_mask = None
            elif linear_attention_mask is not None and torch.all(linear_attention_mask == 1):
                linear_attention_mask = None

        router_token_mask = padding_mask_2d

        # embed positions
        hidden_states = inputs_embeds
        all_hidden_states = () if output_hidden_states else None
        all_self_attns = () if output_attentions else None
        next_decoder_cache = None

        aux_total = None

        for layer_idx, decoder_layer in enumerate(self.layers):
            if output_hidden_states:
                all_hidden_states += (hidden_states,)

            layer_attention_mask = (
                linear_attention_mask if getattr(decoder_layer, "is_linear_attn", False) else full_attention_mask
            )

            layer_outputs = decoder_layer(
                hidden_states,
                attention_mask=layer_attention_mask,
                position_ids=position_ids,
                past_key_value=past_key_values,
                output_attentions=output_attentions,
                use_cache=use_cache,
                sequence_ids=sequence_ids,
                router_token_mask=router_token_mask,
            )

            hidden_states = layer_outputs[0]

            if use_cache:
                next_decoder_cache = layer_outputs[2 if output_attentions else 1]
            if output_attentions:
                all_self_attns += (layer_outputs[1],)

            # 末尾に aux があれば合算
            base_elems = 1 + (1 if output_attentions else 0) + (1 if use_cache else 0)
            if len(layer_outputs) > base_elems:
                layer_aux = layer_outputs[-1]
                if layer_aux is not None:
                    aux_total = layer_aux if aux_total is None else (aux_total + layer_aux)

            # ===== MoE 統計の回収（MoE 層のみ） =====
            mlp = getattr(decoder_layer, "mlp", None)
            if isinstance(mlp, RizeMoE):
                s = getattr(mlp, "_last_moe_stats", None)
                if s:
                    payload = {k: (float(v) if hasattr(v, "item") else v) for k, v in s.items()}
                    payload["layer_idx"] = int(layer_idx)
                    self._last_moe_logging.append(payload)

        hidden_states = self.norm(hidden_states)

        if output_hidden_states:
            all_hidden_states += (hidden_states,)

        next_cache = next_decoder_cache if use_cache else None
        self.router_aux_loss = aux_total  # ForCausalLM で読み出す

        if not return_dict:
            return tuple(v for v in [hidden_states, next_cache, all_hidden_states, all_self_attns] if v is not None)

        return BaseModelOutputWithPast(
            last_hidden_state=hidden_states,
            past_key_values=next_cache,
            hidden_states=all_hidden_states,
            attentions=all_self_attns,
        )

    def reset_global_lbl_buffers(self):
        """Reset per-layer Global LBL count buffers (used across gradient accumulation)."""
        for layer in self.layers:
            mlp = getattr(layer, "mlp", None)
            if isinstance(mlp, RizeMoE):
                mlp.reset_global_lbl_buffer()

    def pop_and_reset_moe_stats(self):
        """直近 step の MoE 統計（層ごと）を返してバッファをクリア"""
        out = list(self._last_moe_logging)
        self._last_moe_logging = []
        return out


class RizeForCausalLM(RizePreTrainedModel, GenerationMixin):
    _tied_weights_keys = ["lm_head.weight"]

    def __init__(self, config):
        super().__init__(config)
        self.model = RizeModel(config)
        self.vocab_size = config.vocab_size
        self.lm_head = nn.Linear(config.hidden_size, config.vocab_size, bias=False)
        use_linear_ce_env = os.getenv("USE_LINEAR_CE")
        if use_linear_ce_env is None:
            self._use_linear_ce = bool(getattr(config, "use_linear_ce", True))
        else:
            self._use_linear_ce = use_linear_ce_env.lower() in ("1", "true", "yes", "y", "on")
        self._linear_ce_impl = getattr(
            config, "linear_ce_impl", os.getenv("LINEAR_CE_IMPL", "cce_exact")
        )

        # Initialize weights and apply final processing
        self.post_init()

        if bool(getattr(self.config, "tie_word_embeddings", False)):
            self.tie_weights()

    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

    # ===== Trainer Callback から呼び出すためのプロキシ =====
    def reset_global_lbl_buffers(self):
        return self.model.reset_global_lbl_buffers()

    def pop_and_reset_moe_stats(self):
        return self.model.pop_and_reset_moe_stats()

    @add_start_docstrings_to_model_forward(Rize_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[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,
        sequence_ids: Optional[torch.LongTensor] = None,
        loss_weights: Optional[torch.Tensor] = None,
    ) -> 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, transformers.,
                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, transformers., config.vocab_size]`.

        Returns:

        Example:

        ```python
        >>> from transformers import AutoTokenizer, RizeForCausalLM

        >>> model = RizeForCausalLM.from_pretrained(PATH_TO_CONVERTED_WEIGHTS)
        >>> tokenizer = AutoTokenizer.from_pretrained(PATH_TO_CONVERTED_TOKENIZER)

        >>> 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?
I'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,
            sequence_ids=sequence_ids,
        )

        hidden_states = outputs[0]

        weighted_loss = labels is not None and loss_weights is not None
        logits = None
        loss = None
        if labels is not None:
            if weighted_loss:
                shift_labels = labels[..., 1:].contiguous()
                shift_weights = loss_weights[..., 1:].contiguous().to(hidden_states.device, dtype=torch.float32)
                valid = shift_labels.ne(-100)
                shift_weights = torch.where(valid, shift_weights, torch.zeros_like(shift_weights))

                loss_logits = self.lm_head(hidden_states[..., :-1, :])
                per_token_loss = F.cross_entropy(
                    loss_logits.reshape(-1, self.config.vocab_size),
                    shift_labels.reshape(-1).to(loss_logits.device),
                    ignore_index=-100,
                    reduction="none",
                ).view(shift_labels.shape)

                denom = shift_weights.sum().clamp_min(1e-12)
                loss = (per_token_loss.to(torch.float32) * shift_weights).sum() / denom
            elif self._use_linear_ce:
                loss = linear_cross_entropy(
                    hidden_states,
                    self.lm_head.weight,
                    labels,
                    shift=1,
                    impl=self._linear_ce_impl,
                    reduction="mean",
                )
            else:
                ignore_index = -100
                if labels.dim() == 2:
                    shift_labels = labels.clone()
                    shift_labels[:, :-1] = labels[:, 1:]
                    shift_labels[:, -1] = ignore_index
                else:
                    shift_labels = labels.clone()
                    shift_labels[..., :-1] = labels[..., 1:]
                    shift_labels[..., -1] = ignore_index

                logits = self.lm_head(hidden_states)
                loss_fct = CrossEntropyLoss(ignore_index=ignore_index)
                loss = loss_fct(
                    logits.view(-1, self.config.vocab_size),
                    shift_labels.view(-1).to(logits.device),
                )

            # MoE 補助損失を加算（config.aux_loss_alpha を使用）
            router_aux = getattr(self.model, "router_aux_loss", None)
            coef = float(getattr(self.config, "aux_loss_alpha", 0.0))
            if router_aux is not None and coef > 0.0:
                loss = loss + coef * router_aux

        # NOTE (DDP find_unused_parameters=True の OOM 回避):
        # find_unused_parameters=True の場合、DDP は forward の戻り値に含まれる Tensor を clone() する。
        # logits=[B,T,V] が巨大だと clone 分だけで数GB増え OOM しやすい。
        # 学習時(labelsあり)のデフォルトでは logits を返さないようにする。
        # 必要なら config.return_logits_in_train=True で復帰できる。
        return_logits = True
        if self.training and labels is not None and not bool(getattr(self.config, "return_logits_in_train", False)):
            return_logits = False

        if labels is None or return_logits:
            if logits is None:
                logits = self.lm_head(hidden_states)

        if not return_dict:
            if return_logits:
                output = (logits,) + outputs[1:]
                return (loss,) + output if loss is not None else output
            return (loss,) if loss is not None else ()

        if not return_logits:
            return CausalLMOutputWithPast(
                loss=loss,
                logits=None,
                past_key_values=None,
                hidden_states=None,
                attentions=None,
            )

        return CausalLMOutputWithPast(
            loss=loss,
            logits=logits,
            past_key_values=outputs.past_key_values,
            hidden_states=outputs.hidden_states,
            attentions=outputs.attentions,
        )

    def prepare_inputs_for_generation(
        self,
        input_ids,
        past_key_values=None,
        attention_mask=None,
        inputs_embeds=None,
        **kwargs,
    ):
        if past_key_values is not None:
            if isinstance(past_key_values, Cache):
                cache_length = past_key_values.get_seq_length()
                past_length = past_key_values.seen_tokens
                max_cache_length = past_key_values.get_seq_length()
            else:
                cache_length = past_length = past_key_values[0][0].shape[2]
                max_cache_length = None

            # Keep only the unprocessed tokens:
            # 1 - If the length of the attention_mask exceeds the length of input_ids, then we are in a setting where
            # some of the inputs are exclusivelly passed as part of the cache (e.g. when passing input_embeds as
            # input)
            if (
                attention_mask is not None
                and attention_mask.shape[1] > input_ids.shape[1]
            ):
                input_ids = input_ids[:, -(attention_mask.shape[1] - past_length) :]
            # 2 - If the past_length is smaller than input_ids', then input_ids holds all input tokens. We can discard
            # input_ids based on the past_length.
            elif past_length < input_ids.shape[1]:
                input_ids = input_ids[:, past_length:]
            # 3 - Otherwise (past_length >= input_ids.shape[1]), let's assume input_ids only has unprocessed tokens.

            # If we are about to go beyond the maximum cache length, we need to crop the input attention mask.
            if (
                max_cache_length is not None
                and attention_mask is not None
                and cache_length + input_ids.shape[1] > max_cache_length
            ):
                attention_mask = attention_mask[:, -max_cache_length:]

        position_ids = kwargs.get("position_ids", None)
        if attention_mask is not None and position_ids is None:
            # create position_ids on the fly for batch generation
            position_ids = attention_mask.long().cumsum(-1) - 1
            position_ids.masked_fill_(attention_mask == 0, 1)
            if past_key_values:
                position_ids = position_ids[:, -input_ids.shape[1] :]

        # if `inputs_embeds` are passed, we only want to use them in the 1st generation step
        if inputs_embeds is not None and past_key_values is None:
            model_inputs = {"inputs_embeds": inputs_embeds}
        else:
            model_inputs = {"input_ids": input_ids}

        model_inputs.update(
            {
                "position_ids": position_ids,
                "past_key_values": past_key_values,
                "use_cache": kwargs.get("use_cache"),
                "attention_mask": attention_mask,
            }
        )
        return model_inputs

    @staticmethod
    def _reorder_cache(past_key_values, beam_idx):
        reordered_past = ()
        for layer_past in past_key_values:
            reordered_past += (
                tuple(
                    past_state.index_select(0, beam_idx.to(past_state.device))
                    for past_state in layer_past
                ),
            )
        return reordered_past


@add_start_docstrings(
    """
    The Rize Model transformer with a sequence classification head on top (linear layer).

    [`RizeForSequenceClassification`] uses the last token in order to do the classification, as other causal models
    (e.g. GPT-2) do.

    Since it does classification on the last token, it requires to know the position of the last token. If a
    `pad_token_id` is defined in the configuration, it finds the last token that is not a padding token in each row. If
    no `pad_token_id` is defined, it simply takes the last value in each row of the batch. Since it cannot guess the
    padding tokens when `inputs_embeds` are passed instead of `input_ids`, it does the same (take the last value in
    each row of the batch).
    """,
    Rize_START_DOCSTRING,
)
class RizeForSequenceClassification(RizePreTrainedModel):
    def __init__(self, config):
        super().__init__(config)
        self.num_labels = config.num_labels
        self.model = RizeModel(config)
        self.score = nn.Linear(config.hidden_size, self.num_labels, 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

    @add_start_docstrings_to_model_forward(Rize_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[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,
    ) -> Union[Tuple, SequenceClassifierOutputWithPast]:
        r"""
        labels (`torch.LongTensor` of shape `(batch_size,)`, *optional*):
            Labels for computing the sequence classification/regression loss. Indices should be in `[0, transformers.,
            config.num_labels - 1]`. If `config.num_labels == 1` a regression loss is computed (Mean-Square loss), If
            `config.num_labels > 1` a classification loss is computed (Cross-Entropy).
        """
        return_dict = (
            return_dict if return_dict is not None else self.config.use_return_dict
        )

        transformer_outputs = self.model(
            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,
        )
        hidden_states = transformer_outputs[0]
        logits = self.score(hidden_states)

        if input_ids is not None:
            batch_size = input_ids.shape[0]
        else:
            batch_size = inputs_embeds.shape[0]

        if self.config.pad_token_id is None and batch_size != 1:
            raise ValueError(
                "Cannot handle batch sizes > 1 if no padding token is defined."
            )
        if self.config.pad_token_id is None:
            sequence_lengths = -1
        else:
            if input_ids is not None:
                sequence_lengths = (
                    torch.eq(input_ids, self.config.pad_token_id).int().argmax(-1) - 1
                ).to(logits.device)
            else:
                sequence_lengths = -1

        pooled_logits = logits[
            torch.arange(batch_size, device=logits.device), sequence_lengths
        ]

        loss = None
        if labels is not None:
            labels = labels.to(logits.device)
            if self.config.problem_type is None:
                if self.num_labels == 1:
                    self.config.problem_type = "regression"
                elif self.num_labels > 1 and (
                    labels.dtype == torch.long or labels.dtype == torch.int
                ):
                    self.config.problem_type = "single_label_classification"
                else:
                    self.config.problem_type = "multi_label_classification"

            if self.config.problem_type == "regression":
                loss_fct = MSELoss()
                if self.num_labels == 1:
                    loss = loss_fct(pooled_logits.squeeze(), labels.squeeze())
                else:
                    loss = loss_fct(pooled_logits, labels)
            elif self.config.problem_type == "single_label_classification":
                loss_fct = CrossEntropyLoss()
                loss = loss_fct(
                    pooled_logits.view(-1, self.num_labels), labels.view(-1)
                )
            elif self.config.problem_type == "multi_label_classification":
                loss_fct = BCEWithLogitsLoss()
                loss = loss_fct(pooled_logits, labels)
        if not return_dict:
            output = (pooled_logits,) + transformer_outputs[1:]
            return ((loss,) + output) if loss is not None else output

        return SequenceClassifierOutputWithPast(
            loss=loss,
            logits=pooled_logits,
            past_key_values=transformer_outputs.past_key_values,
            hidden_states=transformer_outputs.hidden_states,
            attentions=transformer_outputs.attentions,
        )