numba.cuda.cudadrv.error.NvvmError: Failed to verify
Error executing job with overrides: ['model.train_ds.manifest_filepath=data/train.json', 'model.validation_ds.manifest_filepath=data/dev.json', 'model.tokenizer.update_tokenizer=False', 'trainer.devices=-1', 'trainer.accelerator=gpu', 'trainer.max_epochs=50']
Traceback (most recent call last):
  File "/home/wangsang/work/data/NeMo/examples/asr/speech_to_text_finetune.py", line 215, in main
    trainer.fit(asr_model)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/trainer.py", line 532, in fit
    call._call_and_handle_interrupt(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/call.py", line 42, in _call_and_handle_interrupt
    return trainer.strategy.launcher.launch(trainer_fn, *args, trainer=trainer, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/strategies/launchers/subprocess_script.py", line 93, in launch
    return function(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/trainer.py", line 571, in _fit_impl
    self._run(model, ckpt_path=ckpt_path)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/trainer.py", line 980, in _run
    results = self._run_stage()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/trainer.py", line 1023, in _run_stage
    self.fit_loop.run()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/fit_loop.py", line 202, in run
    self.advance()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/fit_loop.py", line 355, in advance
    self.epoch_loop.run(self._data_fetcher)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/training_epoch_loop.py", line 133, in run
    self.advance(data_fetcher)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/training_epoch_loop.py", line 219, in advance
    batch_output = self.automatic_optimization.run(trainer.optimizers[0], kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/optimization/automatic.py", line 188, in run
    self._optimizer_step(kwargs.get("batch_idx", 0), closure)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/optimization/automatic.py", line 266, in _optimizer_step
    call._call_lightning_module_hook(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/call.py", line 145, in _call_lightning_module_hook
    output = fn(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/core/module.py", line 1270, in optimizer_step
    optimizer.step(closure=optimizer_closure)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/core/optimizer.py", line 161, in step
    step_output = self._strategy.optimizer_step(self._optimizer, closure, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/strategies/ddp.py", line 257, in optimizer_step
    optimizer_output = super().optimizer_step(optimizer, closure, model, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/strategies/strategy.py", line 231, in optimizer_step
    return self.precision_plugin.optimizer_step(optimizer, model=model, closure=closure, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/plugins/precision/precision_plugin.py", line 116, in optimizer_step
    return optimizer.step(closure=closure, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/optim/lr_scheduler.py", line 75, in wrapper
    return wrapped(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/optim/optimizer.py", line 76, in _use_grad
    ret = func(self, *args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/optim/adamw.py", line 164, in step
    loss = closure()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/plugins/precision/precision_plugin.py", line 103, in _wrap_closure
    closure_result = closure()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/optimization/automatic.py", line 142, in __call__
    self._result = self.closure(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/utils/_contextlib.py", line 115, in decorate_context
    return func(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/optimization/automatic.py", line 128, in closure
    step_output = self._step_fn()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/loops/optimization/automatic.py", line 315, in _training_step
    training_step_output = call._call_strategy_hook(trainer, "training_step", *kwargs.values())
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/trainer/call.py", line 293, in _call_strategy_hook
    output = fn(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/strategies/ddp.py", line 330, in training_step
    return self.model(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1511, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1520, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/parallel/distributed.py", line 1523, in forward
    else self._run_ddp_forward(*inputs, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/parallel/distributed.py", line 1359, in _run_ddp_forward
    return self.module(*inputs, **kwargs)  # type: ignore[index]
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1511, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1520, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/pytorch_lightning/overrides/base.py", line 90, in forward
    output = self._forward_module.training_step(*inputs, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/utils/model_utils.py", line 381, in wrap_training_step
    output_dict = wrapped(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/models/rnnt_models.py", line 709, in training_step
    loss_value, wer, _, _ = self.joint(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1511, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1520, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/core/classes/common.py", line 1064, in __call__
    outputs = wrapped(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/modules/rnnt.py", line 1452, in forward
    loss_batch = self.loss(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1511, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1520, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/core/classes/common.py", line 1064, in __call__
    outputs = wrapped(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/losses/rnnt.py", line 491, in forward
    loss = self._loss(acts=log_probs, labels=targets, act_lens=input_lengths, label_lens=target_lengths)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1511, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/nn/modules/module.py", line 1520, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/rnnt_pytorch.py", line 435, in forward
    return self.loss(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/torch/autograd/function.py", line 553, in apply
    return super().apply(*args, **kwargs)  # type: ignore[misc]
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/rnnt_pytorch.py", line 62, in forward
    loss_func(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/rnnt.py", line 223, in rnnt_loss_gpu
    status = wrapper.cost_and_grad(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/gpu_rnnt.py", line 252, in cost_and_grad
    return self.compute_cost_and_score(acts, grads, costs, pad_labels, label_lengths, input_lengths)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/gpu_rnnt.py", line 161, in compute_cost_and_score
    self.log_softmax(acts, denom)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/gpu_rnnt.py", line 107, in log_softmax
    reduce.reduce_max(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/reduce.py", line 353, in reduce_max
    return ReduceHelper(
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/nemo/collections/asr/parts/numba/rnnt_loss/utils/cuda_utils/reduce.py", line 294, in ReduceHelper
    _reduce_rows[grid_size, CTA_REDUCE_SIZE, stream, 0](I_opid, R_opid, acts, output, num_rows)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/dispatcher.py", line 539, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/dispatcher.py", line 673, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/dispatcher.py", line 681, in _compile_for_args
    return self.compile(tuple(argtypes))
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/dispatcher.py", line 924, in compile
    kernel = _Kernel(self.py_func, argtypes, **self.targetoptions)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/dispatcher.py", line 103, in __init__
    self.cooperative = 'cudaCGGetIntrinsicHandle' in lib.get_asm_str()
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/codegen.py", line 119, in get_asm_str
    return self._join_ptxes(self._get_ptxes(cc=cc))
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/codegen.py", line 137, in _get_ptxes
    ptxes = [nvvm.llvm_to_ptx(irs, **options)]
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/cudadrv/nvvm.py", line 632, in llvm_to_ptx
    return cu.compile(**opts)
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/cudadrv/nvvm.py", line 270, in compile
    self._try_error(err, 'Failed to verify\n')
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/cudadrv/nvvm.py", line 294, in _try_error
    self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
  File "/home/wangsang/miniconda3/envs/nemo/lib/python3.10/site-packages/numba/cuda/cudadrv/nvvm.py", line 198, in check_error
    raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to verify

<unnamed> (71, 5): parse expected type
NVVM_ERROR_INVALID_IR