code: def test_compute_alphas_kernel(self, dtype):
      numba_utils.skip_numba_cuda_test_if_unsupported(
          __NUMBA_MINIMUM_VERSION__
      )
  
      random = np.random.RandomState(0)
      original_shape = [1, 5, 11, 3]
      B, T, U, V = original_shape
      threshold = 1e-5 if dtype == torch.float32 else 3e-4
  
      # Numpy kernel
      x = random.randn(*original_shape)
      labels = np.array([[1, 1, 1, 2, 2, 2, 1, 2, 2, 1]])  # [1, 10]
      label_len = len(labels[0]) + 1
      blank_idx = 0
  
      x_np = log_softmax(x, axis=-1)
      ground_alphas, ground_log_likelihood = rnnt_numpy.forward_pass(
          x_np[0, :, :label_len, :],
          labels[0, : label_len - 1],
          blank_idx,
      )
  
      # Pytorch kernel
      device = torch.device("cuda")
      if hasattr(cuda, "external_stream"):
          stream = cuda.external_stream(
              torch.cuda.current_stream(device).cuda_stream
          )
      else:
          stream = cuda.default_stream()
  
      x_c = torch.tensor(x, device=device, dtype=dtype)
      labels_c = torch.tensor(labels, device=device, dtype=torch.int64)
  
      # Allocate workspace memory
      denom = torch.zeros(B * T * U, device=device, dtype=x_c.dtype)
      alphas = torch.zeros(B * T * U, device=device, dtype=x_c.dtype)
      llForward = torch.zeros(B, device=device, dtype=x_c.dtype)
      input_lengths = torch.tensor(
          [T], dtype=torch.int64, device=device
      )
      label_lengths = torch.tensor(
          [len(labels[0])], dtype=torch.int64, device=device
      )
  
      # certify input data
      certify_inputs(x_c, labels_c, input_lengths, label_lengths)
  
      # flatten activation tensor (for pointer based indexing)
      x_c = x_c.view([-1])
  
      # call kernel
      # log softmax reduction
      reduce.reduce_max(
          x_c, denom, rows=V, cols=B * T * U, minus=False, stream=stream
      )
      reduce.reduce_exp(
          x_c, denom, rows=V, cols=B * T * U, minus=True, stream=stream
      )
  
      # alpha kernel
      gpu_rnnt_kernel.compute_alphas_kernel[B, U, stream, 0](
          x_c,
          denom,
          alphas,
          llForward,
          input_lengths,
          label_lengths,
          labels_c,
          B,
          T,
          U,
          V,
          blank_idx,
      )
  
      # sync kernel
      stream.synchronize()
  
      # reshape alphas
      alphas = alphas.view([B, T, U])
      diff = ground_alphas - alphas[0].cpu().numpy()
  
      assert np.abs(diff).mean() <= threshold
      assert np.square(diff).mean() <= (threshold**2)
  
      ll_diff = ground_log_likelihood - llForward[0].cpu().numpy()
  
      assert np.abs(ll_diff).mean() <= threshold
      assert np.square(ll_diff).mean() <= (threshold**2)
  
imports: import numpy
  import pytest
  import torch
  from numba import cuda
  from nemo.collections.asr.losses.rnnt_pytorch import MultiblankRNNTLossPytorch, TDTLossPytorch
  from nemo.collections.asr.parts.numba.rnnt_loss import rnnt_numpy
  from nemo.collections.asr.parts.numba.rnnt_loss.rnnt_pytorch import certify_inputs
  from nemo.collections.asr.parts.numba.rnnt_loss.utils.cuda_utils import gpu_rnnt_kernel, reduce
  from nemo.core.utils import numba_utils
  from nemo.core.utils.numba_utils import __NUMBA_MINIMUM_VERSION__
label_desc: Use image node_id 1 for calling the TestRNNTCUDAKernels obj's underlying member method code with example usage: obj.test_compute_alphas_kernel(dtype) without return types