code: def scan_thrust(
      data,
      output_dtype,
      exclusive=True,
      return_reduction=False,
      binop=tvm.tir.generic.add,
  ):
      """Do exclusive or inclusive scan on 1D or multidimensional input, using thrust.
  
      Parameters
      ----------
      data : tvm.te.Tensor
          Input data of any shape. The scan is done over the innermost axis.
  
      output_dtype: string
          The dtype of the output scan tensor.
  
      exclusive: bool, optional
          Whether or not do exclusive or inclusive scan.
  
      return_reduction: bool, optional
          Whether or not return a (N-1)-D tensor storing the reduction of each scan axis.
          Reductions are computed as part of the upsweep pass, so there is no extra cost.
          If False, reductions are ignored. It must be False when exclusive is False.
  
      binop: function, optional
          A binary associative op to use for scan. Since we need to lookup the corresponding
          thrust function, arbitrariy callables are not supported. Currently only
          tvm.tir.generic.add can be passed in.
  
      Returns
      -------
      output : tvm.te.Tensor
          A N-D tensor of the same rank N and shape as the input data.
  
      reduction : tvm.te.Tensor, optional
          (N-1)-D tensor storing the reduction of each scan axis.
          Returned if return_reduction is True.
      """
      data_buf = tvm.tir.decl_buffer(
          data.shape, data.dtype, "data_buf", data_alignment=8
      )
      output_buf = tvm.tir.decl_buffer(
          data.shape, output_dtype, "output_buf", data_alignment=8
      )
  
      output = te.extern(
          [data.shape],
          [data],
          lambda ins, outs: tvm.tir.call_packed(
              _get_thrust_func_name(binop), ins[0], outs[0], exclusive
          ),
          dtype=[output_dtype],
          in_buffers=[data_buf],
          out_buffers=[output_buf],
          name="exclusive_scan_thrust",
          tag="exclusive_scan_thrust_gpu",
      )
  
      if return_reduction:
          assert (
              exclusive
          ), "return_reduction should be False for inclusive scan"
          reduction = get_reduction_from_exclusive_scan(
              data, output, binop
          )
          return output, reduction
  
      return output
  
file_path: tvm/python/tvm/topi/cuda/scan.py
imports: from typing import Callable, Optional, Union
  import tvm
  from tvm import te
  from tvm.contrib.thrust import can_use_rocthrust, can_use_thrust
  from ..None import tag
  from ..math import cast, ceil_log2
  from ..transform import expand_dims, reshape, squeeze, transpose
  from ..utils import ceil_div, get_const_int, prod, swap
  from .injective import schedule_injective_from_existing
label_desc: Use image node_id 4 for calling a global function with example usage: scan_thrust(data, output_dtype, exclusive, return_reduction, binop) and returns: output, output,  reduction