code: def expected(
      A: T.Buffer((128, 128), "float32"), B: T.Buffer(128, "float32")
  ):
      T.func_attr({"target": T.target("cuda", host="llvm")})
      A_flat = T.decl_buffer(16384, data=A.data)
  
      for i in range(128):
          threadIdx_x = T.launch_thread("threadIdx.x", 32)
  
          normal_reduce_data = T.allocate([1], "float32", "local")
          normal_reduce = T.decl_buffer(
              1, data=normal_reduce_data, scope="local"
          )
  
          reduce_data = T.allocate([1], "float32", "local")
          reduce = T.decl_buffer(1, data=reduce_data, scope="local")
  
          normal_reduce[0] = T.float32(0)
          for ko in range(4):
              normal_reduce[0] = (
                  normal_reduce[0]
                  + A_flat[i * 128 + ko * 32 + threadIdx_x]
              )
          with T.attr(
              T.comm_reducer(lambda x, y: x + y, [T.float32(0)]),
              "reduce_scope",
              T.reinterpret("handle", T.uint64(0)),
          ):
              mask_data = T.allocate([1], "uint32", "local")
              mask = T.decl_buffer(
                  1, "uint32", data=mask_data, scope="local"
              )
  
              t0_data = T.allocate([1], "float32", "local")
              t0 = T.decl_buffer(1, data=t0_data, scope="local")
  
              reduce[0] = normal_reduce[0]
              mask[0] = T.tvm_warp_activemask()
  
              t0[0] = T.tvm_warp_shuffle_down(
                  mask[0], reduce[0], 16, 32, 32
              )
              reduce[0] = reduce[0] + t0[0]
              t0[0] = T.tvm_warp_shuffle_down(
                  mask[0], reduce[0], 8, 32, 32
              )
              reduce[0] = reduce[0] + t0[0]
              t0[0] = T.tvm_warp_shuffle_down(
                  mask[0], reduce[0], 4, 32, 32
              )
              reduce[0] = reduce[0] + t0[0]
              t0[0] = T.tvm_warp_shuffle_down(
                  mask[0], reduce[0], 2, 32, 32
              )
              reduce[0] = reduce[0] + t0[0]
              t0[0] = T.tvm_warp_shuffle_down(
                  mask[0], reduce[0], 1, 32, 32
              )
              reduce[0] = reduce[0] + t0[0]
              reduce[0] = T.tvm_warp_shuffle(
                  mask[0], reduce[0], 0, 32, 32
              )
          if threadIdx_x == 0:
              B[i] = reduce[0]
  
imports: import tvm
  import tvm.testing
  from tvm.script import tir
label_desc: Use image node_id 2 for calling the TestReduceSummation obj's underlying member method code with example usage: obj.expected(A, B) without return types