code: def test_cuda_nchw():
      # fmt: off
      @T.prim_func
      def cuda_nchw_0(data: T.Buffer((1, 64, 56, 56), "float32"), weight: T.Buffer((6, 6, 64, 64), "float32"), conv2d_winograd: T.Buffer((1, 64, 56, 56), "float32")) -> None:
          T.func_attr({"global_symbol": "main", "layout_free_buffers": [1], "tir.noalias": T.bool(True)})
          with T.block("root"):
              T.reads()
              T.writes()
              T.block_attr({"meta_schedule.unroll_explicit": 16})
              input_tile_local = T.alloc_buffer((64, 196, 6, 6), scope="local")
              data_pack = T.alloc_buffer((6, 6, 64, 196))
              bgemm = T.alloc_buffer((6, 6, 64, 196))
              inverse_local = T.alloc_buffer((64, 196, 4, 4), scope="local")
              data_pack_local = T.alloc_buffer((6, 6, 64, 196), scope="local")
              bgemm_local = T.alloc_buffer((6, 6, 64, 196), scope="local")
              data_pack_shared = T.alloc_buffer((6, 6, 64, 196), scope="shared")
              weight_shared = T.alloc_buffer((6, 6, 64, 64), scope="shared")
              for ci_p_fused_0 in T.thread_binding(25, thread="blockIdx.x"):
                  for ci_p_fused_1 in T.thread_binding(512, thread="threadIdx.x"):
                      for ax0, ax1, ax2, ax3 in T.grid(1, 1, 6, 6):
                          with T.block("input_tile"):
                              v_ci = T.axis.spatial(64, (ci_p_fused_0 * 512 + ci_p_fused_1) // 196 + ax0)
                              v_p = T.axis.spatial(196, (ci_p_fused_0 * 120 + ci_p_fused_1) % 196 + ax1)
                              v_eps, v_nu = T.axis.remap("SS", [ax2, ax3])
                              T.where(ci_p_fused_0 * 512 + ci_p_fused_1 < 12544)
                              T.reads(data[v_p // 196, v_ci, v_p % 196 // 14 * 4 + v_eps - 1, v_p % 14 * 4 + v_nu - 1])
                              T.writes(input_tile_local[v_ci, v_p, v_eps, v_nu])
                              T.block_attr({"schedule_rule": "None"})
                              input_tile_local[v_ci, v_p, v_eps, v_nu] = T.if_then_else(1 <= v_p % 196 // 14 * 4 + v_eps and v_p % 196 // 14 * 4 + v_eps < 57 and 1 <= v_p % 14 * 4 + v_nu and v_p % 14 * 4 + v_nu < 57, data[v_p // 196, v_ci, v_p % 196 // 14 * 4 + v_eps - 1, v_p % 14 * 4 + v_nu - 1], T.float32(0))
                      for eps in T.unroll(6):
                          for nu in T.unroll(6):
                              for r_a in T.unroll(6):
                                  for r_b in T.unroll(6):
                                      with T.block("data_pack"):
                                          v_eps, v_nu = T.axis.remap("SS", [eps, nu])
                                          v_ci = T.axis.spatial(64, (ci_p_fused_0 * 512 + ci_p_fused_1) // 196)
                                          v_p = T.axis.spatial(196, (ci_p_fused_0 * 512 + ci_p_fused_1) % 196)
                                          v_r_a, v_r_b = T.axis.remap("RR", [r_a, r_b])
                                          T.where(ci_p_fused_0 * 512 + ci_p_fused_1 < 12544)
                                          T.reads(input_tile_local[v_ci, v_p, v_r_a, v_r_b])
                                          T.writes(data_pack_local[v_eps, v_nu, v_ci, v_p])
                                          T.block_attr({"schedule_rule": "conv2d_nchw_winograd_data_pack"})
                                          with T.init():
                                              data_pack_local[v_eps, v_nu, v_ci, v_p] = T.float32(0)
                                          data_pack_local[v_eps, v_nu, v_ci, v_p] = data_pack_local[v_eps, v_nu, v_ci, v_p] + input_tile_local[v_ci, v_p, v_r_a, v_r_b] * T.Select(v_r_a % 6 == 5 and v_eps % 6 == 5, T.float32(1), T.Select(v_r_a % 6 == 5 and v_eps % 6 == 4, T.float32(0), T.Select(v_r_a % 6 == 5 and v_eps % 6 == 3, T.float32(0), T.Select(v_r_a % 6 == 5 and v_eps % 6 == 2, T.float32(0), T.Select(v_r_a % 6 == 5 and v_eps % 6 == 1, T.float32(0), T.Select(v_r_a % 6 == 5 and v_eps % 6 == 0, T.float32(0), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 5, T.float32(1.5), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 4, T.float32(1), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 3, T.float32(1), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 2, T.float32(1), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 1, T.float32(1), T.Select(v_r_a % 6 == 4 and v_eps % 6 == 0, T.float32(1), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 5, T.float32(-2), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 4, T.float32(-0.5), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 3, T.float32(2), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 2, T.float32(2.5), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 1, T.float32(0.5), T.Select(v_r_a % 6 == 3 and v_eps % 6 == 0, T.float32(1.5), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 5, T.float32(-1.5), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 4, T.float32(-1), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 3, T.float32(-1), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 2, T.float32(0.5), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 1, T.float32(-2.5), T.Select(v_r_a % 6 == 2 and v_eps % 6 == 0, T.float32(-2), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 5, T.float32(1), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 4, T.float32(0.5), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 3, T.float32(-2), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 2, T.float32(-1), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 1, T.float32(1), T.Select(v_r_a % 6 == 1 and v_eps % 6 == 0, T.float32(-1.5), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 5, T.float32(0), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 4, T.float32(0), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 3, T.float32(0), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 2, T.float32(0), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 1, T.float32(0), T.Select(v_r_a % 6 == 0 and v_eps % 6 == 0, T.float32(1), T.float32(0))))))))))))))))))))))))))))))))))))) * T.Select(v_r_b % 6 == 5 and v_nu % 6 == 5, T.float32(1), T.Select(v_r_b % 6 == 5 and v_nu % 6 == 4, T.float32(0), T.Select(v_r_b % 6 == 5 and v_nu % 6 == 3, T.float32(0), T.Select(v_r_b % 6 == 5 and v_nu % 6 == 2, T.float32(0), T.Select(v_r_b % 6 == 5 and v_nu % 6 == 1, T.float32(0), T.Select(v_r_b % 6 == 5 and v_nu % 6 == 0, T.float32(0), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 5, T.float32(1.5), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 4, T.float32(1), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 3, T.float32(1), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 2, T.float32(1), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 1, T.float32(1), T.Select(v_r_b % 6 == 4 and v_nu % 6 == 0, T.float32(1), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 5, T.float32(-2), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 4, T.float32(-0.5), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 3, T.float32(2), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 2, T.float32(2.5), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 1, T.float32(0.5), T.Select(v_r_b % 6 == 3 and v_nu % 6 == 0, T.float32(1.5), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 5, T.float32(-1.5), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 4, T.float32(-1), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 3, T.float32(-1), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 2, T.float32(0.5), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 1, T.float32(-2.5), T.Select(v_r_b % 6 == 2 and v_nu % 6 == 0, T.float32(-2), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 5, T.float32(1), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 4, T.float32(0.5), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 3, T.float32(-2), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 2, T.float32(-1), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 1, T.float32(1), T.Select(v_r_b % 6 == 1 and v_nu % 6 == 0, T.float32(-1.5), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 5, T.float32(0), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 4, T.float32(0), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 3, T.float32(0), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 2, T.float32(0), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 1, T.float32(0), T.Select(v_r_b % 6 == 0 and v_nu % 6 == 0, T.float32(1), T.float32(0)))))))))))))))))))))))))))))))))))))
                      for ax0, ax1, ax2, ax3 in T.grid(6, 6, 1, 1):
                          with T.block("data_pack_local"):
                              v0, v1 = T.axis.remap("SS", [ax0, ax1])
                              v2 = T.axis.spatial(64, (ci_p_fused_0 * 512 + ci_p_fused_1) // 196 + ax2)
                              v3 = T.axis.spatial(196, (ci_p_fused_0 * 120 + ci_p_fused_1) % 196 + ax3)
                              T.where(ci_p_fused_0 * 512 + ci_p_fused_1 < 12544)
                              T.reads(data_pack_local[v0, v1, v2, v3])
                              T.writes(data_pack[v0, v1, v2, v3])
                              data_pack[v0, v1, v2, v3] = data_pack_local[v0, v1, v2, v3]
              for eps_0_nu_0_co_0_p_0_fused in T.thread_binding(14, thread="blockIdx.x"):
                  for eps_1_nu_1_co_1_p_1_fused in T.thread_binding(224, thread="vthread.x"):
                      for eps_2_nu_2_co_2_p_2_fused in T.thread_binding(2, thread="threadIdx.x"):
                          for ci_0 in range(2):
                              for ax0_ax1_ax2_ax3_fused in range(32256):
                                  with T.block("data_pack_shared"):
                                      v0 = T.axis.spatial(6, ax0_ax1_ax2_ax3_fused // 5376)
                                      v1 = T.axis.spatial(6, ax0_ax1_ax2_ax3_fused % 5376 // 896)
                                      v2 = T.axis.spatial(64, ci_0 * 32 + ax0_ax1_ax2_ax3_fused % 896 // 28)
                                      v3 = T.axis.spatial(196, eps_0_nu_0_co_0_p_0_fused % 7 * 28 + ax0_ax1_ax2_ax3_fused % 28)
                                      T.reads(data_pack[v0, v1, v2, v3])
                                      T.writes(data_pack_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 4})
                                      data_pack_shared[v0, v1, v2, v3] = data_pack[v0, v1, v2, v3]
                              for ax0_ax1_ax2_ax3_fused in range(36864):
                                  with T.block("weight_shared"):
                                      v0 = T.axis.spatial(6, ax0_ax1_ax2_ax3_fused // 6144)
                                      v1 = T.axis.spatial(6, ax0_ax1_ax2_ax3_fused % 6144 // 1024)
                                      v2 = T.axis.spatial(64, ci_0 * 32 + ax0_ax1_ax2_ax3_fused % 1024 // 32)
                                      v3 = T.axis.spatial(64, eps_0_nu_0_co_0_p_0_fused // 7 * 32 + ax0_ax1_ax2_ax3_fused % 32)
                                      T.reads(weight[v0, v1, v2, v3])
                                      T.writes(weight_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 3})
                                      weight_shared[v0, v1, v2, v3] = weight[v0, v1, v2, v3]
                              for ci_1, eps_3, nu_3, co_3, p_3, ci_2, eps_4, nu_4, co_4, p_4 in T.grid(16, 2, 3, 1, 4, 2, 3, 1, 1, 1):
                                  with T.block("bgemm"):
                                      v_eps = T.axis.spatial(6, eps_3 * 3 + eps_4)
                                      v_nu = T.axis.spatial(6, eps_1_nu_1_co_1_p_1_fused // 112 * 3 + nu_3 + nu_4)
                                      v_co = T.axis.spatial(64, eps_0_nu_0_co_0_p_0_fused // 7 * 32 + eps_1_nu_1_co_1_p_1_fused % 112 // 7 * 2 + eps_2_nu_2_co_2_p_2_fused + co_3 + co_4)
                                      v_p = T.axis.spatial(196, eps_0_nu_0_co_0_p_0_fused % 7 * 28 + eps_1_nu_1_co_1_p_1_fused % 7 * 4 + p_3 + p_4)
                                      v_ci = T.axis.reduce(64, ci_0 * 32 + ci_1 * 2 + ci_2)
                                      T.reads(data_pack_shared[v_eps, v_nu, v_ci, v_p], weight_shared[v_eps, v_nu, v_ci, v_co])
                                      T.writes(bgemm_local[v_eps, v_nu, v_co, v_p])
                                      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():
                                          bgemm_local[v_eps, v_nu, v_co, v_p] = T.float32(0)
                                      bgemm_local[v_eps, v_nu, v_co, v_p] = bgemm_local[v_eps, v_nu, v_co, v_p] + data_pack_shared[v_eps, v_nu, v_ci, v_p] * weight_shared[v_eps, v_nu, v_ci, v_co]
                          for ax0, ax1, ax2, ax3 in T.grid(6, 3, 1, 4):
                              with T.block("bgemm_local"):
                                  v0 = T.axis.spatial(6, ax0)
                                  v1 = T.axis.spatial(6, eps_1_nu_1_co_1_p_1_fused // 112 * 3 + ax1)
                                  v2 = T.axis.spatial(64, eps_0_nu_0_co_0_p_0_fused // 7 * 32 + eps_1_nu_1_co_1_p_1_fused % 112 // 7 * 2 + eps_2_nu_2_co_2_p_2_fused + ax2)
                                  v3 = T.axis.spatial(196, eps_0_nu_0_co_0_p_0_fused % 7 * 28 + eps_1_nu_1_co_1_p_1_fused % 7 * 4 + ax3)
                                  T.reads(bgemm_local[v0, v1, v2, v3])
                                  T.writes(bgemm[v0, v1, v2, v3])
                                  bgemm[v0, v1, v2, v3] = bgemm_local[v0, v1, v2, v3]
              for n_co_h_0_w_0_fused_0 in T.thread_binding(196, thread="blockIdx.x"):
                  for n_co_h_0_w_0_fused_1 in T.thread_binding(64, thread="threadIdx.x"):
                      for ax0, ax1 in T.grid(1, 1):
                          for ax2 in T.unroll(4):
                              for ax3 in T.unroll(4):
                                  for ax4 in T.unroll(6):
                                      for ax5 in T.unroll(6):
                                          with T.block("inverse"):
                                              v_co = T.axis.spatial(64, (n_co_h_0_w_0_fused_0 * 64 + n_co_h_0_w_0_fused_1) // 196 + ax0)
                                              v_p = T.axis.spatial(196, (n_co_h_0_w_0_fused_0 * 64 + n_co_h_0_w_0_fused_1) % 196 + ax1)
                                              v_vh, v_vw, v_r_a, v_r_b = T.axis.remap("SSRR", [ax2, ax3, ax4, ax5])
                                              T.reads(bgemm[v_r_a, v_r_b, v_co, v_p])
                                              T.writes(inverse_local[v_co, v_p, v_vh, v_vw])
                                              T.block_attr({"schedule_rule": "conv2d_nchw_winograd_inverse"})
                                              with T.init():
                                                  inverse_local[v_co, v_p, v_vh, v_vw] = T.float32(0)
                                              inverse_local[v_co, v_p, v_vh, v_vw] = inverse_local[v_co, v_p, v_vh, v_vw] + bgemm[v_r_a, v_r_b, v_co, v_p] * T.Select(v_r_a % 6 == 5 and v_vh % 4 == 3, T.float32(1), T.Select(v_r_a % 6 == 5 and v_vh % 4 == 2, T.float32(0), T.Select(v_r_a % 6 == 5 and v_vh % 4 == 1, T.float32(0), T.Select(v_r_a % 6 == 5 and v_vh % 4 == 0, T.float32(0), T.Select(v_r_a % 6 == 4 and v_vh % 4 == 3, T.float32(-8), T.Select(v_r_a % 6 == 4 and v_vh % 4 == 2, T.float32(4), T.Select(v_r_a % 6 == 4 and v_vh % 4 == 1, T.float32(-2), T.Select(v_r_a % 6 == 4 and v_vh % 4 == 0, T.float32(1), T.Select(v_r_a % 6 == 3 and v_vh % 4 == 3, T.float32(0.125), T.Select(v_r_a % 6 == 3 and v_vh % 4 == 2, T.float32(0.25), T.Select(v_r_a % 6 == 3 and v_vh % 4 == 1, T.float32(0.5), T.Select(v_r_a % 6 == 3 and v_vh % 4 == 0, T.float32(1), T.Select(v_r_a % 6 == 2 and v_vh % 4 == 3, T.float32(1), T.Select(v_r_a % 6 == 2 and v_vh % 4 == 2, T.float32(1), T.Select(v_r_a % 6 == 2 and v_vh % 4 == 1, T.float32(1), T.Select(v_r_a % 6 == 2 and v_vh % 4 == 0, T.float32(1), T.Select(v_r_a % 6 == 1 and v_vh % 4 == 3, T.float32(-1), T.Select(v_r_a % 6 == 1 and v_vh % 4 == 2, T.float32(1), T.Select(v_r_a % 6 == 1 and v_vh % 4 == 1, T.float32(-1), T.Select(v_r_a % 6 == 1 and v_vh % 4 == 0, T.float32(1), T.Select(v_r_a % 6 == 0 and v_vh % 4 == 3, T.float32(0), T.Select(v_r_a % 6 == 0 and v_vh % 4 == 2, T.float32(0), T.Select(v_r_a % 6 == 0 and v_vh % 4 == 1, T.float32(0), T.Select(v_r_a % 6 == 0 and v_vh % 4 == 0, T.float32(1), T.float32(0))))))))))))))))))))))))) * T.Select(v_r_b % 6 == 5 and v_vw % 4 == 3, T.float32(1), T.Select(v_r_b % 6 == 5 and v_vw % 4 == 2, T.float32(0), T.Select(v_r_b % 6 == 5 and v_vw % 4 == 1, T.float32(0), T.Select(v_r_b % 6 == 5 and v_vw % 4 == 0, T.float32(0), T.Select(v_r_b % 6 == 4 and v_vw % 4 == 3, T.float32(-8), T.Select(v_r_b % 6 == 4 and v_vw % 4 == 2, T.float32(4), T.Select(v_r_b % 6 == 4 and v_vw % 4 == 1, T.float32(-2), T.Select(v_r_b % 6 == 4 and v_vw % 4 == 0, T.float32(1), T.Select(v_r_b % 6 == 3 and v_vw % 4 == 3, T.float32(0.125), T.Select(v_r_b % 6 == 3 and v_vw % 4 == 2, T.float32(0.25), T.Select(v_r_b % 6 == 3 and v_vw % 4 == 1, T.float32(0.5), T.Select(v_r_b % 6 == 3 and v_vw % 4 == 0, T.float32(1), T.Select(v_r_b % 6 == 2 and v_vw % 4 == 3, T.float32(1), T.Select(v_r_b % 6 == 2 and v_vw % 4 == 2, T.float32(1), T.Select(v_r_b % 6 == 2 and v_vw % 4 == 1, T.float32(1), T.Select(v_r_b % 6 == 2 and v_vw % 4 == 0, T.float32(1), T.Select(v_r_b % 6 == 1 and v_vw % 4 == 3, T.float32(-1), T.Select(v_r_b % 6 == 1 and v_vw % 4 == 2, T.float32(1), T.Select(v_r_b % 6 == 1 and v_vw % 4 == 1, T.float32(-1), T.Select(v_r_b % 6 == 1 and v_vw % 4 == 0, T.float32(1), T.Select(v_r_b % 6 == 0 and v_vw % 4 == 3, T.float32(0), T.Select(v_r_b % 6 == 0 and v_vw % 4 == 2, T.float32(0), T.Select(v_r_b % 6 == 0 and v_vw % 4 == 1, T.float32(0), T.Select(v_r_b % 6 == 0 and v_vw % 4 == 0, T.float32(1), T.float32(0)))))))))))))))))))))))))
                      for h_1, w_1 in T.grid(4, 4):
                          with T.block("conv2d_winograd"):
                              v_n = T.axis.spatial(1, 0)
                              v_co = T.axis.spatial(64, (n_co_h_0_w_0_fused_0 * 64 + n_co_h_0_w_0_fused_1) // 196)
                              v_h = T.axis.spatial(56, (n_co_h_0_w_0_fused_0 * 64 + n_co_h_0_w_0_fused_1) % 196 // 14 * 4 + h_1)
                              v_w = T.axis.spatial(56, (n_co_h_0_w_0_fused_0 * 64 + n_co_h_0_w_0_fused_1) % 14 * 4 + w_1)
                              T.reads(inverse_local[v_co, v_n * 196 + v_h // 4 * 14 + v_w // 4, v_h % 4, v_w % 4])
                              T.writes(conv2d_winograd[v_n, v_co, v_h, v_w])
                              conv2d_winograd[v_n, v_co, v_h, v_w] = inverse_local[v_co, v_n * 196 + v_h // 4 * 14 + v_w // 4, v_h % 4, v_w % 4]
      # fmt: on
      decision_0 = [
          ("SampleCategorical", 4),
          ("SamplePerfectTile", [1, 1, 1, 2, 3]),
          ("SamplePerfectTile", [1, 2, 1, 3, 1]),
          ("SamplePerfectTile", [2, 16, 2, 1, 1]),
          ("SamplePerfectTile", [7, 7, 1, 4, 1]),
          ("SamplePerfectTile", [2, 16, 2]),
          ("SampleCategorical", 3),
          ("SampleCategorical", 2),
          ("SampleCategorical", 1),
          ("SampleCategorical", 1),
      ]
      with _target():
          mod = create_te_workload("C2D_WIN_NCHW", 0)
      actual = _design_space(mod)
      check_sketches(
          mod,
          sketches=actual,
          expected_mods=[cuda_nchw_0],
          expected_decisions=[decision_0],
          debug_mask=0,
      )
  
file_path: tvm/tests/python/meta_schedule/test_meta_schedule_space_cuda_winograd.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 4 for calling a global function with example usage: test_cuda_nchw() without return types