code: def ReduceHelper(
      I_opid: int,
      R_opid: int,
      acts: torch.Tensor,
      output: torch.Tensor,
      num_rows: int,
      num_cols: int,
      minus: bool,
      stream,
  ):
      """
      CUDA Warp reduction kernel helper which reduces via the R_Op.Add and writes
      the result to `output` according to I_op id.
  
      The result is stored in the blockIdx.
  
      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.
          R_opid: Operator ID for reduction. See R_Op for more information.
          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.
              Represents the number of threads per block.
          num_cols: Flattened shape of activation matrix, without vocabulary dimension
              (B * T * U). Represents number of blocks per grid.
          minus: Bool flag whether to add or subtract as reduction.
              If minus is set; calls _reduce_minus, else calls _reduce_rows kernel.
          stream: CUDA Stream.
      """
  
      if minus:
          grid_size = num_cols
          # call kernel
          _reduce_minus[grid_size, CTA_REDUCE_SIZE, stream, 0](
              I_opid, R_opid, acts, output, num_rows
          )
  
      else:
          grid_size = num_cols
          # call kernel
          _reduce_rows[grid_size, CTA_REDUCE_SIZE, stream, 0](
              I_opid, R_opid, acts, output, num_rows
          )
  
      return True
  
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 4 for calling a global function with example usage: ReduceHelper(I_opid, R_opid, acts, output, num_rows, num_cols, minus, stream) and returns: True