code: def test_num_thread():
      N = 1024
      M = 128
  
      A = te.placeholder((N,), name="A", dtype="float32")
      B = te.compute((N,), lambda i: A[i], name="B")
  
      s = te.create_schedule([B.op])
      o, i = s[B].split(s[B].op.axis[0], M)
  
      s[B].bind(o, te.thread_axis("threadIdx.x"))
      s[B].bind(i, te.thread_axis("threadIdx.y"))
  
      # shared memory usage: 0
      # thread usage: N
  
      for target in ["opencl", "cuda"]:
          if not tvm.testing.device_enabled(target):
              continue
  
          valid = [None]
          with tvm.transform.PassContext(
              config={
                  "tir.add_lower_pass": [
                      (
                          2,
                          get_verify_pass(
                              valid,
                              max_shared_memory_per_block=0,
                              max_threads_per_block=N - 1,
                          ),
                      )
                  ]
              }
          ):
              tvm.build(s, [A, B], target)
          assert not valid[0]
  
          with tvm.transform.PassContext(
              config={
                  "tir.add_lower_pass": [
                      (
                          2,
                          get_verify_pass(
                              valid,
                              max_shared_memory_per_block=0,
                              max_threads_per_block=N,
                          ),
                      )
                  ]
              }
          ):
              tvm.build(s, [A, B], target)
          assert valid[0]
  
          with tvm.transform.PassContext(
              config={
                  "tir.add_lower_pass": [
                      (
                          2,
                          get_verify_pass(
                              valid,
                              max_shared_memory_per_block=0,
                              max_threads_per_block=N,
                              max_thread_y=M - 1,
                          ),
                      )
                  ]
              }
          ):
              tvm.build(s, [A, B], target)
          assert not valid[0]
  
          with tvm.transform.PassContext(
              config={
                  "tir.add_lower_pass": [
                      (
                          2,
                          get_verify_pass(
                              valid,
                              max_shared_memory_per_block=0,
                              max_threads_per_block=N,
                              max_thread_y=M,
                          ),
                      )
                  ]
              }
          ):
              tvm.build(s, [A, B], target)
          assert valid[0]
  
file_path: tvm/tests/python/tir-analysis/test_tir_analysis_verify_gpu_code.py
imports: import tvm
  from tvm import te
  from tvm import topi
  import tvm.testing
  import tvm.topi.testing
label_desc: Use image node_id 4 for calling a global function with example usage: test_num_thread() without return types