code: def test_dense_add_cpu_no_write_cache():
      def apply_trace(sch):
          b0 = sch.get_block(name="T_matmul_NT", func_name="main")
          b1 = sch.get_block(name="root", func_name="main")
          sch.annotate(
              block_or_loop=b0,
              ann_key="meta_schedule.tiling_structure",
              ann_val="SSRSRS",
          )
          l2, l3, l4 = sch.get_loops(block=b0)
          v5, v6, v7, v8 = sch.sample_perfect_tile(
              loop=l2,
              n=4,
              max_innermost_factor=64,
              decision=[4, 4, 4, 2],
          )
          l9, l10, l11, l12 = sch.split(
              loop=l2,
              factors=[v5, v6, v7, v8],
              preserve_unit_iters=True,
          )
          v13, v14, v15, v16 = sch.sample_perfect_tile(
              loop=l3,
              n=4,
              max_innermost_factor=64,
              decision=[1, 1, 4, 32],
          )
          l17, l18, l19, l20 = sch.split(
              loop=l3,
              factors=[v13, v14, v15, v16],
              preserve_unit_iters=True,
          )
          v21, v22 = sch.sample_perfect_tile(
              loop=l4, n=2, max_innermost_factor=64, decision=[8, 16]
          )
          l23, l24 = sch.split(
              loop=l4, factors=[v21, v22], preserve_unit_iters=True
          )
          sch.reorder(l9, l17, l10, l18, l23, l11, l19, l24, l12, l20)
          sch.annotate(
              block_or_loop=b1,
              ann_key="meta_schedule.parallel",
              ann_val=160,
          )
          sch.annotate(
              block_or_loop=b1,
              ann_key="meta_schedule.vectorize",
              ann_val=64,
          )
          v25 = sch.sample_categorical(
              candidates=[0, 16, 64, 512],
              probs=[0.25, 0.25, 0.25, 0.25],
              decision=1,
          )
          sch.annotate(
              block_or_loop=b1,
              ann_key="meta_schedule.unroll_explicit",
              ann_val=v25,
          )
          sch.enter_postproc()
          b26 = sch.get_block(name="root", func_name="main")
          sch.unannotate(
              block_or_loop=b26, ann_key="meta_schedule.parallel"
          )
          sch.unannotate(
              block_or_loop=b26, ann_key="meta_schedule.vectorize"
          )
          sch.unannotate(
              block_or_loop=b26, ann_key="meta_schedule.unroll_explicit"
          )
          (b27,) = sch.get_child_blocks(b26)
          (
              l28,
              l29,
              l30,
              l31,
              l32,
              l33,
              l34,
              l35,
              l36,
              l37,
          ) = sch.get_loops(block=b27)
          l38 = sch.fuse(l28, l29, l30, l31, preserve_unit_iters=True)
          sch.parallel(loop=l38)
          l39 = sch.fuse(l37, preserve_unit_iters=True)
          sch.vectorize(loop=l39)
          sch.annotate(
              block_or_loop=l38,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=16,
          )
          sch.annotate(
              block_or_loop=l38,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          b40 = sch.get_block(name="T_matmul_NT", func_name="main")
          l41, l42, l43, l44, l45, l46, l47 = sch.get_loops(block=b40)
          b48 = sch.decompose_reduction(block=b40, loop=l42)
          b49 = sch.get_block(
              name="T_matmul_NT_update", func_name="main"
          )
          b50 = sch.cache_read(
              block=b49, read_buffer_index=2, storage_scope="global"
          )
          sch.transform_layout(
              block=b49,
              buffer=("read", 2),
              index_map=tvm.tir.IndexMap.from_func(
                  lambda i0, i1: (
                      floordiv(i1, 16),
                      floordiv(i0, 32),
                      floormod(i1, 16),
                      floormod(i0, 32),
                  ),
                  inverse_index_map=lambda i0, i1, i2, i3: (
                      ((i1 * 32) + i3),
                      ((i0 * 16) + i2),
                  ),
                  index_dtype="int32",
              ),
              pad_value=None,
          )
          sch.annotate(
              block_or_loop=b50,
              ann_key="meta_schedule.layout_rewrite_preproc",
              ann_val=1,
          )
  
      verify(
          Dense,
          apply_trace,
          DenseAdd,
          "llvm",
          DenseAdd_cpu_no_write_cache,
      )
  
imports: import pytest
  import tvm
  import tvm.meta_schedule
  import tvm.testing
  from tvm.script import tir
  from tvm.target import Target
  from tvm.target.codegen import llvm_lookup_intrinsic_id
  from tvm.tir import Schedule, floordiv, floormod
  from tvm.tir.tensor_intrin.cuda import *
  from tvm.tir.tensor_intrin.x86 import VNNI_DOT_16x4_INTRIN
label_desc: Use image node_id 4 for calling a global function with example usage: test_dense_add_cpu_no_write_cache() without return types