code: def _reduce_minus(
      I_opid: int, R_opid: int, acts, output, num_rows: int
  ):
      """
      CUDA Warp reduction kernel which reduces via the R_Op.Add
  
      Reduces the input data such that I_Op = Exponential and R_op = Add.
      The result is stored in the blockIdx, and is stored as an exp op.
  
      Note:
          Efficient warp occurs at input shapes of 2 ^ K.
  
      References:
          - Warp Primitives
            [https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/]
  
      Args:
          I_opid: Operator ID for input. See I_Op for more information. For this kernel,
              the Exponential op is chosen in general, and therefore the input
              is reduced in place with scaling.
          R_opid: Operator ID for reduction. See R_Op for more information. For this
              kernel, generally Add op is chosen. It reduces the kernel via summation.
          acts: Flatened activation matrix of shape [B * T * U * (V+1)].
          output: Flatened output matrix of shape [B * T * U * (V+1)].
              Data will be overwritten.
          num_rows: Vocabulary size (including blank token) - V+1.
      """
  
      tid = cuda.threadIdx.x
      idx = tid
      col = cuda.blockIdx.x
  
      # allocate shared thread memory
      storage = cuda.shared.array(
          shape=(CTA_REDUCE_SIZE,), dtype=acts.dtype
      )
  
      max = output[col]
  
      # // Each block works on a column
      if idx < num_rows:
          curr = acts[col * num_rows + idx] - max
          if I_opid == 0:
              curr = rnnt_helper.exponential(curr)
          else:
              curr = rnnt_helper.identity(curr)
  
      idx += CTA_REDUCE_SIZE
  
      while idx < num_rows:
          activation_ = acts[col * num_rows + idx] - max
  
          if I_opid == 0 and R_opid == 0:
              curr = rnnt_helper.add(
                  curr, rnnt_helper.exponential(activation_)
              )
          elif I_opid == 0 and R_opid == 1:
              curr = rnnt_helper.maximum(
                  curr, rnnt_helper.exponential(activation_)
              )
          elif I_opid == 1 and R_opid == 0:
              curr = rnnt_helper.add(
                  curr, rnnt_helper.identity(activation_)
              )
          else:
              curr = rnnt_helper.maximum(
                  curr, rnnt_helper.identity(activation_)
              )
  
          idx += CTA_REDUCE_SIZE
  
      # // Sum thread-totals over the CTA.
      curr = CTAReduce(tid, curr, storage, num_rows, R_opid)
  
      # // Store result in out (inplace, I_op: exponential)
      if tid == 0:
          output[col] = -max - math.log(curr)
  
file_path: espnet/espnet2/asr/transducer/rnnt_multi_blank/utils/cuda_utils/reduce.py
imports: import enum
  import math
  import torch
  from numba import cuda
  from espnet2.asr.transducer.rnnt_multi_blank.utils import global_constants, rnnt_helper
label_desc: Use image node_id 3 for calling a global function with example usage: _reduce_minus(I_opid, R_opid, acts, output, num_rows) without return types