@triton.jit
def _gelu_glu_fwd_kernel(output_ptr, input_ptr, output_row_stride,
    input_row_stride, output_col_stride, input_col_stride,
    output_page_stride, input_page_stride, n_pages, BLOCK_SIZE: tl.constexpr):
    row_idx = tl.program_id(0)
    col_idx = tl.program_id(1)
    x = tl.load(input_ptr + row_idx * input_row_stride + col_idx *
        input_col_stride + tl.arange(0, BLOCK_SIZE // 2) *
        input_page_stride, mask=tl.arange(0, BLOCK_SIZE // 2) < n_pages // 
        2, other=-float('inf'))
    gate = tl.load(input_ptr + row_idx * input_row_stride + col_idx *
        input_col_stride + (tl.arange(0, BLOCK_SIZE // 2) + n_pages // 2) *
        input_page_stride, mask=tl.arange(0, BLOCK_SIZE // 2) < n_pages // 
        2, other=-float('inf'))
    gate_cube = gate * gate * gate
    beta = 0.7978845608028654
    kappa = 0.044715
    inner = beta * (gate + kappa * gate_cube)
    inner_tanh = tanh(inner)
    gate_gelu = 0.5 * gate * (inner_tanh + 1)
    gelu_glu = gate_gelu * x
    tl.store(output_ptr + row_idx * output_row_stride + col_idx *
        output_col_stride + tl.arange(0, BLOCK_SIZE // 2) *
        output_page_stride, gelu_glu, mask=tl.arange(0, BLOCK_SIZE // 2) < 
        n_pages // 2)
