code: def _get_threads(ib, nthread_tx, nthread_bx, nthread_by, nthread_bz):
      tx = te.thread_axis("threadIdx.x")
      bx = te.thread_axis("blockIdx.x")
      ib.scope_attr(tx, "thread_extent", nthread_tx)
      ib.scope_attr(bx, "thread_extent", nthread_bx)
  
      by = te.thread_axis("blockIdx.y")
      bz = te.thread_axis("blockIdx.z")
      ib.scope_attr(by, "thread_extent", nthread_by)
      ib.scope_attr(bz, "thread_extent", nthread_bz)
  
      return tx, bx, by, bz
  
imports: import tvm
  from tvm import te
  from .injective import schedule_injective_from_existing
  from ..transform import strided_slice, transpose
  from ..None import tag
  from ..utils import ceil_div, swap
  from ..math import cast, ceil_log2
label_desc: Use image node_id 2 for calling a global function with example usage: _get_threads(ib, nthread_tx, nthread_bx, nthread_by, nthread_bz) and returns: tx,  bx,  by,  bz