code: def test_dp4a_dense():
      @T.prim_func
      def dp4a_dense_0(
          X: T.Buffer((128, 128), "int8"),
          W: T.Buffer((128, 128), "int8"),
          compute: T.Buffer((128, 128), "int32"),
      ) -> None:
          T.func_attr(
              {"global_symbol": "main", "tir.noalias": T.bool(True)}
          )
          # with T.block("root"):
          compute_local = T.alloc_buffer(
              (128, 128), "int32", scope="local"
          )
          X_shared = T.alloc_buffer((128, 128), "int8", scope="shared")
          W_shared = T.alloc_buffer((128, 128), "int8", scope="shared")
          for i_0_j_0_fused in T.thread_binding(1, thread="blockIdx.x"):
              for i_1_j_1_fused in T.thread_binding(
                  512, thread="vthread.x"
              ):
                  for i_2_j_2_fused in T.thread_binding(
                      2, thread="threadIdx.x"
                  ):
                      for k_0_0 in range(1):
                          for ax0_ax1_fused in range(16384):
                              with T.block("X_shared"):
                                  v0 = T.axis.spatial(
                                      128, ax0_ax1_fused // 128
                                  )
                                  v1 = T.axis.spatial(
                                      128, ax0_ax1_fused % 128
                                  )
                                  T.reads(X[v0, v1])
                                  T.writes(X_shared[v0, v1])
                                  T.block_attr(
                                      {
                                          "meta_schedule.cooperative_fetch": 1
                                      }
                                  )
                                  X_shared[v0, v1] = X[v0, v1]
                          for ax0_ax1_fused in range(16384):
                              with T.block("W_shared"):
                                  v0 = T.axis.spatial(
                                      128, ax0_ax1_fused // 128
                                  )
                                  v1 = T.axis.spatial(
                                      128, ax0_ax1_fused % 128
                                  )
                                  T.reads(W[v0, v1])
                                  T.writes(W_shared[v0, v1])
                                  T.block_attr(
                                      {
                                          "meta_schedule.cooperative_fetch": 1
                                      }
                                  )
                                  W_shared[v0, v1] = W[v0, v1]
                          for (
                              k_0_1,
                              i_3,
                              j_3,
                              k_0_2,
                              i_4,
                              j_4,
                          ) in T.grid(1, 2, 4, 32, 2, 1):
                              with T.block("compute_o"):
                                  v_i = T.axis.spatial(
                                      128,
                                      i_1_j_1_fused // 32 * 8
                                      + i_2_j_2_fused * 4
                                      + i_3 * 2
                                      + i_4,
                                  )
                                  v_j = T.axis.spatial(
                                      128,
                                      i_1_j_1_fused % 32 * 4
                                      + j_3
                                      + j_4,
                                  )
                                  v_k_o = T.axis.reduce(
                                      32,
                                      k_0_0 * 32 + k_0_1 * 32 + k_0_2,
                                  )
                                  T.reads(
                                      X_shared[
                                          v_i, v_k_o * 4 : v_k_o * 4 + 4
                                      ],
                                      W_shared[
                                          v_j, v_k_o * 4 : v_k_o * 4 + 4
                                      ],
                                  )
                                  T.writes(compute_local[v_i, v_j])
                                  T.block_attr(
                                      {
                                          "meta_schedule.auto_tensorize": "dp4a"
                                      }
                                  )
                                  with T.init():
                                      with T.block("compute_init"):
                                          T.reads()
                                          T.writes(
                                              compute_local[v_i, v_j]
                                          )
                                          compute_local[v_i, v_j] = 0
                                  for k_1 in range(4):
                                      with T.block("compute"):
                                          v_k_i = T.axis.reduce(4, k_1)
                                          T.reads(
                                              compute_local[v_i, v_j],
                                              X_shared[
                                                  v_i, v_k_o * 4 + v_k_i
                                              ],
                                              W_shared[
                                                  v_j, v_k_o * 4 + v_k_i
                                              ],
                                          )
                                          T.writes(
                                              compute_local[v_i, v_j]
                                          )
                                          T.block_attr(
                                              {
                                                  "meta_schedule.tiling_structure": "SSSRRSRS"
                                              }
                                          )
                                          compute_local[
                                              v_i, v_j
                                          ] = compute_local[
                                              v_i, v_j
                                          ] + T.Cast(
                                              "int32",
                                              X_shared[
                                                  v_i, v_k_o * 4 + v_k_i
                                              ],
                                          ) * T.Cast(
                                              "int32",
                                              W_shared[
                                                  v_j, v_k_o * 4 + v_k_i
                                              ],
                                          )
                      for ax0, ax1 in T.grid(4, 4):
                          with T.block("compute_local"):
                              v0 = T.axis.spatial(
                                  128,
                                  i_1_j_1_fused // 32 * 8
                                  + i_2_j_2_fused * 4
                                  + ax0,
                              )
                              v1 = T.axis.spatial(
                                  128, i_1_j_1_fused % 32 * 4 + ax1
                              )
                              T.reads(compute_local[v0, v1])
                              T.writes(compute[v0, v1])
                              compute[v0, v1] = compute_local[v0, v1]
  
      decision_0 = [
          ("SamplePerfectTile", [1, 16, 2, 2, 2]),
          ("SamplePerfectTile", [1, 32, 1, 4, 1]),
          ("SamplePerfectTile", [1, 1, 32]),
          ("SampleCategorical", 0),
          ("SampleCategorical", 0),
      ]
      _check_dp4a_dense(
          m=128,
          n=128,
          k=128,
          in_dtype="int8",
          out_dtype="int32",
          expected_mods=[dp4a_dense_0],
          expected_decisions=[decision_0],
      )
  
file_path: tvm/tests/python/meta_schedule/test_meta_schedule_schedule_rule_mlt_intrin.py
imports: from tvm import meta_schedule
  from tvm import te
  from tvm.ir import assert_structural_equal
  from tvm.meta_schedule.testing.space_generation import check_sketches, generate_design_space, print_sketches
  from tvm.script import tir
  from tvm.target import Target
  from tvm.tir.tensor_intrin.arm_cpu import DP4A_INTRIN
  from tvm.tir.tensor_intrin.x86 import AVX512_DOT_16x4_INTRIN
  from tvm.tir.tensor_intrin.x86 import VNNI_DOT_16x4_INTRIN
label_desc: Use image node_id 3 for calling a global function with example usage: test_dp4a_dense() without return types