code: def _calc_first_occurence_ir(
      argsorted_indices, inc_scan, first_occurence
  ):
      """Low level IR to calculate the first occurence of each unique element in the input data.
  
      Parameters
      ----------
      argsorted_indices : Buffer
          A buffer that stores the argsorted indices of the input data.
  
      inc_scan : Buffer
          A buffer that stores the inclusive scan of the binary tir.NE adjacent difference
          of the sorted data.
  
      first_occurence : Buffer
          A buffer that stores the first occurence of each unique element in the input data.
      """
      ib = tir.ir_builder.create()
      argsorted_indices_ptr = ib.buffer_ptr(argsorted_indices)
      inc_scan_ptr = ib.buffer_ptr(inc_scan)
      first_occurence_ptr = ib.buffer_ptr(first_occurence)
      batch_size = argsorted_indices.shape[0]
      max_threads = _get_max_threads(batch_size)
      with ib.new_scope():
          nthread_tx = max_threads
          nthread_bx = ceil_div(batch_size, max_threads)
          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)
          tid = bx * max_threads + tx
          with ib.if_scope(tid < batch_size):
              first_occurence_ptr[tid] = batch_size
      with ib.new_scope():
          nthread_tx = max_threads
          nthread_bx = ceil_div(batch_size, max_threads)
          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)
          tid = bx * max_threads + tx
          with ib.if_scope(tid < batch_size):
              with ib.if_scope(tid == 0):
                  first_occurence_ptr[
                      inc_scan_ptr[tid]
                  ] = argsorted_indices_ptr[tid]
              with ib.else_scope():
                  with ib.if_scope(
                      inc_scan_ptr[tid] != inc_scan_ptr[tid - 1]
                  ):
                      first_occurence_ptr[
                          inc_scan_ptr[tid]
                      ] = argsorted_indices_ptr[tid]
      return ib.get()
  
file_path: tvm/python/tvm/topi/cuda/unique.py
imports: import tvm
  from tvm import te, tir
  from ...te import hybrid
  from .scan import cumsum
  from .sort import sort, argsort
  from ..utils import ceil_div
label_desc: Use image node_id 6 for calling a global function with example usage: _calc_first_occurence_ir(argsorted_indices, inc_scan, first_occurence) and returns: ib