code: def expected(
      A: T.Buffer((1, 1, 2, 128), "float32"),
      B: T.Buffer((1, 1, 2), "float32"),
  ):
      T.func_attr(
          {
              "target": T.target(
                  {
                      "kind": "metal",
                      "max_threads_per_block": 1024,
                      "thread_warp_size": 32,
                      "host": "llvm",
                  }
              ),
          }
      )
      blockIdx_x = T.launch_thread("blockIdx.x", 1)
      red_result = T.allocate([2], "float32", "shared")
      T.attr(red_result, "volatile_scope", 1)
      threadIdx_z = T.launch_thread("threadIdx.z", 1)
      threadIdx_y = T.launch_thread("threadIdx.y", 2)
      threadIdx_x = T.launch_thread("threadIdx.x", 128)
      red_result_1 = T.Buffer((2,), data=red_result, scope="shared")
      with T.attr(
          T.comm_reducer(lambda x0, y0: x0 + y0, [T.float32(0)]),
          "reduce_scope",
          T.reinterpret("handle", T.uint64(0)),
      ):
          red_buf0 = T.decl_buffer([1], "float32", scope="local")
          t0 = T.decl_buffer([1], "float32", scope="local")
          red_buf0_1 = T.decl_buffer([1], "float32", scope="local")
          t0_1 = T.decl_buffer([1], "float32", scope="local")
          red_buf_staging = T.decl_buffer(
              [8], "float32", scope="shared"
          )
          A_1 = T.Buffer((256,), data=A.data)
          red_buf0_1[0] = A_1[threadIdx_y * 128 + threadIdx_x]
          t0_1[0] = T.tvm_warp_shuffle_down(
              0, red_buf0_1[0], 16, 32, 32
          )
          red_buf0_1[0] = red_buf0_1[0] + t0_1[0]
          t0_1[0] = T.tvm_warp_shuffle_down(0, red_buf0_1[0], 8, 32, 32)
          red_buf0_1[0] = red_buf0_1[0] + t0_1[0]
          t0_1[0] = T.tvm_warp_shuffle_down(0, red_buf0_1[0], 4, 32, 32)
          red_buf0_1[0] = red_buf0_1[0] + t0_1[0]
          t0_1[0] = T.tvm_warp_shuffle_down(0, red_buf0_1[0], 2, 32, 32)
          red_buf0_1[0] = red_buf0_1[0] + t0_1[0]
          t0_1[0] = T.tvm_warp_shuffle_down(0, red_buf0_1[0], 1, 32, 32)
          red_buf0_1[0] = red_buf0_1[0] + t0_1[0]
          if threadIdx_x % 32 == 0:
              red_buf_staging[
                  threadIdx_y * 4 + threadIdx_x // 32
              ] = red_buf0_1[0]
          T.tvm_storage_sync("shared")
          if threadIdx_x < 4:
              red_buf0[0] = red_buf_staging[
                  threadIdx_y * 4 + threadIdx_x
              ]
          t0[0] = T.tvm_warp_shuffle_down(0, red_buf0[0], 2, 32, 32)
          red_buf0[0] = red_buf0[0] + t0[0]
          t0[0] = T.tvm_warp_shuffle_down(0, red_buf0[0], 1, 32, 32)
          red_buf0[0] = red_buf0[0] + t0[0]
          if threadIdx_x == 0:
              red_result_1[threadIdx_y] = red_buf0[0]
          T.tvm_storage_sync("shared")
      if threadIdx_x == 0:
          B_1 = T.Buffer((2,), data=B.data)
          B_1[threadIdx_y] = red_result_1[threadIdx_y]
  
imports: import tvm
  import tvm.testing
  from tvm.script import tir
label_desc: Use image node_id 2 for calling the TestMetalNoMask obj's underlying member method code with example usage: obj.expected(A, B) without return types