ERROR_INVALID_IR (6) regression in numba-cuda 0.29 with CCCL transform iterators

[leofang]

Summary

numba-cuda 0.29.0 (with numba 0.64.0) produces invalid NVVM IR when compiling CCCL's transform iterator examples and tests, causing cuda.bindings.nvvm.nvvmError: ERROR_INVALID_IR (6). This is a regression — these tests pass with numba-cuda 0.26.x.

(Note: 0.27.x had a separate, unrelated call_internal breakage that was fixed in 0.28 via #782.)

Failing tests

Consistent across all CI configs (GCC/MSVC, CTK 12.9/13.0/13.1, py3.10/py3.13, L4/H100):

• test_compute_examples_iterator_transform_iterator_basic
• test_compute_examples_iterator_transform_iterator_lambda
• compute/test_iterators.py::test_transform_iterator_with_lambda
• compute/test_iterators.py::test_transform_iterator_with_zip_iterator

Traceback

_, return_type = numba.cuda.compile(...)
  File ".../numba_cuda/numba/cuda/core/compiler_lock.py", line 74, in _acquire_compile_lock
  File ".../numba_cuda/numba/cuda/compiler.py", line 1233, in compile
  File ".../numba_cuda/numba/cuda/codegen.py", line 235, in get_asm_strs
  File ".../numba_cuda/numba/cuda/cudadrv/nvvm.py", line 563, in compile_ir
  File ".../numba_cuda/numba/cuda/cudadrv/nvvm.py", line 236, in verify
cuda.bindings.nvvm.nvvmError: ERROR_INVALID_IR (6)

Installed versions

• numba-cuda 0.29.0
• numba 0.64.0
• cuda-bindings 12.9.6

Links

• CCCL PR where this was discovered: Exclude numba-cuda 0.27.x instead of capping at <0.27.0 cccl#8180
• Example failing GHA job: https://github.com/NVIDIA/cccl/actions/runs/23571427217/job/68635892217

— Leo's bot

[leofang]

cc @gmarkall @shwina

— Leo's bot

[leofang]

cc @NaderAlAwar too (my bot thinks Nader == Graham 😞)