@triton.jit
def _layer_norm_fwd_fused(X_ptr, Y_ptr, W_ptr, B_ptr, Mean_ptr, RSTD_ptr,
    stride, n_cols, eps, BLOCK_SIZE: tl.constexpr):
    row_idx = tl.program_id(0)
    col_offsets = tl.arange(0, BLOCK_SIZE)
    mask = col_offsets < n_cols
    Y_ptr += row_idx * stride
    X_ptr += row_idx * stride
    Mean_ptr += row_idx
    RSTD_ptr += row_idx
    X_row = tl.load(X_ptr + col_offsets, mask=mask, other=0)
    W_row = tl.load(W_ptr + col_offsets, mask=mask, other=0)
    B_row = tl.load(B_ptr + col_offsets, mask=mask, other=0)
    mean = tl.sum(X_row, axis=0) / n_cols
    demeaned = X_row - mean
    var = tl.sum(demeaned * demeaned, axis=0) / n_cols
    rstd = rsqrt(var + eps)
    tl.store(Mean_ptr, mean)
    tl.store(RSTD_ptr, rstd)
    Y_row = tl.fma(demeaned * rstd, W_row, B_row)
    tl.store(Y_ptr + col_offsets, Y_row, mask=mask)
