code: def test_log_sum_exp(self, dtype):
      numba_utils.skip_numba_cuda_test_if_unsupported(
          __NUMBA_MINIMUM_VERSION__
      )
  
      # wrapper kernel for device function that is tested
      @cuda.jit
      def _kernel(x, y):
          x_pos = cuda.grid(1)
          if x_pos < x.shape[0] and x_pos < y.shape[0]:
              x[x_pos] = rnnt_helper.log_sum_exp(x[x_pos], y[x_pos])
  
      x = np.zeros([8]).astype(dtype)  # np.random.rand(8192)
      y = np.ones([8]).astype(dtype)  # np.random.rand(8192)
      threshold = 1e-5 if dtype == np.float32 else 2e-3
  
      stream = cuda.stream()
      x_c = cuda.to_device(x, stream=stream)
      y_c = cuda.to_device(y, stream=stream)
  
      # call kernel
      threads_per_block = global_constants.threads_per_block()
      blocks_per_grid = (
          x.shape[0] + threads_per_block - 1
      ) // threads_per_block
      _kernel[blocks_per_grid, threads_per_block, stream](x_c, y_c)
  
      # sync kernel
      stream.synchronize()
  
      x_new = x_c.copy_to_host(stream=stream)
      del x_c, y_c
  
      assert (x_new.sum() - 10.506093500145782) <= threshold
  
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 1 for calling the TestRNNTHelper obj's underlying member method code with example usage: obj.test_log_sum_exp(dtype) without return types