docstr: Compute both the loss and the gradients.
  
  Args:
      label_acts: A flattened tensor of shape [B, T, U, V] representing the activation matrix for tokens.
      duration_acts: A flattened tensor of shape [B, T, U, D] representing the activation matrix for durations.
      label_grad: A flattented zero tensor of same shape as label_acts.
      duration_grad: A flattented zero tensor of same shape as duration_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",":","label_acts",":","A","flattened","tensor","of","shape","[","B",",","T",",","U",",","V","]","representing","the","activation","matrix","for","tokens",".","duration_acts",":","A","flattened","tensor","of","shape","[","B",",","T",",","U",",","D","]","representing","the","activation","matrix","for","durations",".","label_grad",":","A","flattented","zero","tensor","of","same","shape","as","label_acts",".","duration_grad",":","A","flattented","zero","tensor","of","same","shape","as","duration_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,
      label_acts: torch.Tensor,
      duration_acts: torch.Tensor,
      label_grads: Optional[torch.Tensor],
      duration_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:
          label_acts: A flattened tensor of shape [B, T, U, V] representing the activation matrix for tokens.
          duration_acts: A flattened tensor of shape [B, T, U, D] representing the activation matrix for durations.
          label_grad: A flattented zero tensor of same shape as label_acts.
          duration_grad: A flattented zero tensor of same shape as duration_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 = label_grads is not None
  
      if training:
          label_grads *= 0.0  # zero grads
          duration_grads *= 0.0  # zero grads
  
      _, (
          denom,
          alphas,
          betas,
          llForward,
          llBackward,
          durations,
      ) = self._prepare_workspace()
  
      ######## START EXECUTION ########
      self.log_softmax(label_acts, denom)
  
      r = random.uniform(0, 1)
      if r < self.omega:
          # Compute alphas
          gpu_rnnt_kernel.compute_alphas_kernel[
              self.minibatch_, self.maxU_, self.stream_, 0
          ](
              label_acts,
              denom,
              alphas,
              llForward,
              input_lengths,
              label_lengths,
              labels,
              self.minibatch_,
              self.maxT_,
              self.maxU_,
              self.alphabet_size_,
              self.blank_,
          )
      else:
          # Compute alphas
          gpu_rnnt_kernel.compute_tdt_alphas_kernel[
              self.minibatch_, self.maxU_, self.stream_, 0
          ](
              label_acts,
              duration_acts,
              denom,
              self.sigma,
              alphas,
              llForward,
              input_lengths,
              label_lengths,
              labels,
              self.minibatch_,
              self.maxT_,
              self.maxU_,
              self.alphabet_size_,
              self.blank_,
              durations,
              self.num_durations,
          )
  
      if training:
          # Compute betas
          if r < self.omega:
              gpu_rnnt_kernel.compute_betas_kernel[
                  self.minibatch_, self.maxU_, self.stream_, 0
              ](
                  label_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,
              ](
                  label_grads,
                  label_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_,
              )
          else:
              gpu_rnnt_kernel.compute_tdt_betas_kernel[
                  self.minibatch_, self.maxU_, self.stream_, 0
              ](
                  label_acts,
                  duration_acts,
                  denom,
                  self.sigma,
                  betas,
                  llBackward,
                  input_lengths,
                  label_lengths,
                  labels,
                  self.minibatch_,
                  self.maxT_,
                  self.maxU_,
                  self.alphabet_size_,
                  self.blank_,
                  durations,
                  self.num_durations,
              )
  
              # 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_tdt_grad_kernel[
                  grad_blocks_per_grid,
                  grad_threads_per_block,
                  self.stream_,
                  0,
              ](
                  label_grads,
                  duration_grads,
                  label_acts,
                  duration_acts,
                  denom,
                  self.sigma,
                  alphas,
                  betas,
                  llForward,
                  input_lengths,
                  label_lengths,
                  labels,
                  self.minibatch_,
                  self.maxT_,
                  self.maxU_,
                  self.alphabet_size_,
                  self.blank_,
                  durations,
                  self.num_durations,
                  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 2 for calling the GPUTDT obj's underlying member method code with example usage: obj.compute_cost_and_score(label_acts, duration_acts, label_grads, duration_grads, costs, labels, label_lengths, input_lengths) and returns: global_constants