code: def test_local_memory():
      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])
      AA = s.cache_read(A, "local", [B])
      o, i = s[B].split(s[B].op.axis[0], M)
      s[AA].compute_at(s[B], o)
      s[B].bind(o, te.thread_axis("blockIdx.x"))
  
      # local memory usage: M * 4B
      # thread usage: M
  
      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_local_memory_per_block=4 * M - 1,
                              max_threads_per_block=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_local_memory_per_block=4 * M,
                              max_threads_per_block=1,
                          ),
                      )
                  ]
              }
          ):
              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 3 for calling a global function with example usage: test_local_memory() without return types