code: def test_compute_costs_data(self, batch_size, fastemit_lambda, dtype):
      numba_utils.skip_numba_cuda_test_if_unsupported(
          __NUMBA_MINIMUM_VERSION__
      )
  
      np.random.seed(0)
      x = np.full([batch_size], fill_value=0.0)  # np.random.rand(8192)
      y = np.random.randn(batch_size).astype(
          dtype
      )  # np.random.rand(8192)
      threshold = 1e-5 if dtype == np.float32 else 1e-5
  
      stream = cuda.stream()
      x_c = cuda.to_device(x, stream=stream)
      y_c = cuda.to_device(y, stream=stream)
  
      # call kernel
      threads_per_block = min(x.shape[0], 32)
      blocks_per_grid = (
          x.shape[0] + (threads_per_block - 1)
      ) // threads_per_block
      # Kernel call (source, dest, extra_args_...)
      rnnt_helper.compute_costs_data[
          blocks_per_grid, threads_per_block, stream
      ](y_c, x_c, fastemit_lambda)
  
      # sync kernel
      stream.synchronize()
  
      x_new = x_c.copy_to_host(stream=stream)
      del x_c, y_c
  
      res = -(y.astype(np.float32).copy())
      res *= 1.0 + fastemit_lambda
  
      for i in range(len(x_new)):
          assert abs(x_new[i] - res[i]) < threshold, f"index failed {i}"
  
file_path: NeMo/tests/collections/asr/numba/rnnt_loss/utils/test_rnnt_helper.py
imports: import numpy
  import pytest
  from numba import cuda
  from nemo.collections.asr.parts.numba.rnnt_loss.utils import global_constants, rnnt_helper
  from nemo.core.utils import numba_utils
  from nemo.core.utils.numba_utils import __NUMBA_MINIMUM_VERSION__
label_desc: Use image node_id 10 for calling the TestRNNTHelper obj's underlying member method code with example usage: obj.test_compute_costs_data(batch_size, fastemit_lambda, dtype) without return types