code: def before(
      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)
      cross_thread_B = T.allocate([1], "float32", "local")
      threadIdx_z = T.launch_thread("threadIdx.z", 1)
      threadIdx_y = T.launch_thread("threadIdx.y", 2)
      threadIdx_x = T.launch_thread("threadIdx.x", 128)
      cross_thread_B_1 = T.Buffer(
          (1,), data=cross_thread_B, scope="local"
      )
      with T.attr(
          T.comm_reducer(lambda x0, y0: x0 + y0, [T.float32(0)]),
          "reduce_scope",
          T.reinterpret("handle", T.uint64(0)),
      ):
          A_1 = T.Buffer((256,), data=A.data)
          T.tvm_thread_allreduce(
              T.uint32(1),
              A_1[threadIdx_y * 128 + threadIdx_x],
              T.bool(True),
              cross_thread_B_1[0],
              threadIdx_x,
          )
      if threadIdx_x == 0:
          B_1 = T.Buffer((2,), data=B.data)
          B_1[threadIdx_y] = cross_thread_B_1[0]
  
imports: import tvm
  import tvm.testing
  from tvm.script import tir
label_desc: Use image node_id 1 for calling the TestMetalNoMask obj's underlying member method code with example usage: obj.before(A, B) without return types