def gemm_int8(n, m, l):
    A = te.placeholder((n, l), name="A", dtype="int8")
    B = te.placeholder((m, l), name="B", dtype="int8")

    k = te.reduce_axis((0, l), name="k")
    C = te.compute(
        (n, m),
        lambda i, j: te.sum(A[i, k].astype("int32") * B[j, k].astype("int32"), axis=k),
        name="C",
    )

    cfg = autotvm.get_config()
    s = te.create_schedule(C.op)
    y, x = C.op.axis

    AA = s.cache_read(A, "shared", [C])
    BB = s.cache_read(B, "shared", [C])
    AL = s.cache_read(AA, "local", [C])
    BL = s.cache_read(BB, "local", [C])
    CC = s.cache_write(C, "local")

    k = CC.op.reduce_axis[0]

    cfg.define_split(
        "tile_k",
        cfg.axis(k),
        num_outputs=3,
        filter=lambda entity: entity.size[2] == 4 and entity.size[0] * 2 >= entity.size[1],
    )

    ko, kt, ki = cfg["tile_k"].apply(s, CC, k)

    s[CC].tensorize(ki, intrin_dp4a)

    block_x = te.thread_axis("blockIdx.x")
    block_y = te.thread_axis("blockIdx.y")
    thread_x = te.thread_axis("threadIdx.x")
    thread_y = te.thread_axis("threadIdx.y")

    def block_size_filter(entity):
        return (
            entity.size[0] * 2 >= entity.size[1] * 2
            and entity.size[1] <= 16
            and entity.size[3] <= 4
        )

    cfg.define_split("tile_y", cfg.axis(y), num_outputs=4, filter=block_size_filter)
    cfg.define_split("tile_x", cfg.axis(x), num_outputs=4, filter=block_size_filter)
    by, tyz, ty, yi = cfg["tile_y"].apply(s, C, y)
    bx, txz, tx, xi = cfg["tile_x"].apply(s, C, x)

    s[C].bind(by, block_y)
    s[C].bind(bx, block_x)
    s[C].bind(tyz, te.thread_axis("vthread"))
    s[C].bind(txz, te.thread_axis("vthread"))
    s[C].bind(ty, thread_y)
    s[C].bind(tx, thread_x)
    s[C].reorder(by, bx, tyz, txz, ty, tx, yi, xi)

    s[CC].compute_at(s[C], tx)

    yo, xo = CC.op.axis
    s[CC].reorder(ko, kt, yo, xo, ki)
    s[CC].unroll(kt)

    for stage in [AL, BL]:
        s[stage].compute_at(s[CC], kt)
        _, xi = s[stage].split(stage.op.axis[1], factor=4)
        s[stage].vectorize(xi)
        s[stage].double_buffer()

    cfg.define_knob("storage_align", [16, 48])
    for stage in [AA, BB]:
        s[stage].storage_align(s[stage].op.axis[0], cfg["storage_align"].val, 0)
        s[stage].compute_at(s[CC], ko)

        fused = s[stage].fuse(*s[stage].op.axis)
        ty, tx = s[stage].split(fused, nparts=cfg["tile_y"].size[2])
        tx, xi = s[stage].split(tx, nparts=cfg["tile_x"].size[2])
        _, xi = s[stage].split(xi, factor=16)

        s[stage].bind(ty, thread_y)
        s[stage].bind(tx, thread_x)
        s[stage].vectorize(xi)

    cfg.define_knob("auto_unroll_max_step", [512, 1500])
    s[C].pragma(by, "auto_unroll_max_step", cfg["auto_unroll_max_step"].val)
    s[C].pragma(by, "unroll_explicit", False)

    cfg.add_flop(n * m * l * 2)
    return s, [A, B, C]