Method decode

tests/python/relax/test_codegen_cutlass.py:1267–1297

(
            A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
            B: T.Buffer((T.int64(128),), "float16"),
            decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
        )

    class Module:
        @T.prim_func(s_tir=True)
        def decode(
            A: T.Buffer((T.int64(64), T.int64(64)), "int8"),
            B: T.Buffer((T.int64(128),), "float16"),
            decode_1: T.Buffer((T.int64(64), T.int64(128)), "float16"),
        ):
            T.func_attr({"tirx.noalias": True})
            # with T.sblock("root"):
            for i, j in T.grid(T.int64(64), T.int64(128)):
                with T.sblock("decode"):
                    v_i, v_j = T.axis.remap("SS", [i, j])
                    T.reads(A[v_i, v_j // T.int64(2)], B[v_j])
                    T.writes(decode_1[v_i, v_j])
                    decode_1[v_i, v_j] = (
                        T.Cast(
                            "float16",
                            T.shift_right(
                                T.shift_left(
                                    T.bitwise_and(
                                        T.shift_right(
                                            T.Cast("int32", A[v_i, v_j // T.int64(2)]),
                                            T.Cast("int32", v_j % T.int64(2)) * 4,
                                        ),
                                        15,
                                    ),
                                    28,
                                ),
                                28,
                            ),
                        )
                        * B[v_j]
                    )

        @T.prim_func(s_tir=True)
        def encode(