Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Mostly CUDA: Replace llvmpy API usage with llvmlite APIs #6813

Merged
merged 8 commits into from Mar 15, 2021

Conversation

gmarkall
Copy link
Member

@gmarkall gmarkall commented Mar 10, 2021

This PR replaces all llvmpy API usage with llvmlite API usage in the CUDA target. I started this replacement work because I often want to do something based on existing code in the CUDA target, but I'm loathe to add new code using the compatibility layer, so I felt it would be better to make the entire CUDA target an exemplar of the right API. This spidered out of just the CUDA target because a lot of code is shared between the core and the CUDA target, so the changes are:

  • Remove all llvmpy usage in the CUDA target.
  • Replace all usages of the llvmpy Module class with the llvmlite Module class. I started by replacing only CUDA target modules with the llvmlite Module class, but since it passes through the core so much, with the possibility of a lack of test coverage for some basic functionality (simple mathematical operators, etc, might not be explicitly tested by the CUDA target, with the assumption that simple nopython mode things generally work) I thought the safest route was to make sure that no target uses an llvmpy Module.
  • cgutils now has some extra methods to replace functionality that was in the llvmpy Module class - these functions are now called where the old Module method was used. I didn't replace get_global_variable_named() with anything in cgutils because it was used exactly once.

I'm marking this as medium effort because although it's quite a long PR, the majority of it is very mechanical substitutions.

The llvmpy API exists only for compatibility; this commit replaces all
uses of the llvmpy APIs with llvmlite IR builder API calls in the CUDA
target. Notes:

- Some of the functionality in the `llvmlite.llvmpy.core.Module`
  subclass of `llvmlite.ir.Module` has been moved into new functions in
  cgutils.
- Because CUDA modules are now llvmlite.ir.Module instances, some of the
  core calls to `mod.get_or_insert_function` are replaced with calls to
  the new `cgutils` functions, as the CUDA modules also pass through
  these functions in the core.
@gmarkall gmarkall added 2 - In Progress CUDA CUDA related issue/PR labels Mar 10, 2021
This required a few changes because llvmlite Modules don't have
`get_global_variable_named()` and `add_function()` methods.
@gmarkall gmarkall added 3 - Ready for Review Effort - medium Medium size effort needed and removed 2 - In Progress labels Mar 10, 2021
@gmarkall gmarkall marked this pull request as ready for review March 10, 2021 15:30
@stuartarchibald
Copy link
Contributor

@gmarkall thanks for the patch, please could you resolve the conflict against mainline, many thanks!

@gmarkall
Copy link
Member Author

@stuartarchibald Conflicts resolved!

@stuartarchibald
Copy link
Contributor

@stuartarchibald Conflicts resolved!

Many thanks!

Copy link
Contributor

@stuartarchibald stuartarchibald left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for doing this refactoring. The addition of the few functions to cgutils helps a lot with reducing usage of methods on the module itself. Couple of minor comments else looks good to go. Thanks again!

numba/cpython/randomimpl.py Outdated Show resolved Hide resolved
# Specify alignment to avoid misalignment bug
align = context.get_abi_sizeof(lldtype)
# Alignment is required to be a power of 2 for shared memory. If it is
# not a power of 2 (e.g. for a Record array) then round up accordingly.
gvmem.align = 1 << (align - 1 ).bit_length()

if dynamic_smem:
gvmem.linkage = lc.LINKAGE_EXTERNAL
gvmem.linkage = 'external'
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Wonder if the enum in llvmlite ought to be moved out of llvmpy and into the llvmlite API such that reliance on strings is reduced? (Not for this PR!).

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Maybe... I assumed that the use of strings instead of the enum was a deliberate design decision to permit flexibility in the linkage types for the user of llvmlite, but it's probably flexibility that's turned out to be un-needed (unless it's being used by external non-Numba users).

@stuartarchibald stuartarchibald added 4 - Waiting on author Waiting for author to respond to review and removed 3 - Ready for Review labels Mar 12, 2021
@stuartarchibald stuartarchibald added this to the Numba 0.54 RC milestone Mar 12, 2021
Co-authored-by: stuartarchibald <stuartarchibald@users.noreply.github.com>
@gmarkall gmarkall added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels Mar 12, 2021
Copy link
Contributor

@stuartarchibald stuartarchibald left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the patch and fixes.

@stuartarchibald stuartarchibald added 4 - Waiting on CI Review etc done, waiting for CI to finish Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Mar 12, 2021
@stuartarchibald
Copy link
Contributor

Buildfarm ID: numba_smoketest_cuda_yaml_25.

@esc
Copy link
Member

esc commented Mar 15, 2021

This one fails on CUDA 9.0 and 10.0 tests with:

[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - ERROR: test_kernel_with_debug (numba.cuda.tests.cudapy.test_compiler.TestCompileToPTX)
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_compiler.py", line 88, in test_kernel_with_debug
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - ptx, resty = compile_ptx(f, [], debug=True)
[2021-03-12 20:26:46,270] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 167, in compile_ptx
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - cres.signature.args, debug=debug)
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,271] {docker_operator.py:265} INFO - ERROR: test_debuginfo_in_asm (numba.cuda.tests.cudapy.test_debuginfo.TestCudaDebugInfo)
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 38, in test_debuginfo_in_asm
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - self._check(foo, sig=(types.int32[:],), expect=True)
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 20, in _check
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - asm = self._getasm(fn, sig=sig)
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 16, in _getasm
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - fn.compile(sig)
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,272] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,273] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - ERROR: test_environment_override (numba.cuda.tests.cudapy.test_debuginfo.TestCudaDebugInfo)
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 47, in test_environment_override
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - self._check(foo, sig=(types.int32[:],), expect=True)
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 20, in _check
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - asm = self._getasm(fn, sig=sig)
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 16, in _getasm
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - fn.compile(sig)
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,274] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - ERROR: test_issue_5835 (numba.cuda.tests.cudapy.test_debuginfo.TestCudaDebugInfo)
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,275] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py", line 60, in test_issue_5835
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - @cuda.jit((types.int32[::1],), debug=True, opt=False)
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/decorators.py", line 94, in kernel_jit
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - return Dispatcher(func, [func_or_sig], targetoptions=targetoptions)
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 971, in __init__
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - self.compile(sigs[0])
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,276] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,277] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - ERROR: test_exception (numba.cuda.tests.cudapy.test_exception.TestException)
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_exception.py", line 27, in test_exception
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - safe_foo[1, 3](np.array([0, 1]))
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 894, in __call__
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - self.stream, self.sharedmem)
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1036, in call
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1044, in _compile_for_args
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - return self.compile(tuple(argtypes))
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,278] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,279] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - ERROR: test_user_raise (numba.cuda.tests.cudapy.test_exception.TestException)
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_exception.py", line 36, in test_user_raise
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - foo[1, 1](False)
[2021-03-12 20:26:46,299] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 894, in __call__
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - self.stream, self.sharedmem)
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1036, in call
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1044, in _compile_for_args
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - return self.compile(tuple(argtypes))
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,300] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,301] {docker_operator.py:265} INFO - ERROR: test_divf_exception (numba.cuda.tests.cudapy.test_fastmath.TestFastMathOption)
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_fastmath.py", line 129, in test_divf_exception
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - fastmath=True, debug=True)(f10)
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/decorators.py", line 94, in kernel_jit
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - return Dispatcher(func, [func_or_sig], targetoptions=targetoptions)
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 971, in __init__
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - self.compile(sigs[0])
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,302] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - ======================================================================
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - ERROR: test_user_exception (numba.cuda.tests.cudapy.test_userexc.TestUserExc)
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - Traceback (most recent call last):
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/tests/cudapy/test_userexc.py", line 18, in test_user_exception
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - @cuda.jit("void(int32)", debug=True)
[2021-03-12 20:26:46,303] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/decorators.py", line 94, in kernel_jit
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - return Dispatcher(func, [func_or_sig], targetoptions=targetoptions)
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 971, in __init__
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - self.compile(sigs[0])
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 1166, in compile
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - **self.targetoptions)
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/core/compiler_lock.py", line 35, in _acquire_compile_lock
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - return func(*args, **kwargs)
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/compiler.py", line 586, in __init__
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - debug=self.debug)
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 131, in prepare_cuda_kernel
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - debug=debug)
[2021-03-12 20:26:46,304] {docker_operator.py:265} INFO - File "/opt/conda/envs/testenv_5533e046-3661-440d-be99-dcc96fd2e87b/lib/python3.6/site-packages/numba/cuda/target.py", line 195, in generate_kernel_wrapper
[2021-03-12 20:26:46,305] {docker_operator.py:265} INFO - changed = builder.icmp('==', xchg, old)
[2021-03-12 20:26:46,305] {docker_operator.py:265} INFO - AttributeError: 'IRBuilder' object has no attribute 'icmp'
[2021-03-12 20:26:46,305] {docker_operator.py:265} INFO - 
[2021-03-12 20:26:46,305] {docker_operator.py:265} INFO - ----------------------------------------------------------------------
[2021-03-12 20:26:46,305] {docker_operator.py:265} INFO - Ran 1134 tests in 210.322s

@esc esc added BuildFarm Passed For PRs that have been through the buildfarm and passed and removed Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Mar 15, 2021
@esc
Copy link
Member

esc commented Mar 15, 2021

@gmarkall thanks for the fixes, numba_smoketest_cuda_yaml_25 was all green.

@gmarkall
Copy link
Member Author

@esc many thanks!

Copy link
Contributor

@stuartarchibald stuartarchibald left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the fix, looks good.

@stuartarchibald stuartarchibald added 5 - Ready to merge Review and testing done, is ready to merge and removed 4 - Waiting on CI Review etc done, waiting for CI to finish labels Mar 15, 2021
@sklam sklam merged commit 83b246f into numba:master Mar 15, 2021
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
5 - Ready to merge Review and testing done, is ready to merge BuildFarm Passed For PRs that have been through the buildfarm and passed CUDA CUDA related issue/PR Effort - medium Medium size effort needed
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants