code: def _calc_adjacent_diff_ir(data, output, binop=tir.Sub):
      """Low level IR to calculate adjacent difference in an 1-D array.
  
      Parameters
      ----------
      data : Buffer
          Input 1-D Buffer.
  
      output: Buffer
          A buffer to store adjacent difference, of the same shape as data. The adjacent difference
          is defined as: output[0] = 0, output[i] = binop(data[i], data[i-1])
          where i > 0 and i < len(data).
  
      binop: function, optional
          A binary associative op to use for calculating adjacent difference. The function takes two
          TIR expressions and produce a new TIR expression. By default it uses tvm.tir.Sub to
          compute the adjacent difference.
      """
      ib = tir.ir_builder.create()
      data_ptr = ib.buffer_ptr(data)
      output_ptr = ib.buffer_ptr(output)
      batch_size = data.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):
              with ib.if_scope(tid == 0):
                  output_ptr[tid] = 0
              with ib.else_scope():
                  output_ptr[tid] = tir.Cast(
                      output.dtype,
                      binop(data_ptr[tid], data_ptr[tid - 1]),
                  )
      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 2 for calling a global function with example usage: _calc_adjacent_diff_ir(data, output, binop) and returns: ib