code: def test_reduce_max(self, dtype):
      numba_utils.skip_numba_cuda_test_if_unsupported(
          __NUMBA_MINIMUM_VERSION__
      )
  
      random = np.random.RandomState(0)
      original_shape = [1, 5, 4, 3]
      x = random.randn(*original_shape).reshape([-1]).astype(dtype)
      dx = random.randn(*x.shape).astype(dtype)
  
      stream = cuda.stream()
      x_c = cuda.to_device(x, stream=stream)
      dx_c = cuda.to_device(dx, stream=stream)
  
      # call kernel
      cols = np.prod(original_shape[:3])
      reduce.reduce_max(
          x_c,
          dx_c,
          rows=original_shape[-1],
          cols=cols,
          minus=False,
          stream=stream,
      )
  
      # sync kernel
      stream.synchronize()
  
      dx_result = dx_c.copy_to_host(stream=stream)
      del x_c, dx_c
  
      # collect results in first [B * T * U] values; for all V
      assert np.abs(dx_result[cols:] - dx[cols:]).sum() <= 1e-7
      # make sure dx_result updates the [B * T * U] values
      assert np.abs(dx_result[:cols] - dx[:cols]).sum() > 0
  
file_path: NeMo/tests/collections/asr/numba/rnnt_loss/utils/test_reduce.py
imports: import numpy
  import pytest
  from numba import cuda
  from nemo.collections.asr.parts.numba.rnnt_loss.utils.cuda_utils import 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 TestRNNTCUDAReductions obj's underlying member method code with example usage: obj.test_reduce_max(dtype) without return types