@triton.jit
def _layer_norm_fw(X, Y, W, B, M, V, stride, N, eps, **META):
    """
    Fused layernorm kernel over a 3d tensor.
    The layer norm is applied over the last dimension.

    Compute
        y = (x - E(x))/(sqrt(var(x) + epsilon)) * gamma + beta
    """
    y = _layer_norm_non_affine(X, M, V, stride, N, eps, META)
    y = _affine(W, B, N, y, META)
    _store(y, Y, stride, N, META)
