code: def test_cuda_t2d():
      # fmt: off
      @T.prim_func
      def t2d_0(inputs: T.Buffer((1, 4, 4, 512), "float32"), weight: T.Buffer((4, 4, 512, 256), "float32"), conv2d_transpose_nhwc: T.Buffer((1, 8, 8, 256), "float32")) -> None:
          T.func_attr({"global_symbol": "main", "tir.noalias": T.bool(True)})
          with T.block("root"):
              T.reads()
              T.writes()
              T.block_attr({"meta_schedule.unroll_explicit": 64})
              conv2d_transpose_nhwc_local = T.alloc_buffer((1, 8, 8, 256), scope="local")
              PadInput_shared = T.alloc_buffer((1, 6, 6, 512), scope="shared")
              weight_shared = T.alloc_buffer((4, 4, 512, 256), scope="shared")
              for n_0_h_0_w_0_co_0_fused in T.thread_binding(256, thread="blockIdx.x"):
                  for n_1_h_1_w_1_co_1_fused in T.thread_binding(2, thread="vthread.x"):
                      for n_2_h_2_w_2_co_2_fused in T.thread_binding(1, thread="threadIdx.x"):
                          for rh_0, rw_0, rc_0 in T.grid(4, 1, 16):
                              for ax0_ax1_ax2_ax3_fused in range(rh_0 % 2 * 96 + 96):
                                  with T.block("PadInput_shared"):
                                      v0 = T.axis.spatial(1, 0)
                                      v1 = T.axis.spatial(6, n_0_h_0_w_0_co_0_fused // 64 + rh_0 // 2 + ax0_ax1_ax2_ax3_fused % (96 * (rh_0 % 2 + 1)) // 96)
                                      v2 = T.axis.spatial(6, n_0_h_0_w_0_co_0_fused % 64 // 16 + ax0_ax1_ax2_ax3_fused % 96 // 32)
                                      v3 = T.axis.spatial(512, rc_0 * 32 + ax0_ax1_ax2_ax3_fused % 32)
                                      T.reads(inputs[v0, v1 - 1, v2 - 1, v3])
                                      T.writes(PadInput_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 2})
                                      PadInput_shared[v0, v1, v2, v3] = T.if_then_else(1 <= v1 and v1 < 5 and 1 <= v2 and v2 < 5, inputs[v0, v1 - 1, v2 - 1, v3], T.float32(0))
                              for ax0_ax1_ax2_ax3_fused in range(2048):
                                  with T.block("weight_shared"):
                                      v0 = T.axis.spatial(4, rh_0 * -1 + 3)
                                      v1 = T.axis.spatial(4, ax0_ax1_ax2_ax3_fused // 512)
                                      v2 = T.axis.spatial(512, rc_0 * 32 + ax0_ax1_ax2_ax3_fused % 512 // 16)
                                      v3 = T.axis.spatial(256, n_0_h_0_w_0_co_0_fused % 16 * 16 + ax0_ax1_ax2_ax3_fused % 16)
                                      T.reads(weight[v0, v1, v2, v3])
                                      T.writes(weight_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 4})
                                      weight_shared[v0, v1, v2, v3] = weight[v0, v1, v2, v3]
                              for rh_1, rw_1, rc_1, n_3, h_3, w_3, co_3, rh_2, rw_2, rc_2, n_4, h_4, w_4, co_4 in T.grid(1, 1, 4, 1, 2, 1, 8, 1, 4, 8, 1, 1, 2, 1):
                                  with T.block("conv2d_transpose_nhwc"):
                                      v_n = T.axis.spatial(1, n_3 + n_4)
                                      v_h = T.axis.spatial(8, n_0_h_0_w_0_co_0_fused // 64 * 2 + h_3 + h_4)
                                      v_w = T.axis.spatial(8, n_0_h_0_w_0_co_0_fused % 64 // 16 * 2 + w_3 * 2 + w_4)
                                      v_co = T.axis.spatial(256, n_0_h_0_w_0_co_0_fused % 16 * 16 + n_1_h_1_w_1_co_1_fused * 8 + co_3 + co_4)
                                      v_rh = T.axis.reduce(4, rh_0 + rh_1 + rh_2)
                                      v_rw = T.axis.reduce(4, rw_0 * 4 + rw_1 * 4 + rw_2)
                                      v_rc = T.axis.reduce(512, rc_0 * 32 + rc_1 * 8 + rc_2)
                                      T.reads(PadInput_shared[v_n, (v_h + v_rh) // 2, (v_w + v_rw) // 2, v_rc], weight_shared[3 - v_rh, 3 - v_rw, v_rc, v_co])
                                      T.writes(conv2d_transpose_nhwc_local[v_n, v_h, v_w, v_co])
                                      T.block_attr({"meta_schedule.thread_extent_high_inclusive": 1024, "meta_schedule.thread_extent_low_inclusive": 32, "meta_schedule.tiling_structure": "SSSRRSRS"})
                                      with T.init():
                                          conv2d_transpose_nhwc_local[v_n, v_h, v_w, v_co] = T.float32(0)
                                      conv2d_transpose_nhwc_local[v_n, v_h, v_w, v_co] = conv2d_transpose_nhwc_local[v_n, v_h, v_w, v_co] + T.if_then_else((v_h + v_rh) % 2 == 0 and (v_w + v_rw) % 2 == 0, PadInput_shared[v_n, (v_h + v_rh) // 2, (v_w + v_rw) // 2, v_rc], T.float32(0)) * weight_shared[3 - v_rh, 3 - v_rw, v_rc, v_co]
                          for ax0, ax1, ax2, ax3 in T.grid(1, 2, 2, 8):
                              with T.block("conv2d_transpose_nhwc_local"):
                                  v0 = T.axis.spatial(1, ax0)
                                  v1 = T.axis.spatial(8, n_0_h_0_w_0_co_0_fused // 64 * 2 + ax1)
                                  v2 = T.axis.spatial(8, n_0_h_0_w_0_co_0_fused % 64 // 16 * 2 + ax2)
                                  v3 = T.axis.spatial(256, n_0_h_0_w_0_co_0_fused % 16 * 16 + n_1_h_1_w_1_co_1_fused * 8 + ax3)
                                  T.reads(conv2d_transpose_nhwc_local[v0, v1, v2, v3])
                                  T.writes(conv2d_transpose_nhwc[v0, v1, v2, v3])
                                  conv2d_transpose_nhwc[v0, v1, v2, v3] = conv2d_transpose_nhwc_local[v0, v1, v2, v3]
      # fmt: on
      decision_0 = [
          ("SamplePerfectTile", [1, 1, 1, 1, 1]),
          ("SamplePerfectTile", [4, 1, 1, 2, 1]),
          ("SamplePerfectTile", [4, 1, 1, 1, 2]),
          ("SamplePerfectTile", [16, 2, 1, 8, 1]),
          ("SamplePerfectTile", [4, 1, 1]),
          ("SamplePerfectTile", [1, 1, 4]),
          ("SamplePerfectTile", [16, 4, 8]),
          ("SampleCategorical", 1),
          ("SampleCategorical", 3),
          ("SampleCategorical", 2),
      ]
      mod = create_te_workload("T2D", 0)
      actual = _design_space(mod)
      check_sketches(
          mod,
          sketches=actual,
          expected_mods=[t2d_0],
          expected_decisions=[decision_0],
          debug_mask=0,
      )
  
file_path: tvm/tests/python/meta_schedule/test_meta_schedule_space_cuda.py
imports: from tvm import meta_schedule
  from tvm.meta_schedule.testing.space_generation import check_sketches, generate_design_space, print_sketches
  from tvm.meta_schedule.testing.te_workload import create_te_workload
  from tvm.script import tir
  from tvm.target import Target
label_desc: Use image node_id 11 for calling a global function with example usage: test_cuda_t2d() without return types