docstr: Compute both the loss and the gradients.
  
  Args:
      acts: A flattened tensor of shape [B, T, U, V+1] representing the activation matrix.
      grad: A flattented zero tensor of same shape as acts.
      costs: A zero vector of length B which will be updated inplace with the log probability costs.
      flat_labels: A flattened matrix of labels of shape [B, U]
      label_lengths: A vector of length B that contains the original lengths of the acoustic sequence.
      input_lengths: A vector of length B that contains the original lengths of the target sequence.
  
  Updates:
      This will launch kernels that will update inline the following variables:
      -   grads: Gradients of the activation matrix wrt the costs vector.
      -   costs: Negative log likelihood of the forward variable.
  
  Returns:
      An enum that either represents a successful RNNT operation or failure.
docstr_tok: ["Compute","both","the","loss","and","the","gradients",".","Args",":","acts",":","A","flattened","tensor","of","shape","[","B",",","T",",","U",",","V+1","]","representing","the","activation","matrix",".","grad",":","A","flattented","zero","tensor","of","same","shape","as","acts",".","costs",":","A","zero","vector","of","length","B","which","will","be","updated","inplace","with","the","log","probability","costs",".","flat_labels",":","A","flattened","matrix","of","labels","of","shape","[","B",",","U","]","label_lengths",":","A","vector","of","length","B","that","contains","the","original","lengths","of","the","acoustic","sequence",".","input_lengths",":","A","vector","of","length","B","that","contains","the","original","lengths","of","the","target","sequence",".","Updates",":","This","will","launch","kernels","that","will","update","inline","the","following","variables",":","-","grads",":","Gradients","of","the","activation","matrix","wrt","the","costs","vector",".","-","costs",":","Negative","log","likelihood","of","the","forward","variable",".","Returns",":","An","enum","that","either","represents","a","successful","RNNT","operation","or","failure","."]
code: def compute_cost_and_score(
      self,
      acts: torch.Tensor,
      grads: Optional[torch.Tensor],
      costs: torch.Tensor,
      labels: torch.Tensor,
      label_lengths: torch.Tensor,
      input_lengths: torch.Tensor,
  ) -> global_constants.RNNTStatus:
      """
      Compute both the loss and the gradients.
  
      Args:
          acts: A flattened tensor of shape [B, T, U, V+1] representing the activation matrix.
          grad: A flattented zero tensor of same shape as acts.
          costs: A zero vector of length B which will be updated inplace with the log probability costs.
          flat_labels: A flattened matrix of labels of shape [B, U]
          label_lengths: A vector of length B that contains the original lengths of the acoustic sequence.
          input_lengths: A vector of length B that contains the original lengths of the target sequence.
  
      Updates:
          This will launch kernels that will update inline the following variables:
          -   grads: Gradients of the activation matrix wrt the costs vector.
          -   costs: Negative log likelihood of the forward variable.
  
      Returns:
          An enum that either represents a successful RNNT operation or failure.
      """
      training = grads is not None
  
      if training:
          grads *= 0.0  # zero grads
  
      used_offset, (
          denom,
          alphas,
          betas,
          llForward,
          llBackward,
      ) = self._prepare_workspace()
  
      ######## START EXECUTION ########
      self.log_softmax(acts, denom)
  
      # Compute alphas
      gpu_rnnt_kernel.compute_alphas_kernel[
          self.minibatch_, self.maxU_, self.stream_, 0
      ](
          acts,
          denom,
          alphas,
          llForward,
          input_lengths,
          label_lengths,
          labels,
          self.minibatch_,
          self.maxT_,
          self.maxU_,
          self.alphabet_size_,
          self.blank_,
      )
  
      if training:
          # Compute betas
          gpu_rnnt_kernel.compute_betas_kernel[
              self.minibatch_, self.maxU_, self.stream_, 0
          ](
              acts,
              denom,
              betas,
              llBackward,
              input_lengths,
              label_lengths,
              labels,
              self.minibatch_,
              self.maxT_,
              self.maxU_,
              self.alphabet_size_,
              self.blank_,
          )
  
          # Compute gradient
          grad_blocks_per_grid = (
              self.minibatch_ * self.maxT_ * self.maxU_
          )
          grad_threads_per_block = gpu_rnnt_kernel.GPU_RNNT_THREAD_SIZE
          gpu_rnnt_kernel.compute_grad_kernel[
              grad_blocks_per_grid,
              grad_threads_per_block,
              self.stream_,
              0,
          ](
              grads,
              acts,
              denom,
              alphas,
              betas,
              llForward,
              input_lengths,
              label_lengths,
              labels,
              self.minibatch_,
              self.maxT_,
              self.maxU_,
              self.alphabet_size_,
              self.blank_,
              self.fastemit_lambda_,
              self.clamp_,
          )
  
      # // cost copy, negate (for log likelihood) and update with additional regularizers
      # This needs to be done via CUDA, because we used temporary memory llForward
      # passed to alpha, which was updated with log likelihoods.
      # But copying this data into a pytorch pointer is more difficult (numba api is one way)
      # Therefore launch a pointwise CUDA kernel to update the costs inplace from data of llForward
      # Then negate to compute the loglikelihood.
      threadsperblock = min(costs.shape[0], 32)
      blockspergrid = (
          costs.shape[0] + (threadsperblock - 1)
      ) // threadsperblock
      rnnt_helper.compute_costs_data[
          blockspergrid, threadsperblock, self.stream_, 0
      ](llForward, costs, self.fastemit_lambda_)
      self.stream_.synchronize()
  
      return global_constants.RNNTStatus.RNNT_STATUS_SUCCESS
  
imports: import multiprocessing
  import random
  from typing import Optional, Tuple
  import numba
  import torch
  from numba import cuda
  from nemo.collections.asr.parts.numba.rnnt_loss.utils import global_constants, rnnt_helper
  from nemo.collections.asr.parts.numba.rnnt_loss.utils.cuda_utils import gpu_rnnt_kernel, reduce
label_desc: Use image node_id 3 for calling the GPURNNT obj's underlying member method code with example usage: obj.compute_cost_and_score(acts, grads, costs, labels, label_lengths, input_lengths) and returns: global_constants