code: def test_conv2d_int8_tensorcore():
      def apply_trace(sch):
          b0 = sch.get_block(name="pad_temp", func_name="main")
          b1 = sch.get_block(name="conv2d_nhwc", func_name="main")
          b2 = sch.get_block(name="T_subtract", func_name="main")
          b3 = sch.get_block(name="T_add", func_name="main")
          b4 = sch.get_block(name="T_cast", func_name="main")
          b5 = sch.get_block(name="T_multiply", func_name="main")
          b6 = sch.get_block(name="T_add_1", func_name="main")
          b7 = sch.get_block(name="T_right_shift", func_name="main")
          b8 = sch.get_block(name="T_cast_1", func_name="main")
          b9 = sch.get_block(name="T_add_2", func_name="main")
          b10 = sch.get_block(name="compute", func_name="main")
          b11 = sch.get_block(name="T_cast_2", func_name="main")
          b12 = sch.get_block(name="T_cast_3", func_name="main")
          b13 = sch.get_block(name="T_subtract_1", func_name="main")
          b14 = sch.get_block(name="compute_1", func_name="main")
          b15 = sch.get_block(name="root", func_name="main")
          sch.annotate(
              block_or_loop=b1,
              ann_key="meta_schedule.tiling_structure",
              ann_val="SSSRRSRS",
          )
          b16 = sch.reindex(block=b1, buffer=("write", 0))
          b17 = sch.reindex(block=b1, buffer=("read", 0))
          b18 = sch.reindex(block=b1, buffer=("read", 1))
          sch.transform_layout(
              block=b1,
              buffer=("read", 0),
              index_map=lambda nn, yy, xx, rc: (
                  (((nn * 3136) + (yy * 56)) + xx),
                  rc,
              ),
              pad_value=None,
          )
          sch.transform_layout(
              block=b1,
              buffer=("read", 1),
              index_map=lambda ff, ry, rx, rc: (
                  ry,
                  rx,
                  ff,
                  rc,
              ),
              pad_value=None,
          )
          sch.transform_layout(
              block=b1,
              buffer=("write", 0),
              index_map=lambda nn, yy, xx, ff: (
                  (((nn * 3136) + (yy * 56)) + xx),
                  ff,
              ),
              pad_value=None,
          )
          sch.transform_block_layout(
              block=b16,
              index_map=lambda nn, yy, xx, ff: (
                  (((nn * 3136) + (yy * 56)) + xx),
                  ff,
              ),
          )
          sch.transform_block_layout(
              block=b17,
              index_map=lambda nn, yy, xx, rc: (
                  (((nn * 3136) + (yy * 56)) + xx),
                  rc,
              ),
          )
          sch.transform_block_layout(
              block=b18,
              index_map=lambda ff, ry, rx, rc: (
                  ry,
                  rx,
                  ff,
                  rc,
              ),
          )
          sch.transform_block_layout(
              block=b1,
              index_map=lambda nn, yy, xx, ff, ry, rx, rc: (
                  ry,
                  rx,
                  (((nn * 3136) + (yy * 56)) + xx),
                  ff,
                  rc,
              ),
          )
          l19, l20, l21, l22, l23 = sch.get_loops(block=b1)
          l24, l25 = sch.split(
              loop=l23, factors=[None, 16], preserve_unit_iters=True
          )
          l26, l27 = sch.split(
              loop=l22, factors=[None, 16], preserve_unit_iters=True
          )
          l28, l29 = sch.split(
              loop=l21, factors=[None, 16], preserve_unit_iters=True
          )
          l30, l31, l32, l33, l34, l35, l36, l37 = sch.get_loops(
              block=b1
          )
          sch.reorder(l34, l36, l29, l27, l25)
          b38 = sch.blockize(target=l29)
          sch.annotate(
              block_or_loop=b38,
              ann_key="meta_schedule.auto_tensorize",
              ann_val="wmma_sync_16x16x16_s8s8s32_trans",
          )
          sch.annotate(
              block_or_loop=b38,
              ann_key="meta_schedule.auto_tensorize_init",
              ann_val="wmma_fill_16x16x16_s32",
          )
          sch.annotate(
              block_or_loop=b38, ann_key="warp_execution", ann_val=1
          )
          l39, l40, l41, l42, l43 = sch.get_loops(block=b38)
          v44, v45, v46 = sch.sample_perfect_tile(
              loop=l39, n=3, max_innermost_factor=4, decision=[1, 1, 1]
          )
          l47, l48, l49 = sch.split(
              loop=l39,
              factors=[v44, v45, v46],
              preserve_unit_iters=True,
          )
          v50, v51, v52 = sch.sample_perfect_tile(
              loop=l40, n=3, max_innermost_factor=4, decision=[1, 1, 1]
          )
          l53, l54, l55 = sch.split(
              loop=l40,
              factors=[v50, v51, v52],
              preserve_unit_iters=True,
          )
          v56, v57, v58, v59, v60 = sch.sample_perfect_tile(
              loop=l41,
              n=5,
              max_innermost_factor=4,
              decision=[392, 1, 8, 1, 1],
          )
          l61, l62, l63, l64, l65 = sch.split(
              loop=l41,
              factors=[v56, v57, v58, v59, v60],
              preserve_unit_iters=True,
          )
          v66, v67, v68, v69, v70 = sch.sample_perfect_tile(
              loop=l42,
              n=5,
              max_innermost_factor=4,
              decision=[8, 1, 2, 1, 1],
          )
          l71, l72, l73, l74, l75 = sch.split(
              loop=l42,
              factors=[v66, v67, v68, v69, v70],
              preserve_unit_iters=True,
          )
          v76, v77, v78 = sch.sample_perfect_tile(
              loop=l43, n=3, max_innermost_factor=4, decision=[2, 1, 2]
          )
          l79, l80, l81 = sch.split(
              loop=l43,
              factors=[v76, v77, v78],
              preserve_unit_iters=True,
          )
          sch.reorder(
              l61,
              l71,
              l62,
              l72,
              l63,
              l73,
              l47,
              l53,
              l79,
              l48,
              l54,
              l80,
              l64,
              l74,
              l49,
              l55,
              l81,
              l65,
              l75,
          )
          l82 = sch.fuse(l61, l71, preserve_unit_iters=True)
          sch.bind(loop=l82, thread_axis="blockIdx.x")
          l83 = sch.fuse(l62, l72, preserve_unit_iters=True)
          sch.bind(loop=l83, thread_axis="vthread.x")
          l84 = sch.fuse(l63, l73, preserve_unit_iters=True)
          sch.bind(loop=l84, thread_axis="threadIdx.x")
          sch.annotate(
              block_or_loop=b38,
              ann_key="meta_schedule.thread_extent_low_inclusive",
              ann_val=32,
          )
          sch.annotate(
              block_or_loop=b38,
              ann_key="meta_schedule.thread_extent_high_inclusive",
              ann_val=1024,
          )
          b85 = sch.cache_write(
              block=b38, write_buffer_index=0, storage_scope="shared"
          )
          sch.reverse_compute_at(
              block=b85, loop=l83, preserve_unit_loops=True, index=-1
          )
          b86 = sch.cache_write(
              block=b38,
              write_buffer_index=0,
              storage_scope="wmma.accumulator",
          )
          sch.reverse_compute_at(
              block=b86, loop=l84, preserve_unit_loops=True, index=-1
          )
          v87 = sch.sample_categorical(
              candidates=[1, 2, 3, 4, 8, 16],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=0,
          )
          sch.annotate(
              block_or_loop=b85,
              ann_key="meta_schedule.cooperative_fetch",
              ann_val=v87,
          )
          sch.reverse_compute_inline(block=b16)
          l88, l89, l90, l91, l92 = sch.get_loops(block=b86)
          l93, l94 = sch.split(
              loop=l92, factors=[None, 16], preserve_unit_iters=True
          )
          l95, l96 = sch.split(
              loop=l91, factors=[None, 16], preserve_unit_iters=True
          )
          l97, l98, l99, l100, l101, l102, l103 = sch.get_loops(
              block=b86
          )
          sch.reorder(l102, l96, l94)
          b104 = sch.blockize(target=l96)
          sch.annotate(
              block_or_loop=b104,
              ann_key="meta_schedule.auto_tensorize",
              ann_val="wmma_store_16x16x16_s32_shared",
          )
          b105 = sch.cache_read(
              block=b38,
              read_buffer_index=0,
              storage_scope="shared",
              consumer_blocks=[b38],
          )
          sch.compute_at(
              block=b105, loop=l79, preserve_unit_loops=True, index=-1
          )
          (
              l106,
              l107,
              l108,
              l109,
              l110,
              l111,
              l112,
              l113,
          ) = sch.get_loops(block=b105)
          l114 = sch.fuse(l112, l113, preserve_unit_iters=True)
          v115 = sch.sample_categorical(
              candidates=[1, 2, 3, 4, 8, 16],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=5,
          )
          sch.annotate(
              block_or_loop=b105,
              ann_key="meta_schedule.cooperative_fetch",
              ann_val=v115,
          )
          b116 = sch.cache_read(
              block=b38,
              read_buffer_index=1,
              storage_scope="shared",
              consumer_blocks=[b38],
          )
          sch.compute_at(
              block=b116, loop=l79, preserve_unit_loops=True, index=-1
          )
          (
              l117,
              l118,
              l119,
              l120,
              l121,
              l122,
              l123,
              l124,
              l125,
              l126,
          ) = sch.get_loops(block=b116)
          l127 = sch.fuse(
              l123, l124, l125, l126, preserve_unit_iters=True
          )
          v128 = sch.sample_categorical(
              candidates=[1, 2, 3, 4, 8, 16],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=4,
          )
          sch.annotate(
              block_or_loop=b116,
              ann_key="meta_schedule.cooperative_fetch",
              ann_val=v128,
          )
          b129 = sch.cache_read(
              block=b38,
              read_buffer_index=0,
              storage_scope="wmma.matrix_a",
          )
          sch.compute_at(
              block=b129, loop=l80, preserve_unit_loops=True, index=-1
          )
          (
              l130,
              l131,
              l132,
              l133,
              l134,
              l135,
              l136,
              l137,
              l138,
              l139,
              l140,
          ) = sch.get_loops(block=b129)
          l141, l142 = sch.split(
              loop=l140, factors=[None, 16], preserve_unit_iters=True
          )
          l143, l144 = sch.split(
              loop=l139, factors=[None, 16], preserve_unit_iters=True
          )
          (
              l145,
              l146,
              l147,
              l148,
              l149,
              l150,
              l151,
              l152,
              l153,
              l154,
              l155,
              l156,
              l157,
          ) = sch.get_loops(block=b129)
          sch.reorder(l156, l144, l142)
          b158 = sch.blockize(target=l144)
          sch.annotate(
              block_or_loop=b158,
              ann_key="meta_schedule.auto_tensorize",
              ann_val="wmma_load_16x16x16_s8_a_shared",
          )
          b159 = sch.cache_read(
              block=b38,
              read_buffer_index=1,
              storage_scope="wmma.matrix_b",
          )
          sch.compute_at(
              block=b159, loop=l80, preserve_unit_loops=True, index=-1
          )
          (
              l160,
              l161,
              l162,
              l163,
              l164,
              l165,
              l166,
              l167,
              l168,
              l169,
              l170,
              l171,
              l172,
          ) = sch.get_loops(block=b159)
          l173, l174 = sch.split(
              loop=l172, factors=[None, 16], preserve_unit_iters=True
          )
          l175, l176 = sch.split(
              loop=l171, factors=[None, 16], preserve_unit_iters=True
          )
          (
              l177,
              l178,
              l179,
              l180,
              l181,
              l182,
              l183,
              l184,
              l185,
              l186,
              l187,
              l188,
              l189,
              l190,
              l191,
          ) = sch.get_loops(block=b159)
          sch.reorder(l190, l176, l174)
          b192 = sch.blockize(target=l176)
          sch.annotate(
              block_or_loop=b192,
              ann_key="meta_schedule.auto_tensorize",
              ann_val="wmma_load_16x16x16_s8_b_trans_shared",
          )
          sch.compute_inline(block=b17)
          sch.compute_inline(block=b18)
          sch.storage_align(
              block=b105, buffer_index=0, axis=-2, factor=32, offset=16
          )
          sch.storage_align(
              block=b116, buffer_index=0, axis=-2, factor=32, offset=16
          )
          sch.reverse_compute_inline(block=b14)
          sch.reverse_compute_inline(block=b13)
          sch.reverse_compute_inline(block=b12)
          sch.reverse_compute_inline(block=b11)
          sch.reverse_compute_inline(block=b10)
          sch.reverse_compute_inline(block=b9)
          sch.reverse_compute_inline(block=b8)
          sch.reverse_compute_inline(block=b7)
          sch.reverse_compute_inline(block=b6)
          sch.reverse_compute_inline(block=b5)
          sch.reverse_compute_inline(block=b4)
          sch.reverse_compute_inline(block=b3)
          sch.reverse_compute_inline(block=b2)
          sch.compute_inline(block=b0)
          v193 = sch.sample_categorical(
              candidates=[0, 16, 64, 512, 1024],
              probs=[
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
              ],
              decision=3,
          )
          sch.annotate(
              block_or_loop=b15,
              ann_key="meta_schedule.unroll_explicit",
              ann_val=v193,
          )
          sch.enter_postproc()
          sch.unannotate(
              block_or_loop=b85,
              ann_key="meta_schedule.cooperative_fetch",
          )
          l194, l195, l196, l197 = sch.get_loops(block=b85)
          l198, l199 = sch.split(
              loop=l197, factors=[None, 16], preserve_unit_iters=True
          )
          sch.bind(loop=l199, thread_axis="threadIdx.x")
          sch.unannotate(
              block_or_loop=b105,
              ann_key="meta_schedule.cooperative_fetch",
          )
          l200, l201, l202, l203, l204, l205, l206 = sch.get_loops(
              block=b105
          )
          l207, l208, l209 = sch.split(
              loop=l206,
              factors=[None, 16, 16],
              preserve_unit_iters=True,
          )
          sch.vectorize(loop=l209)
          sch.bind(loop=l208, thread_axis="threadIdx.x")
          sch.unannotate(
              block_or_loop=b116,
              ann_key="meta_schedule.cooperative_fetch",
          )
          l210, l211, l212, l213, l214, l215, l216 = sch.get_loops(
              block=b116
          )
          l217, l218, l219 = sch.split(
              loop=l216, factors=[None, 16, 8], preserve_unit_iters=True
          )
          sch.vectorize(loop=l219)
          sch.bind(loop=l218, thread_axis="threadIdx.x")
          b220 = sch.get_block(name="root", func_name="main")
          sch.unannotate(
              block_or_loop=b220,
              ann_key="meta_schedule.unroll_explicit",
          )
          (
              b221,
              b222,
              b223,
              b224,
              b225,
              b226,
              b227,
          ) = sch.get_child_blocks(b220)
          (
              l228,
              l229,
              l230,
              l231,
              l232,
              l233,
              l234,
              l235,
              l236,
          ) = sch.get_loops(block=b221)
          sch.annotate(
              block_or_loop=l228,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l228,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          (
              l237,
              l238,
              l239,
              l240,
              l241,
              l242,
              l243,
              l244,
              l245,
          ) = sch.get_loops(block=b222)
          sch.annotate(
              block_or_loop=l237,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l237,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          (
              l246,
              l247,
              l248,
              l249,
              l250,
              l251,
              l252,
              l253,
              l254,
              l255,
              l256,
          ) = sch.get_loops(block=b223)
          sch.annotate(
              block_or_loop=l246,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l246,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          (
              l257,
              l258,
              l259,
              l260,
              l261,
              l262,
              l263,
              l264,
              l265,
              l266,
              l267,
              l268,
              l269,
          ) = sch.get_loops(block=b224)
          sch.annotate(
              block_or_loop=l257,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l257,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          (
              l270,
              l271,
              l272,
              l273,
              l274,
              l275,
              l276,
              l277,
              l278,
              l279,
              l280,
              l281,
              l282,
              l283,
              l284,
              l285,
          ) = sch.get_loops(block=b225)
          sch.annotate(
              block_or_loop=l270,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l270,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l286, l287, l288, l289, l290 = sch.get_loops(block=b226)
          sch.annotate(
              block_or_loop=l286,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l286,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l291, l292, l293, l294, l295 = sch.get_loops(block=b227)
          sch.annotate(
              block_or_loop=l291,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=512,
          )
          sch.annotate(
              block_or_loop=l291,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          b296 = sch.get_block(name="conv2d_nhwc_o", func_name="main")
          (
              l297,
              l298,
              l299,
              l300,
              l301,
              l302,
              l303,
              l304,
              l305,
              l306,
              l307,
              l308,
              l309,
              l310,
              l311,
              l312,
          ) = sch.get_loops(block=b296)
          b313 = sch.decompose_reduction(block=b296, loop=l300)
          sch.unannotate(
              block_or_loop=b313, ann_key="meta_schedule.auto_tensorize"
          )
          sch.annotate(
              block_or_loop=b313,
              ann_key="meta_schedule.auto_tensorize",
              ann_val="wmma_fill_16x16x16_s32",
          )
          sch.unannotate(
              block_or_loop=b296,
              ann_key="meta_schedule.auto_tensorize_init",
          )
          sch.unannotate(
              block_or_loop=b313,
              ann_key="meta_schedule.auto_tensorize_init",
          )
          b314 = sch.get_block(
              name="conv2d_nhwc_o_init", func_name="main"
          )
          sch.unannotate(
              block_or_loop=b314, ann_key="meta_schedule.auto_tensorize"
          )
          sch.tensorize(
              block_or_loop=b314, tensor_intrin="wmma_fill_16x16x16_s32"
          )
          b315 = sch.get_block(
              name="pad_temp_reindex_shared_wmma.matrix_a_o",
              func_name="main",
          )
          sch.unannotate(
              block_or_loop=b315, ann_key="meta_schedule.auto_tensorize"
          )
          sch.tensorize(
              block_or_loop=b315,
              tensor_intrin="wmma_load_16x16x16_s8_a_shared",
          )
          b316 = sch.get_block(
              name="p1_reindex_shared_wmma.matrix_b_o", func_name="main"
          )
          sch.unannotate(
              block_or_loop=b316, ann_key="meta_schedule.auto_tensorize"
          )
          sch.tensorize(
              block_or_loop=b316,
              tensor_intrin="wmma_load_16x16x16_s8_b_trans_shared",
          )
          b317 = sch.get_block(
              name="conv2d_nhwc_o_update", func_name="main"
          )
          sch.unannotate(
              block_or_loop=b317, ann_key="meta_schedule.auto_tensorize"
          )
          sch.tensorize(
              block_or_loop=b317,
              tensor_intrin="wmma_sync_16x16x16_s8s8s32_trans",
          )
          b318 = sch.get_block(
              name="conv2d_nhwc_reindex_shared_wmma.accumulator_o",
              func_name="main",
          )
          sch.unannotate(
              block_or_loop=b318, ann_key="meta_schedule.auto_tensorize"
          )
          sch.tensorize(
              block_or_loop=b318,
              tensor_intrin="wmma_store_16x16x16_s32_shared",
          )
  
      verify(
          Conv2dInt8,
          apply_trace,
          Conv2dInt8_target,
          "cuda",
          Conv2dInt8_tensorcore_scheduled,
      )
  
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 6 for calling a global function with example usage: test_conv2d_int8_tensorcore() without return types