code: def CTAReduce(tid: int, x, storage, count: int, R_opid: int):
      """
      CUDA Warp reduction kernel.
  
      It is a device kernel to be called by other kernels.
  
      The data will be read from the right segement recursively, and reduced (ROP) onto the left half.
      Operation continues while warp size is larger than a given offset.
      Beyond this offset, warp reduction is performed via `shfl_down_sync`, which halves the reduction
      space and sums the two halves at each call.
  
      Note:
          Efficient warp occurs at input shapes of 2 ^ K.
  
      References:
          - Warp Primitives [https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/]
  
      Args:
          tid: CUDA thread index
          x: activation. Single float.
          storage: shared memory of size CTA_REDUCE_SIZE used for reduction in parallel threads.
          count: equivalent to num_rows, which is equivalent to alphabet_size (V+1)
          R_opid: Operator ID for reduction. See R_Op for more information.
      """
      storage[tid] = x
  
      cuda.syncthreads()
  
      # Fold the data in half with each pass
      offset = CTA_REDUCE_SIZE // 2
      while offset >= warp_size:
          if (tid + offset) < count and tid < offset:
              # Read from the right half and store to the left half.
              if R_opid == 0:
                  x = rnnt_helper.add(x, storage[offset + tid])
              else:
                  x = rnnt_helper.maximum(x, storage[offset + tid])
  
              storage[tid] = x
  
          cuda.syncthreads()
          offset = offset // 2
  
      offset = warp_size // 2
      while offset > 0:
          # warp reduction and sync
          shuff = cuda.shfl_down_sync(0xFFFFFFFF, x, offset)
  
          if (tid + offset < count) and (tid < offset):
              if R_opid == 0:
                  x = rnnt_helper.add(x, shuff)
              else:
                  x = rnnt_helper.maximum(x, shuff)
  
          offset = offset // 2
  
      return x
  
file_path: NeMo/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/reduce.py
imports: import enum
  import math
  import torch
  from numba import cuda
  from nemo.collections.asr.parts.numba.rnnt_loss.utils import global_constants, rnnt_helper
label_desc: Use image node_id 1 for calling a global function with example usage: CTAReduce(tid, x, storage, count, R_opid) and returns: x