Numba 0.50 depends on llvmlite 0.33, which uses LLVM 9 from the llvmdev package. With the switch to LLVM 9, a patch to LLVM is required to stop it auto-upgrading atomic intrinsics into a form that NVVM doesn't support (see https://github.com/numba/llvmlite/pull/593). However, the llvmdev conda-forge feedstock does not appear to carry this patch, so some CUDA atomics don't work when Numba and llvmlite are installed from conda-forge:
$ python -m numba.runtests numba.cuda.tests.cudapy.test_atomics
...EEEEEEEEEE..........................
----------------------------------------------------------------------
Ran 39 tests in 2.953s
FAILED (errors=10)
The requirement for this patch will be dropped in Numba 0.51 (when numba/numba#6030 is merged) so I don't know whether it's worth doing anything about this issue, or just waiting for Numba 0.51 / llvmlite 0.34 - I've opened the issue here for awareness though.
Full details of the failures:
```
======================================================================
ERROR: test_atomic_add_double (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 259, in test_atomic_add_double
cuda_func = cuda.jit('void(int64[:], float64[:])')(atomic_add_double)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(50, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_double_2 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 272, in test_atomic_add_double_2
cuda_func = cuda.jit('void(float64[:,:])')(atomic_add_double_2)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(62, 18): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_double_3 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 280, in test_atomic_add_double_3
cuda_func = cuda.jit('void(float64[:,:])')(atomic_add_double_3)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(62, 18): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_double_global (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 290, in test_atomic_add_double_global
cuda_func = cuda.jit('void(int64[:], float64[:])')(atomic_add_double_global)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(43, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_double_global_2 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 303, in test_atomic_add_double_global_2
cuda_func = cuda.jit('void(float64[:,:])')(atomic_add_double_global_2)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(35, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_double_global_3 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 311, in test_atomic_add_double_global_3
cuda_func = cuda.jit('void(float64[:,:])')(atomic_add_double_global_3)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(35, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_float (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 215, in test_atomic_add_float
cuda_atomic_add_float = cuda.jit('void(float32[:])')(atomic_add_float)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(50, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_float_2 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 227, in test_atomic_add_float_2
cuda_atomic_add2 = cuda.jit('void(float32[:,:])')(atomic_add_float_2)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(62, 18): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_float_3 (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 234, in test_atomic_add_float_3
cuda_atomic_add3 = cuda.jit('void(float32[:,:])')(atomic_add_float_3)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/decorators.py", line 103, in kernel_jit
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(62, 18): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
======================================================================
ERROR: test_atomic_add_returns_old (numba.cuda.tests.cudapy.test_atomics.TestCudaAtomics)
----------------------------------------------------------------------
Traceback (most recent call last):
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 489, in test_atomic_add_returns_old
self._test_atomic_returns_old(kernel, 10)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/tests/cudapy/test_atomics.py", line 478, in _test_atomic_returns_old
kernel[1, 1](x)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 833, in __call__
kernel = self.specialize(*args)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 844, in specialize
kernel = self.compile(argtypes)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 863, in compile
kernel.bind()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 604, in bind
self._func.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 480, in get
ptx = self.ptx.get()
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/compiler.py", line 451, in get
**self._extra_options)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 515, in llvm_to_ptx
ptx = cu.compile(**opts)
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 232, in compile
self._try_error(err, 'Failed to compile\n')
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 250, in _try_error
self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
File "/home/gmarkall/miniconda3/envs/numba-conda-forge/lib/python3.7/site-packages/numba/cuda/cudadrv/nvvm.py", line 140, in check_error
raise exc
numba.cuda.cudadrv.error.NvvmError: Failed to compile
(25, 17): parse expected binary operation in atomicrmw
NVVM_ERROR_COMPILATION
```
Issue:
Numba 0.50 depends on llvmlite 0.33, which uses LLVM 9 from the llvmdev package. With the switch to LLVM 9, a patch to LLVM is required to stop it auto-upgrading atomic intrinsics into a form that NVVM doesn't support (see https://github.com/numba/llvmlite/pull/593). However, the llvmdev conda-forge feedstock does not appear to carry this patch, so some CUDA atomics don't work when Numba and llvmlite are installed from conda-forge:
The requirement for this patch will be dropped in Numba 0.51 (when numba/numba#6030 is merged) so I don't know whether it's worth doing anything about this issue, or just waiting for Numba 0.51 / llvmlite 0.34 - I've opened the issue here for awareness though.
Full details of the failures:
Environment (
conda list
):Details about
conda
and system (conda info
):