code: def test_winograd_gpu():
      def apply_trace(sch):
          b0 = sch.get_block(name="B", func_name="main")
          b1 = sch.get_block(name="data_pack", func_name="main")
          b2 = sch.get_block(name="bgemm", func_name="main")
          b3 = sch.get_block(name="A", func_name="main")
          b4 = sch.get_block(name="inverse", func_name="main")
          b5 = sch.get_block(name="conv2d_winograd", func_name="main")
          b6 = sch.get_block(name="T_add", func_name="main")
          b7 = sch.get_block(name="T_relu", func_name="main")
          b8 = sch.get_block(name="root", func_name="main")
          sch.compute_inline(block=b0)
          (b9,) = sch.get_producers(block=b1)
          (b10,) = sch.get_producers(block=b9)
          l11, l12, l13, l14, l15, l16 = sch.get_loops(block=b1)
          v17, v18 = sch.sample_perfect_tile(
              loop=l13, n=2, max_innermost_factor=64, decision=[14, 14]
          )
          l19, l20 = sch.split(
              loop=l13, factors=[v17, v18], preserve_unit_iters=True
          )
          v21, v22 = sch.sample_perfect_tile(
              loop=l14, n=2, max_innermost_factor=64, decision=[8, 8]
          )
          l23, l24 = sch.split(
              loop=l14, factors=[v21, v22], preserve_unit_iters=True
          )
          sch.unroll(loop=l11)
          sch.unroll(loop=l12)
          sch.unroll(loop=l15)
          sch.unroll(loop=l16)
          sch.reorder(l19, l23, l20, l24, l11, l12, l15, l16)
          sch.compute_at(
              block=b9, loop=l24, preserve_unit_loops=True, index=-1
          )
          sch.set_scope(block=b9, buffer_index=0, storage_scope="local")
          sch.compute_inline(block=b10)
          l25, l26, l27, l28, l29, l30, l31, l32 = sch.get_loops(
              block=b1
          )
          l33 = sch.fuse(l25, l26, l27, l28, preserve_unit_iters=True)
          v34 = sch.sample_categorical(
              candidates=[32, 64, 128, 256, 512, 1024],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=2,
          )
          l35, l36 = sch.split(
              loop=l33, factors=[None, v34], preserve_unit_iters=True
          )
          sch.bind(loop=l35, thread_axis="blockIdx.x")
          sch.bind(loop=l36, thread_axis="threadIdx.x")
          sch.compute_inline(block=b3)
          l37, l38, l39, l40, l41, l42 = sch.get_loops(block=b4)
          v43, v44 = sch.sample_perfect_tile(
              loop=l39, n=2, max_innermost_factor=64, decision=[28, 7]
          )
          l45, l46 = sch.split(
              loop=l39, factors=[v43, v44], preserve_unit_iters=True
          )
          v47, v48 = sch.sample_perfect_tile(
              loop=l40, n=2, max_innermost_factor=64, decision=[2, 32]
          )
          l49, l50 = sch.split(
              loop=l40, factors=[v47, v48], preserve_unit_iters=True
          )
          sch.unroll(loop=l37)
          sch.unroll(loop=l38)
          sch.unroll(loop=l41)
          sch.unroll(loop=l42)
          sch.reorder(l45, l49, l46, l50, l37, l38, l41, l42)
          l51, l52, l53, l54, l55, l56, l57, l58 = sch.get_loops(
              block=b4
          )
          l59 = sch.fuse(l51, l52, l53, l54, preserve_unit_iters=True)
          v60 = sch.sample_categorical(
              candidates=[32, 64, 128, 256, 512, 1024],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=4,
          )
          l61, l62 = sch.split(
              loop=l59, factors=[None, v60], preserve_unit_iters=True
          )
          sch.bind(loop=l61, thread_axis="blockIdx.x")
          sch.bind(loop=l62, thread_axis="threadIdx.x")
          sch.annotate(
              block_or_loop=b2,
              ann_key="meta_schedule.tiling_structure",
              ann_val="SSSRRSRS",
          )
          l63, l64, l65, l66, l67 = sch.get_loops(block=b2)
          v68, v69, v70, v71, v72 = sch.sample_perfect_tile(
              loop=l63,
              n=5,
              max_innermost_factor=64,
              decision=[1, 2, 3, 1, 1],
          )
          l73, l74, l75, l76, l77 = sch.split(
              loop=l63,
              factors=[v68, v69, v70, v71, v72],
              preserve_unit_iters=True,
          )
          v78, v79, v80, v81, v82 = sch.sample_perfect_tile(
              loop=l64,
              n=5,
              max_innermost_factor=64,
              decision=[6, 1, 1, 1, 1],
          )
          l83, l84, l85, l86, l87 = sch.split(
              loop=l64,
              factors=[v78, v79, v80, v81, v82],
              preserve_unit_iters=True,
          )
          v88, v89, v90, v91, v92 = sch.sample_perfect_tile(
              loop=l65,
              n=5,
              max_innermost_factor=64,
              decision=[7, 2, 1, 14, 1],
          )
          l93, l94, l95, l96, l97 = sch.split(
              loop=l65,
              factors=[v88, v89, v90, v91, v92],
              preserve_unit_iters=True,
          )
          v98, v99, v100, v101, v102 = sch.sample_perfect_tile(
              loop=l66,
              n=5,
              max_innermost_factor=64,
              decision=[4, 1, 16, 1, 1],
          )
          l103, l104, l105, l106, l107 = sch.split(
              loop=l66,
              factors=[v98, v99, v100, v101, v102],
              preserve_unit_iters=True,
          )
          v108, v109, v110 = sch.sample_perfect_tile(
              loop=l67,
              n=3,
              max_innermost_factor=64,
              decision=[2, 2, 16],
          )
          l111, l112, l113 = sch.split(
              loop=l67,
              factors=[v108, v109, v110],
              preserve_unit_iters=True,
          )
          sch.reorder(
              l73,
              l83,
              l93,
              l103,
              l74,
              l84,
              l94,
              l104,
              l75,
              l85,
              l95,
              l105,
              l111,
              l112,
              l76,
              l86,
              l96,
              l106,
              l113,
              l77,
              l87,
              l97,
              l107,
          )
          l114 = sch.fuse(l73, l83, l93, l103, preserve_unit_iters=True)
          sch.bind(loop=l114, thread_axis="blockIdx.x")
          l115 = sch.fuse(l74, l84, l94, l104, preserve_unit_iters=True)
          sch.bind(loop=l115, thread_axis="vthread.x")
          l116 = sch.fuse(l75, l85, l95, l105, preserve_unit_iters=True)
          sch.bind(loop=l116, thread_axis="threadIdx.x")
          sch.annotate(
              block_or_loop=b2,
              ann_key="meta_schedule.thread_extent_low_inclusive",
              ann_val=32,
          )
          sch.annotate(
              block_or_loop=b2,
              ann_key="meta_schedule.thread_extent_high_inclusive",
              ann_val=1024,
          )
          b117 = sch.cache_write(
              block=b2, write_buffer_index=0, storage_scope="local"
          )
          sch.reverse_compute_at(
              block=b117, loop=l116, preserve_unit_loops=True, index=-1
          )
          b118 = sch.cache_read(
              block=b2,
              read_buffer_index=0,
              storage_scope="shared",
              consumer_blocks=[b2],
          )
          sch.compute_at(
              block=b118, loop=l111, preserve_unit_loops=True, index=-1
          )
          (
              l119,
              l120,
              l121,
              l122,
              l123,
              l124,
              l125,
              l126,
          ) = sch.get_loops(block=b118)
          l127 = sch.fuse(
              l123, l124, l125, l126, preserve_unit_iters=True
          )
          v128 = sch.sample_categorical(
              candidates=[1, 2, 3, 4],
              probs=[0.25, 0.25, 0.25, 0.25],
              decision=3,
          )
          sch.annotate(
              block_or_loop=b118,
              ann_key="meta_schedule.cooperative_fetch",
              ann_val=v128,
          )
          b129 = sch.cache_read(
              block=b2,
              read_buffer_index=1,
              storage_scope="shared",
              consumer_blocks=[b2],
          )
          sch.compute_at(
              block=b129, loop=l111, preserve_unit_loops=True, index=-1
          )
          (
              l130,
              l131,
              l132,
              l133,
              l134,
              l135,
              l136,
              l137,
          ) = sch.get_loops(block=b129)
          l138 = sch.fuse(
              l134, l135, l136, l137, preserve_unit_iters=True
          )
          v139 = sch.sample_categorical(
              candidates=[1, 2, 3, 4],
              probs=[0.25, 0.25, 0.25, 0.25],
              decision=3,
          )
          sch.annotate(
              block_or_loop=b129,
              ann_key="meta_schedule.cooperative_fetch",
              ann_val=v139,
          )
          sch.reverse_compute_inline(block=b7)
          sch.reverse_compute_inline(block=b6)
          v140 = sch.sample_categorical(
              candidates=[0, 16, 64, 512, 1024],
              probs=[
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
                  0.20000000000000001,
              ],
              decision=4,
          )
          sch.annotate(
              block_or_loop=b8,
              ann_key="meta_schedule.unroll_explicit",
              ann_val=v140,
          )
          l141, l142, l143, l144 = sch.get_loops(block=b5)
          l145 = sch.fuse(
              l141, l142, l143, l144, preserve_unit_iters=True
          )
          v146 = sch.sample_categorical(
              candidates=[32, 64, 128, 256, 512, 1024],
              probs=[
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
                  0.16666666666666666,
              ],
              decision=2,
          )
          l147, l148 = sch.split(
              loop=l145, factors=[None, v146], preserve_unit_iters=True
          )
          sch.bind(loop=l147, thread_axis="blockIdx.x")
          sch.bind(loop=l148, thread_axis="threadIdx.x")
          sch.enter_postproc()
          sch.unannotate(
              block_or_loop=b118,
              ann_key="meta_schedule.cooperative_fetch",
          )
          l149, l150, l151, l152, l153 = sch.get_loops(block=b118)
          l154, l155, l156 = sch.split(
              loop=l153, factors=[None, 48, 4], preserve_unit_iters=True
          )
          sch.vectorize(loop=l156)
          sch.bind(loop=l155, thread_axis="threadIdx.x")
          sch.unannotate(
              block_or_loop=b129,
              ann_key="meta_schedule.cooperative_fetch",
          )
          l157, l158, l159, l160, l161 = sch.get_loops(block=b129)
          l162, l163, l164 = sch.split(
              loop=l161, factors=[None, 48, 4], preserve_unit_iters=True
          )
          sch.vectorize(loop=l164)
          sch.bind(loop=l163, thread_axis="threadIdx.x")
          b165 = sch.get_block(name="root", func_name="main")
          sch.unannotate(
              block_or_loop=b165,
              ann_key="meta_schedule.unroll_explicit",
          )
          (
              b166,
              b167,
              b168,
              b169,
              b170,
              b171,
              b172,
              b173,
          ) = sch.get_child_blocks(b165)
          l174, l175, l176, l177, l178, l179 = sch.get_loops(block=b166)
          sch.annotate(
              block_or_loop=l174,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l174,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l180, l181, l182, l183, l184, l185 = sch.get_loops(block=b167)
          sch.annotate(
              block_or_loop=l180,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l180,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l186, l187, l188, l189, l190, l191, l192 = sch.get_loops(
              block=b168
          )
          sch.annotate(
              block_or_loop=l186,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l186,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l193, l194, l195, l196, l197, l198, l199 = sch.get_loops(
              block=b169
          )
          sch.annotate(
              block_or_loop=l193,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l193,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          (
              l200,
              l201,
              l202,
              l203,
              l204,
              l205,
              l206,
              l207,
              l208,
              l209,
              l210,
              l211,
              l212,
              l213,
          ) = sch.get_loops(block=b170)
          sch.annotate(
              block_or_loop=l200,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l200,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l214, l215, l216, l217, l218, l219, l220 = sch.get_loops(
              block=b171
          )
          sch.annotate(
              block_or_loop=l214,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l214,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l221, l222, l223, l224, l225, l226 = sch.get_loops(block=b172)
          sch.annotate(
              block_or_loop=l221,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l221,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          l227, l228 = sch.get_loops(block=b173)
          sch.annotate(
              block_or_loop=l227,
              ann_key="pragma_auto_unroll_max_step",
              ann_val=1024,
          )
          sch.annotate(
              block_or_loop=l227,
              ann_key="pragma_unroll_explicit",
              ann_val=1,
          )
          b229 = sch.get_block(name="data_pack", func_name="main")
          l230, l231, l232, l233, l234, l235 = sch.get_loops(block=b229)
          b236 = sch.decompose_reduction(block=b229, loop=l234)
          b237 = sch.get_block(name="bgemm", func_name="main")
          (
              l238,
              l239,
              l240,
              l241,
              l242,
              l243,
              l244,
              l245,
              l246,
              l247,
              l248,
              l249,
              l250,
              l251,
          ) = sch.get_loops(block=b237)
          b252 = sch.decompose_reduction(block=b237, loop=l241)
          b253 = sch.get_block(name="inverse", func_name="main")
          l254, l255, l256, l257, l258, l259 = sch.get_loops(block=b253)
          b260 = sch.decompose_reduction(block=b253, loop=l258)
  
      verify(
          Conv2dWinogradAddRelu,
          apply_trace,
          Conv2dWinogradAddResidualRelu,
          "cuda",
          Conv2dWinogradAddResidualRelu_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 8 for calling a global function with example usage: test_winograd_gpu() without return types