""" 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, )