@triton.jit
def triton_mlp_activation_backward_kernel(grad_output_ptr, grad_input_ptr,
    input_ptr, output_ptr, gated: tl.constexpr, activation_type: tl.
    constexpr, recompute: tl.constexpr, n_cols: tl.constexpr, block_size:
    tl.constexpr):
    row_idx = tl.program_id(0).to(tl.int64)
    columns = tl.program_id(1) * block_size + tl.arange(0, block_size)
    output_offsets = n_cols * row_idx + columns
    input_offsets = 2 * n_cols * row_idx + columns if gated else output_offsets
    input_ptr = input_ptr + input_offsets
    grad_input_ptr = grad_input_ptr + input_offsets
    mask = columns < n_cols
    input_ = tl.load(input_ptr, mask=mask).to(tl.float32)
    output_grad = tl.load(grad_output_ptr + output_offsets, mask=mask).to(tl
        .float32)
    if activation_type == _TritonActivationType.gelu:
        tanh_input = 0.79788456 * input_ * (1 + 0.044715 * input_ * input_)
        tanh = 1 - 2 / (1 + tl.exp(2 * tanh_input))
        grad = 0.5 * input_ * ((1 - tanh * tanh) * (0.79788456 + 
            0.1070322243 * input_ * input_)) + 0.5 * (1 + tanh)
        if gated or recompute:
            out = input_ * 0.5 * (1.0 + tanh)
    elif activation_type == _TritonActivationType.silu:
        exp = tl.exp(-input_)
        sigma = 1 / (1 + exp)
        grad = sigma * sigma + (1 + input_) / (2 + exp + 1 / exp)
        if gated or recompute:
            out = input_ * sigma
    elif activation_type == _TritonActivationType.relu:
        grad = tl.where(input_ > 0, 1, 0)
        if gated or recompute:
            out = tl.where(input_ > 0, input_, 0)
    elif activation_type == _TritonActivationType.squared_relu:
        relu_out = tl.where(input_ > 0, input_, 0)
        grad = 2 * relu_out
        if gated or recompute:
            out = relu_out * relu_out
    else:
        raise NotImplementedError()
    if gated:
        other = tl.load(input_ptr + n_cols, mask=mask)
        tl.store(grad_input_ptr, grad * other * output_grad, mask=mask)
        tl.store(grad_input_ptr + n_cols, out * output_grad, mask=mask)
        out = out * other
    else:
        tl.store(grad_input_ptr, grad * output_grad, mask=mask)
    if recompute:
        tl.store(output_ptr + output_offsets, out, mask=mask)
