code: def test_cuda_mix_threaded_and_normal_reduction():
      def check(device, dtype, m=32, n=32):
          if not tvm.testing.device_enabled(device):
              print("Skipping", device)
              return
          dev = tvm.device(device, 0)
          if dtype == "float16" and not have_fp16(dev.compute_version):
              print("Skip because gpu does not have fp16 support")
              return
  
          a = tvm.te.placeholder((m, n), name="a", dtype=dtype)
          b = topi.sum(a)
          with tvm.target.Target(device):
              sb = tvm.te.create_schedule(b.op)
              i, _ = b.op.reduce_axis
              sb[b].bind(i, tvm.te.thread_axis("threadIdx.x"))
              func = tvm.build(sb, [a, b], device)
              a_np = np.random.uniform(size=(m, n)).astype(a.dtype)
              b_np = np.sum(a_np)
              a_nd = tvm.nd.array(a_np, dev)
              b_nd = tvm.nd.array(
                  np.zeros(b_np.shape, dtype=b_np.dtype), dev
              )
              func(a_nd, b_nd)
              tvm.testing.assert_allclose(b_nd.numpy(), b_np, rtol=1e-3)
  
      check("cuda", "float32")
      check("rocm", "float32")
      check("cuda", "float16")
  
imports: import os
  import re
  import tvm
  from tvm import te
  import numpy
  from tvm import topi
  from tvm.contrib.nvcc import have_fp16, have_int8, have_bf16
  from tvm.contrib import utils
  from tvm.script import tir
  import tvm.testing
  import pytest
label_desc: Use image node_id 15 for calling a global function with example usage: test_cuda_mix_threaded_and_normal_reduction() without return types