code: def cu_kernel_backward(
      log_probs, labels, beta, log_p, T, U, blank, lock
  ):
      """
      Compute backward pass for the forward-backward algorithm using Numba cuda kernel.
      Sequence Transduction with naive implementation : https://arxiv.org/pdf/1211.3711.pdf
  
      Arguments
      ---------
      log_probs : tensor
          4D Tensor of (batch x TimeLength x LabelLength x outputDim) from the Transducer network.
      labels : tensor
          2D Tensor of (batch x MaxSeqLabelLength) containing targets of the batch with zero padding.
      beta : tensor
          3D Tensor of (batch x TimeLength x LabelLength) for backward computation.
      log_p : tensor
          1D Tensor of (batch) for backward cost computation.
      T : tensor
          1D Tensor of (batch) containing TimeLength of each target.
      U : tensor
          1D Tensor of (batch) containing LabelLength of each target.
      blank : int
          Blank indice.
      lock : tensor
          2D Tensor of (batch x LabelLength) containing bool(1-0) lock for parallel computation.
      """
      # parallelize the forward algorithm over batch and target length dim
      b = cuda.blockIdx.x
      u = cuda.threadIdx.x
      t = T[b] - 1
      if u <= U[b]:
          # for each (B,U) Thread
          # wait the unlock of the next computation of beta[b,U+1,:]
          # Do the computation over the whole Time sequence on beta[B,U,:]
          # and then unlock the target U-1 for computation
          while t >= 0:
              if u == U[b]:
                  if t == T[b] - 1:
                      beta[b, t, u] = log_probs[b, t, u, blank]
                  else:
                      beta[b, t, u] = (
                          beta[b, t + 1, u] + log_probs[b, t, u, blank]
                      )
                  cuda.atomic.add(lock, (b, u - 1), -1)
                  t -= 1
              else:
                  if cuda.atomic.add(lock, (b, u), 0) < 0:
                      if t == T[b] - 1:
                          # do logsumexp between log_emit and log_no_emit
                          beta[b, t, u] = (
                              beta[b, t, u + 1]
                              + log_probs[b, t, u, labels[b, u]]
                          )
                      else:
                          # compute emission prob
                          emit = (
                              beta[b, t, u + 1]
                              + log_probs[b, t, u, labels[b, u]]
                          )
                          # compute no_emission prob
                          no_emit = (
                              beta[b, t + 1, u]
                              + log_probs[b, t, u, blank]
                          )
                          # do logsumexp between log_emit and log_no_emit
                          beta[b, t, u] = max(
                              no_emit, emit
                          ) + math.log1p(math.exp(-abs(no_emit - emit)))
                      if u > 0:
                          cuda.atomic.add(lock, (b, u - 1), -1)
                      cuda.atomic.add(lock, (b, u), 1)
                      t -= 1
      if u == 0:
          # for each thread b (utterance)
          # normalize the loss over time
          log_p[b] = beta[b, 0, 0] / T[b]
  
file_path: speechbrain/speechbrain/nnet/loss/transducer_loss.py
imports: import torch
  from torch.autograd import Function
  from torch.nn import Module
  import logging
  import math
  import warnings
label_desc: Use image node_id 2 for calling a global function with example usage: cu_kernel_backward(log_probs, labels, beta, log_p, T, U, blank, lock) without return types