@triton.jit
def triton_normalization_forward_kernel(input_ptr, output_ptr, weight_ptr,
    bias_ptr, inv_var_ptr, n_cols, eps, has_bias: tl.constexpr,
    zero_centered: tl.constexpr, block_size: tl.constexpr):
    row = tl.program_id(0).to(tl.int64)
    cols = tl.arange(0, block_size)
    mask = cols < n_cols
    offsets = row * n_cols + cols
    input_ = tl.load(input_ptr + offsets, mask=mask, other=0.0).to(tl.float32)
    if has_bias:
        mean = tl.sum(input_, axis=0) / n_cols
        input_ = tl.where(mask, input_ - mean, 0.0)
    inv_var = 1 / tl.sqrt(tl.sum(input_ * input_, axis=0) / n_cols + eps)
    tl.store(inv_var_ptr + row, inv_var)
    weight = tl.load(weight_ptr + cols, mask=mask)
    if zero_centered:
        weight += 1
    output = input_ * inv_var * weight
    if has_bias:
        bias = tl.load(bias_ptr + cols, mask=mask)
        output = output + bias
    tl.store(output_ptr + offsets, output, mask=mask)
