code: def test_allreduce_max(dims, target, dev):
      d1, d2, d3 = dims
      _, _, _d1, _d2, _d3 = reduce_max.params
      mod = reduce_max.specialize({_d1: d1, _d2: d2, _d3: d3})
      sch = tvm.tir.Schedule(mod)
      blk = sch.get_block("reduce")
      i, j, k, l = sch.get_loops(blk)
      sch.bind(i, "blockIdx.x")
      sch.bind(j, "threadIdx.z")
      sch.bind(k, "threadIdx.y")
      sch.bind(l, "threadIdx.x")
      f = tvm.build(sch.mod["main"], target=target)
  
      # prepare input and output array
      a_np = -np.random.rand(1, d1, d2, d3).astype("float32")
      b_np = a_np.max(axis=-1).astype("float32")
      a = tvm.nd.array(a_np, dev)
      b = tvm.nd.array(np.zeros_like(b_np), dev)
  
      # launch kernel
      f(a, b)
      tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-6, atol=1e-6)
  
file_path: tvm/tests/python/codegen/test_gpu_codegen_allreduce.py
imports: import tvm
  import tvm.testing
  import numpy
  from tvm.script import tir
  import pytest
label_desc: Use image node_id 7 for calling a global function with example usage: test_allreduce_max(dims, target, dev) without return types