@staticmethod
@util.autotune(configs=linear_configs(), key=['m_size', 'n_size', 'k_size'])
@triton.jit
def forward(output_ptr: tl.tensor, input_ptr: tl.tensor, weight_ptr: tl.
    tensor, bias_ptr: tl.tensor, m_size: tl.int32, n_size: tl.int32, k_size:
    tl.int32, input_batch_stride: tl.int32, input_m_stride: tl.int32,
    input_k_stride: tl.int32, weight_n_stride: tl.int32, weight_k_stride:
    tl.int32, use_accelerator: tl.constexpr, dtype: tl.constexpr,
    m_block_size: tl.constexpr, n_block_size: tl.constexpr, k_block_size:
    tl.constexpr):
    pid = tl.program_id(0)
    num_m_blocks = tl.cdiv(m_size, m_block_size)
    num_n_blocks = tl.cdiv(n_size, n_block_size)
    num_blocks = num_m_blocks * num_n_blocks
    batch = pid // num_blocks
    block = pid % num_blocks
    m_block = block // num_n_blocks
    n_block = block % num_n_blocks
    m_offset = m_block * m_block_size
    n_offset = n_block * n_block_size
    output = language.Linear.forward(input_ptr + batch * input_batch_stride,
        weight_ptr, bias_ptr, m_size, n_size, k_size, input_m_stride,
        input_k_stride, weight_n_stride, weight_k_stride, m_offset,
        n_offset, use_accelerator, m_block_size, n_block_size, k_block_size,
        dtype)
    output_block_ptr = tl.make_block_ptr(output_ptr + batch * m_size *
        n_size, shape=(m_size, n_size), strides=(n_size, 1), offsets=(
        m_offset, n_offset), block_shape=(m_block_size, n_block_size),
        order=(1, 0))
    tl.store(output_block_ptr, output, boundary_check=(0, 1))
