@triton.autotune(configs=autotune_configs, key=['row_dim', 'col_dim',
    'inner_dim'])
@triton.jit
def dense_matmul_kernel(lhs_ptr, rhs_ptr, out_ptr, row_dim: tl.constexpr,
    col_dim: tl.constexpr, inner_dim: tl.constexpr, lhs_stride_row: tl.
    constexpr, lhs_stride_inner: tl.constexpr, rhs_stride_inner: tl.
    constexpr, rhs_stride_col: tl.constexpr, out_stride_row: tl.constexpr,
    out_stride_col: tl.constexpr, accumulate: tl.constexpr, masked: tl.
    constexpr, block_size_row: tl.constexpr, block_size_col: tl.constexpr,
    block_size_inner: tl.constexpr, group_size_row: tl.constexpr):
    if not masked:
        tl.static_assert(row_dim % block_size_row == 0)
        tl.static_assert(col_dim % block_size_col == 0)
        tl.static_assert(inner_dim % block_size_inner == 0)
    pid_row, pid_col = tl.swizzle2d(tl.program_id(axis=0), tl.program_id(
        axis=1), tl.cdiv(row_dim, block_size_row), tl.cdiv(col_dim,
        block_size_col), group_size_row)
    row_offset = pid_row * block_size_row
    col_offset = pid_col * block_size_col
    row_range = tl.arange(0, block_size_row)[:, None] + row_offset
    col_range = tl.arange(0, block_size_col)[None, :] + col_offset
    inner_range = tl.arange(0, block_size_inner)
    lhs_ptr += row_range * lhs_stride_row + inner_range[None, :
        ] * lhs_stride_inner
    rhs_ptr += inner_range[:, None
        ] * rhs_stride_inner + col_range * rhs_stride_col
    out_ptr += row_range * out_stride_row + col_range * out_stride_col
    if masked:
        row_mask = row_range < row_dim
        col_mask = col_range < col_dim
        inner_mask = inner_range < inner_dim
        out = tl.dot(tl.load(lhs_ptr, mask=row_mask * inner_mask[None, :],
            other=0), tl.load(rhs_ptr, mask=inner_mask[:, None] * col_mask,
            other=0), out_dtype=tl.float32)
    else:
        out = tl.dot(tl.load(lhs_ptr), tl.load(rhs_ptr), out_dtype=tl.float32)
    for k in range(1, inner_dim // block_size_inner):
        lhs_ptr += block_size_inner * lhs_stride_inner
        rhs_ptr += block_size_inner * rhs_stride_inner
        if masked:
            inner_range += block_size_inner
            inner_mask = inner_range < inner_dim
            out += tl.dot(tl.load(lhs_ptr, mask=row_mask & inner_mask[None,
                :], other=0), tl.load(rhs_ptr, mask=inner_mask[:, None] &
                col_mask, other=0), out_dtype=tl.float32)
        else:
            out += tl.dot(tl.load(lhs_ptr), tl.load(rhs_ptr))
    if masked:
        out_mask = row_mask & col_mask
        if accumulate:
            out += tl.load(out_ptr, mask=out_mask)
        tl.store(out_ptr, out, mask=out_mask)
    else:
        if accumulate:
            out += tl.load(out_ptr)
        tl.store(out_ptr, out)
