code: def test_cuda_cbr():
      # fmt: off
      @T.prim_func
      def cbr_0(data: T.Buffer((1, 224, 224, 3), "float32"), kernel: T.Buffer((7, 7, 3, 64), "float32"), bias: T.Buffer(64, "float32"), bn_offset: T.Buffer(64, "float32"), bn_scale: T.Buffer(64, "float32"), compute: T.Buffer((1, 112, 112, 64), "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": 512})
              Conv2dOutput_local = T.alloc_buffer((1, 112, 112, 64), scope="local")
              PaddedInput_shared = T.alloc_buffer((1, 230, 230, 3), scope="shared")
              kernel_shared = T.alloc_buffer((7, 7, 3, 64), scope="shared")
              for nn_0_yy_0_xx_0_ff_0_fused in T.thread_binding(14, thread="blockIdx.x"):
                  for nn_1_yy_1_xx_1_ff_1_fused in T.thread_binding(4, thread="vthread.x"):
                      for nn_2_yy_2_xx_2_ff_2_fused in T.thread_binding(128, thread="threadIdx.x"):
                          for ry_0, rx_0, rc_0 in T.grid(7, 1, 3):
                              for ax0_ax1_ax2_ax3_fused in range(8251):
                                  with T.block("PaddedInput_shared"):
                                      v0 = T.axis.spatial(1, 0)
                                      v1 = T.axis.spatial(230, ry_0 + ax0_ax1_ax2_ax3_fused // 37)
                                      v2 = T.axis.spatial(230, nn_0_yy_0_xx_0_ff_0_fused // 2 * 32 + ax0_ax1_ax2_ax3_fused % 37)
                                      v3 = T.axis.spatial(3, rc_0)
                                      T.reads(data[v0, v1 - 3, v2 - 3, v3])
                                      T.writes(PaddedInput_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 1})
                                      PaddedInput_shared[v0, v1, v2, v3] = T.if_then_else(3 <= v1 and v1 < 227 and 3 <= v2 and v2 < 227, data[v0, v1 - 3, v2 - 3, v3], T.float32(0))
                              for ax0_ax1_ax2_ax3_fused in range(224):
                                  with T.block("kernel_shared"):
                                      v0 = T.axis.spatial(7, ry_0)
                                      v1 = T.axis.spatial(7, ax0_ax1_ax2_ax3_fused // 32)
                                      v2 = T.axis.spatial(3, rc_0)
                                      v3 = T.axis.spatial(64, nn_0_yy_0_xx_0_ff_0_fused % 2 * 32 + ax0_ax1_ax2_ax3_fused % 32)
                                      T.reads(kernel[v0, v1, v2, v3])
                                      T.writes(kernel_shared[v0, v1, v2, v3])
                                      T.block_attr({"meta_schedule.cooperative_fetch": 1})
                                      kernel_shared[v0, v1, v2, v3] = kernel[v0, v1, v2, v3]
                              for ry_1, rx_1, rc_1, nn_3, yy_3, xx_3, ff_3, ry_2, rx_2, rc_2, nn_4, yy_4, xx_4, ff_4 in T.grid(1, 1, 1, 1, 1, 1, 2, 1, 7, 1, 1, 7, 1, 8):
                                  with T.block("Conv2dOutput"):
                                      v_nn = T.axis.spatial(1, nn_3 + nn_4)
                                      v_yy = T.axis.spatial(112, nn_1_yy_1_xx_1_ff_1_fused // 2 * 56 + nn_2_yy_2_xx_2_ff_2_fused // 16 * 7 + yy_3 * 7 + yy_4)
                                      v_xx = T.axis.spatial(112, nn_0_yy_0_xx_0_ff_0_fused // 2 * 16 + nn_2_yy_2_xx_2_ff_2_fused % 16 + xx_3 + xx_4)
                                      v_ff = T.axis.spatial(64, nn_0_yy_0_xx_0_ff_0_fused % 2 * 32 + nn_1_yy_1_xx_1_ff_1_fused % 2 * 16 + ff_3 * 8 + ff_4)
                                      v_ry = T.axis.reduce(7, ry_0 + ry_1 + ry_2)
                                      v_rx = T.axis.reduce(7, rx_0 * 7 + rx_1 * 7 + rx_2)
                                      v_rc = T.axis.reduce(3, rc_0 + rc_1 + rc_2)
                                      T.reads(PaddedInput_shared[v_nn, v_yy * 2 + v_ry, v_xx * 2 + v_rx, v_rc], kernel_shared[v_ry, v_rx, v_rc, v_ff])
                                      T.writes(Conv2dOutput_local[v_nn, v_yy, v_xx, v_ff])
                                      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():
                                          Conv2dOutput_local[v_nn, v_yy, v_xx, v_ff] = T.float32(0)
                                      Conv2dOutput_local[v_nn, v_yy, v_xx, v_ff] = Conv2dOutput_local[v_nn, v_yy, v_xx, v_ff] + PaddedInput_shared[v_nn, v_yy * 2 + v_ry, v_xx * 2 + v_rx, v_rc] * kernel_shared[v_ry, v_rx, v_rc, v_ff]
                          for ax0, ax1, ax2, ax3 in T.grid(1, 7, 1, 16):
                              with T.block("Conv2dOutput_local"):
                                  v0 = T.axis.spatial(1, ax0)
                                  v1 = T.axis.spatial(112, nn_1_yy_1_xx_1_ff_1_fused // 2 * 56 + nn_2_yy_2_xx_2_ff_2_fused // 16 * 7 + ax1)
                                  v2 = T.axis.spatial(112, nn_0_yy_0_xx_0_ff_0_fused // 2 * 16 + nn_2_yy_2_xx_2_ff_2_fused % 16 + ax2)
                                  v3 = T.axis.spatial(64, nn_0_yy_0_xx_0_ff_0_fused % 2 * 32 + nn_1_yy_1_xx_1_ff_1_fused % 2 * 16 + ax3)
                                  T.reads(Conv2dOutput_local[v0, v1, v2, v3], bias[v3], bn_scale[v3], bn_offset[v3])
                                  T.writes(compute[v0, v1, v2, v3])
                                  compute[v0, v1, v2, v3] = T.max((Conv2dOutput_local[v0, v1, v2, v3] + bias[v3]) * bn_scale[v3] + bn_offset[v3], T.float32(0))
      # fmt: on
      decision_0 = [
          ("SamplePerfectTile", [1, 1, 1, 1, 1]),
          ("SamplePerfectTile", [1, 2, 8, 1, 7]),
          ("SamplePerfectTile", [7, 1, 16, 1, 1]),
          ("SamplePerfectTile", [2, 2, 1, 2, 8]),
          ("SamplePerfectTile", [7, 1, 1]),
          ("SamplePerfectTile", [1, 1, 7]),
          ("SamplePerfectTile", [3, 1, 1]),
          ("SampleCategorical", 0),
          ("SampleCategorical", 0),
          ("SampleCategorical", 3),
      ]
      mod = create_te_workload("CBR", 0)
      actual = _design_space(mod)
      check_sketches(
          mod,
          sketches=actual,
          expected_mods=[cbr_0],
          expected_decisions=[decision_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 14 for calling a global function with example usage: test_cuda_cbr() without return types