code: def lowered_loop_split(a: T.handle, b: T.handle) -> None:
      A = T.match_buffer(a, [128, 128], dtype="float32")
      B = T.match_buffer(b, [128], dtype="float32")
      reduce_temp0 = T.alloc_buffer(
          [1], dtype="float32", strides=[1], scope="local"
      )
      normal_reduce_temp0 = T.alloc_buffer(
          [1], dtype="float32", strides=[1], scope="local"
      )
      for i in T.serial(0, 128):
          for ki in T.thread_binding(0, 32, thread="threadIdx.x"):
              normal_reduce_temp0[0] = T.float32(0)
              for ko in T.serial(0, 4):
                  with T.block("B_normal_reduction"):
                      vi = T.axis.S(128, i)
                      vk = T.axis.R(128, ko * 32 + ki)
                      T.reads([A[vi, vk], normal_reduce_temp0[0]])
                      T.writes([normal_reduce_temp0[0]])
                      normal_reduce_temp0[0] = (
                          normal_reduce_temp0[0] + A[vi, vk]
                      )
              with T.block("B_cross_thread_reduction"):
                  T.reads([normal_reduce_temp0[0]])
                  T.writes([reduce_temp0[0]])
                  T.attr(
                      T.comm_reducer(
                          lambda x, y: x + y, [T.float32(0)]
                      ),
                      "reduce_scope",
                      T.reinterpret(T.uint64(0), dtype="handle"),
                  )
                  T.evaluate(
                      T.tvm_thread_allreduce(
                          T.uint32(1),
                          normal_reduce_temp0[0],
                          True,
                          reduce_temp0.data,
                          ki,
                          dtype="handle",
                      )
                  )
              with T.block("B_write_back"):
                  vi = T.axis.S(128, i)
                  T.reads([reduce_temp0[0]])
                  T.writes([B[vi]])
                  B[vi] = reduce_temp0[0]
  
imports: from tvm.script import tir
label_desc: Use image node_id 4 for calling a global function with example usage: lowered_loop_split(a, b) without return types