Kernels:

kernels-community
/

liger-kernels

Trusted publisher

Kernel card Files Files and versions

xet

Community

kernels-bot commited on 2 days ago

Commit

20f6e91

verified ·

1 Parent(s): 93fec15

Uploaded using `kernel-builder`.

Browse files

Files changed (46) hide show

build/torch-cuda/__init__.py +8 -2
build/torch-cuda/_ops.py +1 -1
build/torch-cuda/layers.py +172 -344
build/torch-cuda/metadata.json +37 -2
build/torch-cuda/metadata.json.sigstore +1 -0
build/torch-cuda/tiled_mlp.py +31 -8
build/torch-rocm/__init__.py +0 -4
build/torch-rocm/_ops.py +0 -38
build/torch-rocm/cross_entropy.py +0 -558
build/torch-rocm/dyt.py +0 -164
build/torch-rocm/fused_linear_cross_entropy.py +0 -400
build/torch-rocm/geglu.py +0 -143
build/torch-rocm/group_norm.py +0 -311
build/torch-rocm/jsd.py +0 -201
build/torch-rocm/kl_div.py +0 -259
build/torch-rocm/layer_norm.py +0 -320
build/torch-rocm/layers.py +0 -514
build/torch-rocm/liger_kernels/__init__.py +0 -26
build/torch-rocm/metadata.json +0 -10
build/torch-rocm/qwen2vl_mrope.py +0 -222
build/torch-rocm/rms_norm.py +0 -654
build/torch-rocm/rope.py +0 -239
build/torch-rocm/swiglu.py +0 -176
build/torch-rocm/tiled_mlp.py +0 -136
build/torch-rocm/tvd.py +0 -218
build/torch-rocm/utils.py +0 -176
build/torch-xpu/__init__.py +0 -4
build/torch-xpu/_ops.py +0 -38
build/torch-xpu/cross_entropy.py +0 -558
build/torch-xpu/dyt.py +0 -164
build/torch-xpu/fused_linear_cross_entropy.py +0 -400
build/torch-xpu/geglu.py +0 -143
build/torch-xpu/group_norm.py +0 -311
build/torch-xpu/jsd.py +0 -201
build/torch-xpu/kl_div.py +0 -259
build/torch-xpu/layer_norm.py +0 -320
build/torch-xpu/layers.py +0 -514
build/torch-xpu/liger_kernels/__init__.py +0 -26
build/torch-xpu/metadata.json +0 -10
build/torch-xpu/qwen2vl_mrope.py +0 -222
build/torch-xpu/rms_norm.py +0 -654
build/torch-xpu/rope.py +0 -239
build/torch-xpu/swiglu.py +0 -176
build/torch-xpu/tiled_mlp.py +0 -136
build/torch-xpu/tvd.py +0 -218
build/torch-xpu/utils.py +0 -176

build/torch-cuda/__init__.py CHANGED Viewed

@@ -1,4 +1,10 @@
 from . import layers
-from .layers import CrossEntropyOutput, LigerForCausalLMLoss
-__all__ = ["layers", "LigerForCausalLMLoss", "CrossEntropyOutput"]

 from . import layers
+from .layers import LigerForCausalLMLoss, liger_rotary_pos_emb
+__all__ = [
+    "layers",
+    # kept for BC
+    "LigerForCausalLMLoss",
+    "liger_rotary_pos_emb",
+]

build/torch-cuda/_ops.py CHANGED Viewed

@@ -22,7 +22,7 @@ def get_backend() -> str:
 def _find_ops_name() -> str:
     kernel_name = "liger_kernels"
-    unique_id = "ab435e2"
     backend = get_backend()
     return f"_{kernel_name}_{backend}_{unique_id}"

 def _find_ops_name() -> str:
     kernel_name = "liger_kernels"
+    unique_id = "4d9f798"
     backend = get_backend()
     return f"_{kernel_name}_{backend}_{unique_id}"

build/torch-cuda/layers.py CHANGED Viewed

@@ -1,360 +1,142 @@
 import inspect
 from dataclasses import dataclass
-from typing import Optional, Tuple
 import torch
 import torch.nn as nn
-from .cross_entropy import LigerCrossEntropyFunction
-from .dyt import LigerDyTFunction
 from .fused_linear_cross_entropy import LigerFusedLinearCrossEntropyFunction
 from .geglu import LigerGELUMulFunction
-from .group_norm import LigerGroupNormFunction
-from .jsd import LigerJSDFunction
-from .kl_div import LigerKLDivLossFunction
-from .layer_norm import LigerLayerNormFunction
-from .qwen2vl_mrope import LigerQwen2VLMRopeFunction
 from .rms_norm import LigerRMSNormFunction
 from .rope import LigerRopeFunction
 from .swiglu import LigerSiLUMulFunction
 from .tiled_mlp import apply_tiled_mlp
-from .tvd import LigerTVDLossFunction
 class LigerRMSNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        offset: float = 0.0,
-        casting_mode: str = "llama",
-        init_fn: str = "ones",
-        in_place: bool = True,
-        row_mode: Optional[bool] = None,
-        elementwise_affine: bool = True,
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.offset = offset
-        self.casting_mode = casting_mode
-        self.in_place = in_place
-        self.row_mode = row_mode
-        self.elementwise_affine = elementwise_affine
-        if elementwise_affine:
-            init = torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size)
-            self.weight = nn.Parameter(init)
-        else:
-            self.register_parameter("weight", None)
     def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
         return LigerRMSNormFunction.apply(
             hidden_states,
             self.weight,
             self.variance_epsilon,
-            self.offset,
-            self.casting_mode,
-            self.in_place,
-            self.row_mode,
-        )
-    def extra_repr(self) -> str:
-        return (
-            f"{self.hidden_size}, eps={self.variance_epsilon}, offset={self.offset}, "
-            f"in_place={self.in_place}, row_mode={self.row_mode}"
         )
-class LigerLayerNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size))
-        self.bias = nn.Parameter(torch.randn(hidden_size) if bias else torch.zeros(hidden_size))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        return LigerLayerNormFunction.apply(hidden_states, self.weight, self.bias, self.variance_epsilon)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, eps={self.variance_epsilon}"
-class LigerGroupNorm(nn.Module):
-    def __init__(
-        self,
-        num_channels: int,
-        num_groups: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        assert num_channels % num_groups == 0, (
-            f"num_channels ({num_channels}) must be divisible by num_groups ({num_groups})"
-        )
-        self.num_channels = num_channels
-        self.num_groups = num_groups
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(num_channels) if init_fn == "ones" else torch.zeros(num_channels))
-        self.bias = nn.Parameter(torch.randn(num_channels) if bias else torch.zeros(num_channels))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        assert hidden_states.dim() >= 3, f"Input must have at least 3 dimensions, got {hidden_states.dim()}"
-        assert hidden_states.size(1) == self.num_channels, (
-            f"Input must have {self.num_channels} channels, got {hidden_states.size(1)}"
-        )
-        return LigerGroupNormFunction.apply(
-            hidden_states,
-            self.weight,
-            self.bias,
-            self.num_channels,
-            self.num_groups,
-            self.variance_epsilon,
-        )
-    def extra_repr(self) -> str:
-        return f"num_channels={self.num_channels}, num_groups={self.num_groups}, eps={self.variance_epsilon}"
-class LigerDyT(nn.Module):
-    def __init__(self, hidden_size: int, beta: bool = True, init_alpha: float = 0.5):
-        super().__init__()
-        self.hidden_size = hidden_size
-        self.init_alpha = init_alpha
-        self.alpha = nn.Parameter(torch.ones(1) * init_alpha)
-        self.gamma = nn.Parameter(torch.ones(hidden_size))
-        self.beta = nn.Parameter(torch.zeros(hidden_size)) if beta else None
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return LigerDyTFunction.apply(x, self.alpha, self.gamma, self.beta)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, init_alpha={self.init_alpha}, beta={self.beta is not None}"
-class LigerCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.weight = weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-    def forward(self, _input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
-        loss, _, _, _ = LigerCrossEntropyFunction.apply(
-            _input,
-            target,
-            self.weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            False,
-            False,
-        )
-        return loss
-class LigerFusedLinearCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        ce_weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-        accum_dtype: Optional[torch.dtype] = None,
-        use_token_scaling: bool = False,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.ce_weight = ce_weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-        self.accum_dtype = accum_dtype
-        self.use_token_scaling = use_token_scaling
-    def forward(
-        self,
-        lin_weight: torch.Tensor,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        bias: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        loss, _, _, _ = LigerFusedLinearCrossEntropyFunction.apply(
-            _input,
-            lin_weight,
-            target,
-            bias,
-            self.ce_weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            self.accum_dtype,
-            self.use_token_scaling,
-            False,
-            False,
         )
-        return loss
-class LigerJSD(nn.Module):
-    def __init__(self, beta: float = 0.5, ignore_index: int = -100):
-        super().__init__()
-        self.beta = beta
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        log_q: torch.Tensor,
-        log_p: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerJSDFunction.apply(log_q, log_p, shift_labels, self.beta, self.ignore_index)
-class LigerKLDIVLoss(nn.KLDivLoss):
-    def __init__(self, eps: float = 1e-10, *args, **kwargs):
-        super().__init__(*args, **kwargs)
-        self.eps = eps
-    def forward(self, y_pred: torch.Tensor, y_true: torch.Tensor) -> torch.Tensor:
-        return LigerKLDivLossFunction.apply(y_pred, y_true, self.reduction, self.log_target, self.eps)
-class LigerTVDLoss(nn.Module):
-    def __init__(self, reduction: str = "batchmean", ignore_index: int = -100):
-        super().__init__()
-        self.reduction = reduction
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        p: torch.Tensor,
-        q: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerTVDLossFunction.apply(p, q, shift_labels, self.reduction, self.ignore_index)
 class LigerSwiGLUMLP(nn.Module):
-    """SwiGLU MLP block. ``config`` must expose ``hidden_size``, ``intermediate_size``,
-    and ``hidden_act`` (must be ``silu`` or ``swish``)."""
-    def __init__(self, config):
-        super().__init__()
-        if config.hidden_act not in ("silu", "swish"):
-            raise ValueError(f"Activation function {config.hidden_act} not supported.")
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return self.down_proj(LigerSiLUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
 class LigerGEGLUMLP(nn.Module):
-    """GEGLU MLP block. ``config`` must expose ``hidden_size`` and ``intermediate_size``.
-    Uses the tanh approximation of GELU (matches Gemma 1/1.1/2)."""
-    def __init__(self, config):
-        super().__init__()
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return self.down_proj(LigerGELUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
-class LigerTiledGEGLUMLP(nn.Module):
     gate_proj: nn.Linear
     up_proj: nn.Linear
     down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerGELUMulFunction.apply(gate, up))
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         compute_params = [p for p in self.parameters() if p.requires_grad]
         return apply_tiled_mlp(
-            fn=self._mlp_forward,
             mlp_module=self,
             x=x,
-            num_shards=self.num_shards,
             compute_params=compute_params,
         )
-class LigerTiledSwiGLUMLP(nn.Module):
     gate_proj: nn.Linear
     up_proj: nn.Linear
     down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerSiLUMulFunction.apply(gate, up))
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         compute_params = [p for p in self.parameters() if p.requires_grad]
         return apply_tiled_mlp(
-            fn=self._mlp_forward,
             mlp_module=self,
             x=x,
-            num_shards=self.num_shards,
             compute_params=compute_params,
         )
 @dataclass
 class CrossEntropyOutput:
     loss: torch.Tensor
@@ -407,23 +189,33 @@ def liger_fused_linear_cross_entropy(
     )
 def LigerForCausalLMLoss(
-    hidden_states: torch.Tensor,
-    lm_head_weight: torch.Tensor,
     labels: torch.Tensor,
-    hidden_size: int,
     num_items_in_batch: Optional[int] = None,
     ignore_index: int = -100,
     shift_labels: Optional[torch.Tensor] = None,
-    final_logit_softcapping: Optional[float] = None,
-    return_token_accuracy: bool = False,
-    return_predicted_tokens: bool = False,
     **kwargs,
 ):
-    """Drop-in replacement for ``transformers.loss.ForCausalLMLoss`` that fuses the
-    final ``lm_head`` projection with the cross-entropy loss. Returns a scalar
-    ``loss`` by default; returns a :class:`CrossEntropyOutput` when
-    ``return_token_accuracy`` or ``return_predicted_tokens`` is set."""
     applicable_params = inspect.signature(liger_fused_linear_cross_entropy).parameters
     kwargs = {k: v for k, v in kwargs.items() if k in applicable_params}
@@ -435,80 +227,116 @@ def LigerForCausalLMLoss(
     shift_labels = shift_labels.view(-1).to(hidden_states.device)
     reduction = "sum" if num_items_in_batch is not None else "mean"
-    result = liger_fused_linear_cross_entropy(
         hidden_states,
         lm_head_weight,
         shift_labels,
         reduction=reduction,
         ignore_index=ignore_index,
         softcap=final_logit_softcapping,
-        return_token_accuracy=return_token_accuracy,
-        return_predicted_tokens=return_predicted_tokens,
         **kwargs,
     )
-    if isinstance(result, CrossEntropyOutput):
-        loss = result.loss
-        token_accuracy = result.token_accuracy
-        predicted_tokens = result.predicted_tokens
-    else:
-        loss = result
-        token_accuracy = None
-        predicted_tokens = None
     if reduction == "sum":
         loss = loss / num_items_in_batch
-    if return_token_accuracy or return_predicted_tokens:
-        return CrossEntropyOutput(
-            loss=loss,
-            token_accuracy=token_accuracy,
-            predicted_tokens=predicted_tokens,
-        )
     return loss
-def liger_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    position_ids: Optional[torch.Tensor] = None,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply standard rotary positional embedding to ``q`` and ``k``."""
-    return LigerRopeFunction.apply(q, k, cos, sin, position_ids, unsqueeze_dim)
-def liger_multimodal_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    mrope_section,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply Qwen2-VL multimodal rotary positional embedding (M-RoPE) to ``q`` and ``k``."""
-    return LigerQwen2VLMRopeFunction.apply(q, k, cos, sin, mrope_section, unsqueeze_dim)
 __all__ = [
     "LigerRMSNorm",
-    "LigerLayerNorm",
-    "LigerGroupNorm",
-    "LigerDyT",
-    "LigerCrossEntropyLoss",
-    "LigerFusedLinearCrossEntropyLoss",
-    "LigerJSD",
-    "LigerKLDIVLoss",
-    "LigerTVDLoss",
     "LigerSwiGLUMLP",
     "LigerGEGLUMLP",
-    "LigerTiledGEGLUMLP",
     "LigerTiledSwiGLUMLP",
-    "CrossEntropyOutput",
-    "liger_fused_linear_cross_entropy",
-    "LigerForCausalLMLoss",
     "liger_rotary_pos_emb",
-    "liger_multimodal_rotary_pos_emb",
 ]

 import inspect
 from dataclasses import dataclass
+from typing import Optional
 import torch
 import torch.nn as nn
 from .fused_linear_cross_entropy import LigerFusedLinearCrossEntropyFunction
 from .geglu import LigerGELUMulFunction
 from .rms_norm import LigerRMSNormFunction
 from .rope import LigerRopeFunction
 from .swiglu import LigerSiLUMulFunction
 from .tiled_mlp import apply_tiled_mlp
+# NOTE: Not compile-friendly --> large deviations to the original implementation under compile
 class LigerRMSNorm(nn.Module):
+    weight: nn.Parameter
+    variance_epsilon: float
     def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
         return LigerRMSNormFunction.apply(
             hidden_states,
             self.weight,
             self.variance_epsilon,
+            0,
+            "llama",
+            True,
+            None,
         )
+    def extra_repr(self):
+        return f"{tuple(self.weight.shape)}, eps={self.variance_epsilon}"
+class LigerLinear(nn.Module):
+    weight: nn.Parameter
+    bias: nn.Parameter | None
+    can_torch_compile = True
+    def forward(self, input: torch.Tensor) -> torch.Tensor:
+        def _forward(module, x):
+            return nn.functional.linear(x, module.weight, module.bias)
+        compute_params = [p for p in self.parameters() if p.requires_grad]
+        return apply_tiled_mlp(
+            fn=_forward,
+            mlp_module=self,
+            x=input,
+            num_shards=None,
+            compute_params=compute_params,
         )
 class LigerSwiGLUMLP(nn.Module):
+    gate_proj: nn.Linear
+    up_proj: nn.Linear
+    down_proj: nn.Linear
+    can_torch_compile = True
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return self.down_proj(LigerSiLUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
 class LigerGEGLUMLP(nn.Module):
+    gate_proj: nn.Linear
+    up_proj: nn.Linear
+    down_proj: nn.Linear
+    can_torch_compile = True
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return self.down_proj(LigerGELUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
+class LigerTiledSwiGLUMLP(nn.Module):
     gate_proj: nn.Linear
     up_proj: nn.Linear
     down_proj: nn.Linear
+    can_torch_compile = True
     def forward(self, x: torch.Tensor) -> torch.Tensor:
+        def _mlp_forward(module, x):
+            """Internal MLP forward function for tiled computation."""
+            gate = module.gate_proj(x)
+            up = module.up_proj(x)
+            return module.down_proj(LigerSiLUMulFunction.apply(gate, up))
         compute_params = [p for p in self.parameters() if p.requires_grad]
         return apply_tiled_mlp(
+            fn=_mlp_forward,
             mlp_module=self,
             x=x,
+            num_shards=None,
             compute_params=compute_params,
         )
+class LigerTiledGEGLUMLP(nn.Module):
     gate_proj: nn.Linear
     up_proj: nn.Linear
     down_proj: nn.Linear
+    can_torch_compile = True
     def forward(self, x: torch.Tensor) -> torch.Tensor:
+        def _mlp_forward(module, x):
+            """Internal MLP forward function for tiled computation."""
+            gate = module.gate_proj(x)
+            up = module.up_proj(x)
+            return module.down_proj(LigerGELUMulFunction.apply(gate, up))
         compute_params = [p for p in self.parameters() if p.requires_grad]
         return apply_tiled_mlp(
+            fn=_mlp_forward,
             mlp_module=self,
             x=x,
+            num_shards=None,
             compute_params=compute_params,
         )
+def liger_rotary_pos_emb(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    cos: torch.Tensor,
+    sin: torch.Tensor,
+    unsqueeze_dim: int = 1,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Apply standard rotary positional embedding to ``q`` and ``k``."""
+    return LigerRopeFunction.apply(q, k, cos, sin, None, unsqueeze_dim)
 @dataclass
 class CrossEntropyOutput:
     loss: torch.Tensor
     )
+# NOTE: We have this intentional graph break as we encounter issues such as IMAs and Cublas errors.
+#       We know that this is an optimized kernel already so there is less ways to
+#       fuse it either way; we rely on torch compile to go through the base model to optimize.
+@torch.compiler.disable
 def LigerForCausalLMLoss(
+    logits: None,  # to match transformers signature
     labels: torch.Tensor,
+    vocab_size: int,  # to match transformers signature
     num_items_in_batch: Optional[int] = None,
     ignore_index: int = -100,
     shift_labels: Optional[torch.Tensor] = None,
+    hidden_states: torch.Tensor | None = None,
+    lm_head_weight: torch.Tensor | None = None,
+    lm_head_bias: torch.Tensor | None = None,
     **kwargs,
 ):
+    # To match signature we hide these behind the kwargs but we expect a few kwargs to exist
+    hidden_size = kwargs.pop("hidden_size", None)
+    final_logit_softcapping = kwargs.pop("final_logit_softcapping", None)
+    if hidden_size is None or hidden_states is None or lm_head_weight is None:
+        raise ValueError(
+            f"`LigerForCausalLMLoss` requires the LLM's weight (found `{lm_head_weight is not None}`),"
+            f"the last hidden state (found `{hidden_states is not None}`), and the `hidden_size`"
+            f"(found `{hidden_size is not None}`). Please make sure to pass the necessary kwargs."
+        )
     applicable_params = inspect.signature(liger_fused_linear_cross_entropy).parameters
     kwargs = {k: v for k, v in kwargs.items() if k in applicable_params}
     shift_labels = shift_labels.view(-1).to(hidden_states.device)
     reduction = "sum" if num_items_in_batch is not None else "mean"
+    loss = liger_fused_linear_cross_entropy(
         hidden_states,
         lm_head_weight,
         shift_labels,
+        bias=lm_head_bias,
         reduction=reduction,
         ignore_index=ignore_index,
         softcap=final_logit_softcapping,
+        return_token_accuracy=False,
+        return_predicted_tokens=False,
         **kwargs,
     )
     if reduction == "sum":
         loss = loss / num_items_in_batch
     return loss
+# Add torch compile support for functions - in this case it's to allow this to be used
+# but the function itself will not be compiled (see the note at the function)
+liger_rotary_pos_emb.can_torch_compile = True
+LigerForCausalLMLoss.can_torch_compile = True
+class liger_rotary_pos_emb_layer(nn.Module):
+    can_torch_compile = True
+    def forward(
+        self,
+        q: torch.Tensor,
+        k: torch.Tensor,
+        cos: torch.Tensor,
+        sin: torch.Tensor,
+        unsqueeze_dim: int = 1,
+    ) -> tuple[torch.Tensor, torch.Tensor]:
+        """Apply standard rotary positional embedding to ``q`` and ``k``."""
+        return LigerRopeFunction.apply(q, k, cos, sin, None, unsqueeze_dim)
+class LigerForCausalLMLossLayer(nn.Module):
+    can_torch_compile = True
+    # NOTE: We have this intentional graph break as we encounter issues such as IMAs and Cublas errors.
+    #       We know that this is an optimized kernel already so there is less ways to
+    #       fuse it either way; we rely on torch compile to go through the base model to optimize.
+    @torch.compiler.disable
+    def forward(
+        self,
+        logits: None,  # to match transformers signature
+        labels: torch.Tensor,
+        vocab_size: int,  # to match transformers signature
+        num_items_in_batch: Optional[int] = None,
+        ignore_index: int = -100,
+        shift_labels: Optional[torch.Tensor] = None,
+        hidden_states: torch.Tensor | None = None,
+        lm_head_weight: torch.Tensor | None = None,
+        lm_head_bias: torch.Tensor | None = None,
+        **kwargs,
+    ):
+        # To match signature we hide these behind the kwargs but we expect a few kwargs to exist
+        hidden_size = kwargs.pop("hidden_size", None)
+        final_logit_softcapping = kwargs.pop("final_logit_softcapping", None)
+        if hidden_size is None or hidden_states is None or lm_head_weight is None:
+            raise ValueError(
+                f"`LigerForCausalLMLoss` requires the LLM's weight (found `{lm_head_weight is not None}`),"
+                f"the last hidden state (found `{hidden_states is not None}`), and the `hidden_size`"
+                f"(found `{hidden_size is not None}`). Please make sure to pass the necessary kwargs."
+            )
+        applicable_params = inspect.signature(liger_fused_linear_cross_entropy).parameters
+        kwargs = {k: v for k, v in kwargs.items() if k in applicable_params}
+        if shift_labels is None:
+            labels = nn.functional.pad(labels, (0, 1), value=ignore_index)
+            shift_labels = labels[..., 1:].contiguous()
+        hidden_states = hidden_states.view(-1, hidden_size)
+        shift_labels = shift_labels.view(-1).to(hidden_states.device)
+        reduction = "sum" if num_items_in_batch is not None else "mean"
+        loss = liger_fused_linear_cross_entropy(
+            hidden_states,
+            lm_head_weight,
+            shift_labels,
+            bias=lm_head_bias,
+            reduction=reduction,
+            ignore_index=ignore_index,
+            softcap=final_logit_softcapping,
+            return_token_accuracy=False,
+            return_predicted_tokens=False,
+            **kwargs,
+        )
+        if reduction == "sum":
+            loss = loss / num_items_in_batch
+        return loss
 __all__ = [
     "LigerRMSNorm",
+    "LigerLinear",
     "LigerSwiGLUMLP",
     "LigerGEGLUMLP",
     "LigerTiledSwiGLUMLP",
+    "LigerTiledGEGLUMLP",
     "liger_rotary_pos_emb",
+    "liger_rotary_pos_emb_layer",
+    "LigerForCausalLMLoss",
+    "LigerForCausalLMLossLayer",
 ]

build/torch-cuda/metadata.json CHANGED Viewed

@@ -1,10 +1,45 @@
 {
   "name": "liger-kernels",
-  "id": "_liger_kernels_cuda_ab435e2",
-  "version": 1,
   "license": "BSD-2-Clause",
   "python-depends": [],
   "backend": {
     "type": "cuda"
   }
 }

 {
   "name": "liger-kernels",
+  "id": "_liger_kernels_cuda_4d9f798",
+  "version": 3,
   "license": "BSD-2-Clause",
   "python-depends": [],
   "backend": {
     "type": "cuda"
+  },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "__init__.py": "DSZMiK0xOBiMb0JdCc2K9QH4Z6lV7L8D/fTowstf/og=",
+      "_ops.py": "pdeXrc8M9KCvCpe0mtloEEFT3iJWZab9CMNsgT2CK10=",
+      "cross_entropy.py": "LsQN0OI8VHqgUObWF0qqTgNA9z7SYMaZNb2I7FIYnO4=",
+      "dyt.py": "806Tai0T+Yi7Ctf5YGmtwAc87MUjlzjsLky7b1T+oWM=",
+      "fused_linear_cross_entropy.py": "yjtZHWyBcXCBJ57NbbrQ1wMmTr9nARqDBrdjXVYpR74=",
+      "geglu.py": "Lmw0I38IBIHkNROTf9V3bP6PZLktCWCqNAxzVRaipbk=",
+      "group_norm.py": "xdY/tCr4rHOb4PhD827259XB7QDb22z2psUKzTYTKBE=",
+      "jsd.py": "CftyIN36AU63e1GdtuA7hNhs66I3TjS9lttcFriBnGA=",
+      "kl_div.py": "VQG/89mefLifPeVEv+1O8hmsXgGujMLA0JXDwO5QvtU=",
+      "layer_norm.py": "INpxK7Ac27f+eckcMfOOSgZErzDhNF5iGt207eV9aoM=",
+      "layers.py": "+50F9xmnpQXbU1K/lZxdJwbnTotkksJlhoLOWHjo+/M=",
+      "liger_kernels/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
+      "qwen2vl_mrope.py": "3GExhYpLgB4VUtyZyjRk8XjEur3W4EWF6HQ67ML5vBU=",
+      "rms_norm.py": "GLzIZjDP74QxaeKQ5B/VDXKTGxD6Xo9cVaW/iqmsc4c=",
+      "rope.py": "v+7JHRrv+5ImoROkpKfl30WwWI4qTa2tAl7zQeB4ml4=",
+      "swiglu.py": "TLhyvpr7nbVbMqWbx8uz+iU2GxS6q0QrVxWAd1eQCBo=",
+      "tiled_mlp.py": "PJw2R8YdHHxQRSwbFJWjcHxgzssEqsRWqyogkY8pKR8=",
+      "tvd.py": "VbWUtpJ79E2atqjBoVyTSO785aS+idHiQBIrJwifDAw=",
+      "utils.py": "Y8wQsplz6ed7oPmBfX7yGGcX74VB3t/e5VFvYv/X7wU="
+    }
+  },
+  "provenance": {
+    "kernel-builder": {
+      "version": "0.17.0-dev0",
+      "sha": "81580bb92577f2f7228661ca2a221fc052375709",
+      "dirty": false
+    },
+    "kernel": {
+      "sha": "4d9f798b95ed280b770213048b8d3f41509d0c9e",
+      "dirty": false
+    }
   }
 }

build/torch-cuda/metadata.json.sigstore ADDED Viewed

	@@ -0,0 +1 @@

+ {"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json","verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtCgAwIBAgIUIY1Kwpw1ln1fXvDyNdIoHzQ0GEQwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwNzE0MTE1OTA3WhcNMjYwNzE0MTIwOTA3WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAExWsk2F18NXh/g+v39PvCOt83WU6ZcC6CQ1e7DXAbxrYsqgMPsK29ahXv6RnnWwBkrvQRvnq6u/o0tBq27VuaiqOCBe8wggXrMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUMGdfASbZd1xeCeP+w5J3zceZ+oEwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoNGQ5Zjc5OGI5NWVkMjgwYjc3MDIxMzA0OGI4ZDNmNDE1MDlkMGM5ZTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoNGQ5Zjc5OGI5NWVkMjgwYjc3MDIxMzA0OGI4ZDNmNDE1MDlkMGM5ZTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoNGQ5Zjc5OGI5NWVkMjgwYjc3MDIxMzA0OGI4ZDNmNDE1MDlkMGM5ZTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDRkOWY3OThiOTVlZDI4MGI3NzAyMTMwNDhiOGQzZjQxNTA5ZDBjOWUwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMjkzMzA2MjM4NjUvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiQYKKwYBBAHWeQIEAgR7BHkAdwB1AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABn2B+v4EAAAQDAEYwRAIgS0J+dgCFO8GkolLxpLdvSSe2HadB4+nYR3cSvV6oUWkCIAgau6hFzgTpx+2kKwa/ampZ2gDeNYZXD1RoxAsKIyF9MAoGCCqGSM49BAMDA2kAMGYCMQDF7yEhVN8mClCXQDdHjAc7S+w7XOApJlMDaYeyEjQ2J0oEj6Kie8XJSmRQHpuJkO4CMQDU5BtTp5FUKs7zuqkbkvA2+CYey1wQd9vjYOM96KwNwZWxEjHnnKSgBwQteSkBZ7w="},"tlogEntries":[{"logIndex":"2167844577","logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="},"kindVersion":{"kind":"hashedrekord","version":"0.0.1"},"integratedTime":"1784030347","inclusionPromise":{"signedEntryTimestamp":"MEUCIAMiR8my1lRyXLZGGrVKeQBstBWpQZ51qj9FYkzHlKCyAiEA69xGRMRjKD7r1HVWg+fHimAFFTjMkz3PZgNXXl0+AUQ="},"inclusionProof":{"logIndex":"2045940315","rootHash":"myulnHdsTu/va/jsDduWqvHHfP1Bpz/gY5ASWc/vxLc=","treeSize":"2045940321","hashes":["ELPU1S59GnfBMIBqha3CDy8lRp2/UfhI9Ec14GZnsGc=","X+yhpzA31vS3BimZiPaRrAt1+TZCVDiVzX3Ba95EXBI=","VV1YHo726RTeHI4eDhHODulNVcENo4r9jpMRhiHkwAw=","TE+5oEe5ac7OQqO9brT+z6xPst7y8CQWiDgV+zzDbXI=","bHA0KCqwj96QF0bHdBKoFwsVQxc5vBpklhBiXnYJiaQ=","O+GMF6W3u/2gZ2keuQ8kFLbdfn5+iGqBOc3RNs2Cr54=","MUkgUqe8cc9LrfPi6/CeoxIcBig5OluD0n7ts3iQk78=","duwzzfwRIAv1hyLDk7Pw84ajqZB4l/R7mLum9W0K72w=","hluP9JseRAiMY+2slhjSksod6tVtp7e7HP7amn1KEaY=","sMaCrs8erYNQ48jT2FC56F0RAbpTkuSXZ+cs1zcSObw=","zGam2EiRYwS9IjemppqM/acVfU15hw3eAqn9Fb8+8vY=","qQGtNWYnFxNQ1OmqNBR2y92EFrgo5lyaKNQLaE/Dl1c=","/UAkaaq9D8BPMw5qLAxaf+ur4Sfbv3FfLWB2M/py3wM=","56+1wf4OgrBZEzfW9n13hPfm0bRvMRgerPv4TPE+3Ho=","mTSsSZdzVhRgljy9CvDq3GfjxxeOTHrwgpOWK4rRj3Q=","5sxNZDoxEj6DMmQATisX9bQXdFmeRHYfC8BjyDIIgec=","Rq/A2aTC9e54ldNjcpsJ26rX/h8JlN6ZHxgw7yEa6C0=","+/VZ56MsIPxMiyLAodzKXo5TEWdQp36z89qLhpzloAo=","daxmZaajRpZV+JxHiOYZhJBiSKN5ucqjh2WnGbHhirw=","DOCeoSMovIvLExkhIvisow9AuNXgeWs4ECkyR6EcqYU="],"checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n2045940321\nmyulnHdsTu/va/jsDduWqvHHfP1Bpz/gY5ASWc/vxLc=\n\n— rekor.sigstore.dev wNI9ajBFAiB5UZpMuWHHap4flwT77otDC4v44eb3znyf3RKYFvQoywIhAIcGfJ39B+ppsWE0EnZ0LC4Mk/1d7EhtDTrrupVhKN5G\n"}},"canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI5ZGE1ZjdkNGE1ZWRiMDdiYjJmMzJjM2M5NjU1NDFiYWNmZDQ2OWU4Y2Q5N2UyNzZiNWVhNDMwMDM0Mjk2MzA4In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FWUNJUUMvanIxUmY1TisyQi95V1BDdHNCRCsyK0ZZNjNFalFMbWY3c2lWdE43emVnSWhBUFlrYmRaVmxCK3ZTZyt0NmF4V2hHSlRVZjZtUUtHUG1QaTBkYndmWk1UKyIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SRFowRjNTVUpCWjBsVlNWa3hTM2R3ZHpGc2JqRm1XSFpFZVU1a1NXOUllbEV3UjBWUmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDU2UlRCTlZFVXhUMVJCTTFkb1kwNU5hbGwzVG5wRk1FMVVTWGRQVkVFelYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVY0VjNOck1rWXhPRTVZYUM5bkszWXpPVkIyUTA5ME9ETlhWVFphWTBNMlExRXhaVGNLUkZoQlluaHlXWE54WjAxUWMwc3lPV0ZvV0hZMlVtNXVWM2RDYTNKMlVWSjJibkUyZFM5dk1IUkNjVEkzVm5WaGFYRlBRMEpsT0hkbloxaHlUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZOUjJSbUNrRlRZbHBrTVhobFEyVlFLM2MxU2pONlkyVmFLMjlGZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU9SMUUxV21wak5VOUhTVFZPVjFaclRXcG5kMWxxWXpOTlJFbDRDazE2UVRCUFIwazBXa1JPYlU1RVJURk5SR3hyVFVkTk5WcFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDVIVVRWYWFtTTFUMGRKTlU1WFZtdE5hbWQzV1dwak0wMUVTWGhOZWtFd1QwZEpORnBFVG0wS1RrUkZNVTFFYkd0TlIwMDFXbFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDVIVVRWYWFtTTFUMGRKTlU1WFZtdE5hbWQzV1dwak0wMUVTWGdLVFhwQk1FOUhTVFJhUkU1dFRrUkZNVTFFYkd0TlIwMDFXbFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUlNhMDlYV1RNS1QxUm9hVTlVVm14YVJFazBUVWRKTTA1NlFYbE5WRTEzVGtSb2FVOUhVWHBhYWxGNFRsUkJOVnBFUW1wUFYxVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOYW10NlRYcEJNazFxVFRST2FsVjJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwVVZsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTTjBKSWEwRUtaSGRDTVVGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtNHlRaXQyTkVWQlFVRlJSQXBCUlZsM1VrRkpaMU13U2l0a1owTkdUemhIYTI5c1RIaHdUR1IyVTFObE1raGhaRUkwSzI1WlVqTmpVM1pXTm05VlYydERTVUZuWVhVMmFFWjZaMVJ3Q25nck1tdExkMkV2WVcxd1dqSm5SR1ZPV1ZwWVJERlNiM2hCYzB0SmVVWTVUVUZ2UjBORGNVZFRUVFE1UWtGTlJFRXlhMEZOUjFsRFRWRkVSamQ1UldnS1ZrNDRiVU5zUTFoUlJHUklha0ZqTjFNcmR6ZFlUMEZ3U214TlJHRlpaWGxGYWxFeVNqQnZSV28yUzJsbE9GaEtVMjFTVVVod2RVcHJUelJEVFZGRVZRbzFRblJVY0RWR1ZVdHpOM3AxY1d0aWEzWkJNaXREV1dWNU1YZFJaRGwyYWxsUFRUazJTM2RPZDFwWGVFVnFTRzV1UzFOblFuZFJkR1ZUYTBKYU4zYzlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}],"timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQggMo3oHJOCf2rmunccY1LDtaPk2BaAsvKlblZ5XQkUewCFCWorMsrsn3R/LTbLd8stC67Y9fFGA8yMDI2MDcxNDExNTkwN1owAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwNzE0MTE1OTA3WjAvBgkqhkiG9w0BCQQxIgQg1mMvEx741OS29qLX+LwHA1kZBWUUE05jIoWjUwZHRuUwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIxAKnVMkqUShIQdAg6Le33DsY1YfYGVVik24mONsCazZ3A4u2wX0rRnay2rJ0dzH4g8wIwS3H6kXBIDgpF6uGM0CWEDtrF4+ZII4RYd5rFENiTVQmhCVMlpCpH1JFHuVgfzlO3"}]}},"messageSignature":{"messageDigest":{"algorithm":"SHA2_256","digest":"naX31KXtsHuy8yw8llVBus/UaejNl+J2tepDADQpYwg="},"signature":"MEYCIQC/jr1Rf5N+2B/yWPCtsBD+2+FY63EjQLmf7siVtN7zegIhAPYkbdZVlB+vSg+t6axWhGJTUf6mQKGPmPi0dbwfZMT+"}}

build/torch-cuda/tiled_mlp.py CHANGED Viewed

@@ -5,6 +5,7 @@ from typing import List
 from typing import Optional
 import torch
 from .utils import ensure_contiguous
@@ -44,6 +45,7 @@ class LigerTiledMLPFunction(torch.autograd.Function):
         ctx.fn = fn
         ctx.mlp_module = mlp_module
         ctx.shards = shards
         ctx.save_for_backward(x)
         # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
@@ -61,22 +63,33 @@ class LigerTiledMLPFunction(torch.autograd.Function):
         (x,) = ctx.saved_tensors
         mlp_module = ctx.mlp_module
         shards = ctx.shards
         x_requires_grad = x.requires_grad
         x = x.detach()
         # detach() unsets x.requires_grad, so restore it
         x.requires_grad_(x_requires_grad)
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
-        hidden_size = x.shape[-1]
         x_shape_orig = x.shape
         # flatten bs+seqlen to avoid having stride issues when narrowing into seqlen w/ bs>1
-        x = x.view(-1, hidden_size)
-        incoming_grad = grads[0].view(-1, hidden_size)
         x_grad = torch.zeros_like(x)
         x_shards = list(torch.chunk(x, chunks=shards, dim=0))
         for i, x_shard in enumerate(x_shards):
             x_shard.requires_grad_(x_requires_grad)
@@ -86,7 +99,12 @@ class LigerTiledMLPFunction(torch.autograd.Function):
             shard_offset = i * x_shards[0].shape[0]
             x_shard.grad = x_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
-            incoming_grad_shard = incoming_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
             with torch.enable_grad():
                 output = fn(mlp_module, x_shard)
@@ -114,9 +132,6 @@ def apply_tiled_mlp(
         x: the input tensor with shape [bs, seqlen, hidden_size] or [seqlen, hidden_size]
         num_shards: number of shards to use. If None, automatically calculated as ceil(seqlen / hidden_size)
         compute_params: list of parameters for DeepSpeed ZeRO optimization
-    Returns:
-        output tensor with the same shape as input
     """
     if num_shards is None:
         # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size]
@@ -127,6 +142,14 @@ def apply_tiled_mlp(
     # Ensure num_shards is at least 1
     num_shards = max(1, num_shards)
     return LigerTiledMLPFunction.apply(
         fn,
         mlp_module,

 from typing import Optional
 import torch
+import torch.distributed as dist
 from .utils import ensure_contiguous
         ctx.fn = fn
         ctx.mlp_module = mlp_module
         ctx.shards = shards
+        ctx.compute_params = compute_params
         ctx.save_for_backward(x)
         # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
         (x,) = ctx.saved_tensors
         mlp_module = ctx.mlp_module
         shards = ctx.shards
+        compute_params = ctx.compute_params
+        grad_output = grads[0]
         x_requires_grad = x.requires_grad
         x = x.detach()
         # detach() unsets x.requires_grad, so restore it
         x.requires_grad_(x_requires_grad)
+        # x.shape could be [bs, seqlen, in_features] or [seqlen, in_features] (moe experts)
+        in_features = x.shape[-1]
+        out_features = grad_output.shape[-1]
         x_shape_orig = x.shape
         # flatten bs+seqlen to avoid having stride issues when narrowing into seqlen w/ bs>1
+        # NOTE: input and output feature dimensions may differ, e.g.
+        # Linear(in_features, out_features), so flatten x and grad_output with their own last dims.
+        x = x.view(-1, in_features)
+        incoming_grad = grad_output.view(-1, out_features)
         x_grad = torch.zeros_like(x)
         x_shards = list(torch.chunk(x, chunks=shards, dim=0))
+        incoming_grad_shards = list(torch.chunk(incoming_grad, chunks=shards, dim=0))
+        # ZeRO-3 partitioned parameters carry a ds_id; collect them once so the per-shard loop only
+        # flips the ready flag. Parameters on other backends have no ds_id and are left untouched.
+        ds_params = [p for p in compute_params if hasattr(p, "ds_id")] if compute_params else []
         for i, x_shard in enumerate(x_shards):
             x_shard.requires_grad_(x_requires_grad)
             shard_offset = i * x_shards[0].shape[0]
             x_shard.grad = x_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
+            incoming_grad_shard = incoming_grad_shards[i]
+            # Defer DeepSpeed's reduction until the last shard has accumulated into param.grad; the flag
+            # is read by ZeRO's hook during each shard's backward, so it must be set per shard.
+            for param in ds_params:
+                param.ds_grad_is_ready = i + 1 == len(x_shards)
             with torch.enable_grad():
                 output = fn(mlp_module, x_shard)
         x: the input tensor with shape [bs, seqlen, hidden_size] or [seqlen, hidden_size]
         num_shards: number of shards to use. If None, automatically calculated as ceil(seqlen / hidden_size)
         compute_params: list of parameters for DeepSpeed ZeRO optimization
     """
     if num_shards is None:
         # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size]
     # Ensure num_shards is at least 1
     num_shards = max(1, num_shards)
+    # All ranks must run the same number of shards: a sharded-parameter backend (DeepSpeed ZeRO-3, FSDP)
+    # gathers weights inside each shard's recompute, so a rank that runs fewer shards stops participating
+    # in those collectives and deadlocks the others. Harmonize on the per-rank maximum.
+    if dist.is_available() and dist.is_initialized() and dist.get_world_size() > 1:
+        num_shards_tensor = torch.tensor(num_shards, device=x.device)
+        dist.all_reduce(num_shards_tensor, op=dist.ReduceOp.MAX)
+        num_shards = int(num_shards_tensor.item())
     return LigerTiledMLPFunction.apply(
         fn,
         mlp_module,

build/torch-rocm/__init__.py DELETED Viewed

@@ -1,4 +0,0 @@
-from . import layers
-from .layers import CrossEntropyOutput, LigerForCausalLMLoss
-__all__ = ["layers", "LigerForCausalLMLoss", "CrossEntropyOutput"]

build/torch-rocm/_ops.py DELETED Viewed

@@ -1,38 +0,0 @@
-import torch
-def get_backend() -> str:
-    """Detect the backend by inspecting torch."""
-    import torch
-    if hasattr(torch, "neuron"):
-        # Needs to be sorted before specific Torch builds, since Neuron
-        # extension can be loaded into e.g. CUDA Torch builds.
-        return "neuron"
-    elif torch.version.cuda is not None:
-        return "cuda"
-    elif torch.version.hip is not None:
-        return "rocm"
-    elif torch.backends.mps.is_available():
-        return "metal"
-    elif hasattr(torch.version, "xpu") and torch.version.xpu is not None:
-        return "xpu"
-    else:
-        return "cpu"
-def _find_ops_name() -> str:
-    kernel_name = "liger_kernels"
-    unique_id = "ab435e2"
-    backend = get_backend()
-    return f"_{kernel_name}_{backend}_{unique_id}"
-_OPS_NAME = _find_ops_name()
-ops = getattr(torch.ops, _OPS_NAME)
-def add_op_namespace_prefix(op_name: str) -> str:
-    """
-    Prefix op by namespace.
-    """
-    return f"{_OPS_NAME}::{op_name}"

build/torch-rocm/cross_entropy.py DELETED Viewed

@@ -1,558 +0,0 @@
-import operator
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import element_mul_kernel
-from .utils import is_hip
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.jit
-def liger_cross_entropy_kernel(
-    X_ptr,
-    X_stride,
-    Y_ptr,
-    Y_stride,
-    weight_ptr,
-    loss_ptr,
-    z_loss_ptr,
-    loss_stride,
-    token_accuracy_ptr,
-    token_accuracy_stride,
-    predicted_tokens_ptr,
-    predicted_tokens_stride,
-    n_cols,
-    n_non_ignore,
-    sum_non_ignore_weight,
-    weight_sum,
-    ignore_index,
-    lse_square_scale: tl.constexpr,
-    label_smoothing: tl.constexpr,
-    reduction: tl.constexpr,  # set it as constexpr since reduction is always known at compile time
-    softcap,
-    RETURN_Z_LOSS: tl.constexpr,
-    RETURN_TOKEN_ACCURACY: tl.constexpr,
-    RETURN_PREDICTED_TOKENS: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    HAS_WEIGHT: tl.constexpr,
-    HAS_SOFTCAPPING: tl.constexpr,
-    HAS_GRADIENTS: tl.constexpr,
-):
-    """
-    This kernel computes both cross entropy loss and the gradient of the input.
-    We only consider hard label + mean reduction for now. Please refer to https://pytorch.org/docs/stable/generated/torch.nn.CrossEntropyLoss.html for the math.
-    Parameters:
-    X_ptr: Pointer to input tensor.
-    X_stride (int): The stride of the input tensor.
-    Y_ptr: Pointer to target tensor.
-    Y_stride (int): The stride of the target tensor.
-    weight_ptr: Pointer to weight tensor.
-    loss_ptr: Pointer to tensor to store the loss.
-    z_loss_ptr: Pointer to tensor to store the z loss. No operation if RETURN_Z_LOSS is 0.
-    loss_stride (int): The stride of the loss tensor.
-    token_accuracy_ptr: Pointer to tensor to store the per-token accuracy. No operation if RETURN_TOKEN_ACCURACY is 0.
-    token_accuracy_stride (int): The stride of the token accuracy tensor.
-    n_cols (int): The number of columns in the input tensor.
-    n_non_ignore (float): The number of non-ignored elements in the batch.
-    sum_non_ignore_weight (float): The sum of non-ignored target's weights in the batch.
-    weight_sum (float): The sum of weight tensor.
-    ignore_index (int): The index to ignore in the target.
-    label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-    lse_square_scale (float): The scaler of (logsumexp(_input)) ^ 2 adding to the loss for the stability of training.
-    reduction (str): The string for the reduction to apply
-    softcap (float): The upper threshold for scaling logits to the range (-softcap, +softcap).
-    RETURN_Z_LOSS (int): The boolean value to decide whether to store z loss to z_loss_ptr or not. It must be 0 or 1.
-    RETURN_TOKEN_ACCURACY (int): The boolean value to decide whether to store per-token accuracy to token_accuracy_ptr or not. It must be 0 or 1.
-    BLOCK_SIZE (int): The block size for Triton operations.
-    HAS_WEIGHT (bool): The boolean value to determine whether assigning weight to each of the classes.
-    HAS_SOFTCAPPING (bool): The boolean value to determine whether applying soft-capping or not.
-    HAS_GRADIENTS (bool): The boolean value to determine whether calculating gradients in forward pass.
-    """
-    # https://github.com/triton-lang/triton/issues/1058
-    # If B*T*V is too large, program_id * stride will overflow out of int32, so we convert to int64
-    program_id = tl.program_id(0).to(tl.int64)
-    # 1. Load Y_ptr first because if the target is ignore_index, we can return right away
-    Y_ptr += program_id * Y_stride
-    y = tl.load(Y_ptr)
-    # 2. locate the start index
-    X_ptr += program_id * X_stride
-    if y == ignore_index:
-        # set all X_ptr as 0
-        for i in range(0, n_cols, BLOCK_SIZE):
-            X_offsets = i + tl.arange(0, BLOCK_SIZE)
-            tl.store(X_ptr + X_offsets, 0.0, mask=X_offsets < n_cols)
-        # For ignored tokens, set token accuracy to 0
-        if RETURN_TOKEN_ACCURACY:
-            token_accuracy_ptr += program_id * token_accuracy_stride
-            tl.store(token_accuracy_ptr, 0.0)
-        if RETURN_PREDICTED_TOKENS:
-            predicted_tokens_ptr += program_id * predicted_tokens_stride
-            tl.store(predicted_tokens_ptr, -1)
-        return
-    loss_ptr += program_id * loss_stride
-    if RETURN_Z_LOSS:
-        z_loss_ptr += program_id * loss_stride
-    if RETURN_TOKEN_ACCURACY:
-        token_accuracy_ptr += program_id * token_accuracy_stride
-    if RETURN_PREDICTED_TOKENS:
-        predicted_tokens_ptr += program_id * predicted_tokens_stride
-    if HAS_WEIGHT:
-        weight_y = tl.load(weight_ptr + y).cast(tl.float32)
-    # Online softmax: 2 loads + 1 store (compared with 3 loads + 1 store for the safe softmax)
-    # Refer to Algorithm 3 in the paper: https://arxiv.org/pdf/1805.02867
-    # 3. [Online softmax] first pass: find max + sum
-    m = float("-inf")  # m is the max value. use the notation from the paper
-    d = 0.0  # d is the sum. use the notation from the paper
-    argmax_idx = 0  # Track the index of the maximum value for token accuracy / predicted tokens computation
-    ori_X_y = tl.load(X_ptr + y).cast(tl.float32)  # we need to store the original value of X_y for the loss calculation
-    if HAS_SOFTCAPPING:
-        ori_X_y = softcap * tanh(ori_X_y / softcap)
-    # Label smoothing is a general case of normal cross entropy
-    # See the full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issue-2503665310
-    scaled_x_sum = 0.0
-    eps = label_smoothing / n_cols
-    for i in range(0, n_cols, BLOCK_SIZE):
-        X_offsets = i + tl.arange(0, BLOCK_SIZE)
-        X_block = tl.load(
-            X_ptr + X_offsets,
-            mask=X_offsets < n_cols,
-            other=float("-inf"),
-            # Ensure float32 precision for softmax calculation
-        ).cast(tl.float32)
-        if HAS_SOFTCAPPING:
-            X_block = softcap * tanh(X_block / softcap)
-        block_max = tl.max(X_block)
-        # Track argmax for accuracy / predicted tokens computation
-        if RETURN_TOKEN_ACCURACY or RETURN_PREDICTED_TOKENS:
-            # Find the index of the maximum value in this block
-            is_max_mask = X_block == block_max
-            # Mask out invalid indices with a value larger than n_cols
-            masked_offsets = tl.where(is_max_mask, X_offsets, n_cols)
-            # Get the first (smallest) index where max occurs
-            current_block_argmax_idx = tl.min(masked_offsets)
-            is_new_max = block_max > m
-            argmax_idx = tl.where(is_new_max, current_block_argmax_idx, argmax_idx)
-        if label_smoothing > 0:
-            # scale X beforehand to avoid overflow
-            if HAS_WEIGHT:
-                weight_block = tl.load(weight_ptr + X_offsets, mask=X_offsets < n_cols)
-                scaled_x_sum += tl.sum(tl.where(X_offsets < n_cols, -eps * X_block * weight_block, 0.0))
-            else:
-                scaled_x_sum += tl.sum(tl.where(X_offsets < n_cols, -eps * X_block, 0.0))
-        m_new = tl.maximum(m, block_max)
-        d = d * tl.exp(m - m_new) + tl.sum(tl.exp(X_block - m_new))
-        m = m_new
-    # log (sum(e^(X_i))) = log (sum(e ^ (max(X) * e ^ (X_i - max(X)))))
-    #                    = log (e^(max(X)) * sum(e ^ (X_i - max(X))))
-    #                    = max(X) + log (sum(e ^ (X_i - max(X)))) = m + log d
-    lse = m + tl.log(d)
-    # 4. [Online Softmax] Second pass: compute gradients
-    # For 'mean' reduction, gradients are normalized by number of non-ignored elements (N)
-    # dx_y = (softmax(x_y) - 1) / N
-    # dx_i = softmax(x_i) / N, i != y
-    # For label smoothing:
-    # dx_i = (softmax(x_i) - label_smoothing / V) / N, V = n_cols, i != y
-    # dx_y = (softmax(x_y) - label_smoothing / V - (1 - label_smoothing)) / N
-    #      = dx_i - (1 - label_smoothing) / N
-    # With Z loss:
-    # dx_i = ((1 + 2 * lse_square_scale * lse) * softmax(x_i) - label_smoothing / V) / N, i != y
-    # dx_y = dx_i - (1 - label_smoothing) / N
-    # For 'sum' reduction, no normalization is applied:
-    # dx_y = softmax(x_y) - 1
-    # dx_i = softmax(x_i), for i ≠ y
-    if HAS_GRADIENTS:
-        for i in range(0, n_cols, BLOCK_SIZE):
-            X_offsets = i + tl.arange(0, BLOCK_SIZE)
-            X_block = tl.load(
-                X_ptr + X_offsets,
-                mask=X_offsets < n_cols,
-                other=float("-inf"),
-                # Ensure float32 precision for softmax calculation
-            ).cast(tl.float32)
-            if HAS_SOFTCAPPING:
-                intermediate = tanh(X_block / softcap)
-                X_block = softcap * intermediate
-            if not HAS_WEIGHT:
-                # softmax(x_i)
-                X_block = tl.exp(X_block - m) / d
-                # derivative of z-loss: 2 * lse_square_scale * lse * softmax(x_i)
-                X_block += 2 * lse_square_scale * lse * X_block
-                # smoothing term
-                X_block += -eps
-                # special handle dx_y
-                X_block = tl.where(X_offsets != y, X_block, X_block - (1 - label_smoothing))
-                # reduction scale
-                if reduction == "mean":
-                    X_block = X_block / n_non_ignore
-            else:
-                weight_block = tl.load(weight_ptr + X_offsets, mask=X_offsets < n_cols)
-                softmax_X = tl.exp(X_block - m) / d
-                # derivative of original_loss
-                dloss_ori = (1 - label_smoothing) * softmax_X
-                # specially handle dx_y
-                dloss_ori = tl.where(X_offsets != y, dloss_ori, dloss_ori - (1 - label_smoothing))
-                dloss_ori = dloss_ori * weight_y
-                # derivative of smooth_loss
-                dloss_smooth = eps * (-weight_block + softmax_X * weight_sum)
-                # derivative of z-loss
-                dz_loss = 2 * lse_square_scale * lse * softmax_X
-                # reduction scale
-                if reduction == "mean":
-                    dloss_ori = dloss_ori / sum_non_ignore_weight
-                    dloss_smooth = dloss_smooth / sum_non_ignore_weight
-                    # TODO: Implement weighted z_loss. Currently, z_loss is not scaled by weight.
-                    dz_loss = dz_loss / n_non_ignore
-                # derivative of total_loss
-                X_block = dloss_ori + dloss_smooth + dz_loss
-            # chain rule softcapping
-            # d(softcap * tanh(x / softcap)) = (1 - tanh^2(x / softcap))
-            if HAS_SOFTCAPPING:
-                X_block = X_block * (1 - intermediate * intermediate)
-            tl.store(X_ptr + X_offsets, X_block, mask=X_offsets < n_cols)
-    # We need tl.debug_barrier() to ensure the new result of X_ptr is written as mentioned in
-    # https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/ops/cross_entropy.py#L34
-    tl.debug_barrier()
-    # 5. Calculate the loss
-    # loss = log (softmax(X_y)) = log ((e ^ (X_y - max(X)) / sum(e ^ (X - max(X))))
-    #      = (X_y - max(X)) - log(sum(e ^ (X - max(X))))
-    #      = X_y - m - log d = X_y - lse
-    # sum(e ^ (X - max(X))) must >= 1 because the max term is e ^ 0 = 1
-    # So we can safely calculate log (softmax(X_y)) without overflow
-    loss = lse - ori_X_y
-    if HAS_WEIGHT:
-        loss = weight_y * loss
-    # Original loss = H(q, p),  with label smoothing regularization = H(q', p) and (label_smoothing / V) = eps
-    # H(q', p) = (1 - label_smoothing) * H(q, p) + label_smoothing * H(u, p)
-    #          = (1 - label_smoothing) * H(q, p) + eps * sum(logsoftmax(x_i))
-    # By using m (global max of xi) and d (sum of e^(xi-m)), we can simplify as:
-    #          = (1 - label_smoothing) * H(q, p) + (sum(-eps * x_i) + label_smoothing * (m + logd))
-    # Refer to H(q', p) in section 7 of the paper: https://arxiv.org/pdf/1512.00567
-    # pytorch: https://github.com/pytorch/pytorch/blob/2981534f54d49fa3a9755c9b0855e7929c2527f0/aten/src/ATen/native/LossNLL.cpp#L516
-    # See full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issuecomment-2333753087
-    if label_smoothing > 0:
-        if HAS_WEIGHT:
-            smooth_loss = scaled_x_sum + eps * lse * weight_sum
-        else:
-            smooth_loss = scaled_x_sum + label_smoothing * lse
-        loss = loss * (1 - label_smoothing) + smooth_loss
-    # An auxiliary loss, z_loss
-    # Refer to Page14 Loss function section in the paper PaLM: https://www.jmlr.org/papers/v24/22-1144.html
-    z_loss = lse_square_scale * lse * lse
-    # Normalize the loss by the number of non-ignored elements if reduction is "mean"
-    if reduction == "mean":
-        if HAS_WEIGHT:
-            loss = loss / sum_non_ignore_weight
-        else:
-            loss = loss / n_non_ignore
-        # TODO: Implement weighted z_loss. Currently, z_loss is not scaled by weight.
-        z_loss = z_loss / n_non_ignore
-    loss += z_loss
-    tl.store(loss_ptr, loss)
-    if RETURN_Z_LOSS:
-        tl.store(z_loss_ptr, z_loss)
-    if RETURN_TOKEN_ACCURACY:
-        # Store 1.0 if prediction is correct, 0.0 otherwise
-        is_correct = 1.0 if argmax_idx == y else 0.0
-        tl.store(token_accuracy_ptr, is_correct)
-    if RETURN_PREDICTED_TOKENS:
-        tl.store(predicted_tokens_ptr, argmax_idx)
-# The hard limit of TRITON_MAX_TENSOR_NUMEL is 1048576 https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/language/core.py#L19
-# However, setting limit as 65536 as in LayerNorm tutorial is faster because of less register spilling
-# The optimal maximum block size depends on your hardware, your kernel, and your dtype
-# the best size we found by manually tuning on xpu and npu.
-if infer_device() == "xpu":
-    MAX_FUSED_SIZE = 4096
-elif infer_device() == "npu":
-    MAX_FUSED_SIZE = 2048
-else:
-    MAX_FUSED_SIZE = 65536 // 2
-def cross_entropy_forward(
-    _input,
-    target,
-    weight,
-    ignore_index,
-    lse_square_scale,
-    label_smoothing,
-    reduction,
-    softcap,
-    return_z_loss,
-    return_token_accuracy=False,
-    return_predicted_tokens=False,
-):
-    assert isinstance(return_z_loss, bool), f"return_z_loss must be True or False. Got: {return_z_loss}"
-    assert isinstance(return_token_accuracy, bool), (
-        f"return_token_accuracy must be True or False. Got: {return_token_accuracy}"
-    )
-    assert isinstance(return_predicted_tokens, bool), (
-        f"return_predicted_tokens must be True or False. Got: {return_predicted_tokens}"
-    )
-    BT, V = _input.shape
-    n_rows = BT
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    # unreduced loss
-    loss_1d = torch.zeros(n_rows, dtype=_input.dtype, device=_input.device)
-    z_loss_1d = torch.zeros(n_rows, dtype=_input.dtype, device=_input.device) if return_z_loss else None
-    token_accuracy_1d = (
-        torch.zeros(n_rows, dtype=torch.float32, device=_input.device) if return_token_accuracy else None
-    )
-    predicted_tokens_1d = (
-        torch.full((n_rows,), -1, dtype=torch.int64, device=_input.device) if return_predicted_tokens else None
-    )
-    target_mask = target != ignore_index
-    n_non_ignore = target_mask.sum().item()
-    assert (target * target_mask).max() < _input.shape[-1], (
-        f"Target {target.max()} is out of bounds. Expected < {_input.shape[-1]}"
-    )
-    assert (target * target_mask).min() >= 0, f"Target {target.min()} is out of bounds. Expected >= 0"
-    sum_non_ignore_weight = n_non_ignore
-    weight_sum = 0.0
-    if weight is not None:
-        assert weight.shape[0] == V, f"If given, weight has to be a Tensor of size V. Got: {weight.shape}"
-        assert torch.is_floating_point(weight), (
-            f"If given, weight has to be a Tensor of floating point dtype. Got: {weight.dtype}"
-        )
-        sum_non_ignore_weight = torch.gather(weight, dim=0, index=target.masked_select(target_mask)).sum().item()
-        weight_sum = weight.sum().item()
-        # ensure weight is contiguous
-        if weight.stride(-1) != 1:
-            weight = weight.contiguous()
-    # ensure _input and target are contiguous in the last dimension
-    if _input.stride(-1) != 1:
-        _input = _input.contiguous()
-    if target.stride(-1) != 1:
-        target = target.contiguous()
-    # Here we use a trick to store X_ptr gradient in X_ptr so we can save memory
-    liger_cross_entropy_kernel[(n_rows,)](
-        X_ptr=_input,
-        X_stride=_input.stride(-2),
-        Y_ptr=target,
-        Y_stride=target.stride(-1),  # always 1
-        weight_ptr=weight,  # dummy if None
-        loss_ptr=loss_1d,
-        z_loss_ptr=z_loss_1d,
-        loss_stride=loss_1d.stride(-1),  # always 1
-        token_accuracy_ptr=token_accuracy_1d,
-        token_accuracy_stride=token_accuracy_1d.stride(-1)
-        if return_token_accuracy
-        else 0,  # always 1 if accuracy is enabled
-        predicted_tokens_ptr=predicted_tokens_1d,
-        predicted_tokens_stride=predicted_tokens_1d.stride(-1)
-        if return_predicted_tokens
-        else 0,  # always 1 if predicted tokens is enabled
-        n_cols=V,
-        n_non_ignore=n_non_ignore,
-        sum_non_ignore_weight=sum_non_ignore_weight,
-        ignore_index=ignore_index,
-        weight_sum=weight_sum,
-        lse_square_scale=lse_square_scale,
-        label_smoothing=label_smoothing,
-        reduction=reduction,
-        softcap=softcap,
-        RETURN_Z_LOSS=return_z_loss,
-        RETURN_TOKEN_ACCURACY=return_token_accuracy,
-        RETURN_PREDICTED_TOKENS=return_predicted_tokens,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_WEIGHT=True if weight is not None else False,
-        HAS_SOFTCAPPING=True if softcap is not None else False,
-        HAS_GRADIENTS=_input.requires_grad,
-        # TODO: 32 seems to give the best performance
-        # Performance is quite sensitive to num_warps
-        num_warps=32 if not is_hip() else 16,
-    )
-    if reduction == "none":
-        loss = loss_1d
-        z_loss = z_loss_1d if return_z_loss else None
-        token_accuracy = token_accuracy_1d if return_token_accuracy else None
-    else:
-        loss = torch.sum(loss_1d)
-        z_loss = torch.sum(z_loss_1d) if return_z_loss else None
-        # For accuracy, we compute the mean across all non-ignored tokens
-        token_accuracy = torch.sum(token_accuracy_1d) / n_non_ignore if return_token_accuracy else None
-    predicted_tokens = predicted_tokens_1d if return_predicted_tokens else None
-    return loss, z_loss, token_accuracy, predicted_tokens, _input
-def cross_entropy_backward(_input, grad_output):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul to save time
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        pass
-    # If reduction is 'none'
-    elif grad_output.ndim > 0:
-        _input = _input * grad_output.unsqueeze(dim=1)
-    # If reduction is ['mean', 'sum'], grad_output is just a scalar
-    # We use a Triton kernel instead of a PyTorch operation because modifying inputs in-place
-    # for gradient storage and backward multiple times causes anomalies with PyTorch but not with Triton.
-    else:
-        BT, V = _input.shape
-        n_rows = BT
-        BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-        element_mul_kernel[(n_rows,)](
-            _input,
-            _input.stride(-2),
-            grad_output,
-            V,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-    return _input
-class LigerCrossEntropyFunction(torch.autograd.Function):
-    """
-    This class implements a custom autograd function for the Liger Cross Entropy loss.
-    It overrides the forward and backward methods of the torch.autograd.Function class.
-    """
-    @staticmethod
-    def forward(
-        ctx,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        weight: Optional[torch.FloatTensor],
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-        return_z_loss: bool = False,
-        return_token_accuracy: bool = False,
-        return_predicted_tokens: bool = False,
-    ):
-        """
-        The forward pass of the Liger Cross Entropy loss.
-        Parameters:
-        ctx : The context object.
-        _input (tensor): The input tensor of shape (BT, V) where B is batch size, T is sequence length, V is vocab size.
-        target (tensor): The target tensor of shape (BT) where each value is in [0, V-1].
-        weight(Tensor, optional): a manual rescaling weight given to each class. If given, has to be a Tensor of size V and floating point dtype
-        ignore_index (int): The index to ignore in the target.
-        lse_square_scale (float): The scaler of (logsumexp(_input)) ^ 2 adding to the loss for the stability of training.
-        label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-        reduction (str): The reduction to apply to the output: "none" | "mean | "sum".
-        softcap (Optional[float]): The upper threshold for scaling logits to the range (-softcap, +softcap).
-        return_z_loss (bool): When `return_z_loss` is `True`, returns (loss, z_loss, token_accuracy, predicted_tokens) instead of (loss, None, None, None). Default: `False`
-        return_token_accuracy (bool): When `return_token_accuracy` is `True`, computes and returns per-token accuracy without materializing logits. Default: `False`
-        return_predicted_tokens (bool): When `return_predicted_tokens` is `True`, returns per-token predicted class indices (argmax) without materializing logits. Default: `False`
-        Returns:
-        tuple: A tuple with the computed losses, accuracy, and predicted tokens: (loss, z_loss, token_accuracy, predicted_tokens). z_loss, token_accuracy, and predicted_tokens are None if not requested.
-        """
-        input_requires_grad = _input.requires_grad
-        loss, z_loss, token_accuracy, predicted_tokens, _input = cross_entropy_forward(
-            _input,
-            target,
-            weight,
-            ignore_index,
-            lse_square_scale,
-            label_smoothing,
-            reduction,
-            softcap,
-            return_z_loss,
-            return_token_accuracy,
-            return_predicted_tokens,
-        )
-        # TODO: investigation
-        # If we don't detach the _input tensor, the memory will double
-        # Not sure why but seems that there will be a time both grad and value exist but in different location
-        if input_requires_grad:
-            ctx.save_for_backward(_input.detach())
-        ctx.return_z_loss = return_z_loss
-        ctx.return_token_accuracy = return_token_accuracy
-        ctx.return_predicted_tokens = return_predicted_tokens
-        return loss, z_loss, token_accuracy, predicted_tokens
-    @staticmethod
-    def backward(ctx, grad_output, grad_output2, grad_output3, grad_output4):
-        """
-        The backward pass of the Liger Cross Entropy loss.
-        Parameters:
-        ctx : The context object with saved tensors.
-        grad_output (tensor): The tensor containing the gradient of the loss with respect to the output.
-        grad_output2 (tensor): No use. Gradient for z_loss (not used as z_loss is only for logging).
-        grad_output3 (tensor): No use. Gradient for token_accuracy (not used as token_accuracy is only for metrics).
-        grad_output4 (tensor): No use. Gradient for predicted_tokens (not used as predicted_tokens is only for metrics).
-        Returns:
-        tuple: A tuple with the gradients with respect to the inputs. The elements are tensors or None.
-        """
-        if ctx.return_z_loss:
-            del grad_output2  # z_loss is only for logging
-        if ctx.return_token_accuracy:
-            del grad_output3  # token_accuracy is only for metrics
-        if ctx.return_predicted_tokens:
-            del grad_output4  # predicted_tokens is only for metrics
-        (_input,) = ctx.saved_tensors
-        _input = cross_entropy_backward(_input, grad_output)
-        return (
-            _input,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-rocm/dyt.py DELETED Viewed

@@ -1,164 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.autotune(
-    configs=[
-        triton.Config({"BLOCK_N": bn}, num_stages=ns, num_warps=nw)
-        for bn in [1024, 2048, 4096]
-        for ns in [1, 2]
-        for nw in [4, 8, 16]
-    ],
-    key=["N"],
-)
-@triton.jit
-def _dyt_fwd_kernel(X, Y, Alpha, Gamma, Beta, HAVE_BETA: tl.constexpr, N: tl.constexpr, BLOCK_N: tl.constexpr):
-    col = tl.cast(tl.program_id(0), tl.int64) * BLOCK_N + tl.arange(0, BLOCK_N)
-    mask = col < N
-    row_id = tl.cast(tl.program_id(1), tl.int64)
-    X += row_id * N
-    Y += row_id * N
-    alpha = tl.load(Alpha).to(tl.float32)
-    gamma = tl.load(Gamma + col, mask=mask, other=0.0).to(tl.float32)
-    x = tl.load(X + col, mask=mask, other=0.0).to(tl.float32)
-    tanh_x = tanh(alpha * x)
-    y = tanh_x * gamma
-    if HAVE_BETA:
-        beta = tl.load(Beta + col, mask=mask, other=0.0).to(tl.float32)
-        y += beta
-    tl.store(Y + col, y, mask=mask)
-@triton.autotune(
-    configs=[
-        triton.Config({"BLOCK_N": bn}, num_stages=ns, num_warps=nw)
-        for bn in [1024, 2048, 4096]
-        for ns in [1, 2]
-        for nw in [4, 8, 16]
-    ],
-    key=["N"],
-    # DA is indexed by program_id(0), so different BLOCK_N configs write to
-    # different slot counts per SM. Autotune trials don't zero outputs between
-    # runs, so stale slots from a prior trial would leak into da.sum(). Reset
-    # DA between trials to isolate each config's writes.
-    reset_to_zero=["DA"],
-)
-@triton.jit
-def _dyt_bwd_kernel(
-    DY, DX, DA, DG, DB, X, Alpha, Gamma, HAVE_BETA: tl.constexpr, M, N: tl.constexpr, BLOCK_N: tl.constexpr
-):
-    col = tl.cast(tl.program_id(0), tl.int64) * BLOCK_N + tl.arange(0, BLOCK_N)
-    mask = col < N
-    start_row_id = tl.cast(tl.program_id(1), tl.int64)
-    alpha = tl.load(Alpha).to(tl.float32)
-    da = 0.0
-    gamma = tl.load(Gamma + col, mask=mask, other=0.0).to(tl.float32)
-    dg = tl.zeros((BLOCK_N,), dtype=tl.float32)
-    if HAVE_BETA:
-        db = tl.zeros((BLOCK_N,), dtype=tl.float32)
-    for row_id in range(start_row_id, M, tl.num_programs(1)):
-        x = tl.load(X + row_id * N + col, mask=mask, other=0.0).to(tl.float32)
-        dy = tl.load(DY + row_id * N + col, mask=mask, other=0.0).to(tl.float32)
-        tanh_x = tanh(alpha * x)
-        if HAVE_BETA:
-            db += dy
-        dg += dy * tanh_x
-        tmp = (1 - tanh_x * tanh_x) * dy * gamma
-        da += tl.sum(x * tmp, 0)
-        dx = alpha * tmp
-        tl.store(DX + row_id * N + col, dx, mask=mask)
-    tl.store(DG + start_row_id * N + col, dg, mask=mask)
-    if HAVE_BETA:
-        tl.store(DB + start_row_id * N + col, db, mask=mask)
-    tl.store(DA + start_row_id * tl.cdiv(N, 512) + tl.program_id(0), da)
-def liger_dyt_fwd(x, alpha, gamma, beta):
-    assert x.is_contiguous()
-    HAVE_BETA = True if beta is not None else False
-    input_shape = x.shape
-    x = x.view(-1, input_shape[-1])
-    M, N = x.shape
-    y = torch.empty_like(x)
-    grid = lambda meta: (triton.cdiv(N, meta["BLOCK_N"]), M)
-    _dyt_fwd_kernel[grid](
-        x,
-        y,
-        alpha,
-        gamma,
-        beta,
-        HAVE_BETA,
-        N,
-    )
-    return y.view(input_shape)
-def liger_dyt_bwd(dy, x, alpha, gamma, beta):
-    assert dy.is_contiguous()
-    input_shape = x.shape
-    x = x.view(-1, input_shape[-1])
-    M, N = x.shape
-    HAVE_BETA = True if beta is not None else False
-    device = infer_device()
-    if device == "cuda":
-        NUM_SMS = torch.cuda.get_device_properties(x.device).multi_processor_count
-    elif device == "xpu":
-        NUM_SMS = torch.xpu.get_device_properties(x.device).gpu_subslice_count
-    elif device == "npu":
-        NUM_SMS = get_npu_core_count()
-    da = torch.zeros(NUM_SMS, triton.cdiv(N, 512), dtype=torch.float32, device=x.device)
-    dg = torch.empty(NUM_SMS, N, dtype=torch.float32, device=x.device)
-    db = torch.empty(NUM_SMS, N, dtype=torch.float32, device=x.device) if HAVE_BETA else None
-    dx = torch.empty_like(dy)
-    grid = lambda meta: (triton.cdiv(N, meta["BLOCK_N"]), NUM_SMS)
-    _dyt_bwd_kernel[grid](dy, dx, da, dg, db, x, alpha, gamma, HAVE_BETA, M, N)
-    if HAVE_BETA:
-        db = db.sum(0).to(x.dtype)
-    dg = dg.sum(0).to(gamma.dtype)
-    da = da.sum().to(x.dtype).unsqueeze(0)
-    return dx.view(input_shape), da, dg, db
-class LigerDyTFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, x, alpha, gamma, beta):
-        y = liger_dyt_fwd(x, alpha, gamma, beta)
-        ctx.save_for_backward(x, alpha, gamma, beta)
-        return y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dy):
-        x, alpha, gamma, beta = ctx.saved_tensors
-        dx, dalpha, dgamma, dbeta = liger_dyt_bwd(dy, x, alpha, gamma, beta)
-        return dx, dalpha, dgamma, dbeta

build/torch-rocm/fused_linear_cross_entropy.py DELETED Viewed

@@ -1,400 +0,0 @@
-import torch
-import triton
-from .cross_entropy import liger_cross_entropy_kernel
-from .utils import amp_custom_bwd
-from .utils import amp_custom_fwd
-from .utils import element_mul_kernel
-from .utils import is_hip
-from .utils import infer_device
-# The hard limit of TRITON_MAX_TENSOR_NUMEL is 1048576 https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/language/core.py#L19
-# However, setting limit as 65536 as in LayerNorm tutorial is faster because of less register spilling
-# The optimal maximum block size depends on your hardware, your kernel, and your dtype
-MAX_FUSED_SIZE = 2048 if infer_device() == "npu" else 65536 // 2
-def fused_linear_cross_entropy_forward(
-    _input,
-    weight,
-    target,
-    ce_weight=None,
-    bias=None,
-    ignore_index=-100,
-    lse_square_scale=0.0,
-    label_smoothing=0.0,
-    reduction="mean",
-    softcap=None,
-    return_z_loss=False,
-    accum_dtype=None,
-    use_token_scaling=False,
-    return_token_accuracy=False,
-    return_predicted_tokens=False,
-):
-    assert isinstance(return_z_loss, bool), f"return_z_loss must be True or False. Got: {return_z_loss}"
-    assert isinstance(return_token_accuracy, bool), (
-        f"return_token_accuracy must be True or False. Got: {return_token_accuracy}"
-    )
-    assert isinstance(return_predicted_tokens, bool), (
-        f"return_predicted_tokens must be True or False. Got: {return_predicted_tokens}"
-    )
-    device = _input.device
-    input_requires_grad = _input.requires_grad
-    # inputs have shape: BT x H
-    # materialized activations will have shape: BT x V
-    # the increase in memory = BT x V
-    # reduction can be achieved by partitioning the number of tokens BT into smaller chunks.
-    # for ex: if we were to achieve the same memory consumption as BT x H, then the chunk size should be:
-    # inc_factor = (V+H-1)//H, chunk_size = (BT + inc_factor - 1)//inc_factor
-    # for ex: BT = 4096*4, V = 32000, H = 4096 ==> inc_factor = 8, chunk_size = 2048
-    BT, H = _input.shape
-    V = weight.shape[0]
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    inc_factor = triton.cdiv(V, H)  # (V + H - 1) // H
-    chunk_size = triton.next_power_of_2(triton.cdiv(BT, inc_factor))  # (BT + inc_factor - 1) // inc_factor
-    num_chunks = triton.cdiv(BT, chunk_size)  # (BT + chunk_size - 1) // chunk_size
-    grad_input = torch.zeros_like(_input, device=device)
-    # we use fp32 for loss and gradients accumulator
-    if input_requires_grad:
-        if accum_dtype is None:
-            grad_weight = torch.zeros_like(weight, device=device) if weight.requires_grad else None
-            grad_bias = torch.zeros_like(bias, device=device) if bias is not None else None
-        else:
-            grad_weight = torch.zeros_like(weight, dtype=accum_dtype, device=device) if weight.requires_grad else None
-            grad_bias = torch.zeros_like(bias, dtype=accum_dtype, device=device) if bias is not None else None
-    else:
-        grad_weight = None
-        grad_bias = None
-    loss_1d = torch.zeros(BT, dtype=torch.float32, device=device)
-    z_loss_1d = torch.zeros(BT, dtype=_input.dtype, device=_input.device) if return_z_loss else None
-    token_accuracy_1d = torch.zeros(BT, dtype=torch.float32, device=device) if return_token_accuracy else None
-    predicted_tokens_1d = torch.full((BT,), -1, dtype=torch.int64, device=device) if return_predicted_tokens else None
-    # TODO: evaluate how CUDA synchronization caused by .item() affects the speed
-    target_mask = target != ignore_index
-    total_n_non_ignore = target_mask.sum().item()
-    total_sum_non_ignore_ce_weight = total_n_non_ignore
-    ce_weight_sum = 0.0
-    if ce_weight is not None:
-        assert ce_weight.shape[0] == V, f"If given, weight has to be a Tensor of size V. Got: {ce_weight.shape}"
-        assert torch.is_floating_point(ce_weight), (
-            f"If given, weight has to be a Tensor of floating point dtype. Got: {ce_weight.dtype}"
-        )
-        total_sum_non_ignore_ce_weight = (
-            torch.gather(ce_weight, dim=0, index=target.masked_select(target_mask)).sum().item()
-        )
-        ce_weight_sum = ce_weight.sum().item()
-        if ce_weight.stride(-1) != 1:
-            ce_weight = ce_weight.contiguous()
-    for chunk_id in range(num_chunks):
-        start_idx = chunk_id * chunk_size
-        end_idx = min((chunk_id + 1) * chunk_size, BT)
-        _input_chunk = _input[start_idx:end_idx]  # chunk_size x H
-        # when doing matmul, use the original precision
-        logits_chunk = _input_chunk @ weight.t()  # chunk_size x V
-        if bias is not None:
-            logits_chunk = logits_chunk + bias
-        target_chunk = target[start_idx:end_idx]  # chunk_size,
-        n_rows = logits_chunk.shape[0]
-        # Compute predicted probabilities for token scaling if needed
-        if use_token_scaling:
-            # Compute softmax probabilities for scaling
-            # We need to compute this before the cross entropy kernel modifies logits_chunk
-            logits_for_softmax = logits_chunk.detach().clone()  # Detach to avoid gradient flow
-            if softcap is not None:
-                logits_for_softmax = softcap * torch.tanh(logits_for_softmax / softcap)
-            # Compute softmax to get predicted probabilities
-            probs = torch.softmax(logits_for_softmax, dim=-1)
-            # Get predicted probabilities for token scaling, handling ignored targets
-            valid_target_mask = target_chunk != ignore_index
-            valid_targets = target_chunk[valid_target_mask]
-            if len(valid_targets) > 0:
-                # Gather probabilities only for valid targets
-                valid_probs = probs[valid_target_mask]
-                pred_probs_valid = torch.gather(valid_probs, -1, valid_targets.unsqueeze(-1)).squeeze(-1)
-                # Create full tensor with zeros for ignored targets
-                pred_probs = torch.zeros_like(target_chunk, dtype=probs.dtype, device=probs.device)
-                pred_probs[valid_target_mask] = pred_probs_valid
-            else:
-                # All targets are ignored
-                pred_probs = torch.zeros_like(target_chunk, dtype=probs.dtype, device=probs.device)
-            # Store the scaling factors
-            scaling_factors = pred_probs.detach()  # Detach to ensure no gradient flow
-        # unreduced loss
-        loss_1d_slice = loss_1d[start_idx:end_idx]  # chunk_size,
-        z_loss_1d_slice = z_loss_1d[start_idx:end_idx] if return_z_loss else None
-        token_accuracy_1d_slice = token_accuracy_1d[start_idx:end_idx] if return_token_accuracy else None
-        predicted_tokens_1d_slice = predicted_tokens_1d[start_idx:end_idx] if return_predicted_tokens else None
-        # ensure _input and target are contiguous
-        logits_chunk = logits_chunk.contiguous()
-        target_chunk = target_chunk.contiguous()
-        # Here we calculate the gradient of logits_chunk in place so we can save memory.
-        liger_cross_entropy_kernel[(n_rows,)](
-            X_ptr=logits_chunk,
-            X_stride=logits_chunk.stride(-2),
-            Y_ptr=target_chunk,
-            Y_stride=target_chunk.stride(-1),  # always 1
-            weight_ptr=ce_weight,
-            loss_ptr=loss_1d_slice,
-            z_loss_ptr=z_loss_1d_slice,
-            loss_stride=loss_1d_slice.stride(-1),  # always 1
-            token_accuracy_ptr=token_accuracy_1d_slice,
-            token_accuracy_stride=token_accuracy_1d_slice.stride(-1)
-            if return_token_accuracy
-            else 0,  # always 1 if accuracy is enabled
-            predicted_tokens_ptr=predicted_tokens_1d_slice,
-            predicted_tokens_stride=predicted_tokens_1d_slice.stride(-1)
-            if return_predicted_tokens
-            else 0,  # always 1 if predicted tokens is enabled
-            n_cols=V,
-            n_non_ignore=total_n_non_ignore,
-            sum_non_ignore_weight=total_sum_non_ignore_ce_weight,
-            weight_sum=ce_weight_sum,
-            ignore_index=ignore_index,
-            lse_square_scale=lse_square_scale,
-            label_smoothing=label_smoothing,
-            reduction=reduction,
-            softcap=softcap,
-            RETURN_Z_LOSS=return_z_loss,
-            RETURN_TOKEN_ACCURACY=return_token_accuracy,
-            RETURN_PREDICTED_TOKENS=return_predicted_tokens,
-            HAS_WEIGHT=True if ce_weight is not None else False,
-            HAS_SOFTCAPPING=True if softcap is not None else False,
-            HAS_GRADIENTS=input_requires_grad,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-        # Apply token scaling if requested
-        if use_token_scaling:
-            loss_1d_slice = loss_1d_slice * scaling_factors
-            if return_z_loss:
-                z_loss_1d_slice = z_loss_1d_slice * scaling_factors
-        loss_1d[start_idx:end_idx] = loss_1d_slice
-        if return_z_loss:
-            z_loss_1d[start_idx:end_idx] = z_loss_1d_slice
-        if return_token_accuracy:
-            token_accuracy_1d[start_idx:end_idx] = token_accuracy_1d_slice
-        if return_predicted_tokens:
-            predicted_tokens_1d[start_idx:end_idx] = predicted_tokens_1d_slice
-        grad_logits_chunk = logits_chunk  # chunk_size x V
-        # Apply token scaling to gradients if requested
-        if use_token_scaling:
-            # Expand scaling factors to match gradient dimensions
-            scaling_factors_expanded = scaling_factors.unsqueeze(-1)  # chunk_size x 1
-            grad_logits_chunk = grad_logits_chunk * scaling_factors_expanded
-        if input_requires_grad:
-            grad_input[start_idx:end_idx] = grad_logits_chunk @ weight
-        if grad_weight is not None and input_requires_grad:
-            grad_weight += torch.mm(grad_logits_chunk.t(), _input_chunk).float()
-        if bias is not None and input_requires_grad:
-            torch.add(
-                input=grad_bias,
-                other=grad_logits_chunk.sum(dim=0),
-                out=grad_bias,
-                alpha=1.0,
-            )
-    # Need extra calculations for backward if reduction=='none'. Not supporting reduction='none' now.
-    # if reduction == "none":
-    #     loss = loss_1d
-    #     z_loss = z_loss_1d if return_z_loss else None
-    if reduction == "none":
-        # Return per-token losses
-        loss = loss_1d
-        z_loss = z_loss_1d if return_z_loss else None
-        token_accuracy = token_accuracy_1d if return_token_accuracy else None
-    else:
-        loss = torch.sum(loss_1d)
-        z_loss = torch.sum(z_loss_1d) if return_z_loss else None
-        # For accuracy, we compute the mean across all non-ignored tokens
-        token_accuracy = torch.sum(token_accuracy_1d) / total_n_non_ignore if return_token_accuracy else None
-    predicted_tokens = predicted_tokens_1d if return_predicted_tokens else None
-    # Cast back to original dtype
-    grad_weight = grad_weight.to(weight.dtype) if grad_weight is not None else None
-    grad_bias = grad_bias.to(bias.dtype) if grad_bias is not None else None
-    return loss, z_loss, token_accuracy, predicted_tokens, grad_input, grad_weight, grad_bias
-def fused_linear_cross_entropy_backward(grad_output, grad_input, grad_weight, grad_bias):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul to save time
-    if not torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        # We use a Triton kernel instead of a PyTorch operation because modifying inputs in-place
-        # for gradient storage and backward multiple times causes anomalies with PyTorch but not with Triton.
-        BT, H = grad_input.shape
-        n_rows = BT
-        BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(H))
-        element_mul_kernel[(n_rows,)](
-            grad_input,
-            grad_input.stride(-2),
-            grad_output,
-            H,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-        # handle grad_weight
-        if grad_weight is not None:
-            V, H = grad_weight.shape
-            n_rows = V
-            element_mul_kernel[(n_rows,)](
-                grad_weight,
-                grad_weight.stride(-2),
-                grad_output,
-                H,
-                BLOCK_SIZE=BLOCK_SIZE,
-                num_warps=32 if not is_hip() else 16,
-            )
-        if grad_bias is not None:
-            V = grad_bias.shape[0]
-            n_rows = V
-            element_mul_kernel[(n_rows,)](
-                grad_bias,
-                grad_bias.stride(-1),
-                grad_output,
-                1,
-                BLOCK_SIZE=BLOCK_SIZE,
-                num_warps=32 if not is_hip() else 16,
-            )
-    return grad_input, grad_weight, grad_bias
-class LigerFusedLinearCrossEntropyFunction(torch.autograd.Function):
-    @staticmethod
-    @amp_custom_fwd
-    def forward(
-        ctx,
-        _input,
-        weight,
-        target,
-        bias=None,
-        ce_weight=None,
-        ignore_index=-100,
-        lse_square_scale=0.0,
-        label_smoothing=0.0,
-        reduction="mean",
-        softcap=None,
-        return_z_loss: bool = False,
-        accum_dtype=None,
-        use_token_scaling: bool = False,
-        return_token_accuracy: bool = False,
-        return_predicted_tokens: bool = False,
-    ):
-        """
-        Fusing the last linear layer with cross-entropy loss
-            Reference: https://github.com/mgmalek/efficient_cross_entropy
-        Handle the forward and backward pass of the final linear layer via cross-entropy loss by avoiding
-        the materialization of the large logits tensor. Since Cross Entropy Loss is the last layer, we can
-        compute the gradient at the forward pass. By doing so, we don't have to store the _input and target
-        for the backward pass.
-        _input: (B*T, H) where B is batch size, T is sequence length, H is hidden dimension.
-        target: (B*T) where each value is in [0, V-1]
-        weight: (V, H) where V is the number of classes
-        bias: (V) where V is the number of classes
-        ce_weight: a manual rescaling weight given to each class. If given, has to be a Tensor of size V and floating point dtype
-        ignore_index: the index to ignore in the target
-        label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-        reduction: reduction to apply
-        accum_dtype (torch.dtype): the dtype of intermediate result buffers for weight and bias gradient accumulations.
-            Recommended to set `accum_dtype` to higher precision, e.g. `torch.float32`, if the training is unstable with original dtype. Default: `None`, performing accumulations in original dtype
-        use_token_scaling (bool): whether to scale each token's loss by its predicted probability (detached).
-            When True, each token's loss is multiplied by the model's predicted probability for that token's true class.
-            Default: False.
-        return_token_accuracy (bool): When `return_token_accuracy` is `True`, computes and returns per-token accuracy without materializing logits. Default: `False`
-        return_predicted_tokens (bool): When `return_predicted_tokens` is `True`, returns per-token predicted class indices (argmax) without materializing logits. Default: `False`
-        """
-        loss, z_loss, token_accuracy, predicted_tokens, grad_input, grad_weight, grad_bias = (
-            fused_linear_cross_entropy_forward(
-                _input=_input,
-                weight=weight,
-                target=target,
-                bias=bias,
-                ce_weight=ce_weight,
-                ignore_index=ignore_index,
-                lse_square_scale=lse_square_scale,
-                label_smoothing=label_smoothing,
-                reduction=reduction,
-                softcap=softcap,
-                return_z_loss=return_z_loss,
-                accum_dtype=accum_dtype,
-                use_token_scaling=use_token_scaling,
-                return_token_accuracy=return_token_accuracy,
-                return_predicted_tokens=return_predicted_tokens,
-            )
-        )
-        # downcast to dtype and store for backward
-        ctx.save_for_backward(
-            grad_input.detach(),
-            grad_weight.detach() if grad_weight is not None else None,
-            grad_bias.detach() if grad_bias is not None else None,
-        )
-        ctx.return_z_loss = return_z_loss
-        ctx.return_token_accuracy = return_token_accuracy
-        ctx.return_predicted_tokens = return_predicted_tokens
-        return loss, z_loss, token_accuracy, predicted_tokens
-    @staticmethod
-    @amp_custom_bwd
-    def backward(ctx, grad_output, grad_output2, grad_output3, grad_output4):
-        if ctx.return_z_loss:
-            del grad_output2  # z_loss is only for logging
-        if ctx.return_token_accuracy:
-            del grad_output3  # token_accuracy is only for metrics
-        if ctx.return_predicted_tokens:
-            del grad_output4  # predicted_tokens is only for metrics
-        (grad_input, grad_weight, grad_bias) = ctx.saved_tensors
-        grad_input, grad_weight, grad_bias = fused_linear_cross_entropy_backward(
-            grad_output, grad_input, grad_weight, grad_bias
-        )
-        return (
-            grad_input,
-            grad_weight,
-            None,
-            grad_bias,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,  # use_token_scaling
-            None,  # return_token_accuracy
-            None,  # return_predicted_tokens
-        )

build/torch-rocm/geglu.py DELETED Viewed

@@ -1,143 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.jit
-def _geglu_tanh_forward_kernel(a, b, c, stride, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    a += program_id * stride
-    b += program_id * stride
-    c += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    a_row = tl.load(a + col_offsets, mask=mask, other=0).to(tl.float32)
-    b_row = tl.load(b + col_offsets, mask=mask, other=0)
-    # tanh approximation form of GELU is computed with:
-    # 0.5 * a * (1 + tanh(sqrt(2 / pi) * (a + 0.044715 * a^3)))
-    sqrt_2_over_pi = 0.7978845608028654  # sqrt(2 / pi)
-    a_cubed = a_row * a_row * a_row
-    tanh_arg = sqrt_2_over_pi * (a_row + 0.044715 * a_cubed)
-    tanh_result = tanh(tanh_arg)
-    geglu_a = 0.5 * a_row * (1 + tanh_result)
-    c_row = geglu_a.cast(b_row.dtype) * b_row
-    tl.store(c + col_offsets, c_row, mask=mask)
-@triton.jit
-def _geglu_tanh_backward_kernel(dc, a, b, stride, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    dc += program_id * stride
-    a += program_id * stride
-    b += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    dc_row = tl.load(dc + col_offsets, mask=mask, other=0)
-    a_row = tl.load(a + col_offsets, mask=mask, other=0).to(tl.float32)
-    b_row = tl.load(b + col_offsets, mask=mask, other=0)
-    # recomputation to save memory
-    sqrt_2_over_pi = 0.7978845608028654  # sqrt(2 / pi)
-    a_cubed = a_row * a_row * a_row
-    tanh_arg = sqrt_2_over_pi * (a_row + 0.044715 * a_cubed)
-    tanh_result = tanh(tanh_arg)
-    geglu_a = 0.5 * a_row * (1 + tanh_result)
-    geglu_a = geglu_a.to(dc_row.dtype).to(tl.float32)
-    db_row = dc_row.cast(tl.float32) * geglu_a
-    # Gradient w.r.t. a can be computed with:
-    # b * (0.5 * (1 + tanh(z)) + 0.5 * a * (1 - tanh(z)^2) * (sqrt(2/pi) * (1 + 3 * 0.044715 * a^2)))
-    # where z = sqrt(2/pi) * (a + 0.044715 * a^3)
-    term1 = 0.5 * (1 + tanh_result)
-    tanh_sq = tanh_result * tanh_result
-    term2 = 0.5 * a_row * (1 - tanh_sq) * (sqrt_2_over_pi * (1 + 3 * 0.044715 * a_row * a_row))
-    da_row = dc_row * b_row * (term1 + term2)
-    tl.store(a + col_offsets, da_row, mask=mask)
-    tl.store(b + col_offsets, db_row.to(dc_row.dtype), mask=mask)
-def geglu_forward(a, b):
-    ori_shape = a.shape
-    n_cols = ori_shape[-1]
-    a = a.view(-1, n_cols)
-    b = b.view(-1, n_cols)
-    c = torch.empty_like(a)
-    n_rows = a.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _geglu_tanh_forward_kernel[(n_rows,)](
-        a,
-        b,
-        c,
-        c.stride(-2),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a, b, c.view(*ori_shape)
-def geglu_backward(a, b, dc):
-    ori_shape = dc.shape
-    n_cols = ori_shape[-1]
-    dc = dc.view(-1, n_cols)
-    n_rows = dc.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _geglu_tanh_backward_kernel[(n_rows,)](
-        dc,
-        a,
-        b,
-        dc.stride(-2),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a.view(*ori_shape), b.view(*ori_shape)
-class LigerGELUMulFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, a, b):
-        a, b, c = geglu_forward(a, b)
-        ctx.save_for_backward(a, b)
-        return c
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dc):
-        a, b = ctx.saved_tensors
-        a, b = geglu_backward(a, b, dc)
-        return a, b

build/torch-rocm/group_norm.py DELETED Viewed

@@ -1,311 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-if infer_device() == "npu":
-    MAX_FUSED_SIZE = 16384  # 8192
-else:
-    MAX_FUSED_SIZE = 65536
-@triton.jit
-def _group_norm_forward_kernel(
-    Y_ptr,  # pointer to output, shape (n_rows, n_groups, hidden_size)
-    Y_row_stride,  # stride of each row in output
-    Y_col_stride,  # stride of each column in output
-    X_ptr,  # pointer to input, shape (n_rows, n_groups, hidden_size)
-    X_row_stride,  # stride of each row in input
-    X_col_stride,  # stride of each column in input
-    Mean_ptr,  # pointer to mean, shape (n_rows, n_groups)
-    Mean_row_stride,  # stride of each row in mean
-    Mean_col_stride,  # stride of each column in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows, n_groups)
-    RSTD_row_stride,  # stride of each row in rstd
-    RSTD_col_stride,  # stride of each column in rstd
-    W_ptr,  # pointer to W
-    B_ptr,  # pointer to B
-    hidden_size,  # hidden size of X
-    channels_per_group,  # the number of channels per group
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://nn.labml.ai/normalization/group_norm/index.html
-    """
-    batch_idx = tl.program_id(0)
-    group_idx = tl.program_id(1)
-    X_ptr += batch_idx * X_row_stride + group_idx * X_col_stride
-    Y_ptr += batch_idx * Y_row_stride + group_idx * Y_col_stride
-    block_range = tl.arange(0, BLOCK_SIZE)
-    # Compute mean and variance using the online algorithm
-    s = 0.0
-    squared_sum = 0.0
-    for i in tl.range(0, hidden_size, BLOCK_SIZE):
-        hidden_size_offsets = i + block_range
-        mask = hidden_size_offsets < hidden_size
-        X = tl.load(X_ptr + hidden_size_offsets, mask=mask, other=0.0)
-        s += tl.sum(X)
-        # X**2
-        squared_sum += tl.sum(X * X)
-    m = s / hidden_size
-    # variance = E[X**2] - E[X]**2
-    variance = (squared_sum / hidden_size) - (m * m)
-    # 1/std
-    rstd = rsqrt(variance + eps)
-    # Normalize — flat loop over full hidden_size (not per-channel)
-    # This avoids the nested channel × per_channel_hidden loop where
-    # BLOCK_SIZE >> hidden_size_per_channel causes massive padding waste.
-    hidden_size_per_channel = hidden_size // channels_per_group
-    for i in tl.range(0, hidden_size, BLOCK_SIZE):
-        hidden_size_offsets = i + block_range
-        mask = hidden_size_offsets < hidden_size
-        X = tl.load(X_ptr + hidden_size_offsets, mask=mask, other=m)
-        # Determine which channel each element belongs to, then load W/B
-        local_channel = hidden_size_offsets // hidden_size_per_channel
-        global_channel = group_idx * channels_per_group + local_channel
-        W = tl.load(W_ptr + global_channel, mask=mask)
-        B = tl.load(B_ptr + global_channel, mask=mask)
-        Y = (X - m) * rstd * W + B
-        tl.store(Y_ptr + hidden_size_offsets, Y, mask=mask)
-    tl.store(Mean_ptr + batch_idx * Mean_row_stride + group_idx * Mean_col_stride, m)
-    tl.store(RSTD_ptr + batch_idx * RSTD_row_stride + group_idx * RSTD_col_stride, rstd)
-@triton.jit
-def _group_norm_backward_kernel(
-    X_ptr,  # pointer to input, shape (n_rows, n_channels, hidden_size)
-    X_row_stride,  # stride of each row in input
-    X_col_stride,  # stride of each column in input
-    W_ptr,  # pointer to weights, shape (n_channels)
-    Mean_ptr,  # pointer to mean, shape (n_rows, n_groups)
-    Mean_ptr_row_stride,  # stride of each column in mean
-    Mean_ptr_col_stride,  # stride of each column in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows, n_groups)
-    DX_ptr,  # pointer to input grad, shape (n_rows, n_groups, hidden_size)
-    DW_ptr,  # pointer to weights grad, shape (n_channels)
-    DB_ptr,  # pointer to bias grad, shape (n_channels)
-    UPSTREAM_ptr,  # pointer to output grad, shape (n_rows, n_channels, hidden_size)
-    hidden_size: tl.constexpr,  # hidden size
-    channels_per_group: tl.constexpr,  # number of groups in group norm
-    BLOCK_SIZE: tl.constexpr,
-    dtype: tl.constexpr,
-):
-    """
-    References:
-    https://nn.labml.ai/normalization/group_norm/index.html
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    The backprop equations are the same for group_norm and layer_norm
-    the only difference here is that we load the Mean, Rstd corresponding to the
-    group we're computing gradients for and the mean and rstd are computed over n-channels
-    so the total number of elements we compute the mean over is num_channels_per_group * hidden_size
-    We also need to load the Weights corresponding to the current channel to compute the gradients.
-    """
-    batch_idx = tl.program_id(0)
-    group_idx = tl.program_id(1)
-    # Move the pointers to the correct batch
-    X_ptr += batch_idx * X_row_stride
-    DX_ptr += batch_idx * X_row_stride
-    UPSTREAM_ptr += batch_idx * X_row_stride
-    # Mean and rstd are the same shape so have the same strides
-    mean = tl.load(Mean_ptr + batch_idx * Mean_ptr_row_stride + group_idx * Mean_ptr_col_stride)
-    rstd = tl.load(RSTD_ptr + batch_idx * Mean_ptr_row_stride + group_idx * Mean_ptr_col_stride)
-    c1 = 0.0
-    c2 = 0.0
-    block_range = tl.arange(0, BLOCK_SIZE)
-    # We need to compute the sum terms of the backprop equations across all channels in the group
-    for channel_idx in range(group_idx * channels_per_group, (group_idx + 1) * channels_per_group):
-        dW = 0.0
-        dB = 0.0
-        # Move the pointers to the correct channel
-        W = tl.load(W_ptr + channel_idx)
-        for i in tl.range(0, hidden_size, BLOCK_SIZE):
-            hidden_size_offsets = i + block_range
-            mask = hidden_size_offsets < hidden_size
-            X = tl.load(
-                X_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            UPSTREAM_grad = tl.load(
-                UPSTREAM_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            x_hat = (X - mean) * rstd
-            dW += tl.sum(UPSTREAM_grad * x_hat)
-            dB += tl.sum(UPSTREAM_grad)
-            wdy = W * UPSTREAM_grad
-            c1 += tl.sum(x_hat * wdy)
-            c2 += tl.sum(wdy)
-        # Need to ensure additions to the same channel are atomic
-        tl.atomic_add(DW_ptr + channel_idx, dW.to(dtype))
-        tl.atomic_add(DB_ptr + channel_idx, dB.to(dtype))
-    N = hidden_size * channels_per_group
-    c1 = c1 / N
-    c2 = c2 / N
-    for channel_idx in tl.range(group_idx * channels_per_group, (group_idx + 1) * channels_per_group):
-        # Move the pointers to the correct channel
-        W = tl.load(W_ptr + channel_idx)
-        for i in range(0, hidden_size, BLOCK_SIZE):
-            hidden_size_offsets = i + block_range
-            mask = hidden_size_offsets < hidden_size
-            X = tl.load(
-                X_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            UPSTREAM_grad = tl.load(
-                UPSTREAM_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            x_hat = (X - mean) * rstd
-            wdy = W * UPSTREAM_grad
-            dx = (wdy - (x_hat * c1 + c2)) * rstd
-            tl.store(DX_ptr + channel_idx * X_col_stride + hidden_size_offsets, dx, mask=mask)
-def group_norm_forward(X, num_channels, num_groups, W, B, eps):
-    shape = X.shape
-    batch_size = shape[0]
-    channels_per_group = num_channels // num_groups
-    # Reshape X so that the mean and std are computed across the groups
-    X = X.view(batch_size, num_groups, -1).contiguous()
-    hidden_size = X.shape[-1]
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(hidden_size))
-    Y = torch.empty((batch_size, num_groups, hidden_size), dtype=X.dtype, device=X.device)
-    Mean = torch.zeros((batch_size, num_groups), dtype=X.dtype, device=X.device)
-    RSTD = torch.zeros((batch_size, num_groups), dtype=X.dtype, device=X.device)
-    _group_norm_forward_kernel[(batch_size, num_groups)](
-        Y,
-        Y.stride(0),
-        Y.stride(1),
-        X,
-        X.stride(0),
-        X.stride(1),
-        Mean,
-        Mean.stride(0),
-        Mean.stride(1),
-        RSTD,
-        RSTD.stride(0),
-        RSTD.stride(1),
-        W,
-        B,
-        hidden_size,
-        channels_per_group,
-        eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-    )
-    # Return tensors in the original shape
-    return Y.view(*shape), X.view(*shape), Mean, RSTD, BLOCK_SIZE
-def group_norm_backward(dY, X, W, B, Mean, RSTD, num_channels, num_groups):
-    shape = dY.shape
-    batch_size = shape[0]
-    hidden_size = dY.shape[-1]
-    channels_per_group = num_channels // num_groups
-    dY = dY.view(batch_size, num_groups, -1)
-    DX = torch.empty(
-        (batch_size, num_groups, hidden_size * channels_per_group),
-        dtype=X.dtype,
-        device=X.device,
-    )
-    DW = torch.zeros((num_channels), dtype=W.dtype, device=W.device)
-    DB = torch.zeros((num_channels), dtype=B.dtype, device=B.device)
-    triton_dtype = tl.float32 if X.dtype == torch.float32 else tl.bfloat16
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(hidden_size))
-    _group_norm_backward_kernel[(batch_size, num_groups)](
-        X,
-        X.stride(0),
-        X.stride(1),
-        W,
-        Mean,
-        Mean.stride(0),
-        Mean.stride(1),
-        RSTD,
-        DX,
-        DW,
-        DB,
-        dY,
-        hidden_size,
-        channels_per_group,
-        BLOCK_SIZE=BLOCK_SIZE,
-        dtype=triton_dtype,
-    )
-    # Return tensors in the original shape
-    return DX.view(*shape), DW, DB
-class LigerGroupNormFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        X,
-        affine_scaling_weight,
-        affine_shifting_bias,
-        num_channels,
-        num_groups,
-        eps,
-    ):
-        Y, X, Mean, RSTD, BLOCK_SIZE = group_norm_forward(
-            X,
-            num_channels,
-            num_groups,
-            affine_scaling_weight,
-            affine_shifting_bias,
-            eps,
-        )
-        ctx.num_channels = num_channels
-        ctx.num_groups = num_groups
-        ctx.save_for_backward(X, affine_scaling_weight, affine_shifting_bias, Mean, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        X, W, B, Mean, RSTD = ctx.saved_tensors
-        DX, DW, DB = group_norm_backward(dY, X, W, B, Mean, RSTD, ctx.num_channels, ctx.num_groups)
-        return DX, DW, DB, None, None, None

build/torch-rocm/jsd.py DELETED Viewed

@@ -1,201 +0,0 @@
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-from .utils import infer_device
-@triton.jit
-def _jsd_kernel(
-    X_ptr,  # input in logspace, X = log Q
-    X_stride,
-    Y_ptr,  # ground truth in logspace, Y = log P
-    Y_stride,
-    loss_ptr,
-    loss_stride,
-    dX_ptr,
-    dX_stride,
-    label_ptr,
-    beta: tl.constexpr,
-    n_non_ignore: int,
-    ignore_index: tl.constexpr,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-    HAS_LABEL: tl.constexpr,
-):
-    # JSD(P || Q) = (KL(P || M) + KL(Q || M)) / 2, M = (1/2) * (P + Q) = (1/2) * (e ^ Y + e ^ X)
-    #             = sum(P * log P + Q * log Q - 2 * M * log M) / 2
-    #             = sum(e ^ Y * Y + e ^ X * X - 2 * M * log M) / 2
-    # grad_x_i = 0.5 * Q * (X - log_M)
-    pid = tl.program_id(0).to(tl.int64)
-    X_ptr += pid * X_stride
-    dX_ptr += pid * dX_stride
-    Y_ptr += pid * Y_stride
-    loss_ptr += pid * loss_stride
-    label_ptr += pid
-    if HAS_LABEL:
-        label = tl.load(label_ptr)
-        if label == ignore_index:
-            for i in range(0, n_cols, BLOCK_SIZE):
-                offsets = i + tl.arange(0, BLOCK_SIZE)
-                tl.store(dX_ptr + offsets, 0.0, mask=offsets < n_cols)
-            return
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + tl.arange(0, BLOCK_SIZE)
-        mask = offsets < n_cols
-        X = tl.load(X_ptr + offsets, mask=mask, other=float("-inf")).to(tl.float32)
-        Y = tl.load(Y_ptr + offsets, mask=mask, other=float("-inf")).to(tl.float32)
-        if beta == 0.0:  # forward KL
-            Y_max = tl.max(Y, axis=0)
-            Y_shifted = Y - Y_max
-            Y_prob = tl.exp(Y_shifted) * tl.exp(Y_max)  # Compensate for the shift
-            loss = Y_prob * (Y - X)
-            dX = -Y_prob
-        elif beta == 1.0:  # reverse KL
-            X_max = tl.max(X, axis=0)
-            X_shifted = X - X_max
-            X_prob = tl.exp(X_shifted) * tl.exp(X_max)  # Compensate for the shift
-            loss = X_prob * (X - Y)
-            dX = loss + X_prob
-        else:
-            max_val = tl.maximum(tl.max(X, axis=0), tl.max(Y, axis=0))
-            X_shifted = X - max_val
-            Y_shifted = Y - max_val
-            # Pre-compute exp(max_val) since it's used twice
-            exp_max = tl.exp(max_val)
-            # Compute exp terms with compensation
-            Q = tl.exp(X_shifted) * exp_max  # = exp(X)
-            P = tl.exp(Y_shifted) * exp_max  # = exp(Y)
-            # Pre-compute common terms
-            beta_P = beta * P
-            one_minus_beta_Q = (1 - beta) * Q
-            M = beta_P + one_minus_beta_Q
-            log_M = tl.log(M)  # No need to compensate as M is already in original scale
-            loss = beta_P * Y + one_minus_beta_Q * X - M * log_M
-            dX = one_minus_beta_Q * (X - log_M)
-        # Pre-compute scaling factor
-        scale = 1.0 / n_non_ignore
-        loss = loss * scale
-        dX = dX * scale
-        tl.store(loss_ptr + offsets, loss, mask=mask)
-        tl.store(dX_ptr + offsets, dX, mask=mask)
-MAX_FUSED_SIZE = 4096 if infer_device() == "xpu" else 65536
-def jsd_forward(_input, target, shift_labels, beta, ignore_index, has_label):
-    BT, V = _input.shape
-    n_rows = BT
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    # non reduction loss
-    loss = torch.zeros(_input.shape, dtype=torch.float32, device=_input.device)
-    dX = torch.empty_like(_input)
-    if has_label:
-        n_non_ignore = (shift_labels != ignore_index).sum().item()
-    else:
-        n_non_ignore = BT
-    _jsd_kernel[(n_rows,)](
-        X_ptr=_input,  # input in logspace, X = log Q
-        X_stride=_input.stride(-2),
-        Y_ptr=target,  # ground truth in logspace, Y = log P
-        Y_stride=target.stride(-2),
-        loss_ptr=loss,
-        loss_stride=loss.stride(-2),
-        dX_ptr=dX,
-        dX_stride=dX.stride(-2),
-        label_ptr=(shift_labels if has_label else torch.empty(1, device=_input.device)),  # dummy ptr if no label
-        beta=beta,
-        n_non_ignore=n_non_ignore,
-        ignore_index=ignore_index,
-        n_cols=V,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_LABEL=has_label,
-    )
-    loss = torch.sum(loss)
-    return loss.to(_input.dtype), dX
-def jsd_backward(dX, grad_output):
-    # If jsd is the last layer, grad_output is 1.0. Skip the mul to save time
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return dX
-    else:
-        return grad_output * dX
-class LigerJSDFunction(torch.autograd.Function):
-    r"""
-    This class implements the forward and backward pass for the generalized Jensen-Shannon Divergence.
-    .. math::
-        JSD(\beta)(P || Q)
-            = \beta * KLDiv(P || (\beta * P + (1 - \beta) * Q)) + (1 - \beta) * KLDiv(Q || (\beta * P + (1 - \beta) * Q))
-    .. note::
-        As all the other losses in PyTorch, this function expects the first argument,
-        :attr:`_input`, to be the predictions, the output of the student model, in log-space
-        and the second, :attr:`target`, to be the observations, the output of the teacher model, in log-space.
-        This differs from the standard mathematical notation :math:`JSD(P || Q)` where
-        :math:`P` denotes the teacher model and :math:`Q` denotes the student model.
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-        beta: float = 0.5,
-        ignore_index: int = -100,
-    ) -> torch.Tensor:
-        """
-        Args:
-            _input (torch.Tensor): predict values with shape (BT, V) in logspace
-            target (torch.Tensor): ground truth values with shape (BT, V) in logspace
-            shift_labels (Optional[torch.LongTensor]): indicator of next predicted vocab with shape (BT) where each value is in [0, V-1].
-            beta (float): coefficient beta of generalized JSD in the interval [0, 1]. It implements forward/reverse KL when beta equals 0 and 1 respectively. Default: `0.5`
-            ignore_index (int): the index to ignore. Default: -100
-        Returns:
-            loss (torch.Tensor): generalized JSD
-        """
-        has_label = False
-        if shift_labels is not None:
-            assert shift_labels.shape == (_input.shape[0],), (
-                f"the shape of shift_labels must be (BT,). Got: {shift_labels.shape}"
-            )
-            shift_labels = shift_labels.contiguous()
-            has_label = True
-        loss, dX = jsd_forward(_input, target, shift_labels, beta, ignore_index, has_label)
-        ctx.save_for_backward(dX)
-        return loss
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        (dX,) = ctx.saved_tensors
-        dX = jsd_backward(dX, grad_output)
-        return (
-            dX,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-rocm/kl_div.py DELETED Viewed

@@ -1,259 +0,0 @@
-from typing import Literal
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-from .utils import is_hip
-from .utils import infer_device
-def get_num_warps(BLOCK_SIZE):
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32 if not is_hip() else 16
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return num_warps
-if infer_device() == "xpu":
-    MAX_FUSED_SIZE = 8192
-elif infer_device() == "npu":
-    MAX_FUSED_SIZE = 8192
-else:
-    MAX_FUSED_SIZE = 65536 // 4  # 65536 // 4 or 8 works the best
-REDUCTION_LITERAL = Literal["none", "sum", "mean", "batchmean"]
-_REDUCTION_MODE_NONE: tl.constexpr = tl.constexpr(0)
-_REDUCTION_MODE_SUM: tl.constexpr = tl.constexpr(1)
-_REDUCTION_MODE_MEAN: tl.constexpr = tl.constexpr(2)
-_REDUCTION_MODE_BATCHMEAN: tl.constexpr = tl.constexpr(3)
-_str_to_reduction_mode = {
-    "none": _REDUCTION_MODE_NONE.value,
-    "sum": _REDUCTION_MODE_SUM.value,
-    "mean": _REDUCTION_MODE_MEAN.value,
-    "batchmean": _REDUCTION_MODE_BATCHMEAN.value,
-}
-@triton.jit
-def _kldiv_kernel_forward(
-    y_ptr,  # [B, S], prediction ptr, the kernel expects the prediction in log-space
-    y_stride,  # int, prediction stride
-    gt_ptr,  # [B, S], ground truth ptr
-    gt_stride,  # int, ground truth stride
-    loss_ptr,  # [B] or [B, S] if reduction == _REDUCTION_MODE_NONE, output ptr
-    loss_stride,  # int, output stride
-    n_cols,  # int, number of columns in the input tensor
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-    log_target: tl.constexpr = False,
-    reduction: tl.constexpr = _REDUCTION_MODE_BATCHMEAN,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    y_ptr += pid * y_stride
-    gt_ptr += pid * gt_stride
-    loss_ptr += pid * loss_stride
-    base_offsets = tl.arange(0, BLOCK_SIZE)
-    loss_sum = 0.0
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + base_offsets
-        mask = offsets < n_cols
-        y = tl.load(y_ptr + offsets, mask=mask, other=0.0)
-        y_true = tl.load(gt_ptr + offsets, mask=mask, other=0.0)
-        # KL(y_true || y) = y_true * (log(y_true) - log(y))
-        # We compute KL(y_true || y) with y in the log-space
-        if not log_target:
-            loss = y_true * (tl.log(tl.maximum(y_true, eps)) - y)
-        else:
-            loss = tl.exp(y_true) * (y_true - y)
-        if reduction == _REDUCTION_MODE_NONE:
-            tl.store(loss_ptr + offsets, loss, mask=mask)
-        else:
-            loss_sum += tl.sum(loss, axis=0)
-    if reduction != _REDUCTION_MODE_NONE:
-        tl.store(loss_ptr, loss_sum)
-@triton.jit
-def _kldiv_kernel_backward(
-    target_ptr,
-    target_stride,
-    new_grads_ptr,
-    new_grads_stride,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-    log_target: tl.constexpr = False,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    target_ptr += pid * target_stride
-    new_grads_ptr += pid * new_grads_stride
-    offsets = tl.arange(0, BLOCK_SIZE)
-    mask = offsets < n_cols
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + tl.arange(0, BLOCK_SIZE)
-        mask = offsets < n_cols
-        target = tl.load(target_ptr + offsets, mask=mask, other=0.0)
-        if not log_target:
-            res = target * -1
-        else:
-            res = -tl.exp(target)
-        tl.store(new_grads_ptr + offsets, res, mask=mask)
-def kldiv_forward_triton(y_pred, y_true, log_target, reduction, eps):  # [BT, V]
-    BT, V = y_pred.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = 32 if infer_device() == "xpu" else get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    reduction = _str_to_reduction_mode[reduction]
-    out_size = (BT, V) if reduction == _REDUCTION_MODE_NONE.value else (BT,)
-    output_tensor = torch.zeros(out_size, device=y_pred.device, dtype=torch.float32)
-    _kldiv_kernel_forward[grid](
-        y_pred,
-        y_pred.stride(0),
-        y_true,
-        y_true.stride(0),
-        output_tensor,
-        output_tensor.stride(0),
-        V,
-        eps=eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        log_target=log_target,
-        reduction=reduction,
-    )
-    # calculated according to the reduction mode same as in Pytorch. In the later versions, `mean` will be changed to the same behavior as `batchmean`
-    # https://pytorch.org/docs/stable/generated/torch.nn.KLDivLoss.html
-    # https://github.com/pytorch/pytorch/blob/d7b57c4d63edb42e1deeeba9497fcb5f1f748ff2/torch/nn/functional.py#L3372
-    if reduction == _REDUCTION_MODE_BATCHMEAN.value:
-        return output_tensor.sum() / BT
-    elif reduction == _REDUCTION_MODE_SUM.value:
-        return output_tensor.sum(dim=0)
-    elif reduction == _REDUCTION_MODE_MEAN.value:
-        return output_tensor.sum() / (BT * V)
-    else:
-        return output_tensor
-def kldiv_backward_triton(target, grad_output, new_grads, log_target):
-    BT, V = target.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = 32 if infer_device() == "xpu" else get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    # We store the gradients in-place in the input tensor
-    _kldiv_kernel_backward[grid](
-        target,
-        target.stride(0),
-        new_grads,
-        new_grads.stride(0),
-        V,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        log_target=log_target,
-    )
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul then.
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return new_grads
-    return new_grads * grad_output
-class LigerKLDivLossFunction(torch.autograd.Function):
-    """
-    Class implementing the forward and backward pass for the KL Divergence Loss using Triton, as defined by the following formula:
-    ```python
-    if log_target:
-        loss = target.exp() * (target - input)
-    else:
-        loss = target * (target.log() - input)
-    ```,
-    then the loss is reduced according to the `reduction` parameter.
-    as defined in the PyTorch documentation: https://pytorch.org/docs/stable/generated/torch.nn.KLDivLoss.html
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        y_pred: torch.Tensor,
-        y_true: torch.Tensor,
-        reduction: REDUCTION_LITERAL = "batchmean",
-        log_target: bool = False,
-        eps: float = 1e-10,
-    ) -> torch.Tensor:
-        """A forward pass for the KL Divergence Loss.
-        Args:
-            ctx: Torch autograd context
-            y_pred (torch.Tensor): A tensor of shape (BT, V) containing the predicted values, expected to be log-probabilities.
-            y_true (torch.Tensor): A tensor of shape (BT, V) containing the target values, expected to be either probabilities or log-probabilities, depending on the value of `log_target`.
-            reduction (REDUCTION_LITERAL, optional): Reduction to be used. Defaults to "batchmean".
-            log_target (bool, optional): If set to true, expects the ground truth to already be log-probabilities. Defaults to False.
-            eps: (float, optional): A small value to avoid division by zero. Defaults to 1e-10.
-        Returns:
-            torch.Tensor: The computed KL Divergence Loss, with shape (BT, V) if `reduction` is "none", else a scalar.
-        """
-        ctx.save_for_backward(y_true)
-        ctx.reduction = reduction
-        ctx.log_target = log_target
-        return kldiv_forward_triton(y_pred, y_true, log_target=log_target, reduction=reduction, eps=eps)
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        """A backward pass for the KL Divergence Loss.
-        Args:
-            ctx: Torch autograd context
-            grad_output (torch.Tensor): The gradient of the loss with respect to the output.
-        Returns:
-            tuple[torch.Tensor, None, None, None, None]: The gradient of the loss with respect to the inputs and None for the other arguments of the forward method.
-        """
-        (y_true,) = ctx.saved_tensors
-        new_grads = torch.empty_like(y_true)
-        derivative = kldiv_backward_triton(y_true, grad_output, new_grads, ctx.log_target)
-        if ctx.reduction == "batchmean":
-            derivative = derivative / y_true.shape[0]
-        elif ctx.reduction == "sum" or ctx.reduction == "none":
-            pass
-        elif ctx.reduction == "mean":
-            derivative = derivative / (y_true.shape[0] * y_true.shape[1])
-        return (
-            derivative,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-rocm/layer_norm.py DELETED Viewed

@@ -1,320 +0,0 @@
-import math
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import set_large_grf_mode
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-@triton.jit
-def _layer_norm_forward_kernel(
-    Y_ptr,  # pointer to output, shape (n_rows, n_cols)
-    Y_row_stride,  # stride of each row in output
-    X_ptr,  # pointer to input, shape (n_rows, n_cols)
-    X_row_stride,  # stride of each row in input
-    W_ptr,  # pointer to weights, shape (n_cols,)
-    W_row_stride,  # stride of each row in weights
-    B_ptr,  # pointer to bias, shape (n_cols,)
-    B_row_stride,  # stride of each row in bias
-    Mean_ptr,  # pointer to mean, shape (n_rows,)
-    Mean_row_stride,  # stride of each row in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows,)
-    RSTD_row_stride,  # stride of each row in rstd
-    n_cols,
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://arxiv.org/abs/1607.06450
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    """
-    row_idx = tl.program_id(0).to(tl.int64)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    # Pre-load weights and bias in fp32 to avoid repeated conversions
-    W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0.0)
-    B_row = tl.load(B_ptr + col_offsets, mask=mask, other=0.0)
-    W_f32 = W_row.to(tl.float32)
-    B_f32 = B_row.to(tl.float32)
-    # Calculate pointers for this row
-    row_X_ptr = X_ptr + row_idx * X_row_stride
-    row_Y_ptr = Y_ptr + row_idx * Y_row_stride
-    row_Mean_ptr = Mean_ptr + row_idx * Mean_row_stride
-    row_RSTD_ptr = RSTD_ptr + row_idx * RSTD_row_stride
-    # Load input data and convert to fp32 for numerical stability
-    X_row = tl.load(row_X_ptr + col_offsets, mask=mask, other=0.0)
-    X_f32 = X_row.to(tl.float32)
-    # Compute statistics in fp32 for numerical stability
-    mean = tl.sum(X_f32, axis=0) / n_cols
-    X_centered = X_f32 - mean
-    # Apply mask to variance calculation to exclude contributions from masked elements
-    X_centered_masked = tl.where(mask, X_centered, 0.0)
-    var = tl.sum(X_centered_masked * X_centered_masked, axis=0) / n_cols
-    rstd = rsqrt(var + eps)
-    # Store statistics (convert back to original dtype only once)
-    tl.store(row_Mean_ptr, mean.to(X_row.dtype))
-    tl.store(row_RSTD_ptr, rstd.to(X_row.dtype))
-    # Fused normalization and affine transformation
-    # Y = (X - mean) * rstd * W + B = X_centered * rstd * W + B
-    Y_f32 = X_centered * rstd * W_f32 + B_f32
-    # Store output (single conversion back to original dtype)
-    tl.store(row_Y_ptr + col_offsets, Y_f32.to(X_row.dtype), mask=mask)
-@triton.jit
-def _layer_norm_backward_kernel(
-    X_ptr,  # pointer to input, shape (n_rows, n_cols)
-    stride_x,  # stride of each row in input
-    W_ptr,  # pointer to weights, shape (n_cols,)
-    Mean_ptr,  # pointer to mean, shape (n_rows,)
-    stride_mean,  # stride of each row in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows,)
-    stride_rstd,  # stride of each row in rstd
-    DX_ptr,  # pointer to input grad, shape (n_rows, n_cols)
-    stride_dx,  # stride of each row in input grad
-    DW_ptr,  # pointer to weights grad, shape (n_cols,)
-    stride_dw,  # stride of each row in weights grad
-    DB_ptr,  # pointer to bias grad, shape (n_cols,)
-    stride_db,  # stride of each row in bias grad
-    DY_ptr,  # pointer to output grad, shape (n_rows, n_cols)
-    stride_dy,  # stride of each row in output grad
-    n_rows,
-    n_cols,
-    rows_per_program: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://arxiv.org/abs/1607.06450
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    """
-    row_block_id = tl.program_id(0).to(tl.int64)
-    row_start = row_block_id * rows_per_program
-    row_end = min((row_block_id + 1) * rows_per_program, n_rows)
-    cols = tl.arange(0, BLOCK_SIZE)
-    mask = cols < n_cols
-    dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    db_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    # Pre-load weights once (same optimization as forward pass)
-    w = tl.load(W_ptr + cols, mask=mask, other=0.0)
-    w_f32 = w.to(tl.float32)
-    for row_idx in range(row_start, row_end):
-        # Calculate pointers for this specific row
-        row_X_ptr = X_ptr + row_idx * stride_x
-        row_DX_ptr = DX_ptr + row_idx * stride_dx
-        row_DY_ptr = DY_ptr + row_idx * stride_dy
-        row_Mean_ptr = Mean_ptr + row_idx * stride_mean
-        row_RSTD_ptr = RSTD_ptr + row_idx * stride_rstd
-        # Load data for this row
-        x = tl.load(row_X_ptr + cols, mask=mask, other=0.0)
-        dy = tl.load(row_DY_ptr + cols, mask=mask, other=0.0)
-        mean = tl.load(row_Mean_ptr)
-        rstd = tl.load(row_RSTD_ptr)
-        # Convert to fp32 for numerical stability
-        x_f32 = x.to(tl.float32)
-        dy_f32 = dy.to(tl.float32)
-        mean_f32 = mean.to(tl.float32)
-        rstd_f32 = rstd.to(tl.float32)
-        # Compute backward pass for this row
-        x_hat = (x_f32 - mean_f32) * rstd_f32
-        wdy = w_f32 * dy_f32
-        c1 = tl.sum(x_hat * wdy, axis=0) / n_cols
-        c2 = tl.sum(wdy, axis=0) / n_cols
-        dx = (wdy - (x_hat * c1 + c2)) * rstd_f32
-        # Store input gradient
-        tl.store(row_DX_ptr + cols, dx, mask=mask)
-        # Accumulate weight and bias gradients for this thread block's assigned rows
-        dw = dy_f32 * x_hat
-        db = dy_f32
-        dW_row += dw
-        db_row += db
-    tl.store(DW_ptr + row_block_id * stride_dw + cols, dW_row, mask=mask)
-    tl.store(DB_ptr + row_block_id * stride_db + cols, db_row, mask=mask)
-def layer_norm_forward(X, W, B, eps):
-    """
-    Args:
-        X: Input tensor of shape (..., hidden_size)
-        W: Weight tensor of shape (hidden_size,)
-        B: Bias tensor of shape (hidden_size,)
-        eps: Small constant for numerical stability
-    Returns:
-        Tuple of (output, input, mean, rstd, block_size, num_warps)
-    """
-    shape = X.shape
-    dim = shape[-1]
-    X = X.view(-1, dim)
-    n_rows, n_cols = X.shape
-    # Calculate optimal block size and warp configuration
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    # Allocate output tensors
-    Y = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    Mean = torch.empty(n_rows, dtype=X.dtype, device=X.device)
-    RSTD = torch.empty(n_rows, dtype=X.dtype, device=X.device)
-    # Validate input dimensions
-    if X.shape[1] != W.shape[0]:
-        raise ValueError(
-            f"Incompatible dimensions: input feature size (X.shape[1]={X.shape[1]}) "
-            f"must match weight size (W.shape[0]={W.shape[0]})"
-        )
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    # Launch kernel with one thread block per row for optimal performance
-    grid = (n_rows,)
-    _layer_norm_forward_kernel[grid](
-        Y,
-        Y.stride(0),
-        X,
-        X.stride(0),
-        W,
-        W.stride(0),
-        B,
-        B.stride(0),
-        Mean,
-        Mean.stride(0),
-        RSTD,
-        RSTD.stride(0),
-        n_cols,
-        eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        **kernel_args,
-    )
-    return Y.view(*shape), X, Mean, RSTD, BLOCK_SIZE, num_warps
-def layer_norm_backward(dY, X, W, B, Mean, RSTD):
-    """
-    Args:
-        dY: Gradient of output
-        X: Input tensor
-        W: Weight tensor
-        B: Bias tensor
-        Mean: Pre-computed mean
-        RSTD: Pre-computed reciprocal standard deviation
-    Returns:
-        Tuple of (input_grad, weight_grad, bias_grad)
-    """
-    shape = dY.shape
-    dim = shape[-1]
-    dY = dY.view(-1, dim)
-    n_rows, n_cols = dY.shape
-    sm_count = 1
-    if X.device.type == "cuda":
-        sm_count = torch.cuda.get_device_properties(X.device).multi_processor_count
-    elif X.device.type == "xpu":
-        sm_count = torch.xpu.get_device_properties(X.device).gpu_eu_count
-    elif X.device.type == "npu":
-        sm_count = get_npu_core_count()
-    # fp32 for numerical stability especially.
-    _DW = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-    _DB = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-    # Calculate optimal block size and warp configuration
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    if n_cols > BLOCK_SIZE:
-        raise RuntimeError(f"Feature dimension {n_cols} exceeds maximum supported size of {BLOCK_SIZE}.")
-    rows_per_program = math.ceil(n_rows / sm_count)
-    grid = (sm_count,)
-    # Allocate gradient tensors
-    DX = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    kernel_args = {"num_warps": num_warps}
-    # XPU-specific optimization
-    if X.device.type == "xpu":
-        kernel_args.update({"num_warps": 32, "num_stages": 4})
-        set_large_grf_mode(kernel_args)
-    # Launch kernel with one thread block per row for optimal performance
-    _layer_norm_backward_kernel[grid](
-        X,
-        X.stride(0),
-        W,
-        Mean,
-        Mean.stride(0),
-        RSTD,
-        RSTD.stride(0),
-        DX,
-        DX.stride(0),
-        _DW,
-        _DW.stride(0),
-        _DB,
-        _DB.stride(0),
-        dY,
-        dY.stride(0),
-        n_rows,
-        n_cols,
-        rows_per_program=rows_per_program,
-        BLOCK_SIZE=BLOCK_SIZE,
-        **kernel_args,
-    )
-    DX = DX.view(*shape)
-    DW = _DW.sum(dim=0).to(W.dtype)
-    DB = _DB.sum(dim=0).to(B.dtype)
-    return DX, DW, DB
-class LigerLayerNormFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, X, W, B, eps):
-        Y, X, Mean, RSTD, BLOCK_SIZE, num_warps = layer_norm_forward(X, W, B, eps)
-        ctx.save_for_backward(X, W, B, Mean, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        X, W, B, Mean, RSTD = ctx.saved_tensors
-        DX, DW, DB = layer_norm_backward(dY, X, W, B, Mean, RSTD)
-        return DX, DW, DB, None

build/torch-rocm/layers.py DELETED Viewed

@@ -1,514 +0,0 @@
-import inspect
-from dataclasses import dataclass
-from typing import Optional, Tuple
-import torch
-import torch.nn as nn
-from .cross_entropy import LigerCrossEntropyFunction
-from .dyt import LigerDyTFunction
-from .fused_linear_cross_entropy import LigerFusedLinearCrossEntropyFunction
-from .geglu import LigerGELUMulFunction
-from .group_norm import LigerGroupNormFunction
-from .jsd import LigerJSDFunction
-from .kl_div import LigerKLDivLossFunction
-from .layer_norm import LigerLayerNormFunction
-from .qwen2vl_mrope import LigerQwen2VLMRopeFunction
-from .rms_norm import LigerRMSNormFunction
-from .rope import LigerRopeFunction
-from .swiglu import LigerSiLUMulFunction
-from .tiled_mlp import apply_tiled_mlp
-from .tvd import LigerTVDLossFunction
-class LigerRMSNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        offset: float = 0.0,
-        casting_mode: str = "llama",
-        init_fn: str = "ones",
-        in_place: bool = True,
-        row_mode: Optional[bool] = None,
-        elementwise_affine: bool = True,
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.offset = offset
-        self.casting_mode = casting_mode
-        self.in_place = in_place
-        self.row_mode = row_mode
-        self.elementwise_affine = elementwise_affine
-        if elementwise_affine:
-            init = torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size)
-            self.weight = nn.Parameter(init)
-        else:
-            self.register_parameter("weight", None)
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        return LigerRMSNormFunction.apply(
-            hidden_states,
-            self.weight,
-            self.variance_epsilon,
-            self.offset,
-            self.casting_mode,
-            self.in_place,
-            self.row_mode,
-        )
-    def extra_repr(self) -> str:
-        return (
-            f"{self.hidden_size}, eps={self.variance_epsilon}, offset={self.offset}, "
-            f"in_place={self.in_place}, row_mode={self.row_mode}"
-        )
-class LigerLayerNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size))
-        self.bias = nn.Parameter(torch.randn(hidden_size) if bias else torch.zeros(hidden_size))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        return LigerLayerNormFunction.apply(hidden_states, self.weight, self.bias, self.variance_epsilon)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, eps={self.variance_epsilon}"
-class LigerGroupNorm(nn.Module):
-    def __init__(
-        self,
-        num_channels: int,
-        num_groups: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        assert num_channels % num_groups == 0, (
-            f"num_channels ({num_channels}) must be divisible by num_groups ({num_groups})"
-        )
-        self.num_channels = num_channels
-        self.num_groups = num_groups
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(num_channels) if init_fn == "ones" else torch.zeros(num_channels))
-        self.bias = nn.Parameter(torch.randn(num_channels) if bias else torch.zeros(num_channels))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        assert hidden_states.dim() >= 3, f"Input must have at least 3 dimensions, got {hidden_states.dim()}"
-        assert hidden_states.size(1) == self.num_channels, (
-            f"Input must have {self.num_channels} channels, got {hidden_states.size(1)}"
-        )
-        return LigerGroupNormFunction.apply(
-            hidden_states,
-            self.weight,
-            self.bias,
-            self.num_channels,
-            self.num_groups,
-            self.variance_epsilon,
-        )
-    def extra_repr(self) -> str:
-        return f"num_channels={self.num_channels}, num_groups={self.num_groups}, eps={self.variance_epsilon}"
-class LigerDyT(nn.Module):
-    def __init__(self, hidden_size: int, beta: bool = True, init_alpha: float = 0.5):
-        super().__init__()
-        self.hidden_size = hidden_size
-        self.init_alpha = init_alpha
-        self.alpha = nn.Parameter(torch.ones(1) * init_alpha)
-        self.gamma = nn.Parameter(torch.ones(hidden_size))
-        self.beta = nn.Parameter(torch.zeros(hidden_size)) if beta else None
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return LigerDyTFunction.apply(x, self.alpha, self.gamma, self.beta)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, init_alpha={self.init_alpha}, beta={self.beta is not None}"
-class LigerCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.weight = weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-    def forward(self, _input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
-        loss, _, _, _ = LigerCrossEntropyFunction.apply(
-            _input,
-            target,
-            self.weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            False,
-            False,
-        )
-        return loss
-class LigerFusedLinearCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        ce_weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-        accum_dtype: Optional[torch.dtype] = None,
-        use_token_scaling: bool = False,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.ce_weight = ce_weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-        self.accum_dtype = accum_dtype
-        self.use_token_scaling = use_token_scaling
-    def forward(
-        self,
-        lin_weight: torch.Tensor,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        bias: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        loss, _, _, _ = LigerFusedLinearCrossEntropyFunction.apply(
-            _input,
-            lin_weight,
-            target,
-            bias,
-            self.ce_weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            self.accum_dtype,
-            self.use_token_scaling,
-            False,
-            False,
-        )
-        return loss
-class LigerJSD(nn.Module):
-    def __init__(self, beta: float = 0.5, ignore_index: int = -100):
-        super().__init__()
-        self.beta = beta
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        log_q: torch.Tensor,
-        log_p: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerJSDFunction.apply(log_q, log_p, shift_labels, self.beta, self.ignore_index)
-class LigerKLDIVLoss(nn.KLDivLoss):
-    def __init__(self, eps: float = 1e-10, *args, **kwargs):
-        super().__init__(*args, **kwargs)
-        self.eps = eps
-    def forward(self, y_pred: torch.Tensor, y_true: torch.Tensor) -> torch.Tensor:
-        return LigerKLDivLossFunction.apply(y_pred, y_true, self.reduction, self.log_target, self.eps)
-class LigerTVDLoss(nn.Module):
-    def __init__(self, reduction: str = "batchmean", ignore_index: int = -100):
-        super().__init__()
-        self.reduction = reduction
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        p: torch.Tensor,
-        q: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerTVDLossFunction.apply(p, q, shift_labels, self.reduction, self.ignore_index)
-class LigerSwiGLUMLP(nn.Module):
-    """SwiGLU MLP block. ``config`` must expose ``hidden_size``, ``intermediate_size``,
-    and ``hidden_act`` (must be ``silu`` or ``swish``)."""
-    def __init__(self, config):
-        super().__init__()
-        if config.hidden_act not in ("silu", "swish"):
-            raise ValueError(f"Activation function {config.hidden_act} not supported.")
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.down_proj(LigerSiLUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
-class LigerGEGLUMLP(nn.Module):
-    """GEGLU MLP block. ``config`` must expose ``hidden_size`` and ``intermediate_size``.
-    Uses the tanh approximation of GELU (matches Gemma 1/1.1/2)."""
-    def __init__(self, config):
-        super().__init__()
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.down_proj(LigerGELUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
-class LigerTiledGEGLUMLP(nn.Module):
-    gate_proj: nn.Linear
-    up_proj: nn.Linear
-    down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerGELUMulFunction.apply(gate, up))
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        compute_params = [p for p in self.parameters() if p.requires_grad]
-        return apply_tiled_mlp(
-            fn=self._mlp_forward,
-            mlp_module=self,
-            x=x,
-            num_shards=self.num_shards,
-            compute_params=compute_params,
-        )
-class LigerTiledSwiGLUMLP(nn.Module):
-    gate_proj: nn.Linear
-    up_proj: nn.Linear
-    down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerSiLUMulFunction.apply(gate, up))
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        compute_params = [p for p in self.parameters() if p.requires_grad]
-        return apply_tiled_mlp(
-            fn=self._mlp_forward,
-            mlp_module=self,
-            x=x,
-            num_shards=self.num_shards,
-            compute_params=compute_params,
-        )
-@dataclass
-class CrossEntropyOutput:
-    loss: torch.Tensor
-    z_loss: Optional[torch.Tensor] = None
-    token_accuracy: Optional[torch.Tensor] = None
-    predicted_tokens: Optional[torch.Tensor] = None
-def liger_fused_linear_cross_entropy(
-    input: torch.Tensor,
-    weight: torch.Tensor,
-    target: torch.Tensor,
-    bias: Optional[torch.Tensor] = None,
-    ce_weight: Optional[torch.Tensor] = None,
-    ignore_index: int = -100,
-    lse_square_scale: float = 0.0,
-    label_smoothing: float = 0.0,
-    reduction: str = "mean",
-    softcap: Optional[float] = None,
-    return_z_loss: bool = False,
-    accum_dtype: Optional[torch.dtype] = None,
-    use_token_scaling: bool = False,
-    return_token_accuracy: bool = False,
-    return_predicted_tokens: bool = False,
-):
-    loss, z_loss, token_accuracy, predicted_tokens = LigerFusedLinearCrossEntropyFunction.apply(
-        input,
-        weight,
-        target,
-        bias,
-        ce_weight,
-        ignore_index,
-        lse_square_scale,
-        label_smoothing,
-        reduction,
-        softcap,
-        return_z_loss,
-        accum_dtype,
-        use_token_scaling,
-        return_token_accuracy,
-        return_predicted_tokens,
-    )
-    if not return_z_loss and not return_token_accuracy and not return_predicted_tokens:
-        return loss
-    return CrossEntropyOutput(
-        loss=loss,
-        z_loss=z_loss,
-        token_accuracy=token_accuracy,
-        predicted_tokens=predicted_tokens,
-    )
-def LigerForCausalLMLoss(
-    hidden_states: torch.Tensor,
-    lm_head_weight: torch.Tensor,
-    labels: torch.Tensor,
-    hidden_size: int,
-    num_items_in_batch: Optional[int] = None,
-    ignore_index: int = -100,
-    shift_labels: Optional[torch.Tensor] = None,
-    final_logit_softcapping: Optional[float] = None,
-    return_token_accuracy: bool = False,
-    return_predicted_tokens: bool = False,
-    **kwargs,
-):
-    """Drop-in replacement for ``transformers.loss.ForCausalLMLoss`` that fuses the
-    final ``lm_head`` projection with the cross-entropy loss. Returns a scalar
-    ``loss`` by default; returns a :class:`CrossEntropyOutput` when
-    ``return_token_accuracy`` or ``return_predicted_tokens`` is set."""
-    applicable_params = inspect.signature(liger_fused_linear_cross_entropy).parameters
-    kwargs = {k: v for k, v in kwargs.items() if k in applicable_params}
-    if shift_labels is None:
-        labels = nn.functional.pad(labels, (0, 1), value=ignore_index)
-        shift_labels = labels[..., 1:].contiguous()
-    hidden_states = hidden_states.view(-1, hidden_size)
-    shift_labels = shift_labels.view(-1).to(hidden_states.device)
-    reduction = "sum" if num_items_in_batch is not None else "mean"
-    result = liger_fused_linear_cross_entropy(
-        hidden_states,
-        lm_head_weight,
-        shift_labels,
-        reduction=reduction,
-        ignore_index=ignore_index,
-        softcap=final_logit_softcapping,
-        return_token_accuracy=return_token_accuracy,
-        return_predicted_tokens=return_predicted_tokens,
-        **kwargs,
-    )
-    if isinstance(result, CrossEntropyOutput):
-        loss = result.loss
-        token_accuracy = result.token_accuracy
-        predicted_tokens = result.predicted_tokens
-    else:
-        loss = result
-        token_accuracy = None
-        predicted_tokens = None
-    if reduction == "sum":
-        loss = loss / num_items_in_batch
-    if return_token_accuracy or return_predicted_tokens:
-        return CrossEntropyOutput(
-            loss=loss,
-            token_accuracy=token_accuracy,
-            predicted_tokens=predicted_tokens,
-        )
-    return loss
-def liger_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    position_ids: Optional[torch.Tensor] = None,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply standard rotary positional embedding to ``q`` and ``k``."""
-    return LigerRopeFunction.apply(q, k, cos, sin, position_ids, unsqueeze_dim)
-def liger_multimodal_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    mrope_section,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply Qwen2-VL multimodal rotary positional embedding (M-RoPE) to ``q`` and ``k``."""
-    return LigerQwen2VLMRopeFunction.apply(q, k, cos, sin, mrope_section, unsqueeze_dim)
-__all__ = [
-    "LigerRMSNorm",
-    "LigerLayerNorm",
-    "LigerGroupNorm",
-    "LigerDyT",
-    "LigerCrossEntropyLoss",
-    "LigerFusedLinearCrossEntropyLoss",
-    "LigerJSD",
-    "LigerKLDIVLoss",
-    "LigerTVDLoss",
-    "LigerSwiGLUMLP",
-    "LigerGEGLUMLP",
-    "LigerTiledGEGLUMLP",
-    "LigerTiledSwiGLUMLP",
-    "CrossEntropyOutput",
-    "liger_fused_linear_cross_entropy",
-    "LigerForCausalLMLoss",
-    "liger_rotary_pos_emb",
-    "liger_multimodal_rotary_pos_emb",
-]

build/torch-rocm/liger_kernels/__init__.py DELETED Viewed

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch-rocm/metadata.json DELETED Viewed

@@ -1,10 +0,0 @@
-{
-  "name": "liger-kernels",
-  "id": "_liger_kernels_rocm_ab435e2",
-  "version": 1,
-  "license": "BSD-2-Clause",
-  "python-depends": [],
-  "backend": {
-    "type": "rocm"
-  }
-}

build/torch-rocm/qwen2vl_mrope.py DELETED Viewed

@@ -1,222 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-@triton.jit
-def _triton_qwen2vl_mrope(
-    q_ptr,
-    k_ptr,
-    cos,
-    sin,
-    sl,
-    bs: tl.constexpr,
-    n_qh: tl.constexpr,
-    n_kh: tl.constexpr,
-    hd: tl.constexpr,
-    pad_n_qh: tl.constexpr,
-    pad_n_kh: tl.constexpr,
-    pad_hd: tl.constexpr,
-    mrope_section_t: tl.constexpr,
-    mrope_section_h: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BACKWARD_PASS: tl.constexpr = False,
-):
-    pid = tl.program_id(0)
-    # locate start address
-    q_ptr = q_ptr + pid * (n_qh * hd)
-    k_ptr = k_ptr + pid * (n_kh * hd)
-    # ####################################################################
-    # get the cos(mθ_{i...d/2}) and sin(mθ_{i...d/2}) for token position
-    # m of this program instance
-    # ####################################################################
-    # 1. program instances are laid out in a 1D vector of size bsz * seq_len, which
-    # effectively represents a 2D grid of size [bsz, seq_len] with seq_len dimension
-    # being the fastest changing dimension. Thus we can simply do pid // sl to get the batch index
-    # and pid % sl to get the sequence index.
-    # 2. We only need the left half of cos and sin matrix because the right half is just
-    # a clone of the left half.
-    t_end = mrope_section_t
-    h_end = t_end + mrope_section_h
-    t_cos = cos + pid * hd
-    h_cos = t_cos + bs * sl * hd
-    w_cos = h_cos + bs * sl * hd
-    t_sin = sin + pid * hd
-    h_sin = t_sin + bs * sl * hd
-    w_sin = h_sin + bs * sl * hd
-    cos_offsets = tl.arange(0, pad_hd // 2)
-    t_mask = cos_offsets < t_end
-    h_mask = (t_end <= cos_offsets) & (cos_offsets < h_end)
-    w_mask = (h_end <= cos_offsets) & (cos_offsets < hd // 2)
-    t_cos_row = tl.load(t_cos + cos_offsets, mask=t_mask, other=0)
-    h_cos_row = tl.load(h_cos + cos_offsets, mask=h_mask, other=0)
-    w_cos_row = tl.load(w_cos + cos_offsets, mask=w_mask, other=0)
-    t_sin_row = tl.load(t_sin + cos_offsets, mask=t_mask, other=0)
-    h_sin_row = tl.load(h_sin + cos_offsets, mask=h_mask, other=0)
-    w_sin_row = tl.load(w_sin + cos_offsets, mask=w_mask, other=0)
-    cos_row = t_cos_row + h_cos_row + w_cos_row
-    sin_row = t_sin_row + h_sin_row + w_sin_row
-    # ####################################################################
-    # Load the left and right half of q and k for the current
-    # program instance (i.e. for the current token) separately
-    # ####################################################################
-    # left half of the head
-    first_half_q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_half_k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    q_tile_1 = tl.load(q_ptr + first_half_q_offsets, mask=first_q_mask, other=0).to(sin_row.dtype)
-    k_tile_1 = tl.load(k_ptr + first_half_k_offsets, mask=first_k_mask, other=0).to(sin_row.dtype)
-    # right half of the head
-    second_half_q_offsets = first_half_q_offsets + (hd // 2)
-    second_half_k_offsets = first_half_k_offsets + (hd // 2)
-    second_q_mask = first_q_mask
-    second_k_mask = first_k_mask
-    q_tile_2 = tl.load(q_ptr + second_half_q_offsets, mask=second_q_mask, other=0).to(sin_row.dtype)
-    k_tile_2 = tl.load(k_ptr + second_half_k_offsets, mask=second_k_mask, other=0).to(sin_row.dtype)
-    if not BACKWARD_PASS:
-        # y = [x1, x2] * [cos, cos] + [-x2, x1] * [sin, sin]
-        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-    else:
-        # with some math, we can get:
-        # dy = [dx1, dx2] * [cos, cos] + [-dx2, dx1] * [-sin, -sin]
-        new_q_tile_1 = q_tile_1 * cos_row + q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row - q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row + k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row - k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-def qwen2vl_mrope_forward(q, k, cos, sin, mrope_section):
-    # transpose it back to the physical shape because Triton looks at the physical storage
-    # note: q and k are incontiguous before the transformation and will become contiguous after transpose
-    q = q.transpose(1, 2)
-    k = k.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = q.shape
-    n_kv_head = k.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure tensors passed into the kernel are contiguous. It will be no-op if they are already contiguous
-    q = q.contiguous()
-    k = k.contiguous()
-    cos = cos.contiguous()
-    sin = sin.contiguous()
-    _triton_qwen2vl_mrope[(n_row,)](
-        q,
-        k,
-        cos,
-        sin,
-        seq_len,
-        batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        mrope_section[0],
-        mrope_section[1],
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=False,
-    )
-    return q.transpose(1, 2), k.transpose(1, 2), cos, sin
-def qwen2vl_mrope_backward(dq, dk, cos, sin, mrope_section):
-    dq = dq.transpose(1, 2)
-    dk = dk.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = dq.shape
-    n_kv_head = dk.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure dq and dk are contiguous
-    dq = dq.contiguous()
-    dk = dk.contiguous()
-    # backward is similar to forward except swapping few ops
-    _triton_qwen2vl_mrope[(n_row,)](
-        dq,
-        dk,
-        cos,
-        sin,
-        seq_len,
-        batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        mrope_section[0],
-        mrope_section[1],
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=True,
-    )
-    return dq.transpose(1, 2), dk.transpose(1, 2)
-class LigerQwen2VLMRopeFunction(torch.autograd.Function):
-    """
-    Triton implementation of the Qwen2VL Multimodal Rotary Positional Embedding (M-RoPE) operation.
-    Please find the corresponding HuggingFace implementation here:
-    https://github.com/huggingface/transformers/blob/main/src/transformers/models/qwen2_vl/modeling_qwen2_vl.py
-    """
-    @staticmethod
-    def forward(ctx, q, k, cos, sin, mrope_section, unsqueeze_dim=1):
-        """
-        q size: (bsz, n_q_head, seq_len, head_dim)
-        k size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (3, bsz, seq_len, head_dim)
-        sin size: (3, bsz, seq_len, head_dim)
-        """
-        q, k, cos, sin = qwen2vl_mrope_forward(q, k, cos, sin, mrope_section)
-        ctx.save_for_backward(cos, sin)
-        ctx.mrope_section = mrope_section
-        return q, k
-    def backward(ctx, dq, dk):
-        """
-        dq size: (bsz, n_q_head, seq_len, head_dim)
-        dk size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (3, bsz, seq_len, head_dim)
-        sin size: (3, bsz, seq_len, head_dim)
-        """
-        cos, sin = ctx.saved_tensors
-        mrope_section = ctx.mrope_section
-        dq, dk = qwen2vl_mrope_backward(dq, dk, cos, sin, mrope_section)
-        return dq, dk, None, None, None, None

build/torch-rocm/rms_norm.py DELETED Viewed

@@ -1,654 +0,0 @@
-"""
-This file incorporates code from Unsloth licensed under the Apache License, Version 2.0.
-See the original Unsloth repository at https://github.com/unslothai/unsloth.
-The following line
-https://github.com/linkedin/Liger-Kernel/blob/7382a8761f9af679482b968f9348013d933947c7/src/liger_kernel/ops/rms_norm.py#L30
-is based on code from Unsloth, located at:
-https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-Modifications made by Yanning Chen, 2024.
-"""
-import math
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import set_large_grf_mode
-from .utils import torch_to_triton_dtype
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-_CASTING_MODE_NONE: tl.constexpr = tl.constexpr(-1)
-_CASTING_MODE_LLAMA: tl.constexpr = tl.constexpr(0)
-_CASTING_MODE_GEMMA: tl.constexpr = tl.constexpr(1)
-@triton.jit
-def _rms_norm_forward_kernel(
-    Y_ptr,
-    Y_row_stride,
-    X_ptr,
-    X_row_stride,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    n_cols,
-    eps,
-    offset,
-    casting_mode: tl.constexpr,  # constexpr so the `if` blocks can be optimized out
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    y_i = (x_i / (RMS)) * (offset + wi), RMS = sqrt(sum(x_i^2) / N)
-    Reference:
-    1. https://triton-lang.org/main/getting-started/tutorials/05-layer-norm.html
-    2. https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-    3. https://arxiv.org/pdf/1910.07467
-    """
-    row_idx = tl.program_id(0).to(tl.int64)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    y_base = Y_ptr + row_idx * Y_row_stride
-    x_base = X_ptr + row_idx * X_row_stride
-    rstd_base = RSTD_ptr + row_idx * RSTD_row_stride
-    X_row = tl.load(x_base + col_offsets, mask=mask, other=0)
-    X_row_dtype = X_row.dtype
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0)
-    # On Llama, only rstd is computed on fp32
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(tl.float32)
-    # Gemma computes everything on fp32, and then casts back the output to the original dtype
-    if casting_mode == _CASTING_MODE_GEMMA:
-        if elementwise_affine:
-            W_row = W_row.to(tl.float32)
-        X_row = X_row.to(tl.float32)
-    if casting_mode == _CASTING_MODE_NONE:
-        eps = eps.to(X_row_dtype)
-        offset = offset.to(X_row_dtype)
-    mean_square = tl.sum(X_row * X_row, axis=0) / n_cols
-    rstd = rsqrt(mean_square + eps)
-    # We can save time by caching rms with minimal memory overhead
-    # because rms is much smaller compared to X_row, as rms is for each row.
-    # However, on the computation side, it can save 4 operations (*, sum, /, sqrt).
-    tl.store(rstd_base, rstd)
-    X_row = X_row * rstd
-    # On Llama, the multiplication with the weight is done on the original dtype
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(X_row_dtype)
-    if elementwise_affine:
-        Y_row = X_row * (offset + W_row)
-    else:
-        Y_row = X_row
-    if casting_mode == _CASTING_MODE_GEMMA:
-        Y_row = Y_row.to(X_row_dtype)
-    tl.store(y_base + col_offsets, Y_row, mask=mask)
-@triton.jit
-def _rms_norm_backward_kernel(
-    dY_ptr,
-    dY_row_stride,
-    dX_ptr,
-    dX_row_stride,
-    X_ptr,
-    X_row_stride,
-    X_dtype: tl.constexpr,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    dW_ptr,
-    dW_row_stride,
-    n_rows,
-    n_cols,
-    offset,
-    rows_per_program,
-    casting_mode: tl.constexpr,
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    dx = (1 / RMS) * [dy * (w + offset - (1 / N) * (1 / RMS^2) * ((dy * (w + offset)) dot x) * x]. * means element-wise multiplication, whileas dot means dot product
-    dw = sum(dy * (x / RMS)). summation over BxT dimension
-    """
-    row_block_id = tl.program_id(0).to(tl.int64)
-    row_start = row_block_id * rows_per_program
-    row_end = min((row_block_id + 1) * rows_per_program, n_rows)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    if elementwise_affine:
-        dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0.0)
-        W_row = W_row + offset
-    for row_idx in range(row_start, row_end):
-        dy_base = dY_ptr + row_idx * dY_row_stride
-        dx_base = dX_ptr + row_idx * dX_row_stride
-        x_base = X_ptr + row_idx * X_row_stride
-        rstd_base = RSTD_ptr + row_idx * RSTD_row_stride
-        dY_row = tl.load(dy_base + col_offsets, mask=mask, other=0.0)
-        X_row = tl.load(x_base + col_offsets, mask=mask, other=0.0)
-        # Get cached rms
-        rstd_row = tl.load(rstd_base)
-        X_row = X_row.to(tl.float32)
-        # Different bacward graphs for different casting modes
-        if casting_mode == _CASTING_MODE_LLAMA:
-            if elementwise_affine:
-                m = (dY_row * W_row).to(tl.float32)
-            else:
-                m = dY_row.to(tl.float32)
-        elif casting_mode == _CASTING_MODE_GEMMA:
-            dY_row = dY_row.to(tl.float32)
-            if elementwise_affine:
-                m = dY_row * W_row
-            else:
-                m = dY_row
-        else:
-            if elementwise_affine:
-                m = dY_row * W_row
-            else:
-                m = dY_row
-        dX_row = rstd_row * m
-        dX_row += (rstd_row) * (-(1 / n_cols) * rstd_row * rstd_row * tl.sum(m * X_row, axis=0) * X_row)
-        if elementwise_affine:
-            # calculate the gradient of W
-            if casting_mode == _CASTING_MODE_LLAMA:
-                dW_row += dY_row * (X_row * rstd_row).to(X_dtype)
-            else:
-                # here X_row is already in fp32 (see previous if block)
-                dW_row += dY_row * (X_row * rstd_row)
-        tl.store(dx_base + col_offsets, dX_row.to(X_dtype), mask=mask)
-    if elementwise_affine:
-        tl.store(dW_ptr + row_block_id * dW_row_stride + col_offsets, dW_row, mask=mask)
-@triton.jit
-def _block_rms_norm_forward_kernel(
-    Y_ptr,
-    Y_row_stride,
-    X_ptr,
-    X_row_stride,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    n_rows,
-    n_cols,
-    eps,
-    offset,
-    casting_mode: tl.constexpr,  # constexpr so the `if` blocks can be optimized out
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BLOCK_ROW: tl.constexpr,
-):
-    """
-    y_i = (x_i / (RMS)) * (offset + wi), RMS = sqrt(sum(x_i^2) / N)
-    Reference:
-    1. https://triton-lang.org/main/getting-started/tutorials/05-layer-norm.html
-    2. https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-    3. https://arxiv.org/pdf/1910.07467
-    """
-    row_idx = tl.program_id(0) * BLOCK_ROW + tl.arange(0, BLOCK_ROW)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    row_mask = row_idx < n_rows
-    col_mask = col_offsets < n_cols
-    X_row = tl.load(
-        X_ptr + row_idx[:, None] * X_row_stride + col_offsets[None, :],
-        mask=row_mask[:, None] & col_mask[None, :],
-        other=0,
-    )
-    X_row_dtype = X_row.dtype
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=col_mask, other=0)
-    # On Llama, only rstd is computed on fp32
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(tl.float32)
-    # Gemma computes everything on fp32, and then casts back the output to the original dtype
-    if casting_mode == _CASTING_MODE_GEMMA:
-        if elementwise_affine:
-            W_row = W_row.to(tl.float32)
-        X_row = X_row.to(tl.float32)
-    if casting_mode == _CASTING_MODE_NONE:
-        eps = eps.to(X_row_dtype)
-        offset = offset.to(X_row_dtype)
-    mean_square = tl.sum(X_row * X_row, axis=1) / n_cols
-    rstd = rsqrt(mean_square + eps)
-    # We can save time by caching rms with minimal memory overhead
-    # because rms is much smaller compared to X_row, as rms is for each row.
-    # However, on the computation side, it can save 4 operations (*, sum, /, sqrt).
-    tl.store(RSTD_ptr + row_idx * RSTD_row_stride, rstd, row_mask)
-    X_row = X_row * rstd[:, None]
-    # On Llama, the multiplication with the weight is done on the original dtype
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(X_row_dtype)
-    if elementwise_affine:
-        Y_row = X_row * (offset + W_row)[None, :]
-    else:
-        Y_row = X_row
-    if casting_mode == _CASTING_MODE_GEMMA:
-        Y_row = Y_row.to(X_row_dtype)
-    tl.store(
-        Y_ptr + row_idx[:, None] * Y_row_stride + col_offsets[None, :],
-        Y_row,
-        mask=row_mask[:, None] & col_mask[None, :],
-    )
-@triton.jit
-def _block_rms_norm_backward_kernel(
-    dY_ptr,
-    dY_row_stride,
-    dX_ptr,
-    dX_row_stride,
-    X_ptr,
-    X_row_stride,
-    X_dtype: tl.constexpr,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    dW_ptr,
-    dW_row_stride,
-    n_rows,
-    n_cols,
-    offset,
-    casting_mode: tl.constexpr,
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BLOCK_ROW: tl.constexpr,
-):
-    """
-    dx = (1 / RMS) * [dy * (w + offset - (1 / N) * (1 / RMS^2) * ((dy * (w + offset)) dot x) * x]. * means element-wise multiplication, whileas dot means dot product
-    dw = sum(dy * (x / RMS)). summation over BxT dimension
-    """
-    pid = tl.program_id(0).cast(tl.int64)
-    NUM_SMS = tl.num_programs(0)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    col_mask = col_offsets < n_cols
-    if elementwise_affine:
-        dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-        W_row = tl.load(W_ptr + col_offsets, mask=col_mask, other=0.0)
-        W_row = W_row + offset
-    for start in range(pid * BLOCK_ROW, n_rows, NUM_SMS * BLOCK_ROW):
-        row_idx = start + tl.arange(0, BLOCK_ROW)
-        row_mask = row_idx < n_rows
-        dY_row = tl.load(
-            dY_ptr + row_idx[:, None] * dY_row_stride + col_offsets[None, :],
-            mask=row_mask[:, None] & col_mask[None, :],
-            other=0.0,
-        )
-        X_row = tl.load(
-            X_ptr + row_idx[:, None] * X_row_stride + col_offsets[None, :],
-            mask=row_mask[:, None] & col_mask[None, :],
-            other=0.0,
-        )
-        # Get cached rms
-        rstd_row = tl.load(RSTD_ptr + row_idx * RSTD_row_stride, row_mask)
-        X_row = X_row.to(tl.float32)
-        # Different bacward graphs for different casting modes
-        if casting_mode == _CASTING_MODE_LLAMA:
-            if elementwise_affine:
-                m = (dY_row * W_row[None, :]).to(tl.float32)
-            else:
-                m = dY_row.to(tl.float32)
-        elif casting_mode == _CASTING_MODE_GEMMA:
-            dY_row = dY_row.to(tl.float32)
-            if elementwise_affine:
-                m = dY_row * W_row[None, :]
-            else:
-                m = dY_row
-        else:
-            if elementwise_affine:
-                m = dY_row * W_row[None, :]
-            else:
-                m = dY_row
-        dX_row = rstd_row[:, None] * m
-        dX_row += (rstd_row[:, None]) * (
-            -(1 / n_cols) * (rstd_row * rstd_row * tl.sum(m * X_row, axis=1))[:, None] * X_row
-        )
-        if elementwise_affine:
-            if casting_mode == _CASTING_MODE_LLAMA:
-                # TODO(tcc): use tl.sum(..., dtype=tl.float32) once we upgrade to triton>=3.3.0
-                dW_row += tl.sum((dY_row * (X_row * rstd_row[:, None]).to(X_dtype)).to(tl.float32), 0)
-            else:
-                # here X_row is already in fp32 (see previous if block)
-                dW_row += tl.sum(dY_row * (X_row * rstd_row[:, None]), 0)
-        tl.store(
-            dX_ptr + row_idx[:, None] * dX_row_stride + col_offsets[None, :],
-            dX_row,
-            mask=row_mask[:, None] & col_mask[None, :],
-        )
-    if elementwise_affine:
-        tl.store(dW_ptr + pid * dW_row_stride + col_offsets, dW_row, mask=col_mask)
-_str_to_casting_mode = {
-    "llama": _CASTING_MODE_LLAMA.value,
-    "gemma": _CASTING_MODE_GEMMA.value,
-    "none": _CASTING_MODE_NONE.value,
-}
-def rms_norm_forward(X, W, eps, offset, casting_mode, row_mode):
-    if not isinstance(casting_mode, int):
-        assert casting_mode in _str_to_casting_mode, f"Invalid casting mode: {casting_mode}"
-        casting_mode = _str_to_casting_mode[casting_mode]
-    else:
-        assert casting_mode in _str_to_casting_mode.values(), f"Invalid casting mode: {casting_mode}"
-    shape = X.shape
-    dim = shape[-1]
-    X = X.view(-1, dim)
-    n_rows, n_cols = X.shape
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    Y = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    # RSTD is to cache rstd for each row
-    # RSTD is always computed/stored in fp32 if we are using Llama or Gemma casting mode
-    rstd_dtype = torch.float32 if casting_mode in (_CASTING_MODE_LLAMA.value, _CASTING_MODE_GEMMA.value) else X.dtype
-    RSTD = torch.empty(n_rows, dtype=rstd_dtype, device=X.device)
-    if W is not None:
-        # Check constraints.
-        assert X.shape[1] == W.shape[0], (
-            "Incompatible hidden size dimension between tensor1.shape[1] and tensor2.shape[0]"
-        )
-        elementwise_affine = True
-    else:
-        elementwise_affine = False
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    if BLOCK_SIZE > 256 or n_rows < 4096 * 8 or row_mode:
-        _rms_norm_forward_kernel[(n_rows,)](
-            Y,
-            Y.stride(0),
-            X,
-            X.stride(0),
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            n_cols,
-            eps,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    else:
-        BLOCK_ROW = 16
-        kernel_args["BLOCK_ROW"] = BLOCK_ROW
-        _block_rms_norm_forward_kernel[(triton.cdiv(n_rows, BLOCK_ROW),)](
-            Y,
-            Y.stride(0),
-            X,
-            X.stride(0),
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            n_rows,
-            n_cols,
-            eps,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    return Y.view(*shape), X, RSTD, BLOCK_SIZE, num_warps, casting_mode
-def rms_norm_backward(dY, X, W, RSTD, offset, casting_mode, BLOCK_SIZE, num_warps, in_place, row_mode):
-    shape = dY.shape
-    dim = shape[-1]
-    dY = dY.view(-1, dim)
-    n_rows, n_cols = dY.shape
-    sm_count = 1
-    if X.device.type == "cuda":
-        sm_count = torch.cuda.get_device_properties(X.device).multi_processor_count
-    elif X.device.type == "xpu":
-        sm_count = torch.xpu.get_device_properties(X.device).gpu_eu_count
-    elif X.device.type == "npu":
-        sm_count = get_npu_core_count()
-    if W is not None:
-        # fp32 for numerical stability especially.
-        _dW = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-        elementwise_affine = True
-    else:
-        _dW = None
-        elementwise_affine = False
-    if n_cols > BLOCK_SIZE:
-        raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
-    rows_per_program = math.ceil(n_rows / sm_count)
-    grid = (sm_count,)
-    if in_place is True:
-        dX = dY
-    else:
-        dX = torch.zeros_like(dY)
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    if BLOCK_SIZE > 256 or n_rows < 4096 * 8 or row_mode:
-        _rms_norm_backward_kernel[grid](
-            dY,
-            dY.stride(0),
-            dX,
-            dX.stride(0),
-            X,
-            X.stride(0),
-            torch_to_triton_dtype[X.dtype],
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            _dW,
-            _dW.stride(0) if elementwise_affine else 0,
-            n_rows,
-            n_cols,
-            offset,
-            rows_per_program,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    else:
-        BLOCK_ROW = 16
-        kernel_args["BLOCK_ROW"] = BLOCK_ROW
-        _block_rms_norm_backward_kernel[grid](
-            dY,
-            dY.stride(0),
-            dX,
-            dX.stride(0),
-            X,
-            X.stride(0),
-            torch_to_triton_dtype[X.dtype],
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            _dW,
-            _dW.stride(0) if elementwise_affine else 0,
-            n_rows,
-            n_cols,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    dX = dX.view(*shape)
-    if elementwise_affine:
-        dW = _dW.sum(dim=0).to(W.dtype)
-    else:
-        dW = None
-    return dX, dW
-class LigerRMSNormFunction(torch.autograd.Function):
-    """
-    Performs RMSNorm (Root Mean Square Normalization), which normalizes the input tensor `X` using the
-    weight tensor `W`, with an optional offset and casting mode.
-    Some models use an 'offset' to shift the weight tensor `W` by a constant value. For example, Gemma
-    uses an offset of 1.0, so the computation becomes `(X / RMS(X)) * (W + 1.0)` instead of the usual
-    `(X / RMS(X)) * W`. You can pass the offset value as an argument to the forward function.
-    In addition, different models cast their inputs at different places during RMSNorm computation. For
-    example, Gemma casts everything to fp32 nefore starting the computation, while Llama casts only the
-    inverse RMS to fp32. You can specify the casting mode using the `casting_mode` argument. We currently
-    support the following casting modes (they match HuggingFace Transformers' implementations):
-    - 'llama': matches the Llama implementation, where only the inverse RMS is computed on fp32.
-    - 'gemma': matches the Gemma implementation, where everything is cast to fp32, then computed, then cast back to the original dtype.
-    - 'none': no casting is done. The computation is done in the original dtype. This saves memory and is slightly faster, but has more error w.r.t. the original implementation.
-    `in_place` option means whether to in_place modify dY to store dX. This is default to `True` to save memory. However, under certain cases, it can produce incorrect inputs.
-        For example, gemma2 uses two rmsnorm sequentially with residual in between. The resesidual part needs dY so it cannot be modified in-place.
-        Therefore, for the patching of RMSNorm in gemma2, we set `in_place` to `False`
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, X, W, eps, offset=0.0, casting_mode="llama", in_place=True, row_mode=None):
-        """
-        X: (B, T, H) or (BxT, H)
-        W: (H,)
-        """
-        if isinstance(X, torch.distributed.tensor.DTensor):
-            # Input tensor is output of a tensor parallel module and
-            # needs to be gathered to a local tensor to compute
-            # RMSE layer norm on each TP worker.
-            # TODO: support CP.
-            X = X.full_tensor()
-        Y, X, RSTD, BLOCK_SIZE, num_warps, casting_mode = rms_norm_forward(X, W, eps, offset, casting_mode, row_mode)
-        ctx.offset = offset
-        ctx.casting_mode = casting_mode
-        ctx.in_place = in_place
-        ctx.row_mode = row_mode
-        ctx.BLOCK_SIZE = BLOCK_SIZE
-        ctx.num_warps = num_warps
-        ctx.elementwise_affine = W is not None
-        if W is not None:
-            ctx.save_for_backward(X, W, RSTD)
-        else:
-            ctx.save_for_backward(X, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        """
-        Y: (B, T, H) or (BxT, H)
-        """
-        if ctx.elementwise_affine:
-            X, W, RSTD = ctx.saved_tensors
-        else:
-            X, RSTD = ctx.saved_tensors
-            W = None
-        if isinstance(dY, torch.distributed.tensor.DTensor):
-            # Gradients are output of a tensor parallel module and
-            # needs to be gathered to a local tensor for computing RMSE layer.
-            # TODO: support CP.
-            dY = dY.full_tensor()
-        dX, dW = rms_norm_backward(
-            dY, X, W, RSTD, ctx.offset, ctx.casting_mode, ctx.BLOCK_SIZE, ctx.num_warps, ctx.in_place, ctx.row_mode
-        )
-        return dX, dW, None, None, None, None, None

build/torch-rocm/rope.py DELETED Viewed

@@ -1,239 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-@triton.jit
-def _triton_rope(
-    q_ptr,
-    q_row_stride,
-    k_ptr,
-    k_row_stride,
-    cos,
-    cos_row_stride,
-    sin,
-    sin_row_stride,
-    sl,
-    bs: tl.constexpr,
-    cos_bs: tl.constexpr,
-    n_qh: tl.constexpr,
-    n_kh: tl.constexpr,
-    hd: tl.constexpr,
-    pad_n_qh: tl.constexpr,
-    pad_n_kh: tl.constexpr,
-    pad_hd: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BACKWARD_PASS: tl.constexpr = False,
-):
-    # q size: (bsz, seq_len, num_q_heads, head_dim)
-    # q stride: (seq_len * num_q_heads * head_dim, num_q_heads * head_dim, head_dim, 1)
-    # k size: (bsz, seq_len, num_kv_heads, head_dim)
-    # k stride: (seq_len * num_kv_heads * head_dim, num_kv_heads * head_dim, head_dim, 1)
-    # cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-    # stride: (seq_len * head_dim, head_dim, 1)
-    pid = tl.program_id(0).to(tl.int64)
-    # locate start address
-    q_ptr = q_ptr + pid * q_row_stride
-    k_ptr = k_ptr + pid * k_row_stride
-    # ####################################################################
-    # get the cos(mθ_{i...d/2}) and sin(mθ_{i...d/2}) for token position
-    # m of this program instance
-    # ####################################################################
-    # 1. program instances are laid out in a 1D vector of size bsz * seq_len, which
-    # effectively represents a 2D grid of size [bsz, seq_len] with seq_len dimension
-    # being the fastest changing dimension. Thus we can simply do pid // sl to get the batch index
-    # and pid % sl to get the sequence index.
-    # 2. We only need the left half of cos and sin matrix because the right half is just
-    # a clone of the left half.
-    batch_idx = pid // sl
-    cos_row_idx = pid % sl
-    cos = cos + tl.where(
-        cos_bs == 1,
-        cos_row_idx * cos_row_stride,
-        batch_idx * (sl * cos_row_stride) + cos_row_idx * cos_row_stride,
-    )
-    sin = sin + tl.where(
-        cos_bs == 1,
-        cos_row_idx * sin_row_stride,
-        batch_idx * (sl * sin_row_stride) + cos_row_idx * sin_row_stride,
-    )
-    cos_offsets = tl.arange(0, pad_hd // 2)
-    cos_mask = cos_offsets < hd // 2
-    cos_row = tl.load(cos + cos_offsets, mask=cos_mask, other=0)
-    sin_row = tl.load(sin + cos_offsets, mask=cos_mask, other=0)
-    # ####################################################################
-    # Load the left and right half of q and k for the current
-    # program instance (i.e. for the current token) separately
-    # ####################################################################
-    # left half of the head
-    first_half_q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_half_k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    q_tile_1 = tl.load(q_ptr + first_half_q_offsets, mask=first_q_mask, other=0).to(sin_row.dtype)
-    k_tile_1 = tl.load(k_ptr + first_half_k_offsets, mask=first_k_mask, other=0).to(sin_row.dtype)
-    # right half of the head
-    second_half_q_offsets = first_half_q_offsets + (hd // 2)
-    second_half_k_offsets = first_half_k_offsets + (hd // 2)
-    second_q_mask = first_q_mask
-    second_k_mask = first_k_mask
-    q_tile_2 = tl.load(q_ptr + second_half_q_offsets, mask=second_q_mask, other=0).to(sin_row.dtype)
-    k_tile_2 = tl.load(k_ptr + second_half_k_offsets, mask=second_k_mask, other=0).to(sin_row.dtype)
-    if not BACKWARD_PASS:
-        # y = [x1, x2] * [cos, cos] + [-x2, x1] * [sin, sin]
-        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-    else:
-        # with some math, we can get:
-        # dy = [dx1, dx2] * [cos, cos] + [-dx2, dx1] * [-sin, -sin]
-        new_q_tile_1 = q_tile_1 * cos_row + q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row - q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row + k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row - k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-def rope_forward(q, k, cos, sin):
-    # transpose it back to the physical shape because Triton looks at the physical storage
-    # note: q and k are incontiguous before the transformation and will become contiguous after transpose
-    q = q.transpose(1, 2)
-    k = k.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = q.shape
-    n_kv_head = k.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure tensors passed into the kernel are contiguous. It will be no-op if they are already contiguous
-    q = q.contiguous()
-    k = k.contiguous()
-    cos = cos.contiguous()
-    sin = sin.contiguous()
-    cos_batch_size = cos.shape[0]
-    _triton_rope[(n_row,)](
-        q,
-        q.stride(1),
-        k,
-        k.stride(1),
-        cos,
-        cos.stride(-2),
-        sin,
-        sin.stride(-2),
-        seq_len,
-        batch_size,
-        cos_batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=False,
-    )
-    return q.transpose(1, 2), k.transpose(1, 2), cos, sin
-def rope_backward(dq, dk, cos, sin):
-    dq = dq.transpose(1, 2)
-    dk = dk.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = dq.shape
-    cos_batch_size = cos.shape[0]
-    n_kv_head = dk.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure dq and dk are contiguous
-    dq = dq.contiguous()
-    dk = dk.contiguous()
-    # backward is similar to forward except swapping few ops
-    _triton_rope[(n_row,)](
-        dq,
-        dq.stride(1),
-        dk,
-        dk.stride(1),
-        cos,
-        cos.stride(-2),
-        sin,
-        sin.stride(-2),
-        seq_len,
-        batch_size,
-        cos_batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=True,
-    )
-    return dq.transpose(1, 2), dk.transpose(1, 2)
-class LigerRopeFunction(torch.autograd.Function):
-    """
-    Triton implementation of the Rotary Positional Embedding (RoPE) operation. Please note that
-    this implements the HuggingFace Llama & Mistral version, whose rotation matrix is slightly different
-    than the original RoPE paper.
-    Please find the corresponding HuggingFace implementation here:
-    https://github.com/huggingface/transformers/blob/v4.40.2/src/transformers/models/llama/modeling_llama.py#L184
-    For more details about the rotation matrix used here, please refer to:
-    https://discuss.huggingface.co/t/is-llama-rotary-embedding-implementation-correct/44509/2
-    """
-    @staticmethod
-    def forward(ctx, q, k, cos, sin, position_ids=None, unsqueeze_dim=1):
-        """
-        q size: (bsz, n_q_head, seq_len, head_dim)
-        k size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        sin size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        """
-        q, k, cos, sin = rope_forward(q, k, cos, sin)
-        ctx.save_for_backward(cos, sin)
-        return q, k
-    def backward(ctx, dq, dk):
-        """
-        dq size: (bsz, n_q_head, seq_len, head_dim)
-        dk size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        sin size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        """
-        cos, sin = ctx.saved_tensors
-        dq, dk = rope_backward(dq, dk, cos, sin)
-        return dq, dk, None, None, None, None

build/torch-rocm/swiglu.py DELETED Viewed

@@ -1,176 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import ensure_contiguous
-@triton.jit
-def silu(x):
-    return x * tl.sigmoid(x)
-@triton.jit
-def _swiglu_forward_kernel(
-    a_ptr, b_ptr, c_ptr, stride, gate_multiplier, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr
-):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    a_ptr += program_id * stride
-    b_ptr += program_id * stride
-    c_ptr += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    # sigmoid requires type float32
-    a_row = tl.load(a_ptr + col_offsets, mask=mask, other=0).to(tl.float32) * gate_multiplier
-    b_row = tl.load(b_ptr + col_offsets, mask=mask, other=0)
-    c_row = silu(a_row).cast(b_row.dtype) * b_row
-    tl.store(c_ptr + col_offsets, c_row, mask=mask)
-@triton.jit
-def _swiglu_backward_kernel(
-    dc_ptr, a_ptr, b_ptr, stride, gate_multiplier, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr
-):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    dc_ptr += program_id * stride
-    a_ptr += program_id * stride
-    b_ptr += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    dc_row = tl.load(dc_ptr + col_offsets, mask=mask, other=0)
-    # sigmoid requires type float32
-    a_row = tl.load(a_ptr + col_offsets, mask=mask, other=0).to(tl.float32) * gate_multiplier
-    b_row = tl.load(b_ptr + col_offsets, mask=mask, other=0)
-    # recomputation to save memory. a_row already holds a * gate_multiplier.
-    sig_a = tl.sigmoid(a_row)
-    silu_a = a_row * sig_a
-    db_row = dc_row * silu_a
-    # chain rule pulls an extra factor of gate_multiplier through the pre-activation scaling
-    da_row = dc_row * (silu_a * (1 - sig_a) + sig_a) * b_row * gate_multiplier
-    tl.store(a_ptr + col_offsets, da_row, mask=mask)
-    tl.store(b_ptr + col_offsets, db_row, mask=mask)
-def swiglu_forward(a, b, gate_multiplier: float = 1.0):
-    ori_shape = a.shape
-    n_cols = ori_shape[-1]
-    a = a.view(-1, n_cols)
-    b = b.view(-1, n_cols)
-    c = torch.empty_like(a)
-    n_rows = a.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _swiglu_forward_kernel[(n_rows,)](
-        a,
-        b,
-        c,
-        c.stride(-2),
-        float(gate_multiplier),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a, b, c.view(*ori_shape)
-def swiglu_backward(a, b, dc, gate_multiplier: float = 1.0):
-    ori_shape = dc.shape
-    n_cols = ori_shape[-1]
-    dc = dc.view(-1, n_cols)
-    n_rows = dc.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _swiglu_backward_kernel[(n_rows,)](
-        dc,
-        a,
-        b,
-        dc.stride(-2),
-        float(gate_multiplier),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a.view(*ori_shape), b.view(*ori_shape)
-class LigerSiLUMulFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, a, b, gate_multiplier: float = 1.0, down_multiplier: float = 1.0):
-        gate_multiplier = float(gate_multiplier)
-        down_multiplier = float(down_multiplier)
-        ctx.gate_multiplier = gate_multiplier
-        ctx.down_multiplier = down_multiplier
-        if isinstance(a, torch.distributed.tensor.DTensor) or isinstance(b, torch.distributed.tensor.DTensor):
-            device_mesh, placements = (
-                (a.device_mesh, a.placements)
-                if isinstance(a, torch.distributed.tensor.DTensor)
-                else (b.device_mesh, b.placements)
-            )
-            # Assume that full tensors are gathered before and identical across
-            # the associated process groups.
-            if not isinstance(a, torch.distributed.tensor.DTensor):
-                a = torch.distributed.tensor.distribute_tensor(a, device_mesh=device_mesh, placements=placements)
-            if not isinstance(b, torch.distributed.tensor.DTensor):
-                b = torch.distributed.tensor.distribute_tensor(b, device_mesh=device_mesh, placements=placements)
-            a_local, b_local, c_local = swiglu_forward(a.to_local(), b.to_local(), gate_multiplier)
-            if down_multiplier != 1.0:
-                c_local = c_local * down_multiplier
-            ctx.save_for_backward(a_local, b_local)
-            ctx.dtensor_metadata = (device_mesh, placements)
-            return torch.distributed.tensor.DTensor.from_local(c_local, device_mesh, placements)
-        else:
-            a, b, c = swiglu_forward(a, b, gate_multiplier)
-            if down_multiplier != 1.0:
-                c = c * down_multiplier
-            ctx.save_for_backward(a, b)
-            ctx.dtensor_metadata = None
-            return c
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dc):
-        a, b = ctx.saved_tensors
-        gate_multiplier = ctx.gate_multiplier
-        down_multiplier = ctx.down_multiplier
-        if ctx.dtensor_metadata is not None:
-            device_mesh, placements = ctx.dtensor_metadata
-            # Assume that full tensors are gathered before and identical across
-            # the associated process groups.
-            dc_local = (
-                dc.to_local()
-                if isinstance(dc, torch.distributed.tensor.DTensor)
-                else torch.distributed.tensor.distribute_tensor(dc, device_mesh=device_mesh, placements=placements)
-            )
-            if down_multiplier != 1.0:
-                dc_local = dc_local * down_multiplier
-            a_local, b_local = swiglu_backward(a, b, dc_local, gate_multiplier)
-            return (
-                torch.distributed.tensor.DTensor.from_local(a_local, device_mesh, placements),
-                torch.distributed.tensor.DTensor.from_local(b_local, device_mesh, placements),
-                None,
-                None,
-            )
-        if down_multiplier != 1.0:
-            dc = dc * down_multiplier
-        a, b = swiglu_backward(a, b, dc, gate_multiplier)
-        return a, b, None, None

build/torch-rocm/tiled_mlp.py DELETED Viewed

@@ -1,136 +0,0 @@
-import math
-from typing import Callable
-from typing import List
-from typing import Optional
-import torch
-from .utils import ensure_contiguous
-class LigerTiledMLPFunction(torch.autograd.Function):
-    """
-    Based on DeepSpeed's TiledMLP:
-    https://github.com/deepspeedai/DeepSpeed/blob/v0.18.2/deepspeed/runtime/sequence_parallel/ulysses_sp.py#L838
-    Perform a tiled MLP computation to massively reduce memory usage needed to compute MLP
-    when using very long sequence lengths.
-    This module re-computes `forward` in the `backward`. So the `forward` occurs twice each iteration.
-    And if you're using activation checkpointing it then occurs thrice.
-    Args:
-        fn: the function to call on sharded inputs (e.g., mlp.forward)
-        mlp_module: the MLP nn.Module object
-        x: the input to MLP.forward (hidden_states)
-        shards: how many shards to use
-        compute_params: a list of weights engaged in the compute
-    Returns:
-        the computed hidden_states
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        fn: Callable,
-        mlp_module: torch.nn.Module,
-        x: torch.Tensor,
-        shards: int,
-        compute_params: Optional[List[torch.nn.Parameter]] = None,
-    ) -> torch.Tensor:
-        ctx.fn = fn
-        ctx.mlp_module = mlp_module
-        ctx.shards = shards
-        ctx.save_for_backward(x)
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
-        x_shards = list(torch.chunk(x, chunks=shards, dim=-2))
-        with torch.no_grad():
-            output_shards = [fn(mlp_module, x_shard) for x_shard in x_shards]
-        output_unsharded = torch.cat(output_shards, dim=-2)
-        return output_unsharded
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, *grads) -> tuple:
-        fn = ctx.fn
-        (x,) = ctx.saved_tensors
-        mlp_module = ctx.mlp_module
-        shards = ctx.shards
-        x_requires_grad = x.requires_grad
-        x = x.detach()
-        # detach() unsets x.requires_grad, so restore it
-        x.requires_grad_(x_requires_grad)
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
-        hidden_size = x.shape[-1]
-        x_shape_orig = x.shape
-        # flatten bs+seqlen to avoid having stride issues when narrowing into seqlen w/ bs>1
-        x = x.view(-1, hidden_size)
-        incoming_grad = grads[0].view(-1, hidden_size)
-        x_grad = torch.zeros_like(x)
-        x_shards = list(torch.chunk(x, chunks=shards, dim=0))
-        for i, x_shard in enumerate(x_shards):
-            x_shard.requires_grad_(x_requires_grad)
-            # if seqlen is not exactly divisible by shards the last step will be shorter than shard_step
-            shard_step = x_shards[i].shape[0]
-            shard_offset = i * x_shards[0].shape[0]
-            x_shard.grad = x_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
-            incoming_grad_shard = incoming_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
-            with torch.enable_grad():
-                output = fn(mlp_module, x_shard)
-            torch.autograd.backward(output, incoming_grad_shard)
-        # unflatten
-        x_grad = x_grad.view(x_shape_orig)
-        return (None, None, x_grad, None, None)
-def apply_tiled_mlp(
-    fn: Callable,
-    mlp_module: torch.nn.Module,
-    x: torch.Tensor,
-    num_shards: Optional[int] = None,
-    compute_params: Optional[List[torch.nn.Parameter]] = None,
-) -> torch.Tensor:
-    """
-    Apply tiled MLP computation for memory efficiency.
-    Args:
-        fn: the function to call on sharded inputs (e.g., lambda module, x: module(x))
-        mlp_module: the MLP nn.Module object
-        x: the input tensor with shape [bs, seqlen, hidden_size] or [seqlen, hidden_size]
-        num_shards: number of shards to use. If None, automatically calculated as ceil(seqlen / hidden_size)
-        compute_params: list of parameters for DeepSpeed ZeRO optimization
-    Returns:
-        output tensor with the same shape as input
-    """
-    if num_shards is None:
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size]
-        hidden_size = x.shape[-1]
-        seqlen = x.shape[-2]
-        num_shards = math.ceil(seqlen / hidden_size)
-    # Ensure num_shards is at least 1
-    num_shards = max(1, num_shards)
-    return LigerTiledMLPFunction.apply(
-        fn,
-        mlp_module,
-        x,
-        num_shards,
-        compute_params,
-    )

build/torch-rocm/tvd.py DELETED Viewed

@@ -1,218 +0,0 @@
-from typing import Literal
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-MAX_FUSED_SIZE = 65536 // 4
-REDUCTION_LITERAL = Literal["none", "sum", "mean", "batchmean"]
-_REDUCTION_MODE_NONE = tl.constexpr(0)
-_REDUCTION_MODE_SUM = tl.constexpr(1)
-_REDUCTION_MODE_MEAN = tl.constexpr(2)
-_REDUCTION_MODE_BATCHMEAN = tl.constexpr(3)
-_str_to_reduction_mode = {
-    "none": _REDUCTION_MODE_NONE.value,
-    "sum": _REDUCTION_MODE_SUM.value,
-    "mean": _REDUCTION_MODE_MEAN.value,
-    "batchmean": _REDUCTION_MODE_BATCHMEAN.value,
-}
-def get_num_warps(BLOCK_SIZE):
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return num_warps
-@triton.jit
-def _tv_distance_kernel(
-    p_ptr,
-    p_stride,
-    q_ptr,
-    q_stride,
-    loss_ptr,
-    loss_stride,
-    grads_ptr,
-    grads_stride,
-    label_ptr,
-    ignore_index: tl.constexpr,
-    n_cols,
-    scale,  # pre-computed reduction scale for gradients (fused into kernel)
-    BLOCK_SIZE: tl.constexpr,
-    HAS_LABEL: tl.constexpr,
-    reduction: tl.constexpr = _REDUCTION_MODE_BATCHMEAN,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    p_ptr += pid * p_stride
-    q_ptr += pid * q_stride
-    loss_ptr += pid * loss_stride
-    grads_ptr += pid * grads_stride
-    label_ptr += pid
-    base_offsets = tl.arange(0, BLOCK_SIZE)
-    if HAS_LABEL:
-        label = tl.load(label_ptr)
-        if label == ignore_index:
-            for i in range(0, n_cols, BLOCK_SIZE):
-                offsets = i + base_offsets
-                mask = offsets < n_cols
-                tl.store(grads_ptr + offsets, 0.0, mask=mask)
-                if reduction == _REDUCTION_MODE_NONE:
-                    tl.store(loss_ptr + offsets, 0.0, mask=mask)
-            return
-    loss_sum = 0.0
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + base_offsets
-        mask = offsets < n_cols
-        p = tl.load(p_ptr + offsets, mask=mask, other=0.0)
-        q = tl.load(q_ptr + offsets, mask=mask, other=0.0)
-        # TVD(P || Q) = 0.5 * |P - Q|
-        tv_loss = 0.5 * tl.abs(p - q)
-        # Fuse reduction scaling into gradient computation (eliminates separate Python division)
-        grad_res = tl.where(p > q, 0.5 * scale, -0.5 * scale)
-        tl.store(grads_ptr + offsets, grad_res, mask=mask)
-        if reduction == _REDUCTION_MODE_NONE:
-            tl.store(loss_ptr + offsets, tv_loss, mask=mask)
-        else:
-            loss_sum += tl.sum(tv_loss, axis=0)
-    if reduction != _REDUCTION_MODE_NONE:
-        # Fuse reduction scaling into loss (same scale as gradients; avoids Python division)
-        tl.store(loss_ptr, loss_sum * scale)
-def tv_distance_forward_triton(p, q, shift_labels, reduction, ignore_index, has_label):
-    BT, V = p.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    reduction = _str_to_reduction_mode[reduction]
-    out_size = (BT, V) if reduction == _REDUCTION_MODE_NONE.value else (BT,)
-    output_tensor = torch.zeros(out_size, device=p.device, dtype=torch.float32)
-    grads = torch.empty_like(p)
-    n_non_ignore = (shift_labels != ignore_index).sum().item() if has_label else BT
-    # Pre-compute gradient scale factor (fused into kernel to avoid separate division)
-    if reduction == _REDUCTION_MODE_BATCHMEAN.value:
-        scale = 1.0 / n_non_ignore
-    elif reduction == _REDUCTION_MODE_MEAN.value:
-        scale = 1.0 / (n_non_ignore * V)
-    else:
-        scale = 1.0
-    _tv_distance_kernel[grid](
-        p,
-        p.stride(0),
-        q,
-        q.stride(0),
-        output_tensor,
-        output_tensor.stride(0),
-        grads,
-        grads.stride(0),
-        shift_labels if has_label else torch.empty(1, device=p.device),
-        ignore_index,
-        V,
-        scale,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_LABEL=has_label,
-        num_warps=num_warps,
-        reduction=reduction,
-    )
-    # Loss and gradients are already scaled inside the kernel — no separate division needed
-    if reduction in (_REDUCTION_MODE_BATCHMEAN.value, _REDUCTION_MODE_MEAN.value):
-        return output_tensor.sum(), grads
-    elif reduction == _REDUCTION_MODE_SUM.value:
-        return output_tensor.sum(dim=0), grads
-    else:
-        return output_tensor, grads
-def tvd_backward_triton(grad_output, grads):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul then.
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return grads
-    return grads * grad_output
-class LigerTVDLossFunction(torch.autograd.Function):
-    """
-    Class implementing the forward and backward pass for the Total Variation Distance Loss using Triton.
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        p: torch.Tensor,
-        q: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-        reduction: REDUCTION_LITERAL = "batchmean",
-        ignore_index: int = -100,
-    ) -> torch.Tensor:
-        """A forward pass for the Total Variation Distance Loss.
-        Args:
-            ctx: Torch autograd context
-            p (torch.Tensor): A tensor of shape (BT, V) containing the first distribution.
-            q (torch.Tensor): A tensor of shape (BT, V) containing the second distribution.
-            shift_labels (Optional[torch.Tensor]): A tensor of shape (BT,) containing the labels.
-            reduction (REDUCTION_LITERAL, optional): The reduction method to be applied. Defaults to "batchmean".
-            ignore_index (int, optional): The index to ignore during loss calculation. Defaults to -100.
-        Returns:
-            torch.Tensor: The computed Total Variation Distance Loss.
-        """
-        has_label = False
-        if shift_labels is not None:
-            assert shift_labels.shape == (p.shape[0],), (
-                f"the shape of shift_labels must be (BT,). Got: {shift_labels.shape}"
-            )
-            shift_labels = shift_labels.contiguous()
-            has_label = True
-        loss, grads = tv_distance_forward_triton(p, q, shift_labels, reduction, ignore_index, has_label)
-        ctx.save_for_backward(grads)
-        return loss
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        """A backward pass for the Total Variation Distance Loss.
-        Args:
-            ctx: Torch autograd context
-            grad_output (torch.Tensor): The gradient of the loss with respect to the output.
-        Returns:
-            tuple[torch.Tensor, None, None, None, None]: The gradient of the loss with respect to the inputs.
-        """
-        (grads,) = ctx.saved_tensors
-        grads = tvd_backward_triton(grad_output, grads)
-        return grads, None, None, None, None

build/torch-rocm/utils.py DELETED Viewed

@@ -1,176 +0,0 @@
-"""
-This file incorporates code from Unsloth licensed under the Apache License, Version 2.0.
-See the original Unsloth repository at https://github.com/unslothai/unsloth.
-The following line
-https://github.com/linkedin/Liger-Kernel/blob/7382a8761f9af679482b968f9348013d933947c7/src/liger_kernel/ops/utils.py#L23
-is based on code from Unsloth, located at:
-https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/utils.py#L43
-Modifications made by Yanning Chen, 2024.
-"""
-import functools
-import importlib
-import operator
-from typing import Callable
-import torch
-import triton
-import triton.language as tl
-from packaging.version import Version
-def is_npu_available() -> bool:
-    """Detect Ascend NPU availability."""
-    try:
-        from transformers.utils import is_torch_npu_available
-        return is_torch_npu_available()
-    except Exception:
-        return False
-def infer_device():
-    """
-    Get current device name based on available devices
-    """
-    if torch.cuda.is_available():  # Works for both Nvidia and AMD
-        return "cuda"
-    # Use Ascend NPU if available (torch.npu)
-    elif is_npu_available():
-        return "npu"
-    # XPU (Intel) if available
-    elif torch.xpu.is_available():
-        return "xpu"
-    else:
-        return "cpu"
-def is_hip() -> bool:
-    return torch.version.hip is not None
-def ensure_contiguous(fn):
-    @functools.wraps(fn)
-    def wrapper(ctx, *args, **kwargs):
-        def maybe_to_contiguous(x):
-            return x.contiguous() if isinstance(x, torch.Tensor) else x
-        args = [maybe_to_contiguous(arg) for arg in args]
-        kwargs = {k: maybe_to_contiguous(v) for k, v in kwargs.items()}
-        return fn(ctx, *args, **kwargs)
-    return wrapper
-def calculate_settings(n):
-    # reference: https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/utils.py#L43
-    MAX_FUSED_SIZE = 65536
-    BLOCK_SIZE = triton.next_power_of_2(n)
-    if BLOCK_SIZE > MAX_FUSED_SIZE:
-        raise RuntimeError(
-            f"Cannot launch Triton kernel since n = {n} exceeds the recommended Triton blocksize = {MAX_FUSED_SIZE}."
-        )
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32 if not is_hip() else 16
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return BLOCK_SIZE, num_warps
-def compare_version(package: str, operator: Callable, target: str):
-    try:
-        pkg = importlib.import_module(package)
-    except ImportError:
-        return False
-    pkg_version = Version(pkg.__version__)
-    return operator(pkg_version, Version(target))
-def get_amp_custom_fwd_bwd() -> Callable:
-    device = infer_device()
-    if compare_version("torch", operator.ge, "2.4.0"):
-        return (
-            functools.partial(torch.amp.custom_fwd, device_type=device),
-            functools.partial(torch.amp.custom_bwd, device_type=device),
-        )
-    if hasattr(torch, "npu") and getattr(torch.npu, "amp", None) is not None:
-        return torch.npu.amp.custom_fwd, torch.npu.amp.custom_bwd
-    return torch.cuda.amp.custom_fwd, torch.cuda.amp.custom_bwd
-amp_custom_fwd, amp_custom_bwd = get_amp_custom_fwd_bwd()
-torch_to_triton_dtype = {
-    torch.float32: tl.float32,
-    torch.float16: tl.float16,
-    torch.bfloat16: tl.bfloat16,
-}
-@triton.jit
-def element_mul_kernel(
-    X_ptr,
-    X_stride,
-    grad_output_ptr,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    This function multiplies each element of the tensor pointed by X_ptr with the value pointed by grad_output_ptr.
-    The multiplication is performed in-place on the tensor pointed by X_ptr.
-    Parameters:
-    X_ptr: Pointer to the input tensor.
-    X_stride (int): The stride of the input tensor.
-    grad_output_ptr: Pointer to the gradient output value.
-    n_cols (int): The number of columns in the input tensor.
-    BLOCK_SIZE (int): The block size for Triton operations.
-    """
-    # Get the program ID and convert it to int64 to avoid overflow
-    program_id = tl.program_id(0).to(tl.int64)
-    # Locate the start index
-    X_ptr += program_id * X_stride
-    # Load the gradient output value
-    grad_output = tl.load(grad_output_ptr)
-    # Perform the element-wise multiplication
-    for i in range(0, n_cols, BLOCK_SIZE):
-        X_offsets = i + tl.arange(0, BLOCK_SIZE)
-        X_block = tl.load(X_ptr + X_offsets, mask=X_offsets < n_cols)
-        tl.store(X_ptr + X_offsets, X_block * grad_output, mask=X_offsets < n_cols)
-def get_npu_core_count(default: int = 20) -> int:
-    """Return NPU vector core count.
-    Fallback to `default` if Triton runtime or NPU device is unavailable.
-    """
-    try:
-        utils = triton.runtime.driver.active.utils
-        props = utils.get_device_properties(0)
-        return int(props.get("num_vectorcore", default))
-    except Exception:
-        return default
-def set_large_grf_mode(kernel_args: dict):
-    """Set large GRF mode for XPU devices."""
-    # On XPU triton installed along with pytorch-xpu will be called `pytorch-triton-xpu`,
-    # triton XPU installed from source will be called `triton`.
-    if compare_version("pytorch-triton-xpu", operator.ge, "3.6.0") or compare_version("triton", operator.ge, "3.6.0"):
-        kernel_args["grf_mode"] = "256"
-    else:
-        # API was changed in https://github.com/intel/intel-xpu-backend-for-triton/pull/5430
-        kernel_args["grf_mode"] = "large"

build/torch-xpu/__init__.py DELETED Viewed

@@ -1,4 +0,0 @@
-from . import layers
-from .layers import CrossEntropyOutput, LigerForCausalLMLoss
-__all__ = ["layers", "LigerForCausalLMLoss", "CrossEntropyOutput"]

build/torch-xpu/_ops.py DELETED Viewed

@@ -1,38 +0,0 @@
-import torch
-def get_backend() -> str:
-    """Detect the backend by inspecting torch."""
-    import torch
-    if hasattr(torch, "neuron"):
-        # Needs to be sorted before specific Torch builds, since Neuron
-        # extension can be loaded into e.g. CUDA Torch builds.
-        return "neuron"
-    elif torch.version.cuda is not None:
-        return "cuda"
-    elif torch.version.hip is not None:
-        return "rocm"
-    elif torch.backends.mps.is_available():
-        return "metal"
-    elif hasattr(torch.version, "xpu") and torch.version.xpu is not None:
-        return "xpu"
-    else:
-        return "cpu"
-def _find_ops_name() -> str:
-    kernel_name = "liger_kernels"
-    unique_id = "ab435e2"
-    backend = get_backend()
-    return f"_{kernel_name}_{backend}_{unique_id}"
-_OPS_NAME = _find_ops_name()
-ops = getattr(torch.ops, _OPS_NAME)
-def add_op_namespace_prefix(op_name: str) -> str:
-    """
-    Prefix op by namespace.
-    """
-    return f"{_OPS_NAME}::{op_name}"

build/torch-xpu/cross_entropy.py DELETED Viewed

@@ -1,558 +0,0 @@
-import operator
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import element_mul_kernel
-from .utils import is_hip
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.jit
-def liger_cross_entropy_kernel(
-    X_ptr,
-    X_stride,
-    Y_ptr,
-    Y_stride,
-    weight_ptr,
-    loss_ptr,
-    z_loss_ptr,
-    loss_stride,
-    token_accuracy_ptr,
-    token_accuracy_stride,
-    predicted_tokens_ptr,
-    predicted_tokens_stride,
-    n_cols,
-    n_non_ignore,
-    sum_non_ignore_weight,
-    weight_sum,
-    ignore_index,
-    lse_square_scale: tl.constexpr,
-    label_smoothing: tl.constexpr,
-    reduction: tl.constexpr,  # set it as constexpr since reduction is always known at compile time
-    softcap,
-    RETURN_Z_LOSS: tl.constexpr,
-    RETURN_TOKEN_ACCURACY: tl.constexpr,
-    RETURN_PREDICTED_TOKENS: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    HAS_WEIGHT: tl.constexpr,
-    HAS_SOFTCAPPING: tl.constexpr,
-    HAS_GRADIENTS: tl.constexpr,
-):
-    """
-    This kernel computes both cross entropy loss and the gradient of the input.
-    We only consider hard label + mean reduction for now. Please refer to https://pytorch.org/docs/stable/generated/torch.nn.CrossEntropyLoss.html for the math.
-    Parameters:
-    X_ptr: Pointer to input tensor.
-    X_stride (int): The stride of the input tensor.
-    Y_ptr: Pointer to target tensor.
-    Y_stride (int): The stride of the target tensor.
-    weight_ptr: Pointer to weight tensor.
-    loss_ptr: Pointer to tensor to store the loss.
-    z_loss_ptr: Pointer to tensor to store the z loss. No operation if RETURN_Z_LOSS is 0.
-    loss_stride (int): The stride of the loss tensor.
-    token_accuracy_ptr: Pointer to tensor to store the per-token accuracy. No operation if RETURN_TOKEN_ACCURACY is 0.
-    token_accuracy_stride (int): The stride of the token accuracy tensor.
-    n_cols (int): The number of columns in the input tensor.
-    n_non_ignore (float): The number of non-ignored elements in the batch.
-    sum_non_ignore_weight (float): The sum of non-ignored target's weights in the batch.
-    weight_sum (float): The sum of weight tensor.
-    ignore_index (int): The index to ignore in the target.
-    label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-    lse_square_scale (float): The scaler of (logsumexp(_input)) ^ 2 adding to the loss for the stability of training.
-    reduction (str): The string for the reduction to apply
-    softcap (float): The upper threshold for scaling logits to the range (-softcap, +softcap).
-    RETURN_Z_LOSS (int): The boolean value to decide whether to store z loss to z_loss_ptr or not. It must be 0 or 1.
-    RETURN_TOKEN_ACCURACY (int): The boolean value to decide whether to store per-token accuracy to token_accuracy_ptr or not. It must be 0 or 1.
-    BLOCK_SIZE (int): The block size for Triton operations.
-    HAS_WEIGHT (bool): The boolean value to determine whether assigning weight to each of the classes.
-    HAS_SOFTCAPPING (bool): The boolean value to determine whether applying soft-capping or not.
-    HAS_GRADIENTS (bool): The boolean value to determine whether calculating gradients in forward pass.
-    """
-    # https://github.com/triton-lang/triton/issues/1058
-    # If B*T*V is too large, program_id * stride will overflow out of int32, so we convert to int64
-    program_id = tl.program_id(0).to(tl.int64)
-    # 1. Load Y_ptr first because if the target is ignore_index, we can return right away
-    Y_ptr += program_id * Y_stride
-    y = tl.load(Y_ptr)
-    # 2. locate the start index
-    X_ptr += program_id * X_stride
-    if y == ignore_index:
-        # set all X_ptr as 0
-        for i in range(0, n_cols, BLOCK_SIZE):
-            X_offsets = i + tl.arange(0, BLOCK_SIZE)
-            tl.store(X_ptr + X_offsets, 0.0, mask=X_offsets < n_cols)
-        # For ignored tokens, set token accuracy to 0
-        if RETURN_TOKEN_ACCURACY:
-            token_accuracy_ptr += program_id * token_accuracy_stride
-            tl.store(token_accuracy_ptr, 0.0)
-        if RETURN_PREDICTED_TOKENS:
-            predicted_tokens_ptr += program_id * predicted_tokens_stride
-            tl.store(predicted_tokens_ptr, -1)
-        return
-    loss_ptr += program_id * loss_stride
-    if RETURN_Z_LOSS:
-        z_loss_ptr += program_id * loss_stride
-    if RETURN_TOKEN_ACCURACY:
-        token_accuracy_ptr += program_id * token_accuracy_stride
-    if RETURN_PREDICTED_TOKENS:
-        predicted_tokens_ptr += program_id * predicted_tokens_stride
-    if HAS_WEIGHT:
-        weight_y = tl.load(weight_ptr + y).cast(tl.float32)
-    # Online softmax: 2 loads + 1 store (compared with 3 loads + 1 store for the safe softmax)
-    # Refer to Algorithm 3 in the paper: https://arxiv.org/pdf/1805.02867
-    # 3. [Online softmax] first pass: find max + sum
-    m = float("-inf")  # m is the max value. use the notation from the paper
-    d = 0.0  # d is the sum. use the notation from the paper
-    argmax_idx = 0  # Track the index of the maximum value for token accuracy / predicted tokens computation
-    ori_X_y = tl.load(X_ptr + y).cast(tl.float32)  # we need to store the original value of X_y for the loss calculation
-    if HAS_SOFTCAPPING:
-        ori_X_y = softcap * tanh(ori_X_y / softcap)
-    # Label smoothing is a general case of normal cross entropy
-    # See the full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issue-2503665310
-    scaled_x_sum = 0.0
-    eps = label_smoothing / n_cols
-    for i in range(0, n_cols, BLOCK_SIZE):
-        X_offsets = i + tl.arange(0, BLOCK_SIZE)
-        X_block = tl.load(
-            X_ptr + X_offsets,
-            mask=X_offsets < n_cols,
-            other=float("-inf"),
-            # Ensure float32 precision for softmax calculation
-        ).cast(tl.float32)
-        if HAS_SOFTCAPPING:
-            X_block = softcap * tanh(X_block / softcap)
-        block_max = tl.max(X_block)
-        # Track argmax for accuracy / predicted tokens computation
-        if RETURN_TOKEN_ACCURACY or RETURN_PREDICTED_TOKENS:
-            # Find the index of the maximum value in this block
-            is_max_mask = X_block == block_max
-            # Mask out invalid indices with a value larger than n_cols
-            masked_offsets = tl.where(is_max_mask, X_offsets, n_cols)
-            # Get the first (smallest) index where max occurs
-            current_block_argmax_idx = tl.min(masked_offsets)
-            is_new_max = block_max > m
-            argmax_idx = tl.where(is_new_max, current_block_argmax_idx, argmax_idx)
-        if label_smoothing > 0:
-            # scale X beforehand to avoid overflow
-            if HAS_WEIGHT:
-                weight_block = tl.load(weight_ptr + X_offsets, mask=X_offsets < n_cols)
-                scaled_x_sum += tl.sum(tl.where(X_offsets < n_cols, -eps * X_block * weight_block, 0.0))
-            else:
-                scaled_x_sum += tl.sum(tl.where(X_offsets < n_cols, -eps * X_block, 0.0))
-        m_new = tl.maximum(m, block_max)
-        d = d * tl.exp(m - m_new) + tl.sum(tl.exp(X_block - m_new))
-        m = m_new
-    # log (sum(e^(X_i))) = log (sum(e ^ (max(X) * e ^ (X_i - max(X)))))
-    #                    = log (e^(max(X)) * sum(e ^ (X_i - max(X))))
-    #                    = max(X) + log (sum(e ^ (X_i - max(X)))) = m + log d
-    lse = m + tl.log(d)
-    # 4. [Online Softmax] Second pass: compute gradients
-    # For 'mean' reduction, gradients are normalized by number of non-ignored elements (N)
-    # dx_y = (softmax(x_y) - 1) / N
-    # dx_i = softmax(x_i) / N, i != y
-    # For label smoothing:
-    # dx_i = (softmax(x_i) - label_smoothing / V) / N, V = n_cols, i != y
-    # dx_y = (softmax(x_y) - label_smoothing / V - (1 - label_smoothing)) / N
-    #      = dx_i - (1 - label_smoothing) / N
-    # With Z loss:
-    # dx_i = ((1 + 2 * lse_square_scale * lse) * softmax(x_i) - label_smoothing / V) / N, i != y
-    # dx_y = dx_i - (1 - label_smoothing) / N
-    # For 'sum' reduction, no normalization is applied:
-    # dx_y = softmax(x_y) - 1
-    # dx_i = softmax(x_i), for i ≠ y
-    if HAS_GRADIENTS:
-        for i in range(0, n_cols, BLOCK_SIZE):
-            X_offsets = i + tl.arange(0, BLOCK_SIZE)
-            X_block = tl.load(
-                X_ptr + X_offsets,
-                mask=X_offsets < n_cols,
-                other=float("-inf"),
-                # Ensure float32 precision for softmax calculation
-            ).cast(tl.float32)
-            if HAS_SOFTCAPPING:
-                intermediate = tanh(X_block / softcap)
-                X_block = softcap * intermediate
-            if not HAS_WEIGHT:
-                # softmax(x_i)
-                X_block = tl.exp(X_block - m) / d
-                # derivative of z-loss: 2 * lse_square_scale * lse * softmax(x_i)
-                X_block += 2 * lse_square_scale * lse * X_block
-                # smoothing term
-                X_block += -eps
-                # special handle dx_y
-                X_block = tl.where(X_offsets != y, X_block, X_block - (1 - label_smoothing))
-                # reduction scale
-                if reduction == "mean":
-                    X_block = X_block / n_non_ignore
-            else:
-                weight_block = tl.load(weight_ptr + X_offsets, mask=X_offsets < n_cols)
-                softmax_X = tl.exp(X_block - m) / d
-                # derivative of original_loss
-                dloss_ori = (1 - label_smoothing) * softmax_X
-                # specially handle dx_y
-                dloss_ori = tl.where(X_offsets != y, dloss_ori, dloss_ori - (1 - label_smoothing))
-                dloss_ori = dloss_ori * weight_y
-                # derivative of smooth_loss
-                dloss_smooth = eps * (-weight_block + softmax_X * weight_sum)
-                # derivative of z-loss
-                dz_loss = 2 * lse_square_scale * lse * softmax_X
-                # reduction scale
-                if reduction == "mean":
-                    dloss_ori = dloss_ori / sum_non_ignore_weight
-                    dloss_smooth = dloss_smooth / sum_non_ignore_weight
-                    # TODO: Implement weighted z_loss. Currently, z_loss is not scaled by weight.
-                    dz_loss = dz_loss / n_non_ignore
-                # derivative of total_loss
-                X_block = dloss_ori + dloss_smooth + dz_loss
-            # chain rule softcapping
-            # d(softcap * tanh(x / softcap)) = (1 - tanh^2(x / softcap))
-            if HAS_SOFTCAPPING:
-                X_block = X_block * (1 - intermediate * intermediate)
-            tl.store(X_ptr + X_offsets, X_block, mask=X_offsets < n_cols)
-    # We need tl.debug_barrier() to ensure the new result of X_ptr is written as mentioned in
-    # https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/ops/cross_entropy.py#L34
-    tl.debug_barrier()
-    # 5. Calculate the loss
-    # loss = log (softmax(X_y)) = log ((e ^ (X_y - max(X)) / sum(e ^ (X - max(X))))
-    #      = (X_y - max(X)) - log(sum(e ^ (X - max(X))))
-    #      = X_y - m - log d = X_y - lse
-    # sum(e ^ (X - max(X))) must >= 1 because the max term is e ^ 0 = 1
-    # So we can safely calculate log (softmax(X_y)) without overflow
-    loss = lse - ori_X_y
-    if HAS_WEIGHT:
-        loss = weight_y * loss
-    # Original loss = H(q, p),  with label smoothing regularization = H(q', p) and (label_smoothing / V) = eps
-    # H(q', p) = (1 - label_smoothing) * H(q, p) + label_smoothing * H(u, p)
-    #          = (1 - label_smoothing) * H(q, p) + eps * sum(logsoftmax(x_i))
-    # By using m (global max of xi) and d (sum of e^(xi-m)), we can simplify as:
-    #          = (1 - label_smoothing) * H(q, p) + (sum(-eps * x_i) + label_smoothing * (m + logd))
-    # Refer to H(q', p) in section 7 of the paper: https://arxiv.org/pdf/1512.00567
-    # pytorch: https://github.com/pytorch/pytorch/blob/2981534f54d49fa3a9755c9b0855e7929c2527f0/aten/src/ATen/native/LossNLL.cpp#L516
-    # See full derivation at https://github.com/linkedin/Liger-Kernel/pull/198#issuecomment-2333753087
-    if label_smoothing > 0:
-        if HAS_WEIGHT:
-            smooth_loss = scaled_x_sum + eps * lse * weight_sum
-        else:
-            smooth_loss = scaled_x_sum + label_smoothing * lse
-        loss = loss * (1 - label_smoothing) + smooth_loss
-    # An auxiliary loss, z_loss
-    # Refer to Page14 Loss function section in the paper PaLM: https://www.jmlr.org/papers/v24/22-1144.html
-    z_loss = lse_square_scale * lse * lse
-    # Normalize the loss by the number of non-ignored elements if reduction is "mean"
-    if reduction == "mean":
-        if HAS_WEIGHT:
-            loss = loss / sum_non_ignore_weight
-        else:
-            loss = loss / n_non_ignore
-        # TODO: Implement weighted z_loss. Currently, z_loss is not scaled by weight.
-        z_loss = z_loss / n_non_ignore
-    loss += z_loss
-    tl.store(loss_ptr, loss)
-    if RETURN_Z_LOSS:
-        tl.store(z_loss_ptr, z_loss)
-    if RETURN_TOKEN_ACCURACY:
-        # Store 1.0 if prediction is correct, 0.0 otherwise
-        is_correct = 1.0 if argmax_idx == y else 0.0
-        tl.store(token_accuracy_ptr, is_correct)
-    if RETURN_PREDICTED_TOKENS:
-        tl.store(predicted_tokens_ptr, argmax_idx)
-# The hard limit of TRITON_MAX_TENSOR_NUMEL is 1048576 https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/language/core.py#L19
-# However, setting limit as 65536 as in LayerNorm tutorial is faster because of less register spilling
-# The optimal maximum block size depends on your hardware, your kernel, and your dtype
-# the best size we found by manually tuning on xpu and npu.
-if infer_device() == "xpu":
-    MAX_FUSED_SIZE = 4096
-elif infer_device() == "npu":
-    MAX_FUSED_SIZE = 2048
-else:
-    MAX_FUSED_SIZE = 65536 // 2
-def cross_entropy_forward(
-    _input,
-    target,
-    weight,
-    ignore_index,
-    lse_square_scale,
-    label_smoothing,
-    reduction,
-    softcap,
-    return_z_loss,
-    return_token_accuracy=False,
-    return_predicted_tokens=False,
-):
-    assert isinstance(return_z_loss, bool), f"return_z_loss must be True or False. Got: {return_z_loss}"
-    assert isinstance(return_token_accuracy, bool), (
-        f"return_token_accuracy must be True or False. Got: {return_token_accuracy}"
-    )
-    assert isinstance(return_predicted_tokens, bool), (
-        f"return_predicted_tokens must be True or False. Got: {return_predicted_tokens}"
-    )
-    BT, V = _input.shape
-    n_rows = BT
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    # unreduced loss
-    loss_1d = torch.zeros(n_rows, dtype=_input.dtype, device=_input.device)
-    z_loss_1d = torch.zeros(n_rows, dtype=_input.dtype, device=_input.device) if return_z_loss else None
-    token_accuracy_1d = (
-        torch.zeros(n_rows, dtype=torch.float32, device=_input.device) if return_token_accuracy else None
-    )
-    predicted_tokens_1d = (
-        torch.full((n_rows,), -1, dtype=torch.int64, device=_input.device) if return_predicted_tokens else None
-    )
-    target_mask = target != ignore_index
-    n_non_ignore = target_mask.sum().item()
-    assert (target * target_mask).max() < _input.shape[-1], (
-        f"Target {target.max()} is out of bounds. Expected < {_input.shape[-1]}"
-    )
-    assert (target * target_mask).min() >= 0, f"Target {target.min()} is out of bounds. Expected >= 0"
-    sum_non_ignore_weight = n_non_ignore
-    weight_sum = 0.0
-    if weight is not None:
-        assert weight.shape[0] == V, f"If given, weight has to be a Tensor of size V. Got: {weight.shape}"
-        assert torch.is_floating_point(weight), (
-            f"If given, weight has to be a Tensor of floating point dtype. Got: {weight.dtype}"
-        )
-        sum_non_ignore_weight = torch.gather(weight, dim=0, index=target.masked_select(target_mask)).sum().item()
-        weight_sum = weight.sum().item()
-        # ensure weight is contiguous
-        if weight.stride(-1) != 1:
-            weight = weight.contiguous()
-    # ensure _input and target are contiguous in the last dimension
-    if _input.stride(-1) != 1:
-        _input = _input.contiguous()
-    if target.stride(-1) != 1:
-        target = target.contiguous()
-    # Here we use a trick to store X_ptr gradient in X_ptr so we can save memory
-    liger_cross_entropy_kernel[(n_rows,)](
-        X_ptr=_input,
-        X_stride=_input.stride(-2),
-        Y_ptr=target,
-        Y_stride=target.stride(-1),  # always 1
-        weight_ptr=weight,  # dummy if None
-        loss_ptr=loss_1d,
-        z_loss_ptr=z_loss_1d,
-        loss_stride=loss_1d.stride(-1),  # always 1
-        token_accuracy_ptr=token_accuracy_1d,
-        token_accuracy_stride=token_accuracy_1d.stride(-1)
-        if return_token_accuracy
-        else 0,  # always 1 if accuracy is enabled
-        predicted_tokens_ptr=predicted_tokens_1d,
-        predicted_tokens_stride=predicted_tokens_1d.stride(-1)
-        if return_predicted_tokens
-        else 0,  # always 1 if predicted tokens is enabled
-        n_cols=V,
-        n_non_ignore=n_non_ignore,
-        sum_non_ignore_weight=sum_non_ignore_weight,
-        ignore_index=ignore_index,
-        weight_sum=weight_sum,
-        lse_square_scale=lse_square_scale,
-        label_smoothing=label_smoothing,
-        reduction=reduction,
-        softcap=softcap,
-        RETURN_Z_LOSS=return_z_loss,
-        RETURN_TOKEN_ACCURACY=return_token_accuracy,
-        RETURN_PREDICTED_TOKENS=return_predicted_tokens,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_WEIGHT=True if weight is not None else False,
-        HAS_SOFTCAPPING=True if softcap is not None else False,
-        HAS_GRADIENTS=_input.requires_grad,
-        # TODO: 32 seems to give the best performance
-        # Performance is quite sensitive to num_warps
-        num_warps=32 if not is_hip() else 16,
-    )
-    if reduction == "none":
-        loss = loss_1d
-        z_loss = z_loss_1d if return_z_loss else None
-        token_accuracy = token_accuracy_1d if return_token_accuracy else None
-    else:
-        loss = torch.sum(loss_1d)
-        z_loss = torch.sum(z_loss_1d) if return_z_loss else None
-        # For accuracy, we compute the mean across all non-ignored tokens
-        token_accuracy = torch.sum(token_accuracy_1d) / n_non_ignore if return_token_accuracy else None
-    predicted_tokens = predicted_tokens_1d if return_predicted_tokens else None
-    return loss, z_loss, token_accuracy, predicted_tokens, _input
-def cross_entropy_backward(_input, grad_output):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul to save time
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        pass
-    # If reduction is 'none'
-    elif grad_output.ndim > 0:
-        _input = _input * grad_output.unsqueeze(dim=1)
-    # If reduction is ['mean', 'sum'], grad_output is just a scalar
-    # We use a Triton kernel instead of a PyTorch operation because modifying inputs in-place
-    # for gradient storage and backward multiple times causes anomalies with PyTorch but not with Triton.
-    else:
-        BT, V = _input.shape
-        n_rows = BT
-        BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-        element_mul_kernel[(n_rows,)](
-            _input,
-            _input.stride(-2),
-            grad_output,
-            V,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-    return _input
-class LigerCrossEntropyFunction(torch.autograd.Function):
-    """
-    This class implements a custom autograd function for the Liger Cross Entropy loss.
-    It overrides the forward and backward methods of the torch.autograd.Function class.
-    """
-    @staticmethod
-    def forward(
-        ctx,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        weight: Optional[torch.FloatTensor],
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-        return_z_loss: bool = False,
-        return_token_accuracy: bool = False,
-        return_predicted_tokens: bool = False,
-    ):
-        """
-        The forward pass of the Liger Cross Entropy loss.
-        Parameters:
-        ctx : The context object.
-        _input (tensor): The input tensor of shape (BT, V) where B is batch size, T is sequence length, V is vocab size.
-        target (tensor): The target tensor of shape (BT) where each value is in [0, V-1].
-        weight(Tensor, optional): a manual rescaling weight given to each class. If given, has to be a Tensor of size V and floating point dtype
-        ignore_index (int): The index to ignore in the target.
-        lse_square_scale (float): The scaler of (logsumexp(_input)) ^ 2 adding to the loss for the stability of training.
-        label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-        reduction (str): The reduction to apply to the output: "none" | "mean | "sum".
-        softcap (Optional[float]): The upper threshold for scaling logits to the range (-softcap, +softcap).
-        return_z_loss (bool): When `return_z_loss` is `True`, returns (loss, z_loss, token_accuracy, predicted_tokens) instead of (loss, None, None, None). Default: `False`
-        return_token_accuracy (bool): When `return_token_accuracy` is `True`, computes and returns per-token accuracy without materializing logits. Default: `False`
-        return_predicted_tokens (bool): When `return_predicted_tokens` is `True`, returns per-token predicted class indices (argmax) without materializing logits. Default: `False`
-        Returns:
-        tuple: A tuple with the computed losses, accuracy, and predicted tokens: (loss, z_loss, token_accuracy, predicted_tokens). z_loss, token_accuracy, and predicted_tokens are None if not requested.
-        """
-        input_requires_grad = _input.requires_grad
-        loss, z_loss, token_accuracy, predicted_tokens, _input = cross_entropy_forward(
-            _input,
-            target,
-            weight,
-            ignore_index,
-            lse_square_scale,
-            label_smoothing,
-            reduction,
-            softcap,
-            return_z_loss,
-            return_token_accuracy,
-            return_predicted_tokens,
-        )
-        # TODO: investigation
-        # If we don't detach the _input tensor, the memory will double
-        # Not sure why but seems that there will be a time both grad and value exist but in different location
-        if input_requires_grad:
-            ctx.save_for_backward(_input.detach())
-        ctx.return_z_loss = return_z_loss
-        ctx.return_token_accuracy = return_token_accuracy
-        ctx.return_predicted_tokens = return_predicted_tokens
-        return loss, z_loss, token_accuracy, predicted_tokens
-    @staticmethod
-    def backward(ctx, grad_output, grad_output2, grad_output3, grad_output4):
-        """
-        The backward pass of the Liger Cross Entropy loss.
-        Parameters:
-        ctx : The context object with saved tensors.
-        grad_output (tensor): The tensor containing the gradient of the loss with respect to the output.
-        grad_output2 (tensor): No use. Gradient for z_loss (not used as z_loss is only for logging).
-        grad_output3 (tensor): No use. Gradient for token_accuracy (not used as token_accuracy is only for metrics).
-        grad_output4 (tensor): No use. Gradient for predicted_tokens (not used as predicted_tokens is only for metrics).
-        Returns:
-        tuple: A tuple with the gradients with respect to the inputs. The elements are tensors or None.
-        """
-        if ctx.return_z_loss:
-            del grad_output2  # z_loss is only for logging
-        if ctx.return_token_accuracy:
-            del grad_output3  # token_accuracy is only for metrics
-        if ctx.return_predicted_tokens:
-            del grad_output4  # predicted_tokens is only for metrics
-        (_input,) = ctx.saved_tensors
-        _input = cross_entropy_backward(_input, grad_output)
-        return (
-            _input,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-xpu/dyt.py DELETED Viewed

@@ -1,164 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.autotune(
-    configs=[
-        triton.Config({"BLOCK_N": bn}, num_stages=ns, num_warps=nw)
-        for bn in [1024, 2048, 4096]
-        for ns in [1, 2]
-        for nw in [4, 8, 16]
-    ],
-    key=["N"],
-)
-@triton.jit
-def _dyt_fwd_kernel(X, Y, Alpha, Gamma, Beta, HAVE_BETA: tl.constexpr, N: tl.constexpr, BLOCK_N: tl.constexpr):
-    col = tl.cast(tl.program_id(0), tl.int64) * BLOCK_N + tl.arange(0, BLOCK_N)
-    mask = col < N
-    row_id = tl.cast(tl.program_id(1), tl.int64)
-    X += row_id * N
-    Y += row_id * N
-    alpha = tl.load(Alpha).to(tl.float32)
-    gamma = tl.load(Gamma + col, mask=mask, other=0.0).to(tl.float32)
-    x = tl.load(X + col, mask=mask, other=0.0).to(tl.float32)
-    tanh_x = tanh(alpha * x)
-    y = tanh_x * gamma
-    if HAVE_BETA:
-        beta = tl.load(Beta + col, mask=mask, other=0.0).to(tl.float32)
-        y += beta
-    tl.store(Y + col, y, mask=mask)
-@triton.autotune(
-    configs=[
-        triton.Config({"BLOCK_N": bn}, num_stages=ns, num_warps=nw)
-        for bn in [1024, 2048, 4096]
-        for ns in [1, 2]
-        for nw in [4, 8, 16]
-    ],
-    key=["N"],
-    # DA is indexed by program_id(0), so different BLOCK_N configs write to
-    # different slot counts per SM. Autotune trials don't zero outputs between
-    # runs, so stale slots from a prior trial would leak into da.sum(). Reset
-    # DA between trials to isolate each config's writes.
-    reset_to_zero=["DA"],
-)
-@triton.jit
-def _dyt_bwd_kernel(
-    DY, DX, DA, DG, DB, X, Alpha, Gamma, HAVE_BETA: tl.constexpr, M, N: tl.constexpr, BLOCK_N: tl.constexpr
-):
-    col = tl.cast(tl.program_id(0), tl.int64) * BLOCK_N + tl.arange(0, BLOCK_N)
-    mask = col < N
-    start_row_id = tl.cast(tl.program_id(1), tl.int64)
-    alpha = tl.load(Alpha).to(tl.float32)
-    da = 0.0
-    gamma = tl.load(Gamma + col, mask=mask, other=0.0).to(tl.float32)
-    dg = tl.zeros((BLOCK_N,), dtype=tl.float32)
-    if HAVE_BETA:
-        db = tl.zeros((BLOCK_N,), dtype=tl.float32)
-    for row_id in range(start_row_id, M, tl.num_programs(1)):
-        x = tl.load(X + row_id * N + col, mask=mask, other=0.0).to(tl.float32)
-        dy = tl.load(DY + row_id * N + col, mask=mask, other=0.0).to(tl.float32)
-        tanh_x = tanh(alpha * x)
-        if HAVE_BETA:
-            db += dy
-        dg += dy * tanh_x
-        tmp = (1 - tanh_x * tanh_x) * dy * gamma
-        da += tl.sum(x * tmp, 0)
-        dx = alpha * tmp
-        tl.store(DX + row_id * N + col, dx, mask=mask)
-    tl.store(DG + start_row_id * N + col, dg, mask=mask)
-    if HAVE_BETA:
-        tl.store(DB + start_row_id * N + col, db, mask=mask)
-    tl.store(DA + start_row_id * tl.cdiv(N, 512) + tl.program_id(0), da)
-def liger_dyt_fwd(x, alpha, gamma, beta):
-    assert x.is_contiguous()
-    HAVE_BETA = True if beta is not None else False
-    input_shape = x.shape
-    x = x.view(-1, input_shape[-1])
-    M, N = x.shape
-    y = torch.empty_like(x)
-    grid = lambda meta: (triton.cdiv(N, meta["BLOCK_N"]), M)
-    _dyt_fwd_kernel[grid](
-        x,
-        y,
-        alpha,
-        gamma,
-        beta,
-        HAVE_BETA,
-        N,
-    )
-    return y.view(input_shape)
-def liger_dyt_bwd(dy, x, alpha, gamma, beta):
-    assert dy.is_contiguous()
-    input_shape = x.shape
-    x = x.view(-1, input_shape[-1])
-    M, N = x.shape
-    HAVE_BETA = True if beta is not None else False
-    device = infer_device()
-    if device == "cuda":
-        NUM_SMS = torch.cuda.get_device_properties(x.device).multi_processor_count
-    elif device == "xpu":
-        NUM_SMS = torch.xpu.get_device_properties(x.device).gpu_subslice_count
-    elif device == "npu":
-        NUM_SMS = get_npu_core_count()
-    da = torch.zeros(NUM_SMS, triton.cdiv(N, 512), dtype=torch.float32, device=x.device)
-    dg = torch.empty(NUM_SMS, N, dtype=torch.float32, device=x.device)
-    db = torch.empty(NUM_SMS, N, dtype=torch.float32, device=x.device) if HAVE_BETA else None
-    dx = torch.empty_like(dy)
-    grid = lambda meta: (triton.cdiv(N, meta["BLOCK_N"]), NUM_SMS)
-    _dyt_bwd_kernel[grid](dy, dx, da, dg, db, x, alpha, gamma, HAVE_BETA, M, N)
-    if HAVE_BETA:
-        db = db.sum(0).to(x.dtype)
-    dg = dg.sum(0).to(gamma.dtype)
-    da = da.sum().to(x.dtype).unsqueeze(0)
-    return dx.view(input_shape), da, dg, db
-class LigerDyTFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, x, alpha, gamma, beta):
-        y = liger_dyt_fwd(x, alpha, gamma, beta)
-        ctx.save_for_backward(x, alpha, gamma, beta)
-        return y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dy):
-        x, alpha, gamma, beta = ctx.saved_tensors
-        dx, dalpha, dgamma, dbeta = liger_dyt_bwd(dy, x, alpha, gamma, beta)
-        return dx, dalpha, dgamma, dbeta

build/torch-xpu/fused_linear_cross_entropy.py DELETED Viewed

@@ -1,400 +0,0 @@
-import torch
-import triton
-from .cross_entropy import liger_cross_entropy_kernel
-from .utils import amp_custom_bwd
-from .utils import amp_custom_fwd
-from .utils import element_mul_kernel
-from .utils import is_hip
-from .utils import infer_device
-# The hard limit of TRITON_MAX_TENSOR_NUMEL is 1048576 https://github.com/triton-lang/triton/blob/ba42a5c68fd0505f8c42f4202d53be0f8d9a5fe0/python/triton/language/core.py#L19
-# However, setting limit as 65536 as in LayerNorm tutorial is faster because of less register spilling
-# The optimal maximum block size depends on your hardware, your kernel, and your dtype
-MAX_FUSED_SIZE = 2048 if infer_device() == "npu" else 65536 // 2
-def fused_linear_cross_entropy_forward(
-    _input,
-    weight,
-    target,
-    ce_weight=None,
-    bias=None,
-    ignore_index=-100,
-    lse_square_scale=0.0,
-    label_smoothing=0.0,
-    reduction="mean",
-    softcap=None,
-    return_z_loss=False,
-    accum_dtype=None,
-    use_token_scaling=False,
-    return_token_accuracy=False,
-    return_predicted_tokens=False,
-):
-    assert isinstance(return_z_loss, bool), f"return_z_loss must be True or False. Got: {return_z_loss}"
-    assert isinstance(return_token_accuracy, bool), (
-        f"return_token_accuracy must be True or False. Got: {return_token_accuracy}"
-    )
-    assert isinstance(return_predicted_tokens, bool), (
-        f"return_predicted_tokens must be True or False. Got: {return_predicted_tokens}"
-    )
-    device = _input.device
-    input_requires_grad = _input.requires_grad
-    # inputs have shape: BT x H
-    # materialized activations will have shape: BT x V
-    # the increase in memory = BT x V
-    # reduction can be achieved by partitioning the number of tokens BT into smaller chunks.
-    # for ex: if we were to achieve the same memory consumption as BT x H, then the chunk size should be:
-    # inc_factor = (V+H-1)//H, chunk_size = (BT + inc_factor - 1)//inc_factor
-    # for ex: BT = 4096*4, V = 32000, H = 4096 ==> inc_factor = 8, chunk_size = 2048
-    BT, H = _input.shape
-    V = weight.shape[0]
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    inc_factor = triton.cdiv(V, H)  # (V + H - 1) // H
-    chunk_size = triton.next_power_of_2(triton.cdiv(BT, inc_factor))  # (BT + inc_factor - 1) // inc_factor
-    num_chunks = triton.cdiv(BT, chunk_size)  # (BT + chunk_size - 1) // chunk_size
-    grad_input = torch.zeros_like(_input, device=device)
-    # we use fp32 for loss and gradients accumulator
-    if input_requires_grad:
-        if accum_dtype is None:
-            grad_weight = torch.zeros_like(weight, device=device) if weight.requires_grad else None
-            grad_bias = torch.zeros_like(bias, device=device) if bias is not None else None
-        else:
-            grad_weight = torch.zeros_like(weight, dtype=accum_dtype, device=device) if weight.requires_grad else None
-            grad_bias = torch.zeros_like(bias, dtype=accum_dtype, device=device) if bias is not None else None
-    else:
-        grad_weight = None
-        grad_bias = None
-    loss_1d = torch.zeros(BT, dtype=torch.float32, device=device)
-    z_loss_1d = torch.zeros(BT, dtype=_input.dtype, device=_input.device) if return_z_loss else None
-    token_accuracy_1d = torch.zeros(BT, dtype=torch.float32, device=device) if return_token_accuracy else None
-    predicted_tokens_1d = torch.full((BT,), -1, dtype=torch.int64, device=device) if return_predicted_tokens else None
-    # TODO: evaluate how CUDA synchronization caused by .item() affects the speed
-    target_mask = target != ignore_index
-    total_n_non_ignore = target_mask.sum().item()
-    total_sum_non_ignore_ce_weight = total_n_non_ignore
-    ce_weight_sum = 0.0
-    if ce_weight is not None:
-        assert ce_weight.shape[0] == V, f"If given, weight has to be a Tensor of size V. Got: {ce_weight.shape}"
-        assert torch.is_floating_point(ce_weight), (
-            f"If given, weight has to be a Tensor of floating point dtype. Got: {ce_weight.dtype}"
-        )
-        total_sum_non_ignore_ce_weight = (
-            torch.gather(ce_weight, dim=0, index=target.masked_select(target_mask)).sum().item()
-        )
-        ce_weight_sum = ce_weight.sum().item()
-        if ce_weight.stride(-1) != 1:
-            ce_weight = ce_weight.contiguous()
-    for chunk_id in range(num_chunks):
-        start_idx = chunk_id * chunk_size
-        end_idx = min((chunk_id + 1) * chunk_size, BT)
-        _input_chunk = _input[start_idx:end_idx]  # chunk_size x H
-        # when doing matmul, use the original precision
-        logits_chunk = _input_chunk @ weight.t()  # chunk_size x V
-        if bias is not None:
-            logits_chunk = logits_chunk + bias
-        target_chunk = target[start_idx:end_idx]  # chunk_size,
-        n_rows = logits_chunk.shape[0]
-        # Compute predicted probabilities for token scaling if needed
-        if use_token_scaling:
-            # Compute softmax probabilities for scaling
-            # We need to compute this before the cross entropy kernel modifies logits_chunk
-            logits_for_softmax = logits_chunk.detach().clone()  # Detach to avoid gradient flow
-            if softcap is not None:
-                logits_for_softmax = softcap * torch.tanh(logits_for_softmax / softcap)
-            # Compute softmax to get predicted probabilities
-            probs = torch.softmax(logits_for_softmax, dim=-1)
-            # Get predicted probabilities for token scaling, handling ignored targets
-            valid_target_mask = target_chunk != ignore_index
-            valid_targets = target_chunk[valid_target_mask]
-            if len(valid_targets) > 0:
-                # Gather probabilities only for valid targets
-                valid_probs = probs[valid_target_mask]
-                pred_probs_valid = torch.gather(valid_probs, -1, valid_targets.unsqueeze(-1)).squeeze(-1)
-                # Create full tensor with zeros for ignored targets
-                pred_probs = torch.zeros_like(target_chunk, dtype=probs.dtype, device=probs.device)
-                pred_probs[valid_target_mask] = pred_probs_valid
-            else:
-                # All targets are ignored
-                pred_probs = torch.zeros_like(target_chunk, dtype=probs.dtype, device=probs.device)
-            # Store the scaling factors
-            scaling_factors = pred_probs.detach()  # Detach to ensure no gradient flow
-        # unreduced loss
-        loss_1d_slice = loss_1d[start_idx:end_idx]  # chunk_size,
-        z_loss_1d_slice = z_loss_1d[start_idx:end_idx] if return_z_loss else None
-        token_accuracy_1d_slice = token_accuracy_1d[start_idx:end_idx] if return_token_accuracy else None
-        predicted_tokens_1d_slice = predicted_tokens_1d[start_idx:end_idx] if return_predicted_tokens else None
-        # ensure _input and target are contiguous
-        logits_chunk = logits_chunk.contiguous()
-        target_chunk = target_chunk.contiguous()
-        # Here we calculate the gradient of logits_chunk in place so we can save memory.
-        liger_cross_entropy_kernel[(n_rows,)](
-            X_ptr=logits_chunk,
-            X_stride=logits_chunk.stride(-2),
-            Y_ptr=target_chunk,
-            Y_stride=target_chunk.stride(-1),  # always 1
-            weight_ptr=ce_weight,
-            loss_ptr=loss_1d_slice,
-            z_loss_ptr=z_loss_1d_slice,
-            loss_stride=loss_1d_slice.stride(-1),  # always 1
-            token_accuracy_ptr=token_accuracy_1d_slice,
-            token_accuracy_stride=token_accuracy_1d_slice.stride(-1)
-            if return_token_accuracy
-            else 0,  # always 1 if accuracy is enabled
-            predicted_tokens_ptr=predicted_tokens_1d_slice,
-            predicted_tokens_stride=predicted_tokens_1d_slice.stride(-1)
-            if return_predicted_tokens
-            else 0,  # always 1 if predicted tokens is enabled
-            n_cols=V,
-            n_non_ignore=total_n_non_ignore,
-            sum_non_ignore_weight=total_sum_non_ignore_ce_weight,
-            weight_sum=ce_weight_sum,
-            ignore_index=ignore_index,
-            lse_square_scale=lse_square_scale,
-            label_smoothing=label_smoothing,
-            reduction=reduction,
-            softcap=softcap,
-            RETURN_Z_LOSS=return_z_loss,
-            RETURN_TOKEN_ACCURACY=return_token_accuracy,
-            RETURN_PREDICTED_TOKENS=return_predicted_tokens,
-            HAS_WEIGHT=True if ce_weight is not None else False,
-            HAS_SOFTCAPPING=True if softcap is not None else False,
-            HAS_GRADIENTS=input_requires_grad,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-        # Apply token scaling if requested
-        if use_token_scaling:
-            loss_1d_slice = loss_1d_slice * scaling_factors
-            if return_z_loss:
-                z_loss_1d_slice = z_loss_1d_slice * scaling_factors
-        loss_1d[start_idx:end_idx] = loss_1d_slice
-        if return_z_loss:
-            z_loss_1d[start_idx:end_idx] = z_loss_1d_slice
-        if return_token_accuracy:
-            token_accuracy_1d[start_idx:end_idx] = token_accuracy_1d_slice
-        if return_predicted_tokens:
-            predicted_tokens_1d[start_idx:end_idx] = predicted_tokens_1d_slice
-        grad_logits_chunk = logits_chunk  # chunk_size x V
-        # Apply token scaling to gradients if requested
-        if use_token_scaling:
-            # Expand scaling factors to match gradient dimensions
-            scaling_factors_expanded = scaling_factors.unsqueeze(-1)  # chunk_size x 1
-            grad_logits_chunk = grad_logits_chunk * scaling_factors_expanded
-        if input_requires_grad:
-            grad_input[start_idx:end_idx] = grad_logits_chunk @ weight
-        if grad_weight is not None and input_requires_grad:
-            grad_weight += torch.mm(grad_logits_chunk.t(), _input_chunk).float()
-        if bias is not None and input_requires_grad:
-            torch.add(
-                input=grad_bias,
-                other=grad_logits_chunk.sum(dim=0),
-                out=grad_bias,
-                alpha=1.0,
-            )
-    # Need extra calculations for backward if reduction=='none'. Not supporting reduction='none' now.
-    # if reduction == "none":
-    #     loss = loss_1d
-    #     z_loss = z_loss_1d if return_z_loss else None
-    if reduction == "none":
-        # Return per-token losses
-        loss = loss_1d
-        z_loss = z_loss_1d if return_z_loss else None
-        token_accuracy = token_accuracy_1d if return_token_accuracy else None
-    else:
-        loss = torch.sum(loss_1d)
-        z_loss = torch.sum(z_loss_1d) if return_z_loss else None
-        # For accuracy, we compute the mean across all non-ignored tokens
-        token_accuracy = torch.sum(token_accuracy_1d) / total_n_non_ignore if return_token_accuracy else None
-    predicted_tokens = predicted_tokens_1d if return_predicted_tokens else None
-    # Cast back to original dtype
-    grad_weight = grad_weight.to(weight.dtype) if grad_weight is not None else None
-    grad_bias = grad_bias.to(bias.dtype) if grad_bias is not None else None
-    return loss, z_loss, token_accuracy, predicted_tokens, grad_input, grad_weight, grad_bias
-def fused_linear_cross_entropy_backward(grad_output, grad_input, grad_weight, grad_bias):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul to save time
-    if not torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        # We use a Triton kernel instead of a PyTorch operation because modifying inputs in-place
-        # for gradient storage and backward multiple times causes anomalies with PyTorch but not with Triton.
-        BT, H = grad_input.shape
-        n_rows = BT
-        BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(H))
-        element_mul_kernel[(n_rows,)](
-            grad_input,
-            grad_input.stride(-2),
-            grad_output,
-            H,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=32 if not is_hip() else 16,
-        )
-        # handle grad_weight
-        if grad_weight is not None:
-            V, H = grad_weight.shape
-            n_rows = V
-            element_mul_kernel[(n_rows,)](
-                grad_weight,
-                grad_weight.stride(-2),
-                grad_output,
-                H,
-                BLOCK_SIZE=BLOCK_SIZE,
-                num_warps=32 if not is_hip() else 16,
-            )
-        if grad_bias is not None:
-            V = grad_bias.shape[0]
-            n_rows = V
-            element_mul_kernel[(n_rows,)](
-                grad_bias,
-                grad_bias.stride(-1),
-                grad_output,
-                1,
-                BLOCK_SIZE=BLOCK_SIZE,
-                num_warps=32 if not is_hip() else 16,
-            )
-    return grad_input, grad_weight, grad_bias
-class LigerFusedLinearCrossEntropyFunction(torch.autograd.Function):
-    @staticmethod
-    @amp_custom_fwd
-    def forward(
-        ctx,
-        _input,
-        weight,
-        target,
-        bias=None,
-        ce_weight=None,
-        ignore_index=-100,
-        lse_square_scale=0.0,
-        label_smoothing=0.0,
-        reduction="mean",
-        softcap=None,
-        return_z_loss: bool = False,
-        accum_dtype=None,
-        use_token_scaling: bool = False,
-        return_token_accuracy: bool = False,
-        return_predicted_tokens: bool = False,
-    ):
-        """
-        Fusing the last linear layer with cross-entropy loss
-            Reference: https://github.com/mgmalek/efficient_cross_entropy
-        Handle the forward and backward pass of the final linear layer via cross-entropy loss by avoiding
-        the materialization of the large logits tensor. Since Cross Entropy Loss is the last layer, we can
-        compute the gradient at the forward pass. By doing so, we don't have to store the _input and target
-        for the backward pass.
-        _input: (B*T, H) where B is batch size, T is sequence length, H is hidden dimension.
-        target: (B*T) where each value is in [0, V-1]
-        weight: (V, H) where V is the number of classes
-        bias: (V) where V is the number of classes
-        ce_weight: a manual rescaling weight given to each class. If given, has to be a Tensor of size V and floating point dtype
-        ignore_index: the index to ignore in the target
-        label_smoothing (float): The amount of smoothing when computing the loss, where 0.0 means no smoothing.
-        reduction: reduction to apply
-        accum_dtype (torch.dtype): the dtype of intermediate result buffers for weight and bias gradient accumulations.
-            Recommended to set `accum_dtype` to higher precision, e.g. `torch.float32`, if the training is unstable with original dtype. Default: `None`, performing accumulations in original dtype
-        use_token_scaling (bool): whether to scale each token's loss by its predicted probability (detached).
-            When True, each token's loss is multiplied by the model's predicted probability for that token's true class.
-            Default: False.
-        return_token_accuracy (bool): When `return_token_accuracy` is `True`, computes and returns per-token accuracy without materializing logits. Default: `False`
-        return_predicted_tokens (bool): When `return_predicted_tokens` is `True`, returns per-token predicted class indices (argmax) without materializing logits. Default: `False`
-        """
-        loss, z_loss, token_accuracy, predicted_tokens, grad_input, grad_weight, grad_bias = (
-            fused_linear_cross_entropy_forward(
-                _input=_input,
-                weight=weight,
-                target=target,
-                bias=bias,
-                ce_weight=ce_weight,
-                ignore_index=ignore_index,
-                lse_square_scale=lse_square_scale,
-                label_smoothing=label_smoothing,
-                reduction=reduction,
-                softcap=softcap,
-                return_z_loss=return_z_loss,
-                accum_dtype=accum_dtype,
-                use_token_scaling=use_token_scaling,
-                return_token_accuracy=return_token_accuracy,
-                return_predicted_tokens=return_predicted_tokens,
-            )
-        )
-        # downcast to dtype and store for backward
-        ctx.save_for_backward(
-            grad_input.detach(),
-            grad_weight.detach() if grad_weight is not None else None,
-            grad_bias.detach() if grad_bias is not None else None,
-        )
-        ctx.return_z_loss = return_z_loss
-        ctx.return_token_accuracy = return_token_accuracy
-        ctx.return_predicted_tokens = return_predicted_tokens
-        return loss, z_loss, token_accuracy, predicted_tokens
-    @staticmethod
-    @amp_custom_bwd
-    def backward(ctx, grad_output, grad_output2, grad_output3, grad_output4):
-        if ctx.return_z_loss:
-            del grad_output2  # z_loss is only for logging
-        if ctx.return_token_accuracy:
-            del grad_output3  # token_accuracy is only for metrics
-        if ctx.return_predicted_tokens:
-            del grad_output4  # predicted_tokens is only for metrics
-        (grad_input, grad_weight, grad_bias) = ctx.saved_tensors
-        grad_input, grad_weight, grad_bias = fused_linear_cross_entropy_backward(
-            grad_output, grad_input, grad_weight, grad_bias
-        )
-        return (
-            grad_input,
-            grad_weight,
-            None,
-            grad_bias,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,
-            None,  # use_token_scaling
-            None,  # return_token_accuracy
-            None,  # return_predicted_tokens
-        )

build/torch-xpu/geglu.py DELETED Viewed

@@ -1,143 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import tanh
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import tanh
-else:
-    from triton.language.math import tanh
-@triton.jit
-def _geglu_tanh_forward_kernel(a, b, c, stride, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    a += program_id * stride
-    b += program_id * stride
-    c += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    a_row = tl.load(a + col_offsets, mask=mask, other=0).to(tl.float32)
-    b_row = tl.load(b + col_offsets, mask=mask, other=0)
-    # tanh approximation form of GELU is computed with:
-    # 0.5 * a * (1 + tanh(sqrt(2 / pi) * (a + 0.044715 * a^3)))
-    sqrt_2_over_pi = 0.7978845608028654  # sqrt(2 / pi)
-    a_cubed = a_row * a_row * a_row
-    tanh_arg = sqrt_2_over_pi * (a_row + 0.044715 * a_cubed)
-    tanh_result = tanh(tanh_arg)
-    geglu_a = 0.5 * a_row * (1 + tanh_result)
-    c_row = geglu_a.cast(b_row.dtype) * b_row
-    tl.store(c + col_offsets, c_row, mask=mask)
-@triton.jit
-def _geglu_tanh_backward_kernel(dc, a, b, stride, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    dc += program_id * stride
-    a += program_id * stride
-    b += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    dc_row = tl.load(dc + col_offsets, mask=mask, other=0)
-    a_row = tl.load(a + col_offsets, mask=mask, other=0).to(tl.float32)
-    b_row = tl.load(b + col_offsets, mask=mask, other=0)
-    # recomputation to save memory
-    sqrt_2_over_pi = 0.7978845608028654  # sqrt(2 / pi)
-    a_cubed = a_row * a_row * a_row
-    tanh_arg = sqrt_2_over_pi * (a_row + 0.044715 * a_cubed)
-    tanh_result = tanh(tanh_arg)
-    geglu_a = 0.5 * a_row * (1 + tanh_result)
-    geglu_a = geglu_a.to(dc_row.dtype).to(tl.float32)
-    db_row = dc_row.cast(tl.float32) * geglu_a
-    # Gradient w.r.t. a can be computed with:
-    # b * (0.5 * (1 + tanh(z)) + 0.5 * a * (1 - tanh(z)^2) * (sqrt(2/pi) * (1 + 3 * 0.044715 * a^2)))
-    # where z = sqrt(2/pi) * (a + 0.044715 * a^3)
-    term1 = 0.5 * (1 + tanh_result)
-    tanh_sq = tanh_result * tanh_result
-    term2 = 0.5 * a_row * (1 - tanh_sq) * (sqrt_2_over_pi * (1 + 3 * 0.044715 * a_row * a_row))
-    da_row = dc_row * b_row * (term1 + term2)
-    tl.store(a + col_offsets, da_row, mask=mask)
-    tl.store(b + col_offsets, db_row.to(dc_row.dtype), mask=mask)
-def geglu_forward(a, b):
-    ori_shape = a.shape
-    n_cols = ori_shape[-1]
-    a = a.view(-1, n_cols)
-    b = b.view(-1, n_cols)
-    c = torch.empty_like(a)
-    n_rows = a.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _geglu_tanh_forward_kernel[(n_rows,)](
-        a,
-        b,
-        c,
-        c.stride(-2),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a, b, c.view(*ori_shape)
-def geglu_backward(a, b, dc):
-    ori_shape = dc.shape
-    n_cols = ori_shape[-1]
-    dc = dc.view(-1, n_cols)
-    n_rows = dc.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _geglu_tanh_backward_kernel[(n_rows,)](
-        dc,
-        a,
-        b,
-        dc.stride(-2),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a.view(*ori_shape), b.view(*ori_shape)
-class LigerGELUMulFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, a, b):
-        a, b, c = geglu_forward(a, b)
-        ctx.save_for_backward(a, b)
-        return c
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dc):
-        a, b = ctx.saved_tensors
-        a, b = geglu_backward(a, b, dc)
-        return a, b

build/torch-xpu/group_norm.py DELETED Viewed

@@ -1,311 +0,0 @@
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import infer_device
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-if infer_device() == "npu":
-    MAX_FUSED_SIZE = 16384  # 8192
-else:
-    MAX_FUSED_SIZE = 65536
-@triton.jit
-def _group_norm_forward_kernel(
-    Y_ptr,  # pointer to output, shape (n_rows, n_groups, hidden_size)
-    Y_row_stride,  # stride of each row in output
-    Y_col_stride,  # stride of each column in output
-    X_ptr,  # pointer to input, shape (n_rows, n_groups, hidden_size)
-    X_row_stride,  # stride of each row in input
-    X_col_stride,  # stride of each column in input
-    Mean_ptr,  # pointer to mean, shape (n_rows, n_groups)
-    Mean_row_stride,  # stride of each row in mean
-    Mean_col_stride,  # stride of each column in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows, n_groups)
-    RSTD_row_stride,  # stride of each row in rstd
-    RSTD_col_stride,  # stride of each column in rstd
-    W_ptr,  # pointer to W
-    B_ptr,  # pointer to B
-    hidden_size,  # hidden size of X
-    channels_per_group,  # the number of channels per group
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://nn.labml.ai/normalization/group_norm/index.html
-    """
-    batch_idx = tl.program_id(0)
-    group_idx = tl.program_id(1)
-    X_ptr += batch_idx * X_row_stride + group_idx * X_col_stride
-    Y_ptr += batch_idx * Y_row_stride + group_idx * Y_col_stride
-    block_range = tl.arange(0, BLOCK_SIZE)
-    # Compute mean and variance using the online algorithm
-    s = 0.0
-    squared_sum = 0.0
-    for i in tl.range(0, hidden_size, BLOCK_SIZE):
-        hidden_size_offsets = i + block_range
-        mask = hidden_size_offsets < hidden_size
-        X = tl.load(X_ptr + hidden_size_offsets, mask=mask, other=0.0)
-        s += tl.sum(X)
-        # X**2
-        squared_sum += tl.sum(X * X)
-    m = s / hidden_size
-    # variance = E[X**2] - E[X]**2
-    variance = (squared_sum / hidden_size) - (m * m)
-    # 1/std
-    rstd = rsqrt(variance + eps)
-    # Normalize — flat loop over full hidden_size (not per-channel)
-    # This avoids the nested channel × per_channel_hidden loop where
-    # BLOCK_SIZE >> hidden_size_per_channel causes massive padding waste.
-    hidden_size_per_channel = hidden_size // channels_per_group
-    for i in tl.range(0, hidden_size, BLOCK_SIZE):
-        hidden_size_offsets = i + block_range
-        mask = hidden_size_offsets < hidden_size
-        X = tl.load(X_ptr + hidden_size_offsets, mask=mask, other=m)
-        # Determine which channel each element belongs to, then load W/B
-        local_channel = hidden_size_offsets // hidden_size_per_channel
-        global_channel = group_idx * channels_per_group + local_channel
-        W = tl.load(W_ptr + global_channel, mask=mask)
-        B = tl.load(B_ptr + global_channel, mask=mask)
-        Y = (X - m) * rstd * W + B
-        tl.store(Y_ptr + hidden_size_offsets, Y, mask=mask)
-    tl.store(Mean_ptr + batch_idx * Mean_row_stride + group_idx * Mean_col_stride, m)
-    tl.store(RSTD_ptr + batch_idx * RSTD_row_stride + group_idx * RSTD_col_stride, rstd)
-@triton.jit
-def _group_norm_backward_kernel(
-    X_ptr,  # pointer to input, shape (n_rows, n_channels, hidden_size)
-    X_row_stride,  # stride of each row in input
-    X_col_stride,  # stride of each column in input
-    W_ptr,  # pointer to weights, shape (n_channels)
-    Mean_ptr,  # pointer to mean, shape (n_rows, n_groups)
-    Mean_ptr_row_stride,  # stride of each column in mean
-    Mean_ptr_col_stride,  # stride of each column in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows, n_groups)
-    DX_ptr,  # pointer to input grad, shape (n_rows, n_groups, hidden_size)
-    DW_ptr,  # pointer to weights grad, shape (n_channels)
-    DB_ptr,  # pointer to bias grad, shape (n_channels)
-    UPSTREAM_ptr,  # pointer to output grad, shape (n_rows, n_channels, hidden_size)
-    hidden_size: tl.constexpr,  # hidden size
-    channels_per_group: tl.constexpr,  # number of groups in group norm
-    BLOCK_SIZE: tl.constexpr,
-    dtype: tl.constexpr,
-):
-    """
-    References:
-    https://nn.labml.ai/normalization/group_norm/index.html
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    The backprop equations are the same for group_norm and layer_norm
-    the only difference here is that we load the Mean, Rstd corresponding to the
-    group we're computing gradients for and the mean and rstd are computed over n-channels
-    so the total number of elements we compute the mean over is num_channels_per_group * hidden_size
-    We also need to load the Weights corresponding to the current channel to compute the gradients.
-    """
-    batch_idx = tl.program_id(0)
-    group_idx = tl.program_id(1)
-    # Move the pointers to the correct batch
-    X_ptr += batch_idx * X_row_stride
-    DX_ptr += batch_idx * X_row_stride
-    UPSTREAM_ptr += batch_idx * X_row_stride
-    # Mean and rstd are the same shape so have the same strides
-    mean = tl.load(Mean_ptr + batch_idx * Mean_ptr_row_stride + group_idx * Mean_ptr_col_stride)
-    rstd = tl.load(RSTD_ptr + batch_idx * Mean_ptr_row_stride + group_idx * Mean_ptr_col_stride)
-    c1 = 0.0
-    c2 = 0.0
-    block_range = tl.arange(0, BLOCK_SIZE)
-    # We need to compute the sum terms of the backprop equations across all channels in the group
-    for channel_idx in range(group_idx * channels_per_group, (group_idx + 1) * channels_per_group):
-        dW = 0.0
-        dB = 0.0
-        # Move the pointers to the correct channel
-        W = tl.load(W_ptr + channel_idx)
-        for i in tl.range(0, hidden_size, BLOCK_SIZE):
-            hidden_size_offsets = i + block_range
-            mask = hidden_size_offsets < hidden_size
-            X = tl.load(
-                X_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            UPSTREAM_grad = tl.load(
-                UPSTREAM_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            x_hat = (X - mean) * rstd
-            dW += tl.sum(UPSTREAM_grad * x_hat)
-            dB += tl.sum(UPSTREAM_grad)
-            wdy = W * UPSTREAM_grad
-            c1 += tl.sum(x_hat * wdy)
-            c2 += tl.sum(wdy)
-        # Need to ensure additions to the same channel are atomic
-        tl.atomic_add(DW_ptr + channel_idx, dW.to(dtype))
-        tl.atomic_add(DB_ptr + channel_idx, dB.to(dtype))
-    N = hidden_size * channels_per_group
-    c1 = c1 / N
-    c2 = c2 / N
-    for channel_idx in tl.range(group_idx * channels_per_group, (group_idx + 1) * channels_per_group):
-        # Move the pointers to the correct channel
-        W = tl.load(W_ptr + channel_idx)
-        for i in range(0, hidden_size, BLOCK_SIZE):
-            hidden_size_offsets = i + block_range
-            mask = hidden_size_offsets < hidden_size
-            X = tl.load(
-                X_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            UPSTREAM_grad = tl.load(
-                UPSTREAM_ptr + channel_idx * X_col_stride + hidden_size_offsets,
-                mask=mask,
-                other=0.0,
-            )
-            x_hat = (X - mean) * rstd
-            wdy = W * UPSTREAM_grad
-            dx = (wdy - (x_hat * c1 + c2)) * rstd
-            tl.store(DX_ptr + channel_idx * X_col_stride + hidden_size_offsets, dx, mask=mask)
-def group_norm_forward(X, num_channels, num_groups, W, B, eps):
-    shape = X.shape
-    batch_size = shape[0]
-    channels_per_group = num_channels // num_groups
-    # Reshape X so that the mean and std are computed across the groups
-    X = X.view(batch_size, num_groups, -1).contiguous()
-    hidden_size = X.shape[-1]
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(hidden_size))
-    Y = torch.empty((batch_size, num_groups, hidden_size), dtype=X.dtype, device=X.device)
-    Mean = torch.zeros((batch_size, num_groups), dtype=X.dtype, device=X.device)
-    RSTD = torch.zeros((batch_size, num_groups), dtype=X.dtype, device=X.device)
-    _group_norm_forward_kernel[(batch_size, num_groups)](
-        Y,
-        Y.stride(0),
-        Y.stride(1),
-        X,
-        X.stride(0),
-        X.stride(1),
-        Mean,
-        Mean.stride(0),
-        Mean.stride(1),
-        RSTD,
-        RSTD.stride(0),
-        RSTD.stride(1),
-        W,
-        B,
-        hidden_size,
-        channels_per_group,
-        eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-    )
-    # Return tensors in the original shape
-    return Y.view(*shape), X.view(*shape), Mean, RSTD, BLOCK_SIZE
-def group_norm_backward(dY, X, W, B, Mean, RSTD, num_channels, num_groups):
-    shape = dY.shape
-    batch_size = shape[0]
-    hidden_size = dY.shape[-1]
-    channels_per_group = num_channels // num_groups
-    dY = dY.view(batch_size, num_groups, -1)
-    DX = torch.empty(
-        (batch_size, num_groups, hidden_size * channels_per_group),
-        dtype=X.dtype,
-        device=X.device,
-    )
-    DW = torch.zeros((num_channels), dtype=W.dtype, device=W.device)
-    DB = torch.zeros((num_channels), dtype=B.dtype, device=B.device)
-    triton_dtype = tl.float32 if X.dtype == torch.float32 else tl.bfloat16
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(hidden_size))
-    _group_norm_backward_kernel[(batch_size, num_groups)](
-        X,
-        X.stride(0),
-        X.stride(1),
-        W,
-        Mean,
-        Mean.stride(0),
-        Mean.stride(1),
-        RSTD,
-        DX,
-        DW,
-        DB,
-        dY,
-        hidden_size,
-        channels_per_group,
-        BLOCK_SIZE=BLOCK_SIZE,
-        dtype=triton_dtype,
-    )
-    # Return tensors in the original shape
-    return DX.view(*shape), DW, DB
-class LigerGroupNormFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        X,
-        affine_scaling_weight,
-        affine_shifting_bias,
-        num_channels,
-        num_groups,
-        eps,
-    ):
-        Y, X, Mean, RSTD, BLOCK_SIZE = group_norm_forward(
-            X,
-            num_channels,
-            num_groups,
-            affine_scaling_weight,
-            affine_shifting_bias,
-            eps,
-        )
-        ctx.num_channels = num_channels
-        ctx.num_groups = num_groups
-        ctx.save_for_backward(X, affine_scaling_weight, affine_shifting_bias, Mean, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        X, W, B, Mean, RSTD = ctx.saved_tensors
-        DX, DW, DB = group_norm_backward(dY, X, W, B, Mean, RSTD, ctx.num_channels, ctx.num_groups)
-        return DX, DW, DB, None, None, None

build/torch-xpu/jsd.py DELETED Viewed

@@ -1,201 +0,0 @@
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-from .utils import infer_device
-@triton.jit
-def _jsd_kernel(
-    X_ptr,  # input in logspace, X = log Q
-    X_stride,
-    Y_ptr,  # ground truth in logspace, Y = log P
-    Y_stride,
-    loss_ptr,
-    loss_stride,
-    dX_ptr,
-    dX_stride,
-    label_ptr,
-    beta: tl.constexpr,
-    n_non_ignore: int,
-    ignore_index: tl.constexpr,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-    HAS_LABEL: tl.constexpr,
-):
-    # JSD(P || Q) = (KL(P || M) + KL(Q || M)) / 2, M = (1/2) * (P + Q) = (1/2) * (e ^ Y + e ^ X)
-    #             = sum(P * log P + Q * log Q - 2 * M * log M) / 2
-    #             = sum(e ^ Y * Y + e ^ X * X - 2 * M * log M) / 2
-    # grad_x_i = 0.5 * Q * (X - log_M)
-    pid = tl.program_id(0).to(tl.int64)
-    X_ptr += pid * X_stride
-    dX_ptr += pid * dX_stride
-    Y_ptr += pid * Y_stride
-    loss_ptr += pid * loss_stride
-    label_ptr += pid
-    if HAS_LABEL:
-        label = tl.load(label_ptr)
-        if label == ignore_index:
-            for i in range(0, n_cols, BLOCK_SIZE):
-                offsets = i + tl.arange(0, BLOCK_SIZE)
-                tl.store(dX_ptr + offsets, 0.0, mask=offsets < n_cols)
-            return
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + tl.arange(0, BLOCK_SIZE)
-        mask = offsets < n_cols
-        X = tl.load(X_ptr + offsets, mask=mask, other=float("-inf")).to(tl.float32)
-        Y = tl.load(Y_ptr + offsets, mask=mask, other=float("-inf")).to(tl.float32)
-        if beta == 0.0:  # forward KL
-            Y_max = tl.max(Y, axis=0)
-            Y_shifted = Y - Y_max
-            Y_prob = tl.exp(Y_shifted) * tl.exp(Y_max)  # Compensate for the shift
-            loss = Y_prob * (Y - X)
-            dX = -Y_prob
-        elif beta == 1.0:  # reverse KL
-            X_max = tl.max(X, axis=0)
-            X_shifted = X - X_max
-            X_prob = tl.exp(X_shifted) * tl.exp(X_max)  # Compensate for the shift
-            loss = X_prob * (X - Y)
-            dX = loss + X_prob
-        else:
-            max_val = tl.maximum(tl.max(X, axis=0), tl.max(Y, axis=0))
-            X_shifted = X - max_val
-            Y_shifted = Y - max_val
-            # Pre-compute exp(max_val) since it's used twice
-            exp_max = tl.exp(max_val)
-            # Compute exp terms with compensation
-            Q = tl.exp(X_shifted) * exp_max  # = exp(X)
-            P = tl.exp(Y_shifted) * exp_max  # = exp(Y)
-            # Pre-compute common terms
-            beta_P = beta * P
-            one_minus_beta_Q = (1 - beta) * Q
-            M = beta_P + one_minus_beta_Q
-            log_M = tl.log(M)  # No need to compensate as M is already in original scale
-            loss = beta_P * Y + one_minus_beta_Q * X - M * log_M
-            dX = one_minus_beta_Q * (X - log_M)
-        # Pre-compute scaling factor
-        scale = 1.0 / n_non_ignore
-        loss = loss * scale
-        dX = dX * scale
-        tl.store(loss_ptr + offsets, loss, mask=mask)
-        tl.store(dX_ptr + offsets, dX, mask=mask)
-MAX_FUSED_SIZE = 4096 if infer_device() == "xpu" else 65536
-def jsd_forward(_input, target, shift_labels, beta, ignore_index, has_label):
-    BT, V = _input.shape
-    n_rows = BT
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    # non reduction loss
-    loss = torch.zeros(_input.shape, dtype=torch.float32, device=_input.device)
-    dX = torch.empty_like(_input)
-    if has_label:
-        n_non_ignore = (shift_labels != ignore_index).sum().item()
-    else:
-        n_non_ignore = BT
-    _jsd_kernel[(n_rows,)](
-        X_ptr=_input,  # input in logspace, X = log Q
-        X_stride=_input.stride(-2),
-        Y_ptr=target,  # ground truth in logspace, Y = log P
-        Y_stride=target.stride(-2),
-        loss_ptr=loss,
-        loss_stride=loss.stride(-2),
-        dX_ptr=dX,
-        dX_stride=dX.stride(-2),
-        label_ptr=(shift_labels if has_label else torch.empty(1, device=_input.device)),  # dummy ptr if no label
-        beta=beta,
-        n_non_ignore=n_non_ignore,
-        ignore_index=ignore_index,
-        n_cols=V,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_LABEL=has_label,
-    )
-    loss = torch.sum(loss)
-    return loss.to(_input.dtype), dX
-def jsd_backward(dX, grad_output):
-    # If jsd is the last layer, grad_output is 1.0. Skip the mul to save time
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return dX
-    else:
-        return grad_output * dX
-class LigerJSDFunction(torch.autograd.Function):
-    r"""
-    This class implements the forward and backward pass for the generalized Jensen-Shannon Divergence.
-    .. math::
-        JSD(\beta)(P || Q)
-            = \beta * KLDiv(P || (\beta * P + (1 - \beta) * Q)) + (1 - \beta) * KLDiv(Q || (\beta * P + (1 - \beta) * Q))
-    .. note::
-        As all the other losses in PyTorch, this function expects the first argument,
-        :attr:`_input`, to be the predictions, the output of the student model, in log-space
-        and the second, :attr:`target`, to be the observations, the output of the teacher model, in log-space.
-        This differs from the standard mathematical notation :math:`JSD(P || Q)` where
-        :math:`P` denotes the teacher model and :math:`Q` denotes the student model.
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-        beta: float = 0.5,
-        ignore_index: int = -100,
-    ) -> torch.Tensor:
-        """
-        Args:
-            _input (torch.Tensor): predict values with shape (BT, V) in logspace
-            target (torch.Tensor): ground truth values with shape (BT, V) in logspace
-            shift_labels (Optional[torch.LongTensor]): indicator of next predicted vocab with shape (BT) where each value is in [0, V-1].
-            beta (float): coefficient beta of generalized JSD in the interval [0, 1]. It implements forward/reverse KL when beta equals 0 and 1 respectively. Default: `0.5`
-            ignore_index (int): the index to ignore. Default: -100
-        Returns:
-            loss (torch.Tensor): generalized JSD
-        """
-        has_label = False
-        if shift_labels is not None:
-            assert shift_labels.shape == (_input.shape[0],), (
-                f"the shape of shift_labels must be (BT,). Got: {shift_labels.shape}"
-            )
-            shift_labels = shift_labels.contiguous()
-            has_label = True
-        loss, dX = jsd_forward(_input, target, shift_labels, beta, ignore_index, has_label)
-        ctx.save_for_backward(dX)
-        return loss
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        (dX,) = ctx.saved_tensors
-        dX = jsd_backward(dX, grad_output)
-        return (
-            dX,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-xpu/kl_div.py DELETED Viewed

@@ -1,259 +0,0 @@
-from typing import Literal
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-from .utils import is_hip
-from .utils import infer_device
-def get_num_warps(BLOCK_SIZE):
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32 if not is_hip() else 16
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return num_warps
-if infer_device() == "xpu":
-    MAX_FUSED_SIZE = 8192
-elif infer_device() == "npu":
-    MAX_FUSED_SIZE = 8192
-else:
-    MAX_FUSED_SIZE = 65536 // 4  # 65536 // 4 or 8 works the best
-REDUCTION_LITERAL = Literal["none", "sum", "mean", "batchmean"]
-_REDUCTION_MODE_NONE: tl.constexpr = tl.constexpr(0)
-_REDUCTION_MODE_SUM: tl.constexpr = tl.constexpr(1)
-_REDUCTION_MODE_MEAN: tl.constexpr = tl.constexpr(2)
-_REDUCTION_MODE_BATCHMEAN: tl.constexpr = tl.constexpr(3)
-_str_to_reduction_mode = {
-    "none": _REDUCTION_MODE_NONE.value,
-    "sum": _REDUCTION_MODE_SUM.value,
-    "mean": _REDUCTION_MODE_MEAN.value,
-    "batchmean": _REDUCTION_MODE_BATCHMEAN.value,
-}
-@triton.jit
-def _kldiv_kernel_forward(
-    y_ptr,  # [B, S], prediction ptr, the kernel expects the prediction in log-space
-    y_stride,  # int, prediction stride
-    gt_ptr,  # [B, S], ground truth ptr
-    gt_stride,  # int, ground truth stride
-    loss_ptr,  # [B] or [B, S] if reduction == _REDUCTION_MODE_NONE, output ptr
-    loss_stride,  # int, output stride
-    n_cols,  # int, number of columns in the input tensor
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-    log_target: tl.constexpr = False,
-    reduction: tl.constexpr = _REDUCTION_MODE_BATCHMEAN,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    y_ptr += pid * y_stride
-    gt_ptr += pid * gt_stride
-    loss_ptr += pid * loss_stride
-    base_offsets = tl.arange(0, BLOCK_SIZE)
-    loss_sum = 0.0
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + base_offsets
-        mask = offsets < n_cols
-        y = tl.load(y_ptr + offsets, mask=mask, other=0.0)
-        y_true = tl.load(gt_ptr + offsets, mask=mask, other=0.0)
-        # KL(y_true || y) = y_true * (log(y_true) - log(y))
-        # We compute KL(y_true || y) with y in the log-space
-        if not log_target:
-            loss = y_true * (tl.log(tl.maximum(y_true, eps)) - y)
-        else:
-            loss = tl.exp(y_true) * (y_true - y)
-        if reduction == _REDUCTION_MODE_NONE:
-            tl.store(loss_ptr + offsets, loss, mask=mask)
-        else:
-            loss_sum += tl.sum(loss, axis=0)
-    if reduction != _REDUCTION_MODE_NONE:
-        tl.store(loss_ptr, loss_sum)
-@triton.jit
-def _kldiv_kernel_backward(
-    target_ptr,
-    target_stride,
-    new_grads_ptr,
-    new_grads_stride,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-    log_target: tl.constexpr = False,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    target_ptr += pid * target_stride
-    new_grads_ptr += pid * new_grads_stride
-    offsets = tl.arange(0, BLOCK_SIZE)
-    mask = offsets < n_cols
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + tl.arange(0, BLOCK_SIZE)
-        mask = offsets < n_cols
-        target = tl.load(target_ptr + offsets, mask=mask, other=0.0)
-        if not log_target:
-            res = target * -1
-        else:
-            res = -tl.exp(target)
-        tl.store(new_grads_ptr + offsets, res, mask=mask)
-def kldiv_forward_triton(y_pred, y_true, log_target, reduction, eps):  # [BT, V]
-    BT, V = y_pred.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = 32 if infer_device() == "xpu" else get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    reduction = _str_to_reduction_mode[reduction]
-    out_size = (BT, V) if reduction == _REDUCTION_MODE_NONE.value else (BT,)
-    output_tensor = torch.zeros(out_size, device=y_pred.device, dtype=torch.float32)
-    _kldiv_kernel_forward[grid](
-        y_pred,
-        y_pred.stride(0),
-        y_true,
-        y_true.stride(0),
-        output_tensor,
-        output_tensor.stride(0),
-        V,
-        eps=eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        log_target=log_target,
-        reduction=reduction,
-    )
-    # calculated according to the reduction mode same as in Pytorch. In the later versions, `mean` will be changed to the same behavior as `batchmean`
-    # https://pytorch.org/docs/stable/generated/torch.nn.KLDivLoss.html
-    # https://github.com/pytorch/pytorch/blob/d7b57c4d63edb42e1deeeba9497fcb5f1f748ff2/torch/nn/functional.py#L3372
-    if reduction == _REDUCTION_MODE_BATCHMEAN.value:
-        return output_tensor.sum() / BT
-    elif reduction == _REDUCTION_MODE_SUM.value:
-        return output_tensor.sum(dim=0)
-    elif reduction == _REDUCTION_MODE_MEAN.value:
-        return output_tensor.sum() / (BT * V)
-    else:
-        return output_tensor
-def kldiv_backward_triton(target, grad_output, new_grads, log_target):
-    BT, V = target.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = 32 if infer_device() == "xpu" else get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    # We store the gradients in-place in the input tensor
-    _kldiv_kernel_backward[grid](
-        target,
-        target.stride(0),
-        new_grads,
-        new_grads.stride(0),
-        V,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        log_target=log_target,
-    )
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul then.
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return new_grads
-    return new_grads * grad_output
-class LigerKLDivLossFunction(torch.autograd.Function):
-    """
-    Class implementing the forward and backward pass for the KL Divergence Loss using Triton, as defined by the following formula:
-    ```python
-    if log_target:
-        loss = target.exp() * (target - input)
-    else:
-        loss = target * (target.log() - input)
-    ```,
-    then the loss is reduced according to the `reduction` parameter.
-    as defined in the PyTorch documentation: https://pytorch.org/docs/stable/generated/torch.nn.KLDivLoss.html
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        y_pred: torch.Tensor,
-        y_true: torch.Tensor,
-        reduction: REDUCTION_LITERAL = "batchmean",
-        log_target: bool = False,
-        eps: float = 1e-10,
-    ) -> torch.Tensor:
-        """A forward pass for the KL Divergence Loss.
-        Args:
-            ctx: Torch autograd context
-            y_pred (torch.Tensor): A tensor of shape (BT, V) containing the predicted values, expected to be log-probabilities.
-            y_true (torch.Tensor): A tensor of shape (BT, V) containing the target values, expected to be either probabilities or log-probabilities, depending on the value of `log_target`.
-            reduction (REDUCTION_LITERAL, optional): Reduction to be used. Defaults to "batchmean".
-            log_target (bool, optional): If set to true, expects the ground truth to already be log-probabilities. Defaults to False.
-            eps: (float, optional): A small value to avoid division by zero. Defaults to 1e-10.
-        Returns:
-            torch.Tensor: The computed KL Divergence Loss, with shape (BT, V) if `reduction` is "none", else a scalar.
-        """
-        ctx.save_for_backward(y_true)
-        ctx.reduction = reduction
-        ctx.log_target = log_target
-        return kldiv_forward_triton(y_pred, y_true, log_target=log_target, reduction=reduction, eps=eps)
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        """A backward pass for the KL Divergence Loss.
-        Args:
-            ctx: Torch autograd context
-            grad_output (torch.Tensor): The gradient of the loss with respect to the output.
-        Returns:
-            tuple[torch.Tensor, None, None, None, None]: The gradient of the loss with respect to the inputs and None for the other arguments of the forward method.
-        """
-        (y_true,) = ctx.saved_tensors
-        new_grads = torch.empty_like(y_true)
-        derivative = kldiv_backward_triton(y_true, grad_output, new_grads, ctx.log_target)
-        if ctx.reduction == "batchmean":
-            derivative = derivative / y_true.shape[0]
-        elif ctx.reduction == "sum" or ctx.reduction == "none":
-            pass
-        elif ctx.reduction == "mean":
-            derivative = derivative / (y_true.shape[0] * y_true.shape[1])
-        return (
-            derivative,
-            None,
-            None,
-            None,
-            None,
-        )

build/torch-xpu/layer_norm.py DELETED Viewed

@@ -1,320 +0,0 @@
-import math
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import set_large_grf_mode
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-@triton.jit
-def _layer_norm_forward_kernel(
-    Y_ptr,  # pointer to output, shape (n_rows, n_cols)
-    Y_row_stride,  # stride of each row in output
-    X_ptr,  # pointer to input, shape (n_rows, n_cols)
-    X_row_stride,  # stride of each row in input
-    W_ptr,  # pointer to weights, shape (n_cols,)
-    W_row_stride,  # stride of each row in weights
-    B_ptr,  # pointer to bias, shape (n_cols,)
-    B_row_stride,  # stride of each row in bias
-    Mean_ptr,  # pointer to mean, shape (n_rows,)
-    Mean_row_stride,  # stride of each row in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows,)
-    RSTD_row_stride,  # stride of each row in rstd
-    n_cols,
-    eps,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://arxiv.org/abs/1607.06450
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    """
-    row_idx = tl.program_id(0).to(tl.int64)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    # Pre-load weights and bias in fp32 to avoid repeated conversions
-    W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0.0)
-    B_row = tl.load(B_ptr + col_offsets, mask=mask, other=0.0)
-    W_f32 = W_row.to(tl.float32)
-    B_f32 = B_row.to(tl.float32)
-    # Calculate pointers for this row
-    row_X_ptr = X_ptr + row_idx * X_row_stride
-    row_Y_ptr = Y_ptr + row_idx * Y_row_stride
-    row_Mean_ptr = Mean_ptr + row_idx * Mean_row_stride
-    row_RSTD_ptr = RSTD_ptr + row_idx * RSTD_row_stride
-    # Load input data and convert to fp32 for numerical stability
-    X_row = tl.load(row_X_ptr + col_offsets, mask=mask, other=0.0)
-    X_f32 = X_row.to(tl.float32)
-    # Compute statistics in fp32 for numerical stability
-    mean = tl.sum(X_f32, axis=0) / n_cols
-    X_centered = X_f32 - mean
-    # Apply mask to variance calculation to exclude contributions from masked elements
-    X_centered_masked = tl.where(mask, X_centered, 0.0)
-    var = tl.sum(X_centered_masked * X_centered_masked, axis=0) / n_cols
-    rstd = rsqrt(var + eps)
-    # Store statistics (convert back to original dtype only once)
-    tl.store(row_Mean_ptr, mean.to(X_row.dtype))
-    tl.store(row_RSTD_ptr, rstd.to(X_row.dtype))
-    # Fused normalization and affine transformation
-    # Y = (X - mean) * rstd * W + B = X_centered * rstd * W + B
-    Y_f32 = X_centered * rstd * W_f32 + B_f32
-    # Store output (single conversion back to original dtype)
-    tl.store(row_Y_ptr + col_offsets, Y_f32.to(X_row.dtype), mask=mask)
-@triton.jit
-def _layer_norm_backward_kernel(
-    X_ptr,  # pointer to input, shape (n_rows, n_cols)
-    stride_x,  # stride of each row in input
-    W_ptr,  # pointer to weights, shape (n_cols,)
-    Mean_ptr,  # pointer to mean, shape (n_rows,)
-    stride_mean,  # stride of each row in mean
-    RSTD_ptr,  # pointer to rstd, shape (n_rows,)
-    stride_rstd,  # stride of each row in rstd
-    DX_ptr,  # pointer to input grad, shape (n_rows, n_cols)
-    stride_dx,  # stride of each row in input grad
-    DW_ptr,  # pointer to weights grad, shape (n_cols,)
-    stride_dw,  # stride of each row in weights grad
-    DB_ptr,  # pointer to bias grad, shape (n_cols,)
-    stride_db,  # stride of each row in bias grad
-    DY_ptr,  # pointer to output grad, shape (n_rows, n_cols)
-    stride_dy,  # stride of each row in output grad
-    n_rows,
-    n_cols,
-    rows_per_program: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    References:
-    https://arxiv.org/abs/1607.06450
-    https://github.com/karpathy/llm.c/blob/master/doc/layernorm/layernorm.md
-    """
-    row_block_id = tl.program_id(0).to(tl.int64)
-    row_start = row_block_id * rows_per_program
-    row_end = min((row_block_id + 1) * rows_per_program, n_rows)
-    cols = tl.arange(0, BLOCK_SIZE)
-    mask = cols < n_cols
-    dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    db_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    # Pre-load weights once (same optimization as forward pass)
-    w = tl.load(W_ptr + cols, mask=mask, other=0.0)
-    w_f32 = w.to(tl.float32)
-    for row_idx in range(row_start, row_end):
-        # Calculate pointers for this specific row
-        row_X_ptr = X_ptr + row_idx * stride_x
-        row_DX_ptr = DX_ptr + row_idx * stride_dx
-        row_DY_ptr = DY_ptr + row_idx * stride_dy
-        row_Mean_ptr = Mean_ptr + row_idx * stride_mean
-        row_RSTD_ptr = RSTD_ptr + row_idx * stride_rstd
-        # Load data for this row
-        x = tl.load(row_X_ptr + cols, mask=mask, other=0.0)
-        dy = tl.load(row_DY_ptr + cols, mask=mask, other=0.0)
-        mean = tl.load(row_Mean_ptr)
-        rstd = tl.load(row_RSTD_ptr)
-        # Convert to fp32 for numerical stability
-        x_f32 = x.to(tl.float32)
-        dy_f32 = dy.to(tl.float32)
-        mean_f32 = mean.to(tl.float32)
-        rstd_f32 = rstd.to(tl.float32)
-        # Compute backward pass for this row
-        x_hat = (x_f32 - mean_f32) * rstd_f32
-        wdy = w_f32 * dy_f32
-        c1 = tl.sum(x_hat * wdy, axis=0) / n_cols
-        c2 = tl.sum(wdy, axis=0) / n_cols
-        dx = (wdy - (x_hat * c1 + c2)) * rstd_f32
-        # Store input gradient
-        tl.store(row_DX_ptr + cols, dx, mask=mask)
-        # Accumulate weight and bias gradients for this thread block's assigned rows
-        dw = dy_f32 * x_hat
-        db = dy_f32
-        dW_row += dw
-        db_row += db
-    tl.store(DW_ptr + row_block_id * stride_dw + cols, dW_row, mask=mask)
-    tl.store(DB_ptr + row_block_id * stride_db + cols, db_row, mask=mask)
-def layer_norm_forward(X, W, B, eps):
-    """
-    Args:
-        X: Input tensor of shape (..., hidden_size)
-        W: Weight tensor of shape (hidden_size,)
-        B: Bias tensor of shape (hidden_size,)
-        eps: Small constant for numerical stability
-    Returns:
-        Tuple of (output, input, mean, rstd, block_size, num_warps)
-    """
-    shape = X.shape
-    dim = shape[-1]
-    X = X.view(-1, dim)
-    n_rows, n_cols = X.shape
-    # Calculate optimal block size and warp configuration
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    # Allocate output tensors
-    Y = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    Mean = torch.empty(n_rows, dtype=X.dtype, device=X.device)
-    RSTD = torch.empty(n_rows, dtype=X.dtype, device=X.device)
-    # Validate input dimensions
-    if X.shape[1] != W.shape[0]:
-        raise ValueError(
-            f"Incompatible dimensions: input feature size (X.shape[1]={X.shape[1]}) "
-            f"must match weight size (W.shape[0]={W.shape[0]})"
-        )
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    # Launch kernel with one thread block per row for optimal performance
-    grid = (n_rows,)
-    _layer_norm_forward_kernel[grid](
-        Y,
-        Y.stride(0),
-        X,
-        X.stride(0),
-        W,
-        W.stride(0),
-        B,
-        B.stride(0),
-        Mean,
-        Mean.stride(0),
-        RSTD,
-        RSTD.stride(0),
-        n_cols,
-        eps,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-        **kernel_args,
-    )
-    return Y.view(*shape), X, Mean, RSTD, BLOCK_SIZE, num_warps
-def layer_norm_backward(dY, X, W, B, Mean, RSTD):
-    """
-    Args:
-        dY: Gradient of output
-        X: Input tensor
-        W: Weight tensor
-        B: Bias tensor
-        Mean: Pre-computed mean
-        RSTD: Pre-computed reciprocal standard deviation
-    Returns:
-        Tuple of (input_grad, weight_grad, bias_grad)
-    """
-    shape = dY.shape
-    dim = shape[-1]
-    dY = dY.view(-1, dim)
-    n_rows, n_cols = dY.shape
-    sm_count = 1
-    if X.device.type == "cuda":
-        sm_count = torch.cuda.get_device_properties(X.device).multi_processor_count
-    elif X.device.type == "xpu":
-        sm_count = torch.xpu.get_device_properties(X.device).gpu_eu_count
-    elif X.device.type == "npu":
-        sm_count = get_npu_core_count()
-    # fp32 for numerical stability especially.
-    _DW = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-    _DB = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-    # Calculate optimal block size and warp configuration
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    if n_cols > BLOCK_SIZE:
-        raise RuntimeError(f"Feature dimension {n_cols} exceeds maximum supported size of {BLOCK_SIZE}.")
-    rows_per_program = math.ceil(n_rows / sm_count)
-    grid = (sm_count,)
-    # Allocate gradient tensors
-    DX = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    kernel_args = {"num_warps": num_warps}
-    # XPU-specific optimization
-    if X.device.type == "xpu":
-        kernel_args.update({"num_warps": 32, "num_stages": 4})
-        set_large_grf_mode(kernel_args)
-    # Launch kernel with one thread block per row for optimal performance
-    _layer_norm_backward_kernel[grid](
-        X,
-        X.stride(0),
-        W,
-        Mean,
-        Mean.stride(0),
-        RSTD,
-        RSTD.stride(0),
-        DX,
-        DX.stride(0),
-        _DW,
-        _DW.stride(0),
-        _DB,
-        _DB.stride(0),
-        dY,
-        dY.stride(0),
-        n_rows,
-        n_cols,
-        rows_per_program=rows_per_program,
-        BLOCK_SIZE=BLOCK_SIZE,
-        **kernel_args,
-    )
-    DX = DX.view(*shape)
-    DW = _DW.sum(dim=0).to(W.dtype)
-    DB = _DB.sum(dim=0).to(B.dtype)
-    return DX, DW, DB
-class LigerLayerNormFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, X, W, B, eps):
-        Y, X, Mean, RSTD, BLOCK_SIZE, num_warps = layer_norm_forward(X, W, B, eps)
-        ctx.save_for_backward(X, W, B, Mean, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        X, W, B, Mean, RSTD = ctx.saved_tensors
-        DX, DW, DB = layer_norm_backward(dY, X, W, B, Mean, RSTD)
-        return DX, DW, DB, None

build/torch-xpu/layers.py DELETED Viewed

@@ -1,514 +0,0 @@
-import inspect
-from dataclasses import dataclass
-from typing import Optional, Tuple
-import torch
-import torch.nn as nn
-from .cross_entropy import LigerCrossEntropyFunction
-from .dyt import LigerDyTFunction
-from .fused_linear_cross_entropy import LigerFusedLinearCrossEntropyFunction
-from .geglu import LigerGELUMulFunction
-from .group_norm import LigerGroupNormFunction
-from .jsd import LigerJSDFunction
-from .kl_div import LigerKLDivLossFunction
-from .layer_norm import LigerLayerNormFunction
-from .qwen2vl_mrope import LigerQwen2VLMRopeFunction
-from .rms_norm import LigerRMSNormFunction
-from .rope import LigerRopeFunction
-from .swiglu import LigerSiLUMulFunction
-from .tiled_mlp import apply_tiled_mlp
-from .tvd import LigerTVDLossFunction
-class LigerRMSNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        offset: float = 0.0,
-        casting_mode: str = "llama",
-        init_fn: str = "ones",
-        in_place: bool = True,
-        row_mode: Optional[bool] = None,
-        elementwise_affine: bool = True,
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.offset = offset
-        self.casting_mode = casting_mode
-        self.in_place = in_place
-        self.row_mode = row_mode
-        self.elementwise_affine = elementwise_affine
-        if elementwise_affine:
-            init = torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size)
-            self.weight = nn.Parameter(init)
-        else:
-            self.register_parameter("weight", None)
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        return LigerRMSNormFunction.apply(
-            hidden_states,
-            self.weight,
-            self.variance_epsilon,
-            self.offset,
-            self.casting_mode,
-            self.in_place,
-            self.row_mode,
-        )
-    def extra_repr(self) -> str:
-        return (
-            f"{self.hidden_size}, eps={self.variance_epsilon}, offset={self.offset}, "
-            f"in_place={self.in_place}, row_mode={self.row_mode}"
-        )
-class LigerLayerNorm(nn.Module):
-    def __init__(
-        self,
-        hidden_size: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        self.hidden_size = hidden_size
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(hidden_size) if init_fn == "ones" else torch.zeros(hidden_size))
-        self.bias = nn.Parameter(torch.randn(hidden_size) if bias else torch.zeros(hidden_size))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        return LigerLayerNormFunction.apply(hidden_states, self.weight, self.bias, self.variance_epsilon)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, eps={self.variance_epsilon}"
-class LigerGroupNorm(nn.Module):
-    def __init__(
-        self,
-        num_channels: int,
-        num_groups: int,
-        eps: float = 1e-6,
-        bias: bool = False,
-        init_fn: str = "ones",
-    ):
-        super().__init__()
-        assert init_fn in ("ones", "zeros"), f"init_fn must be 'ones' or 'zeros', got {init_fn}"
-        assert num_channels % num_groups == 0, (
-            f"num_channels ({num_channels}) must be divisible by num_groups ({num_groups})"
-        )
-        self.num_channels = num_channels
-        self.num_groups = num_groups
-        self.variance_epsilon = eps
-        self.weight = nn.Parameter(torch.ones(num_channels) if init_fn == "ones" else torch.zeros(num_channels))
-        self.bias = nn.Parameter(torch.randn(num_channels) if bias else torch.zeros(num_channels))
-    def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
-        assert hidden_states.dim() >= 3, f"Input must have at least 3 dimensions, got {hidden_states.dim()}"
-        assert hidden_states.size(1) == self.num_channels, (
-            f"Input must have {self.num_channels} channels, got {hidden_states.size(1)}"
-        )
-        return LigerGroupNormFunction.apply(
-            hidden_states,
-            self.weight,
-            self.bias,
-            self.num_channels,
-            self.num_groups,
-            self.variance_epsilon,
-        )
-    def extra_repr(self) -> str:
-        return f"num_channels={self.num_channels}, num_groups={self.num_groups}, eps={self.variance_epsilon}"
-class LigerDyT(nn.Module):
-    def __init__(self, hidden_size: int, beta: bool = True, init_alpha: float = 0.5):
-        super().__init__()
-        self.hidden_size = hidden_size
-        self.init_alpha = init_alpha
-        self.alpha = nn.Parameter(torch.ones(1) * init_alpha)
-        self.gamma = nn.Parameter(torch.ones(hidden_size))
-        self.beta = nn.Parameter(torch.zeros(hidden_size)) if beta else None
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return LigerDyTFunction.apply(x, self.alpha, self.gamma, self.beta)
-    def extra_repr(self) -> str:
-        return f"{self.hidden_size}, init_alpha={self.init_alpha}, beta={self.beta is not None}"
-class LigerCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.weight = weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-    def forward(self, _input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
-        loss, _, _, _ = LigerCrossEntropyFunction.apply(
-            _input,
-            target,
-            self.weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            False,
-            False,
-        )
-        return loss
-class LigerFusedLinearCrossEntropyLoss(nn.Module):
-    def __init__(
-        self,
-        ce_weight: Optional[torch.Tensor] = None,
-        ignore_index: int = -100,
-        lse_square_scale: float = 0.0,
-        label_smoothing: float = 0.0,
-        reduction: str = "mean",
-        softcap: Optional[float] = None,
-        accum_dtype: Optional[torch.dtype] = None,
-        use_token_scaling: bool = False,
-    ):
-        super().__init__()
-        assert 0.0 <= label_smoothing <= 1.0, f"label_smoothing must be in [0, 1], got {label_smoothing}"
-        assert reduction in ("mean", "sum", "none"), f"reduction must be 'mean', 'sum', or 'none', got {reduction}"
-        assert softcap is None or softcap > 0, f"softcap must be > 0 or None, got {softcap}"
-        self.ce_weight = ce_weight
-        self.ignore_index = ignore_index
-        self.lse_square_scale = lse_square_scale
-        self.label_smoothing = label_smoothing
-        self.reduction = reduction
-        self.softcap = softcap
-        self.accum_dtype = accum_dtype
-        self.use_token_scaling = use_token_scaling
-    def forward(
-        self,
-        lin_weight: torch.Tensor,
-        _input: torch.Tensor,
-        target: torch.Tensor,
-        bias: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        loss, _, _, _ = LigerFusedLinearCrossEntropyFunction.apply(
-            _input,
-            lin_weight,
-            target,
-            bias,
-            self.ce_weight,
-            self.ignore_index,
-            self.lse_square_scale,
-            self.label_smoothing,
-            self.reduction,
-            self.softcap,
-            False,
-            self.accum_dtype,
-            self.use_token_scaling,
-            False,
-            False,
-        )
-        return loss
-class LigerJSD(nn.Module):
-    def __init__(self, beta: float = 0.5, ignore_index: int = -100):
-        super().__init__()
-        self.beta = beta
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        log_q: torch.Tensor,
-        log_p: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerJSDFunction.apply(log_q, log_p, shift_labels, self.beta, self.ignore_index)
-class LigerKLDIVLoss(nn.KLDivLoss):
-    def __init__(self, eps: float = 1e-10, *args, **kwargs):
-        super().__init__(*args, **kwargs)
-        self.eps = eps
-    def forward(self, y_pred: torch.Tensor, y_true: torch.Tensor) -> torch.Tensor:
-        return LigerKLDivLossFunction.apply(y_pred, y_true, self.reduction, self.log_target, self.eps)
-class LigerTVDLoss(nn.Module):
-    def __init__(self, reduction: str = "batchmean", ignore_index: int = -100):
-        super().__init__()
-        self.reduction = reduction
-        self.ignore_index = ignore_index
-    def forward(
-        self,
-        p: torch.Tensor,
-        q: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-    ) -> torch.Tensor:
-        return LigerTVDLossFunction.apply(p, q, shift_labels, self.reduction, self.ignore_index)
-class LigerSwiGLUMLP(nn.Module):
-    """SwiGLU MLP block. ``config`` must expose ``hidden_size``, ``intermediate_size``,
-    and ``hidden_act`` (must be ``silu`` or ``swish``)."""
-    def __init__(self, config):
-        super().__init__()
-        if config.hidden_act not in ("silu", "swish"):
-            raise ValueError(f"Activation function {config.hidden_act} not supported.")
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.down_proj(LigerSiLUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
-class LigerGEGLUMLP(nn.Module):
-    """GEGLU MLP block. ``config`` must expose ``hidden_size`` and ``intermediate_size``.
-    Uses the tanh approximation of GELU (matches Gemma 1/1.1/2)."""
-    def __init__(self, config):
-        super().__init__()
-        self.config = config
-        self.hidden_size = config.hidden_size
-        self.intermediate_size = config.intermediate_size
-        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
-        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.down_proj(LigerGELUMulFunction.apply(self.gate_proj(x), self.up_proj(x)))
-class LigerTiledGEGLUMLP(nn.Module):
-    gate_proj: nn.Linear
-    up_proj: nn.Linear
-    down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerGELUMulFunction.apply(gate, up))
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        compute_params = [p for p in self.parameters() if p.requires_grad]
-        return apply_tiled_mlp(
-            fn=self._mlp_forward,
-            mlp_module=self,
-            x=x,
-            num_shards=self.num_shards,
-            compute_params=compute_params,
-        )
-class LigerTiledSwiGLUMLP(nn.Module):
-    gate_proj: nn.Linear
-    up_proj: nn.Linear
-    down_proj: nn.Linear
-    num_shards: int
-    def _mlp_forward(self, module, x):
-        """Internal MLP forward function for tiled computation."""
-        gate = module.gate_proj(x)
-        up = module.up_proj(x)
-        return module.down_proj(LigerSiLUMulFunction.apply(gate, up))
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        compute_params = [p for p in self.parameters() if p.requires_grad]
-        return apply_tiled_mlp(
-            fn=self._mlp_forward,
-            mlp_module=self,
-            x=x,
-            num_shards=self.num_shards,
-            compute_params=compute_params,
-        )
-@dataclass
-class CrossEntropyOutput:
-    loss: torch.Tensor
-    z_loss: Optional[torch.Tensor] = None
-    token_accuracy: Optional[torch.Tensor] = None
-    predicted_tokens: Optional[torch.Tensor] = None
-def liger_fused_linear_cross_entropy(
-    input: torch.Tensor,
-    weight: torch.Tensor,
-    target: torch.Tensor,
-    bias: Optional[torch.Tensor] = None,
-    ce_weight: Optional[torch.Tensor] = None,
-    ignore_index: int = -100,
-    lse_square_scale: float = 0.0,
-    label_smoothing: float = 0.0,
-    reduction: str = "mean",
-    softcap: Optional[float] = None,
-    return_z_loss: bool = False,
-    accum_dtype: Optional[torch.dtype] = None,
-    use_token_scaling: bool = False,
-    return_token_accuracy: bool = False,
-    return_predicted_tokens: bool = False,
-):
-    loss, z_loss, token_accuracy, predicted_tokens = LigerFusedLinearCrossEntropyFunction.apply(
-        input,
-        weight,
-        target,
-        bias,
-        ce_weight,
-        ignore_index,
-        lse_square_scale,
-        label_smoothing,
-        reduction,
-        softcap,
-        return_z_loss,
-        accum_dtype,
-        use_token_scaling,
-        return_token_accuracy,
-        return_predicted_tokens,
-    )
-    if not return_z_loss and not return_token_accuracy and not return_predicted_tokens:
-        return loss
-    return CrossEntropyOutput(
-        loss=loss,
-        z_loss=z_loss,
-        token_accuracy=token_accuracy,
-        predicted_tokens=predicted_tokens,
-    )
-def LigerForCausalLMLoss(
-    hidden_states: torch.Tensor,
-    lm_head_weight: torch.Tensor,
-    labels: torch.Tensor,
-    hidden_size: int,
-    num_items_in_batch: Optional[int] = None,
-    ignore_index: int = -100,
-    shift_labels: Optional[torch.Tensor] = None,
-    final_logit_softcapping: Optional[float] = None,
-    return_token_accuracy: bool = False,
-    return_predicted_tokens: bool = False,
-    **kwargs,
-):
-    """Drop-in replacement for ``transformers.loss.ForCausalLMLoss`` that fuses the
-    final ``lm_head`` projection with the cross-entropy loss. Returns a scalar
-    ``loss`` by default; returns a :class:`CrossEntropyOutput` when
-    ``return_token_accuracy`` or ``return_predicted_tokens`` is set."""
-    applicable_params = inspect.signature(liger_fused_linear_cross_entropy).parameters
-    kwargs = {k: v for k, v in kwargs.items() if k in applicable_params}
-    if shift_labels is None:
-        labels = nn.functional.pad(labels, (0, 1), value=ignore_index)
-        shift_labels = labels[..., 1:].contiguous()
-    hidden_states = hidden_states.view(-1, hidden_size)
-    shift_labels = shift_labels.view(-1).to(hidden_states.device)
-    reduction = "sum" if num_items_in_batch is not None else "mean"
-    result = liger_fused_linear_cross_entropy(
-        hidden_states,
-        lm_head_weight,
-        shift_labels,
-        reduction=reduction,
-        ignore_index=ignore_index,
-        softcap=final_logit_softcapping,
-        return_token_accuracy=return_token_accuracy,
-        return_predicted_tokens=return_predicted_tokens,
-        **kwargs,
-    )
-    if isinstance(result, CrossEntropyOutput):
-        loss = result.loss
-        token_accuracy = result.token_accuracy
-        predicted_tokens = result.predicted_tokens
-    else:
-        loss = result
-        token_accuracy = None
-        predicted_tokens = None
-    if reduction == "sum":
-        loss = loss / num_items_in_batch
-    if return_token_accuracy or return_predicted_tokens:
-        return CrossEntropyOutput(
-            loss=loss,
-            token_accuracy=token_accuracy,
-            predicted_tokens=predicted_tokens,
-        )
-    return loss
-def liger_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    position_ids: Optional[torch.Tensor] = None,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply standard rotary positional embedding to ``q`` and ``k``."""
-    return LigerRopeFunction.apply(q, k, cos, sin, position_ids, unsqueeze_dim)
-def liger_multimodal_rotary_pos_emb(
-    q: torch.Tensor,
-    k: torch.Tensor,
-    cos: torch.Tensor,
-    sin: torch.Tensor,
-    mrope_section,
-    unsqueeze_dim: int = 1,
-) -> Tuple[torch.Tensor, torch.Tensor]:
-    """Apply Qwen2-VL multimodal rotary positional embedding (M-RoPE) to ``q`` and ``k``."""
-    return LigerQwen2VLMRopeFunction.apply(q, k, cos, sin, mrope_section, unsqueeze_dim)
-__all__ = [
-    "LigerRMSNorm",
-    "LigerLayerNorm",
-    "LigerGroupNorm",
-    "LigerDyT",
-    "LigerCrossEntropyLoss",
-    "LigerFusedLinearCrossEntropyLoss",
-    "LigerJSD",
-    "LigerKLDIVLoss",
-    "LigerTVDLoss",
-    "LigerSwiGLUMLP",
-    "LigerGEGLUMLP",
-    "LigerTiledGEGLUMLP",
-    "LigerTiledSwiGLUMLP",
-    "CrossEntropyOutput",
-    "liger_fused_linear_cross_entropy",
-    "LigerForCausalLMLoss",
-    "liger_rotary_pos_emb",
-    "liger_multimodal_rotary_pos_emb",
-]

build/torch-xpu/liger_kernels/__init__.py DELETED Viewed

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch-xpu/metadata.json DELETED Viewed

@@ -1,10 +0,0 @@
-{
-  "name": "liger-kernels",
-  "id": "_liger_kernels_xpu_ab435e2",
-  "version": 1,
-  "license": "BSD-2-Clause",
-  "python-depends": [],
-  "backend": {
-    "type": "xpu"
-  }
-}

build/torch-xpu/qwen2vl_mrope.py DELETED Viewed

@@ -1,222 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-@triton.jit
-def _triton_qwen2vl_mrope(
-    q_ptr,
-    k_ptr,
-    cos,
-    sin,
-    sl,
-    bs: tl.constexpr,
-    n_qh: tl.constexpr,
-    n_kh: tl.constexpr,
-    hd: tl.constexpr,
-    pad_n_qh: tl.constexpr,
-    pad_n_kh: tl.constexpr,
-    pad_hd: tl.constexpr,
-    mrope_section_t: tl.constexpr,
-    mrope_section_h: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BACKWARD_PASS: tl.constexpr = False,
-):
-    pid = tl.program_id(0)
-    # locate start address
-    q_ptr = q_ptr + pid * (n_qh * hd)
-    k_ptr = k_ptr + pid * (n_kh * hd)
-    # ####################################################################
-    # get the cos(mθ_{i...d/2}) and sin(mθ_{i...d/2}) for token position
-    # m of this program instance
-    # ####################################################################
-    # 1. program instances are laid out in a 1D vector of size bsz * seq_len, which
-    # effectively represents a 2D grid of size [bsz, seq_len] with seq_len dimension
-    # being the fastest changing dimension. Thus we can simply do pid // sl to get the batch index
-    # and pid % sl to get the sequence index.
-    # 2. We only need the left half of cos and sin matrix because the right half is just
-    # a clone of the left half.
-    t_end = mrope_section_t
-    h_end = t_end + mrope_section_h
-    t_cos = cos + pid * hd
-    h_cos = t_cos + bs * sl * hd
-    w_cos = h_cos + bs * sl * hd
-    t_sin = sin + pid * hd
-    h_sin = t_sin + bs * sl * hd
-    w_sin = h_sin + bs * sl * hd
-    cos_offsets = tl.arange(0, pad_hd // 2)
-    t_mask = cos_offsets < t_end
-    h_mask = (t_end <= cos_offsets) & (cos_offsets < h_end)
-    w_mask = (h_end <= cos_offsets) & (cos_offsets < hd // 2)
-    t_cos_row = tl.load(t_cos + cos_offsets, mask=t_mask, other=0)
-    h_cos_row = tl.load(h_cos + cos_offsets, mask=h_mask, other=0)
-    w_cos_row = tl.load(w_cos + cos_offsets, mask=w_mask, other=0)
-    t_sin_row = tl.load(t_sin + cos_offsets, mask=t_mask, other=0)
-    h_sin_row = tl.load(h_sin + cos_offsets, mask=h_mask, other=0)
-    w_sin_row = tl.load(w_sin + cos_offsets, mask=w_mask, other=0)
-    cos_row = t_cos_row + h_cos_row + w_cos_row
-    sin_row = t_sin_row + h_sin_row + w_sin_row
-    # ####################################################################
-    # Load the left and right half of q and k for the current
-    # program instance (i.e. for the current token) separately
-    # ####################################################################
-    # left half of the head
-    first_half_q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_half_k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    q_tile_1 = tl.load(q_ptr + first_half_q_offsets, mask=first_q_mask, other=0).to(sin_row.dtype)
-    k_tile_1 = tl.load(k_ptr + first_half_k_offsets, mask=first_k_mask, other=0).to(sin_row.dtype)
-    # right half of the head
-    second_half_q_offsets = first_half_q_offsets + (hd // 2)
-    second_half_k_offsets = first_half_k_offsets + (hd // 2)
-    second_q_mask = first_q_mask
-    second_k_mask = first_k_mask
-    q_tile_2 = tl.load(q_ptr + second_half_q_offsets, mask=second_q_mask, other=0).to(sin_row.dtype)
-    k_tile_2 = tl.load(k_ptr + second_half_k_offsets, mask=second_k_mask, other=0).to(sin_row.dtype)
-    if not BACKWARD_PASS:
-        # y = [x1, x2] * [cos, cos] + [-x2, x1] * [sin, sin]
-        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-    else:
-        # with some math, we can get:
-        # dy = [dx1, dx2] * [cos, cos] + [-dx2, dx1] * [-sin, -sin]
-        new_q_tile_1 = q_tile_1 * cos_row + q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row - q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row + k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row - k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-def qwen2vl_mrope_forward(q, k, cos, sin, mrope_section):
-    # transpose it back to the physical shape because Triton looks at the physical storage
-    # note: q and k are incontiguous before the transformation and will become contiguous after transpose
-    q = q.transpose(1, 2)
-    k = k.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = q.shape
-    n_kv_head = k.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure tensors passed into the kernel are contiguous. It will be no-op if they are already contiguous
-    q = q.contiguous()
-    k = k.contiguous()
-    cos = cos.contiguous()
-    sin = sin.contiguous()
-    _triton_qwen2vl_mrope[(n_row,)](
-        q,
-        k,
-        cos,
-        sin,
-        seq_len,
-        batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        mrope_section[0],
-        mrope_section[1],
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=False,
-    )
-    return q.transpose(1, 2), k.transpose(1, 2), cos, sin
-def qwen2vl_mrope_backward(dq, dk, cos, sin, mrope_section):
-    dq = dq.transpose(1, 2)
-    dk = dk.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = dq.shape
-    n_kv_head = dk.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure dq and dk are contiguous
-    dq = dq.contiguous()
-    dk = dk.contiguous()
-    # backward is similar to forward except swapping few ops
-    _triton_qwen2vl_mrope[(n_row,)](
-        dq,
-        dk,
-        cos,
-        sin,
-        seq_len,
-        batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        mrope_section[0],
-        mrope_section[1],
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=True,
-    )
-    return dq.transpose(1, 2), dk.transpose(1, 2)
-class LigerQwen2VLMRopeFunction(torch.autograd.Function):
-    """
-    Triton implementation of the Qwen2VL Multimodal Rotary Positional Embedding (M-RoPE) operation.
-    Please find the corresponding HuggingFace implementation here:
-    https://github.com/huggingface/transformers/blob/main/src/transformers/models/qwen2_vl/modeling_qwen2_vl.py
-    """
-    @staticmethod
-    def forward(ctx, q, k, cos, sin, mrope_section, unsqueeze_dim=1):
-        """
-        q size: (bsz, n_q_head, seq_len, head_dim)
-        k size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (3, bsz, seq_len, head_dim)
-        sin size: (3, bsz, seq_len, head_dim)
-        """
-        q, k, cos, sin = qwen2vl_mrope_forward(q, k, cos, sin, mrope_section)
-        ctx.save_for_backward(cos, sin)
-        ctx.mrope_section = mrope_section
-        return q, k
-    def backward(ctx, dq, dk):
-        """
-        dq size: (bsz, n_q_head, seq_len, head_dim)
-        dk size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (3, bsz, seq_len, head_dim)
-        sin size: (3, bsz, seq_len, head_dim)
-        """
-        cos, sin = ctx.saved_tensors
-        mrope_section = ctx.mrope_section
-        dq, dk = qwen2vl_mrope_backward(dq, dk, cos, sin, mrope_section)
-        return dq, dk, None, None, None, None

build/torch-xpu/rms_norm.py DELETED Viewed

@@ -1,654 +0,0 @@
-"""
-This file incorporates code from Unsloth licensed under the Apache License, Version 2.0.
-See the original Unsloth repository at https://github.com/unslothai/unsloth.
-The following line
-https://github.com/linkedin/Liger-Kernel/blob/7382a8761f9af679482b968f9348013d933947c7/src/liger_kernel/ops/rms_norm.py#L30
-is based on code from Unsloth, located at:
-https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-Modifications made by Yanning Chen, 2024.
-"""
-import math
-import operator
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import compare_version
-from .utils import ensure_contiguous
-from .utils import get_npu_core_count
-from .utils import set_large_grf_mode
-from .utils import torch_to_triton_dtype
-from .utils import is_npu_available
-if compare_version("triton", operator.ge, "3.0.0") and not is_npu_available():
-    try:
-        # typical import path with dispatch available
-        from triton.language.extra.libdevice import rsqrt
-    except ModuleNotFoundError:
-        # for working with NGC containers
-        from triton.language.extra.cuda.libdevice import rsqrt
-else:
-    from triton.language.math import rsqrt
-_CASTING_MODE_NONE: tl.constexpr = tl.constexpr(-1)
-_CASTING_MODE_LLAMA: tl.constexpr = tl.constexpr(0)
-_CASTING_MODE_GEMMA: tl.constexpr = tl.constexpr(1)
-@triton.jit
-def _rms_norm_forward_kernel(
-    Y_ptr,
-    Y_row_stride,
-    X_ptr,
-    X_row_stride,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    n_cols,
-    eps,
-    offset,
-    casting_mode: tl.constexpr,  # constexpr so the `if` blocks can be optimized out
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    y_i = (x_i / (RMS)) * (offset + wi), RMS = sqrt(sum(x_i^2) / N)
-    Reference:
-    1. https://triton-lang.org/main/getting-started/tutorials/05-layer-norm.html
-    2. https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-    3. https://arxiv.org/pdf/1910.07467
-    """
-    row_idx = tl.program_id(0).to(tl.int64)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    y_base = Y_ptr + row_idx * Y_row_stride
-    x_base = X_ptr + row_idx * X_row_stride
-    rstd_base = RSTD_ptr + row_idx * RSTD_row_stride
-    X_row = tl.load(x_base + col_offsets, mask=mask, other=0)
-    X_row_dtype = X_row.dtype
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0)
-    # On Llama, only rstd is computed on fp32
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(tl.float32)
-    # Gemma computes everything on fp32, and then casts back the output to the original dtype
-    if casting_mode == _CASTING_MODE_GEMMA:
-        if elementwise_affine:
-            W_row = W_row.to(tl.float32)
-        X_row = X_row.to(tl.float32)
-    if casting_mode == _CASTING_MODE_NONE:
-        eps = eps.to(X_row_dtype)
-        offset = offset.to(X_row_dtype)
-    mean_square = tl.sum(X_row * X_row, axis=0) / n_cols
-    rstd = rsqrt(mean_square + eps)
-    # We can save time by caching rms with minimal memory overhead
-    # because rms is much smaller compared to X_row, as rms is for each row.
-    # However, on the computation side, it can save 4 operations (*, sum, /, sqrt).
-    tl.store(rstd_base, rstd)
-    X_row = X_row * rstd
-    # On Llama, the multiplication with the weight is done on the original dtype
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(X_row_dtype)
-    if elementwise_affine:
-        Y_row = X_row * (offset + W_row)
-    else:
-        Y_row = X_row
-    if casting_mode == _CASTING_MODE_GEMMA:
-        Y_row = Y_row.to(X_row_dtype)
-    tl.store(y_base + col_offsets, Y_row, mask=mask)
-@triton.jit
-def _rms_norm_backward_kernel(
-    dY_ptr,
-    dY_row_stride,
-    dX_ptr,
-    dX_row_stride,
-    X_ptr,
-    X_row_stride,
-    X_dtype: tl.constexpr,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    dW_ptr,
-    dW_row_stride,
-    n_rows,
-    n_cols,
-    offset,
-    rows_per_program,
-    casting_mode: tl.constexpr,
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    dx = (1 / RMS) * [dy * (w + offset - (1 / N) * (1 / RMS^2) * ((dy * (w + offset)) dot x) * x]. * means element-wise multiplication, whileas dot means dot product
-    dw = sum(dy * (x / RMS)). summation over BxT dimension
-    """
-    row_block_id = tl.program_id(0).to(tl.int64)
-    row_start = row_block_id * rows_per_program
-    row_end = min((row_block_id + 1) * rows_per_program, n_rows)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    if elementwise_affine:
-        dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0.0)
-        W_row = W_row + offset
-    for row_idx in range(row_start, row_end):
-        dy_base = dY_ptr + row_idx * dY_row_stride
-        dx_base = dX_ptr + row_idx * dX_row_stride
-        x_base = X_ptr + row_idx * X_row_stride
-        rstd_base = RSTD_ptr + row_idx * RSTD_row_stride
-        dY_row = tl.load(dy_base + col_offsets, mask=mask, other=0.0)
-        X_row = tl.load(x_base + col_offsets, mask=mask, other=0.0)
-        # Get cached rms
-        rstd_row = tl.load(rstd_base)
-        X_row = X_row.to(tl.float32)
-        # Different bacward graphs for different casting modes
-        if casting_mode == _CASTING_MODE_LLAMA:
-            if elementwise_affine:
-                m = (dY_row * W_row).to(tl.float32)
-            else:
-                m = dY_row.to(tl.float32)
-        elif casting_mode == _CASTING_MODE_GEMMA:
-            dY_row = dY_row.to(tl.float32)
-            if elementwise_affine:
-                m = dY_row * W_row
-            else:
-                m = dY_row
-        else:
-            if elementwise_affine:
-                m = dY_row * W_row
-            else:
-                m = dY_row
-        dX_row = rstd_row * m
-        dX_row += (rstd_row) * (-(1 / n_cols) * rstd_row * rstd_row * tl.sum(m * X_row, axis=0) * X_row)
-        if elementwise_affine:
-            # calculate the gradient of W
-            if casting_mode == _CASTING_MODE_LLAMA:
-                dW_row += dY_row * (X_row * rstd_row).to(X_dtype)
-            else:
-                # here X_row is already in fp32 (see previous if block)
-                dW_row += dY_row * (X_row * rstd_row)
-        tl.store(dx_base + col_offsets, dX_row.to(X_dtype), mask=mask)
-    if elementwise_affine:
-        tl.store(dW_ptr + row_block_id * dW_row_stride + col_offsets, dW_row, mask=mask)
-@triton.jit
-def _block_rms_norm_forward_kernel(
-    Y_ptr,
-    Y_row_stride,
-    X_ptr,
-    X_row_stride,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    n_rows,
-    n_cols,
-    eps,
-    offset,
-    casting_mode: tl.constexpr,  # constexpr so the `if` blocks can be optimized out
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BLOCK_ROW: tl.constexpr,
-):
-    """
-    y_i = (x_i / (RMS)) * (offset + wi), RMS = sqrt(sum(x_i^2) / N)
-    Reference:
-    1. https://triton-lang.org/main/getting-started/tutorials/05-layer-norm.html
-    2. https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/rms_layernorm.py#L22
-    3. https://arxiv.org/pdf/1910.07467
-    """
-    row_idx = tl.program_id(0) * BLOCK_ROW + tl.arange(0, BLOCK_ROW)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    row_mask = row_idx < n_rows
-    col_mask = col_offsets < n_cols
-    X_row = tl.load(
-        X_ptr + row_idx[:, None] * X_row_stride + col_offsets[None, :],
-        mask=row_mask[:, None] & col_mask[None, :],
-        other=0,
-    )
-    X_row_dtype = X_row.dtype
-    if elementwise_affine:
-        W_row = tl.load(W_ptr + col_offsets, mask=col_mask, other=0)
-    # On Llama, only rstd is computed on fp32
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(tl.float32)
-    # Gemma computes everything on fp32, and then casts back the output to the original dtype
-    if casting_mode == _CASTING_MODE_GEMMA:
-        if elementwise_affine:
-            W_row = W_row.to(tl.float32)
-        X_row = X_row.to(tl.float32)
-    if casting_mode == _CASTING_MODE_NONE:
-        eps = eps.to(X_row_dtype)
-        offset = offset.to(X_row_dtype)
-    mean_square = tl.sum(X_row * X_row, axis=1) / n_cols
-    rstd = rsqrt(mean_square + eps)
-    # We can save time by caching rms with minimal memory overhead
-    # because rms is much smaller compared to X_row, as rms is for each row.
-    # However, on the computation side, it can save 4 operations (*, sum, /, sqrt).
-    tl.store(RSTD_ptr + row_idx * RSTD_row_stride, rstd, row_mask)
-    X_row = X_row * rstd[:, None]
-    # On Llama, the multiplication with the weight is done on the original dtype
-    if casting_mode == _CASTING_MODE_LLAMA:
-        X_row = X_row.to(X_row_dtype)
-    if elementwise_affine:
-        Y_row = X_row * (offset + W_row)[None, :]
-    else:
-        Y_row = X_row
-    if casting_mode == _CASTING_MODE_GEMMA:
-        Y_row = Y_row.to(X_row_dtype)
-    tl.store(
-        Y_ptr + row_idx[:, None] * Y_row_stride + col_offsets[None, :],
-        Y_row,
-        mask=row_mask[:, None] & col_mask[None, :],
-    )
-@triton.jit
-def _block_rms_norm_backward_kernel(
-    dY_ptr,
-    dY_row_stride,
-    dX_ptr,
-    dX_row_stride,
-    X_ptr,
-    X_row_stride,
-    X_dtype: tl.constexpr,
-    W_ptr,
-    W_row_stride,
-    RSTD_ptr,
-    RSTD_row_stride,
-    dW_ptr,
-    dW_row_stride,
-    n_rows,
-    n_cols,
-    offset,
-    casting_mode: tl.constexpr,
-    elementwise_affine: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BLOCK_ROW: tl.constexpr,
-):
-    """
-    dx = (1 / RMS) * [dy * (w + offset - (1 / N) * (1 / RMS^2) * ((dy * (w + offset)) dot x) * x]. * means element-wise multiplication, whileas dot means dot product
-    dw = sum(dy * (x / RMS)). summation over BxT dimension
-    """
-    pid = tl.program_id(0).cast(tl.int64)
-    NUM_SMS = tl.num_programs(0)
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    col_mask = col_offsets < n_cols
-    if elementwise_affine:
-        dW_row = tl.zeros((BLOCK_SIZE,), dtype=tl.float32)
-        W_row = tl.load(W_ptr + col_offsets, mask=col_mask, other=0.0)
-        W_row = W_row + offset
-    for start in range(pid * BLOCK_ROW, n_rows, NUM_SMS * BLOCK_ROW):
-        row_idx = start + tl.arange(0, BLOCK_ROW)
-        row_mask = row_idx < n_rows
-        dY_row = tl.load(
-            dY_ptr + row_idx[:, None] * dY_row_stride + col_offsets[None, :],
-            mask=row_mask[:, None] & col_mask[None, :],
-            other=0.0,
-        )
-        X_row = tl.load(
-            X_ptr + row_idx[:, None] * X_row_stride + col_offsets[None, :],
-            mask=row_mask[:, None] & col_mask[None, :],
-            other=0.0,
-        )
-        # Get cached rms
-        rstd_row = tl.load(RSTD_ptr + row_idx * RSTD_row_stride, row_mask)
-        X_row = X_row.to(tl.float32)
-        # Different bacward graphs for different casting modes
-        if casting_mode == _CASTING_MODE_LLAMA:
-            if elementwise_affine:
-                m = (dY_row * W_row[None, :]).to(tl.float32)
-            else:
-                m = dY_row.to(tl.float32)
-        elif casting_mode == _CASTING_MODE_GEMMA:
-            dY_row = dY_row.to(tl.float32)
-            if elementwise_affine:
-                m = dY_row * W_row[None, :]
-            else:
-                m = dY_row
-        else:
-            if elementwise_affine:
-                m = dY_row * W_row[None, :]
-            else:
-                m = dY_row
-        dX_row = rstd_row[:, None] * m
-        dX_row += (rstd_row[:, None]) * (
-            -(1 / n_cols) * (rstd_row * rstd_row * tl.sum(m * X_row, axis=1))[:, None] * X_row
-        )
-        if elementwise_affine:
-            if casting_mode == _CASTING_MODE_LLAMA:
-                # TODO(tcc): use tl.sum(..., dtype=tl.float32) once we upgrade to triton>=3.3.0
-                dW_row += tl.sum((dY_row * (X_row * rstd_row[:, None]).to(X_dtype)).to(tl.float32), 0)
-            else:
-                # here X_row is already in fp32 (see previous if block)
-                dW_row += tl.sum(dY_row * (X_row * rstd_row[:, None]), 0)
-        tl.store(
-            dX_ptr + row_idx[:, None] * dX_row_stride + col_offsets[None, :],
-            dX_row,
-            mask=row_mask[:, None] & col_mask[None, :],
-        )
-    if elementwise_affine:
-        tl.store(dW_ptr + pid * dW_row_stride + col_offsets, dW_row, mask=col_mask)
-_str_to_casting_mode = {
-    "llama": _CASTING_MODE_LLAMA.value,
-    "gemma": _CASTING_MODE_GEMMA.value,
-    "none": _CASTING_MODE_NONE.value,
-}
-def rms_norm_forward(X, W, eps, offset, casting_mode, row_mode):
-    if not isinstance(casting_mode, int):
-        assert casting_mode in _str_to_casting_mode, f"Invalid casting mode: {casting_mode}"
-        casting_mode = _str_to_casting_mode[casting_mode]
-    else:
-        assert casting_mode in _str_to_casting_mode.values(), f"Invalid casting mode: {casting_mode}"
-    shape = X.shape
-    dim = shape[-1]
-    X = X.view(-1, dim)
-    n_rows, n_cols = X.shape
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    Y = torch.empty((n_rows, n_cols), dtype=X.dtype, device=X.device)
-    # RSTD is to cache rstd for each row
-    # RSTD is always computed/stored in fp32 if we are using Llama or Gemma casting mode
-    rstd_dtype = torch.float32 if casting_mode in (_CASTING_MODE_LLAMA.value, _CASTING_MODE_GEMMA.value) else X.dtype
-    RSTD = torch.empty(n_rows, dtype=rstd_dtype, device=X.device)
-    if W is not None:
-        # Check constraints.
-        assert X.shape[1] == W.shape[0], (
-            "Incompatible hidden size dimension between tensor1.shape[1] and tensor2.shape[0]"
-        )
-        elementwise_affine = True
-    else:
-        elementwise_affine = False
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    if BLOCK_SIZE > 256 or n_rows < 4096 * 8 or row_mode:
-        _rms_norm_forward_kernel[(n_rows,)](
-            Y,
-            Y.stride(0),
-            X,
-            X.stride(0),
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            n_cols,
-            eps,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    else:
-        BLOCK_ROW = 16
-        kernel_args["BLOCK_ROW"] = BLOCK_ROW
-        _block_rms_norm_forward_kernel[(triton.cdiv(n_rows, BLOCK_ROW),)](
-            Y,
-            Y.stride(0),
-            X,
-            X.stride(0),
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            n_rows,
-            n_cols,
-            eps,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    return Y.view(*shape), X, RSTD, BLOCK_SIZE, num_warps, casting_mode
-def rms_norm_backward(dY, X, W, RSTD, offset, casting_mode, BLOCK_SIZE, num_warps, in_place, row_mode):
-    shape = dY.shape
-    dim = shape[-1]
-    dY = dY.view(-1, dim)
-    n_rows, n_cols = dY.shape
-    sm_count = 1
-    if X.device.type == "cuda":
-        sm_count = torch.cuda.get_device_properties(X.device).multi_processor_count
-    elif X.device.type == "xpu":
-        sm_count = torch.xpu.get_device_properties(X.device).gpu_eu_count
-    elif X.device.type == "npu":
-        sm_count = get_npu_core_count()
-    if W is not None:
-        # fp32 for numerical stability especially.
-        _dW = torch.empty((sm_count, n_cols), dtype=torch.float32, device=W.device)
-        elementwise_affine = True
-    else:
-        _dW = None
-        elementwise_affine = False
-    if n_cols > BLOCK_SIZE:
-        raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
-    rows_per_program = math.ceil(n_rows / sm_count)
-    grid = (sm_count,)
-    if in_place is True:
-        dX = dY
-    else:
-        dX = torch.zeros_like(dY)
-    # XPU-specific optimization
-    kernel_args = {}
-    if X.device.type == "xpu":
-        set_large_grf_mode(kernel_args)
-    if BLOCK_SIZE > 256 or n_rows < 4096 * 8 or row_mode:
-        _rms_norm_backward_kernel[grid](
-            dY,
-            dY.stride(0),
-            dX,
-            dX.stride(0),
-            X,
-            X.stride(0),
-            torch_to_triton_dtype[X.dtype],
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            _dW,
-            _dW.stride(0) if elementwise_affine else 0,
-            n_rows,
-            n_cols,
-            offset,
-            rows_per_program,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    else:
-        BLOCK_ROW = 16
-        kernel_args["BLOCK_ROW"] = BLOCK_ROW
-        _block_rms_norm_backward_kernel[grid](
-            dY,
-            dY.stride(0),
-            dX,
-            dX.stride(0),
-            X,
-            X.stride(0),
-            torch_to_triton_dtype[X.dtype],
-            W,
-            W.stride(0) if elementwise_affine else 0,
-            RSTD,
-            RSTD.stride(0),
-            _dW,
-            _dW.stride(0) if elementwise_affine else 0,
-            n_rows,
-            n_cols,
-            offset,
-            casting_mode,
-            elementwise_affine=elementwise_affine,
-            BLOCK_SIZE=BLOCK_SIZE,
-            num_warps=num_warps,
-            **kernel_args,  # XPU-specific optimization
-        )
-    dX = dX.view(*shape)
-    if elementwise_affine:
-        dW = _dW.sum(dim=0).to(W.dtype)
-    else:
-        dW = None
-    return dX, dW
-class LigerRMSNormFunction(torch.autograd.Function):
-    """
-    Performs RMSNorm (Root Mean Square Normalization), which normalizes the input tensor `X` using the
-    weight tensor `W`, with an optional offset and casting mode.
-    Some models use an 'offset' to shift the weight tensor `W` by a constant value. For example, Gemma
-    uses an offset of 1.0, so the computation becomes `(X / RMS(X)) * (W + 1.0)` instead of the usual
-    `(X / RMS(X)) * W`. You can pass the offset value as an argument to the forward function.
-    In addition, different models cast their inputs at different places during RMSNorm computation. For
-    example, Gemma casts everything to fp32 nefore starting the computation, while Llama casts only the
-    inverse RMS to fp32. You can specify the casting mode using the `casting_mode` argument. We currently
-    support the following casting modes (they match HuggingFace Transformers' implementations):
-    - 'llama': matches the Llama implementation, where only the inverse RMS is computed on fp32.
-    - 'gemma': matches the Gemma implementation, where everything is cast to fp32, then computed, then cast back to the original dtype.
-    - 'none': no casting is done. The computation is done in the original dtype. This saves memory and is slightly faster, but has more error w.r.t. the original implementation.
-    `in_place` option means whether to in_place modify dY to store dX. This is default to `True` to save memory. However, under certain cases, it can produce incorrect inputs.
-        For example, gemma2 uses two rmsnorm sequentially with residual in between. The resesidual part needs dY so it cannot be modified in-place.
-        Therefore, for the patching of RMSNorm in gemma2, we set `in_place` to `False`
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, X, W, eps, offset=0.0, casting_mode="llama", in_place=True, row_mode=None):
-        """
-        X: (B, T, H) or (BxT, H)
-        W: (H,)
-        """
-        if isinstance(X, torch.distributed.tensor.DTensor):
-            # Input tensor is output of a tensor parallel module and
-            # needs to be gathered to a local tensor to compute
-            # RMSE layer norm on each TP worker.
-            # TODO: support CP.
-            X = X.full_tensor()
-        Y, X, RSTD, BLOCK_SIZE, num_warps, casting_mode = rms_norm_forward(X, W, eps, offset, casting_mode, row_mode)
-        ctx.offset = offset
-        ctx.casting_mode = casting_mode
-        ctx.in_place = in_place
-        ctx.row_mode = row_mode
-        ctx.BLOCK_SIZE = BLOCK_SIZE
-        ctx.num_warps = num_warps
-        ctx.elementwise_affine = W is not None
-        if W is not None:
-            ctx.save_for_backward(X, W, RSTD)
-        else:
-            ctx.save_for_backward(X, RSTD)
-        return Y
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dY):
-        """
-        Y: (B, T, H) or (BxT, H)
-        """
-        if ctx.elementwise_affine:
-            X, W, RSTD = ctx.saved_tensors
-        else:
-            X, RSTD = ctx.saved_tensors
-            W = None
-        if isinstance(dY, torch.distributed.tensor.DTensor):
-            # Gradients are output of a tensor parallel module and
-            # needs to be gathered to a local tensor for computing RMSE layer.
-            # TODO: support CP.
-            dY = dY.full_tensor()
-        dX, dW = rms_norm_backward(
-            dY, X, W, RSTD, ctx.offset, ctx.casting_mode, ctx.BLOCK_SIZE, ctx.num_warps, ctx.in_place, ctx.row_mode
-        )
-        return dX, dW, None, None, None, None, None

build/torch-xpu/rope.py DELETED Viewed

@@ -1,239 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-@triton.jit
-def _triton_rope(
-    q_ptr,
-    q_row_stride,
-    k_ptr,
-    k_row_stride,
-    cos,
-    cos_row_stride,
-    sin,
-    sin_row_stride,
-    sl,
-    bs: tl.constexpr,
-    cos_bs: tl.constexpr,
-    n_qh: tl.constexpr,
-    n_kh: tl.constexpr,
-    hd: tl.constexpr,
-    pad_n_qh: tl.constexpr,
-    pad_n_kh: tl.constexpr,
-    pad_hd: tl.constexpr,
-    BLOCK_SIZE: tl.constexpr,
-    BACKWARD_PASS: tl.constexpr = False,
-):
-    # q size: (bsz, seq_len, num_q_heads, head_dim)
-    # q stride: (seq_len * num_q_heads * head_dim, num_q_heads * head_dim, head_dim, 1)
-    # k size: (bsz, seq_len, num_kv_heads, head_dim)
-    # k stride: (seq_len * num_kv_heads * head_dim, num_kv_heads * head_dim, head_dim, 1)
-    # cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-    # stride: (seq_len * head_dim, head_dim, 1)
-    pid = tl.program_id(0).to(tl.int64)
-    # locate start address
-    q_ptr = q_ptr + pid * q_row_stride
-    k_ptr = k_ptr + pid * k_row_stride
-    # ####################################################################
-    # get the cos(mθ_{i...d/2}) and sin(mθ_{i...d/2}) for token position
-    # m of this program instance
-    # ####################################################################
-    # 1. program instances are laid out in a 1D vector of size bsz * seq_len, which
-    # effectively represents a 2D grid of size [bsz, seq_len] with seq_len dimension
-    # being the fastest changing dimension. Thus we can simply do pid // sl to get the batch index
-    # and pid % sl to get the sequence index.
-    # 2. We only need the left half of cos and sin matrix because the right half is just
-    # a clone of the left half.
-    batch_idx = pid // sl
-    cos_row_idx = pid % sl
-    cos = cos + tl.where(
-        cos_bs == 1,
-        cos_row_idx * cos_row_stride,
-        batch_idx * (sl * cos_row_stride) + cos_row_idx * cos_row_stride,
-    )
-    sin = sin + tl.where(
-        cos_bs == 1,
-        cos_row_idx * sin_row_stride,
-        batch_idx * (sl * sin_row_stride) + cos_row_idx * sin_row_stride,
-    )
-    cos_offsets = tl.arange(0, pad_hd // 2)
-    cos_mask = cos_offsets < hd // 2
-    cos_row = tl.load(cos + cos_offsets, mask=cos_mask, other=0)
-    sin_row = tl.load(sin + cos_offsets, mask=cos_mask, other=0)
-    # ####################################################################
-    # Load the left and right half of q and k for the current
-    # program instance (i.e. for the current token) separately
-    # ####################################################################
-    # left half of the head
-    first_half_q_offsets = tl.arange(0, pad_n_qh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_half_k_offsets = tl.arange(0, pad_n_kh)[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
-    first_q_mask = (tl.arange(0, pad_n_qh)[:, None] < n_qh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    first_k_mask = (tl.arange(0, pad_n_kh)[:, None] < n_kh) & (tl.arange(0, pad_hd // 2)[None, :] < hd // 2)
-    q_tile_1 = tl.load(q_ptr + first_half_q_offsets, mask=first_q_mask, other=0).to(sin_row.dtype)
-    k_tile_1 = tl.load(k_ptr + first_half_k_offsets, mask=first_k_mask, other=0).to(sin_row.dtype)
-    # right half of the head
-    second_half_q_offsets = first_half_q_offsets + (hd // 2)
-    second_half_k_offsets = first_half_k_offsets + (hd // 2)
-    second_q_mask = first_q_mask
-    second_k_mask = first_k_mask
-    q_tile_2 = tl.load(q_ptr + second_half_q_offsets, mask=second_q_mask, other=0).to(sin_row.dtype)
-    k_tile_2 = tl.load(k_ptr + second_half_k_offsets, mask=second_k_mask, other=0).to(sin_row.dtype)
-    if not BACKWARD_PASS:
-        # y = [x1, x2] * [cos, cos] + [-x2, x1] * [sin, sin]
-        new_q_tile_1 = q_tile_1 * cos_row - q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row + q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row - k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row + k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-    else:
-        # with some math, we can get:
-        # dy = [dx1, dx2] * [cos, cos] + [-dx2, dx1] * [-sin, -sin]
-        new_q_tile_1 = q_tile_1 * cos_row + q_tile_2 * sin_row
-        tl.store(q_ptr + first_half_q_offsets, new_q_tile_1, mask=first_q_mask)
-        new_q_tile_2 = q_tile_2 * cos_row - q_tile_1 * sin_row
-        tl.store(q_ptr + second_half_q_offsets, new_q_tile_2, mask=second_q_mask)
-        new_k_tile_1 = k_tile_1 * cos_row + k_tile_2 * sin_row
-        tl.store(k_ptr + first_half_k_offsets, new_k_tile_1, mask=first_k_mask)
-        new_k_tile_2 = k_tile_2 * cos_row - k_tile_1 * sin_row
-        tl.store(k_ptr + second_half_k_offsets, new_k_tile_2, mask=second_k_mask)
-def rope_forward(q, k, cos, sin):
-    # transpose it back to the physical shape because Triton looks at the physical storage
-    # note: q and k are incontiguous before the transformation and will become contiguous after transpose
-    q = q.transpose(1, 2)
-    k = k.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = q.shape
-    n_kv_head = k.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure tensors passed into the kernel are contiguous. It will be no-op if they are already contiguous
-    q = q.contiguous()
-    k = k.contiguous()
-    cos = cos.contiguous()
-    sin = sin.contiguous()
-    cos_batch_size = cos.shape[0]
-    _triton_rope[(n_row,)](
-        q,
-        q.stride(1),
-        k,
-        k.stride(1),
-        cos,
-        cos.stride(-2),
-        sin,
-        sin.stride(-2),
-        seq_len,
-        batch_size,
-        cos_batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=False,
-    )
-    return q.transpose(1, 2), k.transpose(1, 2), cos, sin
-def rope_backward(dq, dk, cos, sin):
-    dq = dq.transpose(1, 2)
-    dk = dk.transpose(1, 2)
-    batch_size, seq_len, n_q_head, head_dim = dq.shape
-    cos_batch_size = cos.shape[0]
-    n_kv_head = dk.shape[2]
-    pad_hd = triton.next_power_of_2(head_dim)
-    pad_n_q_head = triton.next_power_of_2(n_q_head)
-    pad_n_kv_head = triton.next_power_of_2(n_kv_head)
-    BLOCK_SIZE = max(pad_n_q_head, pad_n_kv_head)
-    n_row = batch_size * seq_len
-    # ensure dq and dk are contiguous
-    dq = dq.contiguous()
-    dk = dk.contiguous()
-    # backward is similar to forward except swapping few ops
-    _triton_rope[(n_row,)](
-        dq,
-        dq.stride(1),
-        dk,
-        dk.stride(1),
-        cos,
-        cos.stride(-2),
-        sin,
-        sin.stride(-2),
-        seq_len,
-        batch_size,
-        cos_batch_size,
-        n_q_head,
-        n_kv_head,
-        head_dim,
-        pad_n_q_head,
-        pad_n_kv_head,
-        pad_hd,
-        BLOCK_SIZE=BLOCK_SIZE,
-        BACKWARD_PASS=True,
-    )
-    return dq.transpose(1, 2), dk.transpose(1, 2)
-class LigerRopeFunction(torch.autograd.Function):
-    """
-    Triton implementation of the Rotary Positional Embedding (RoPE) operation. Please note that
-    this implements the HuggingFace Llama & Mistral version, whose rotation matrix is slightly different
-    than the original RoPE paper.
-    Please find the corresponding HuggingFace implementation here:
-    https://github.com/huggingface/transformers/blob/v4.40.2/src/transformers/models/llama/modeling_llama.py#L184
-    For more details about the rotation matrix used here, please refer to:
-    https://discuss.huggingface.co/t/is-llama-rotary-embedding-implementation-correct/44509/2
-    """
-    @staticmethod
-    def forward(ctx, q, k, cos, sin, position_ids=None, unsqueeze_dim=1):
-        """
-        q size: (bsz, n_q_head, seq_len, head_dim)
-        k size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        sin size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        """
-        q, k, cos, sin = rope_forward(q, k, cos, sin)
-        ctx.save_for_backward(cos, sin)
-        return q, k
-    def backward(ctx, dq, dk):
-        """
-        dq size: (bsz, n_q_head, seq_len, head_dim)
-        dk size: (bsz, n_kv_head, seq_len, head_dim)
-        cos size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        sin size: (1, seq_len, head_dim) or (bsz, seq_len, head_dim)
-        """
-        cos, sin = ctx.saved_tensors
-        dq, dk = rope_backward(dq, dk, cos, sin)
-        return dq, dk, None, None, None, None

build/torch-xpu/swiglu.py DELETED Viewed

@@ -1,176 +0,0 @@
-import torch
-import triton
-import triton.language as tl
-from .utils import calculate_settings
-from .utils import ensure_contiguous
-@triton.jit
-def silu(x):
-    return x * tl.sigmoid(x)
-@triton.jit
-def _swiglu_forward_kernel(
-    a_ptr, b_ptr, c_ptr, stride, gate_multiplier, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr
-):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    a_ptr += program_id * stride
-    b_ptr += program_id * stride
-    c_ptr += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    # sigmoid requires type float32
-    a_row = tl.load(a_ptr + col_offsets, mask=mask, other=0).to(tl.float32) * gate_multiplier
-    b_row = tl.load(b_ptr + col_offsets, mask=mask, other=0)
-    c_row = silu(a_row).cast(b_row.dtype) * b_row
-    tl.store(c_ptr + col_offsets, c_row, mask=mask)
-@triton.jit
-def _swiglu_backward_kernel(
-    dc_ptr, a_ptr, b_ptr, stride, gate_multiplier, n_cols: tl.constexpr, BLOCK_SIZE: tl.constexpr
-):
-    program_id = tl.program_id(0).to(tl.int64)
-    # locate start index
-    dc_ptr += program_id * stride
-    a_ptr += program_id * stride
-    b_ptr += program_id * stride
-    col_offsets = tl.arange(0, BLOCK_SIZE)
-    mask = col_offsets < n_cols
-    dc_row = tl.load(dc_ptr + col_offsets, mask=mask, other=0)
-    # sigmoid requires type float32
-    a_row = tl.load(a_ptr + col_offsets, mask=mask, other=0).to(tl.float32) * gate_multiplier
-    b_row = tl.load(b_ptr + col_offsets, mask=mask, other=0)
-    # recomputation to save memory. a_row already holds a * gate_multiplier.
-    sig_a = tl.sigmoid(a_row)
-    silu_a = a_row * sig_a
-    db_row = dc_row * silu_a
-    # chain rule pulls an extra factor of gate_multiplier through the pre-activation scaling
-    da_row = dc_row * (silu_a * (1 - sig_a) + sig_a) * b_row * gate_multiplier
-    tl.store(a_ptr + col_offsets, da_row, mask=mask)
-    tl.store(b_ptr + col_offsets, db_row, mask=mask)
-def swiglu_forward(a, b, gate_multiplier: float = 1.0):
-    ori_shape = a.shape
-    n_cols = ori_shape[-1]
-    a = a.view(-1, n_cols)
-    b = b.view(-1, n_cols)
-    c = torch.empty_like(a)
-    n_rows = a.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _swiglu_forward_kernel[(n_rows,)](
-        a,
-        b,
-        c,
-        c.stride(-2),
-        float(gate_multiplier),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a, b, c.view(*ori_shape)
-def swiglu_backward(a, b, dc, gate_multiplier: float = 1.0):
-    ori_shape = dc.shape
-    n_cols = ori_shape[-1]
-    dc = dc.view(-1, n_cols)
-    n_rows = dc.shape[0]
-    BLOCK_SIZE, num_warps = calculate_settings(n_cols)
-    _swiglu_backward_kernel[(n_rows,)](
-        dc,
-        a,
-        b,
-        dc.stride(-2),
-        float(gate_multiplier),
-        n_cols=n_cols,
-        BLOCK_SIZE=BLOCK_SIZE,
-        num_warps=num_warps,
-    )
-    return a.view(*ori_shape), b.view(*ori_shape)
-class LigerSiLUMulFunction(torch.autograd.Function):
-    @staticmethod
-    @ensure_contiguous
-    def forward(ctx, a, b, gate_multiplier: float = 1.0, down_multiplier: float = 1.0):
-        gate_multiplier = float(gate_multiplier)
-        down_multiplier = float(down_multiplier)
-        ctx.gate_multiplier = gate_multiplier
-        ctx.down_multiplier = down_multiplier
-        if isinstance(a, torch.distributed.tensor.DTensor) or isinstance(b, torch.distributed.tensor.DTensor):
-            device_mesh, placements = (
-                (a.device_mesh, a.placements)
-                if isinstance(a, torch.distributed.tensor.DTensor)
-                else (b.device_mesh, b.placements)
-            )
-            # Assume that full tensors are gathered before and identical across
-            # the associated process groups.
-            if not isinstance(a, torch.distributed.tensor.DTensor):
-                a = torch.distributed.tensor.distribute_tensor(a, device_mesh=device_mesh, placements=placements)
-            if not isinstance(b, torch.distributed.tensor.DTensor):
-                b = torch.distributed.tensor.distribute_tensor(b, device_mesh=device_mesh, placements=placements)
-            a_local, b_local, c_local = swiglu_forward(a.to_local(), b.to_local(), gate_multiplier)
-            if down_multiplier != 1.0:
-                c_local = c_local * down_multiplier
-            ctx.save_for_backward(a_local, b_local)
-            ctx.dtensor_metadata = (device_mesh, placements)
-            return torch.distributed.tensor.DTensor.from_local(c_local, device_mesh, placements)
-        else:
-            a, b, c = swiglu_forward(a, b, gate_multiplier)
-            if down_multiplier != 1.0:
-                c = c * down_multiplier
-            ctx.save_for_backward(a, b)
-            ctx.dtensor_metadata = None
-            return c
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, dc):
-        a, b = ctx.saved_tensors
-        gate_multiplier = ctx.gate_multiplier
-        down_multiplier = ctx.down_multiplier
-        if ctx.dtensor_metadata is not None:
-            device_mesh, placements = ctx.dtensor_metadata
-            # Assume that full tensors are gathered before and identical across
-            # the associated process groups.
-            dc_local = (
-                dc.to_local()
-                if isinstance(dc, torch.distributed.tensor.DTensor)
-                else torch.distributed.tensor.distribute_tensor(dc, device_mesh=device_mesh, placements=placements)
-            )
-            if down_multiplier != 1.0:
-                dc_local = dc_local * down_multiplier
-            a_local, b_local = swiglu_backward(a, b, dc_local, gate_multiplier)
-            return (
-                torch.distributed.tensor.DTensor.from_local(a_local, device_mesh, placements),
-                torch.distributed.tensor.DTensor.from_local(b_local, device_mesh, placements),
-                None,
-                None,
-            )
-        if down_multiplier != 1.0:
-            dc = dc * down_multiplier
-        a, b = swiglu_backward(a, b, dc, gate_multiplier)
-        return a, b, None, None

build/torch-xpu/tiled_mlp.py DELETED Viewed

@@ -1,136 +0,0 @@
-import math
-from typing import Callable
-from typing import List
-from typing import Optional
-import torch
-from .utils import ensure_contiguous
-class LigerTiledMLPFunction(torch.autograd.Function):
-    """
-    Based on DeepSpeed's TiledMLP:
-    https://github.com/deepspeedai/DeepSpeed/blob/v0.18.2/deepspeed/runtime/sequence_parallel/ulysses_sp.py#L838
-    Perform a tiled MLP computation to massively reduce memory usage needed to compute MLP
-    when using very long sequence lengths.
-    This module re-computes `forward` in the `backward`. So the `forward` occurs twice each iteration.
-    And if you're using activation checkpointing it then occurs thrice.
-    Args:
-        fn: the function to call on sharded inputs (e.g., mlp.forward)
-        mlp_module: the MLP nn.Module object
-        x: the input to MLP.forward (hidden_states)
-        shards: how many shards to use
-        compute_params: a list of weights engaged in the compute
-    Returns:
-        the computed hidden_states
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        fn: Callable,
-        mlp_module: torch.nn.Module,
-        x: torch.Tensor,
-        shards: int,
-        compute_params: Optional[List[torch.nn.Parameter]] = None,
-    ) -> torch.Tensor:
-        ctx.fn = fn
-        ctx.mlp_module = mlp_module
-        ctx.shards = shards
-        ctx.save_for_backward(x)
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
-        x_shards = list(torch.chunk(x, chunks=shards, dim=-2))
-        with torch.no_grad():
-            output_shards = [fn(mlp_module, x_shard) for x_shard in x_shards]
-        output_unsharded = torch.cat(output_shards, dim=-2)
-        return output_unsharded
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, *grads) -> tuple:
-        fn = ctx.fn
-        (x,) = ctx.saved_tensors
-        mlp_module = ctx.mlp_module
-        shards = ctx.shards
-        x_requires_grad = x.requires_grad
-        x = x.detach()
-        # detach() unsets x.requires_grad, so restore it
-        x.requires_grad_(x_requires_grad)
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size] (moe experts)
-        hidden_size = x.shape[-1]
-        x_shape_orig = x.shape
-        # flatten bs+seqlen to avoid having stride issues when narrowing into seqlen w/ bs>1
-        x = x.view(-1, hidden_size)
-        incoming_grad = grads[0].view(-1, hidden_size)
-        x_grad = torch.zeros_like(x)
-        x_shards = list(torch.chunk(x, chunks=shards, dim=0))
-        for i, x_shard in enumerate(x_shards):
-            x_shard.requires_grad_(x_requires_grad)
-            # if seqlen is not exactly divisible by shards the last step will be shorter than shard_step
-            shard_step = x_shards[i].shape[0]
-            shard_offset = i * x_shards[0].shape[0]
-            x_shard.grad = x_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
-            incoming_grad_shard = incoming_grad.narrow(0, shard_offset, shard_step).view_as(x_shard)
-            with torch.enable_grad():
-                output = fn(mlp_module, x_shard)
-            torch.autograd.backward(output, incoming_grad_shard)
-        # unflatten
-        x_grad = x_grad.view(x_shape_orig)
-        return (None, None, x_grad, None, None)
-def apply_tiled_mlp(
-    fn: Callable,
-    mlp_module: torch.nn.Module,
-    x: torch.Tensor,
-    num_shards: Optional[int] = None,
-    compute_params: Optional[List[torch.nn.Parameter]] = None,
-) -> torch.Tensor:
-    """
-    Apply tiled MLP computation for memory efficiency.
-    Args:
-        fn: the function to call on sharded inputs (e.g., lambda module, x: module(x))
-        mlp_module: the MLP nn.Module object
-        x: the input tensor with shape [bs, seqlen, hidden_size] or [seqlen, hidden_size]
-        num_shards: number of shards to use. If None, automatically calculated as ceil(seqlen / hidden_size)
-        compute_params: list of parameters for DeepSpeed ZeRO optimization
-    Returns:
-        output tensor with the same shape as input
-    """
-    if num_shards is None:
-        # x.shape could be [bs, seqlen, hidden_size] or [seqlen, hidden_size]
-        hidden_size = x.shape[-1]
-        seqlen = x.shape[-2]
-        num_shards = math.ceil(seqlen / hidden_size)
-    # Ensure num_shards is at least 1
-    num_shards = max(1, num_shards)
-    return LigerTiledMLPFunction.apply(
-        fn,
-        mlp_module,
-        x,
-        num_shards,
-        compute_params,
-    )

build/torch-xpu/tvd.py DELETED Viewed

@@ -1,218 +0,0 @@
-from typing import Literal
-from typing import Optional
-import torch
-import triton
-import triton.language as tl
-from .utils import ensure_contiguous
-MAX_FUSED_SIZE = 65536 // 4
-REDUCTION_LITERAL = Literal["none", "sum", "mean", "batchmean"]
-_REDUCTION_MODE_NONE = tl.constexpr(0)
-_REDUCTION_MODE_SUM = tl.constexpr(1)
-_REDUCTION_MODE_MEAN = tl.constexpr(2)
-_REDUCTION_MODE_BATCHMEAN = tl.constexpr(3)
-_str_to_reduction_mode = {
-    "none": _REDUCTION_MODE_NONE.value,
-    "sum": _REDUCTION_MODE_SUM.value,
-    "mean": _REDUCTION_MODE_MEAN.value,
-    "batchmean": _REDUCTION_MODE_BATCHMEAN.value,
-}
-def get_num_warps(BLOCK_SIZE):
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return num_warps
-@triton.jit
-def _tv_distance_kernel(
-    p_ptr,
-    p_stride,
-    q_ptr,
-    q_stride,
-    loss_ptr,
-    loss_stride,
-    grads_ptr,
-    grads_stride,
-    label_ptr,
-    ignore_index: tl.constexpr,
-    n_cols,
-    scale,  # pre-computed reduction scale for gradients (fused into kernel)
-    BLOCK_SIZE: tl.constexpr,
-    HAS_LABEL: tl.constexpr,
-    reduction: tl.constexpr = _REDUCTION_MODE_BATCHMEAN,
-):
-    pid = tl.program_id(0).to(tl.int64)
-    p_ptr += pid * p_stride
-    q_ptr += pid * q_stride
-    loss_ptr += pid * loss_stride
-    grads_ptr += pid * grads_stride
-    label_ptr += pid
-    base_offsets = tl.arange(0, BLOCK_SIZE)
-    if HAS_LABEL:
-        label = tl.load(label_ptr)
-        if label == ignore_index:
-            for i in range(0, n_cols, BLOCK_SIZE):
-                offsets = i + base_offsets
-                mask = offsets < n_cols
-                tl.store(grads_ptr + offsets, 0.0, mask=mask)
-                if reduction == _REDUCTION_MODE_NONE:
-                    tl.store(loss_ptr + offsets, 0.0, mask=mask)
-            return
-    loss_sum = 0.0
-    for i in range(0, n_cols, BLOCK_SIZE):
-        offsets = i + base_offsets
-        mask = offsets < n_cols
-        p = tl.load(p_ptr + offsets, mask=mask, other=0.0)
-        q = tl.load(q_ptr + offsets, mask=mask, other=0.0)
-        # TVD(P || Q) = 0.5 * |P - Q|
-        tv_loss = 0.5 * tl.abs(p - q)
-        # Fuse reduction scaling into gradient computation (eliminates separate Python division)
-        grad_res = tl.where(p > q, 0.5 * scale, -0.5 * scale)
-        tl.store(grads_ptr + offsets, grad_res, mask=mask)
-        if reduction == _REDUCTION_MODE_NONE:
-            tl.store(loss_ptr + offsets, tv_loss, mask=mask)
-        else:
-            loss_sum += tl.sum(tv_loss, axis=0)
-    if reduction != _REDUCTION_MODE_NONE:
-        # Fuse reduction scaling into loss (same scale as gradients; avoids Python division)
-        tl.store(loss_ptr, loss_sum * scale)
-def tv_distance_forward_triton(p, q, shift_labels, reduction, ignore_index, has_label):
-    BT, V = p.shape
-    BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(V))
-    num_warps = get_num_warps(BLOCK_SIZE)
-    grid = (BT,)
-    reduction = _str_to_reduction_mode[reduction]
-    out_size = (BT, V) if reduction == _REDUCTION_MODE_NONE.value else (BT,)
-    output_tensor = torch.zeros(out_size, device=p.device, dtype=torch.float32)
-    grads = torch.empty_like(p)
-    n_non_ignore = (shift_labels != ignore_index).sum().item() if has_label else BT
-    # Pre-compute gradient scale factor (fused into kernel to avoid separate division)
-    if reduction == _REDUCTION_MODE_BATCHMEAN.value:
-        scale = 1.0 / n_non_ignore
-    elif reduction == _REDUCTION_MODE_MEAN.value:
-        scale = 1.0 / (n_non_ignore * V)
-    else:
-        scale = 1.0
-    _tv_distance_kernel[grid](
-        p,
-        p.stride(0),
-        q,
-        q.stride(0),
-        output_tensor,
-        output_tensor.stride(0),
-        grads,
-        grads.stride(0),
-        shift_labels if has_label else torch.empty(1, device=p.device),
-        ignore_index,
-        V,
-        scale,
-        BLOCK_SIZE=BLOCK_SIZE,
-        HAS_LABEL=has_label,
-        num_warps=num_warps,
-        reduction=reduction,
-    )
-    # Loss and gradients are already scaled inside the kernel — no separate division needed
-    if reduction in (_REDUCTION_MODE_BATCHMEAN.value, _REDUCTION_MODE_MEAN.value):
-        return output_tensor.sum(), grads
-    elif reduction == _REDUCTION_MODE_SUM.value:
-        return output_tensor.sum(dim=0), grads
-    else:
-        return output_tensor, grads
-def tvd_backward_triton(grad_output, grads):
-    # If cross entropy is the last layer, grad_output is 1.0. Skip the mul then.
-    if torch.equal(grad_output, torch.tensor(1.0, device=grad_output.device)):
-        return grads
-    return grads * grad_output
-class LigerTVDLossFunction(torch.autograd.Function):
-    """
-    Class implementing the forward and backward pass for the Total Variation Distance Loss using Triton.
-    """
-    @staticmethod
-    @ensure_contiguous
-    def forward(
-        ctx,
-        p: torch.Tensor,
-        q: torch.Tensor,
-        shift_labels: Optional[torch.Tensor] = None,
-        reduction: REDUCTION_LITERAL = "batchmean",
-        ignore_index: int = -100,
-    ) -> torch.Tensor:
-        """A forward pass for the Total Variation Distance Loss.
-        Args:
-            ctx: Torch autograd context
-            p (torch.Tensor): A tensor of shape (BT, V) containing the first distribution.
-            q (torch.Tensor): A tensor of shape (BT, V) containing the second distribution.
-            shift_labels (Optional[torch.Tensor]): A tensor of shape (BT,) containing the labels.
-            reduction (REDUCTION_LITERAL, optional): The reduction method to be applied. Defaults to "batchmean".
-            ignore_index (int, optional): The index to ignore during loss calculation. Defaults to -100.
-        Returns:
-            torch.Tensor: The computed Total Variation Distance Loss.
-        """
-        has_label = False
-        if shift_labels is not None:
-            assert shift_labels.shape == (p.shape[0],), (
-                f"the shape of shift_labels must be (BT,). Got: {shift_labels.shape}"
-            )
-            shift_labels = shift_labels.contiguous()
-            has_label = True
-        loss, grads = tv_distance_forward_triton(p, q, shift_labels, reduction, ignore_index, has_label)
-        ctx.save_for_backward(grads)
-        return loss
-    @staticmethod
-    @ensure_contiguous
-    def backward(ctx, grad_output: torch.Tensor) -> torch.Tensor:
-        """A backward pass for the Total Variation Distance Loss.
-        Args:
-            ctx: Torch autograd context
-            grad_output (torch.Tensor): The gradient of the loss with respect to the output.
-        Returns:
-            tuple[torch.Tensor, None, None, None, None]: The gradient of the loss with respect to the inputs.
-        """
-        (grads,) = ctx.saved_tensors
-        grads = tvd_backward_triton(grad_output, grads)
-        return grads, None, None, None, None

build/torch-xpu/utils.py DELETED Viewed

@@ -1,176 +0,0 @@
-"""
-This file incorporates code from Unsloth licensed under the Apache License, Version 2.0.
-See the original Unsloth repository at https://github.com/unslothai/unsloth.
-The following line
-https://github.com/linkedin/Liger-Kernel/blob/7382a8761f9af679482b968f9348013d933947c7/src/liger_kernel/ops/utils.py#L23
-is based on code from Unsloth, located at:
-https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/utils.py#L43
-Modifications made by Yanning Chen, 2024.
-"""
-import functools
-import importlib
-import operator
-from typing import Callable
-import torch
-import triton
-import triton.language as tl
-from packaging.version import Version
-def is_npu_available() -> bool:
-    """Detect Ascend NPU availability."""
-    try:
-        from transformers.utils import is_torch_npu_available
-        return is_torch_npu_available()
-    except Exception:
-        return False
-def infer_device():
-    """
-    Get current device name based on available devices
-    """
-    if torch.cuda.is_available():  # Works for both Nvidia and AMD
-        return "cuda"
-    # Use Ascend NPU if available (torch.npu)
-    elif is_npu_available():
-        return "npu"
-    # XPU (Intel) if available
-    elif torch.xpu.is_available():
-        return "xpu"
-    else:
-        return "cpu"
-def is_hip() -> bool:
-    return torch.version.hip is not None
-def ensure_contiguous(fn):
-    @functools.wraps(fn)
-    def wrapper(ctx, *args, **kwargs):
-        def maybe_to_contiguous(x):
-            return x.contiguous() if isinstance(x, torch.Tensor) else x
-        args = [maybe_to_contiguous(arg) for arg in args]
-        kwargs = {k: maybe_to_contiguous(v) for k, v in kwargs.items()}
-        return fn(ctx, *args, **kwargs)
-    return wrapper
-def calculate_settings(n):
-    # reference: https://github.com/unslothai/unsloth/blob/fd753fed99ed5f10ef8a9b7139588d9de9ddecfb/unsloth/kernels/utils.py#L43
-    MAX_FUSED_SIZE = 65536
-    BLOCK_SIZE = triton.next_power_of_2(n)
-    if BLOCK_SIZE > MAX_FUSED_SIZE:
-        raise RuntimeError(
-            f"Cannot launch Triton kernel since n = {n} exceeds the recommended Triton blocksize = {MAX_FUSED_SIZE}."
-        )
-    num_warps = 4
-    if BLOCK_SIZE >= 32768:
-        num_warps = 32 if not is_hip() else 16
-    elif BLOCK_SIZE >= 8192:
-        num_warps = 16
-    elif BLOCK_SIZE >= 2048:
-        num_warps = 8
-    return BLOCK_SIZE, num_warps
-def compare_version(package: str, operator: Callable, target: str):
-    try:
-        pkg = importlib.import_module(package)
-    except ImportError:
-        return False
-    pkg_version = Version(pkg.__version__)
-    return operator(pkg_version, Version(target))
-def get_amp_custom_fwd_bwd() -> Callable:
-    device = infer_device()
-    if compare_version("torch", operator.ge, "2.4.0"):
-        return (
-            functools.partial(torch.amp.custom_fwd, device_type=device),
-            functools.partial(torch.amp.custom_bwd, device_type=device),
-        )
-    if hasattr(torch, "npu") and getattr(torch.npu, "amp", None) is not None:
-        return torch.npu.amp.custom_fwd, torch.npu.amp.custom_bwd
-    return torch.cuda.amp.custom_fwd, torch.cuda.amp.custom_bwd
-amp_custom_fwd, amp_custom_bwd = get_amp_custom_fwd_bwd()
-torch_to_triton_dtype = {
-    torch.float32: tl.float32,
-    torch.float16: tl.float16,
-    torch.bfloat16: tl.bfloat16,
-}
-@triton.jit
-def element_mul_kernel(
-    X_ptr,
-    X_stride,
-    grad_output_ptr,
-    n_cols,
-    BLOCK_SIZE: tl.constexpr,
-):
-    """
-    This function multiplies each element of the tensor pointed by X_ptr with the value pointed by grad_output_ptr.
-    The multiplication is performed in-place on the tensor pointed by X_ptr.
-    Parameters:
-    X_ptr: Pointer to the input tensor.
-    X_stride (int): The stride of the input tensor.
-    grad_output_ptr: Pointer to the gradient output value.
-    n_cols (int): The number of columns in the input tensor.
-    BLOCK_SIZE (int): The block size for Triton operations.
-    """
-    # Get the program ID and convert it to int64 to avoid overflow
-    program_id = tl.program_id(0).to(tl.int64)
-    # Locate the start index
-    X_ptr += program_id * X_stride
-    # Load the gradient output value
-    grad_output = tl.load(grad_output_ptr)
-    # Perform the element-wise multiplication
-    for i in range(0, n_cols, BLOCK_SIZE):
-        X_offsets = i + tl.arange(0, BLOCK_SIZE)
-        X_block = tl.load(X_ptr + X_offsets, mask=X_offsets < n_cols)
-        tl.store(X_ptr + X_offsets, X_block * grad_output, mask=X_offsets < n_cols)
-def get_npu_core_count(default: int = 20) -> int:
-    """Return NPU vector core count.
-    Fallback to `default` if Triton runtime or NPU device is unavailable.
-    """
-    try:
-        utils = triton.runtime.driver.active.utils
-        props = utils.get_device_properties(0)
-        return int(props.get("num_vectorcore", default))
-    except Exception:
-        return default
-def set_large_grf_mode(kernel_args: dict):
-    """Set large GRF mode for XPU devices."""
-    # On XPU triton installed along with pytorch-xpu will be called `pytorch-triton-xpu`,
-    # triton XPU installed from source will be called `triton`.
-    if compare_version("pytorch-triton-xpu", operator.ge, "3.6.0") or compare_version("triton", operator.ge, "3.6.0"):
-        kernel_args["grf_mode"] = "256"
-    else:
-        # API was changed in https://github.com/intel/intel-xpu-backend-for-triton/pull/5430
-        kernel_args["grf_mode"] = "large"