@triton.jit
def triton_mm(x_ptr, y_ptr, out_ptr, n: tl.constexpr, m: tl.constexpr, p:
    tl.constexpr, BLOCK_SIZE: tl.constexpr):
    pid0 = tl.program_id(axis=0)
    pid1 = tl.program_id(axis=1)
    x_row = pid0 * BLOCK_SIZE * m + tl.arange(0, m)
    x_col = tl.arange(0, BLOCK_SIZE) * m
    x_offset = x_row[None, :] + x_col[:, None]
    x_mask = tl.core.full((1, m), True, dtype=tl.int1) and (pid0 *
        BLOCK_SIZE + tl.arange(0, BLOCK_SIZE))[:, None] < n
    y_row = pid1 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    y_col = tl.arange(0, m) * p
    y_offset = y_row[None, :] + y_col[:, None]
    y_mask = (pid1 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE))[None, :
        ] < p and tl.core.full((m, 1), True, dtype=tl.int1)
    x = tl.load(x_ptr + x_offset, mask=x_mask, other=0.0)
    y = tl.load(y_ptr + y_offset, mask=y_mask, other=0.0)
    out = tl.dot(x, y, allow_tf32=False)
    out_row = pid0 * BLOCK_SIZE * p + pid1 * BLOCK_SIZE + tl.arange(0,
        BLOCK_SIZE)
    out_col = tl.arange(0, BLOCK_SIZE) * p
    out_offset = out_row[None, :] + out_col[:, None]
    out_mask = (pid1 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE))[None, :
        ] < p and (pid0 * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE))[:, None] < n
    tl.store(out_ptr + out_offset, out, mask=out_mask)
