[WIP]Vendor in _dispatcher, _devicearray, mviewbuf C extension for CUDA-specific customization #451
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…p level numba_cuda
This is not sufficient. We also need: - One build for each Python version for both of amd64 and aarch64 - The docs build to be changed to download the py313-amd64 wheels - Maybe other things I missed?
|
Auto-sync is disabled for ready for review pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
Contributor
Author
|
/ok to test 3c48a3d |
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
Contributor
Author
|
/ok to test 618a885 |
Contributor
|
/ok to test |
Contributor
|
/ok to test |
gmarkall
added
3 - Ready for Review
VijayKandiah
changed the title
gmarkall
reviewed
Sep 5, 2025
| extern "C" { | ||
| #endif | ||
|
|
||
| #define NUMBA_DEVICEARRAY_IMPORT_NAME "numba_cuda._devicearray" |
Contributor
There was a problem hiding this comment.
I suspect this should match the module name:
Suggested change
| #define NUMBA_DEVICEARRAY_IMPORT_NAME "numba_cuda._devicearray" | |
| #define NUMBA_DEVICEARRAY_IMPORT_NAME "numba.cuda.cext._devicearray" |
Contributor
There was a problem hiding this comment.
I'm not sure how this was working when it was imported in import_devicearray in _dispatcher.cpp, but maybe it's something to do with the code in cext.__init__.py - will update here once I understand better.
gmarkall
reviewed
Sep 5, 2025
gmarkall
reviewed
Sep 5, 2025
gmarkall
reviewed
Sep 5, 2025
gmarkall
reviewed
Sep 5, 2025
Contributor
|
/ok to test |
Co-authored-by: Graham Markall <[email protected]>
Contributor
|
/ok to test |
The package is `numba.cuda.cext`, not `numba_cuda.numba.cuda.cext` - the redirector should ensure that this should always be the case.
Noticing the following traceback when trying to do
```
from numba import cuda
```
gives:
```
AttributeError: cannot access submodule 'cuda' of module 'numba' (most likely due to a circular import)
Traceback (most recent call last):
File "<string>", line 1, in <module>
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/__init__.py", line 73, in <module>
from .device_init import *
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/device_init.py", line 66, in <module>
from .decorators import jit, declare_device
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/decorators.py", line 9, in <module>
from numba.cuda.dispatcher import CUDADispatcher
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 50, in <module>
from numba.cuda.cext import _dispatcher
ImportError: numba.cuda.cext._devicearray failed to import
```
let's instrument the initialization of `_devicearray` to see what's
happening.
Observe that the "failed to import" message actually comes from
`_dispatcher.cpp`, in its `MOD_INIT()` function:
```
PyErr_SetString(PyExc_ImportError, NUMBA_DEVICEARRAY_IMPORT_NAME " failed to import");
```
This is produced if the `import_devicearray()` function fails. That
function does two kinds of imports - a module import, and a capsule
import.
This commit adds debug prints to `import_devicearray()` so that we can
see which kind of import is failing.
With this commit, we get:
```
In init
About to create capsule
About to check if type ready
About to add object
About to get module dict
About to set item
Success
Imported device array
Capsule did not import.
AttributeError: cannot access submodule 'cuda' of module 'numba' (most likely due to a circular import)
Traceback (most recent call last):
File "<string>", line 1, in <module>
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/__init__.py", line 73, in <module>
from .device_init import *
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/device_init.py", line 66, in <module>
from .decorators import jit, declare_device
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/decorators.py", line 9, in <module>
from numba.cuda.dispatcher import CUDADispatcher
File "/home/gmarkall/numbadev/numba-cuda/numba_cuda/numba/cuda/dispatcher.py", line 50, in <module>
from numba.cuda.cext import _dispatcher
ImportError: numba.cuda.cext._devicearray failed to import
```
so we can see that the `_devicearray` module does import, but the
capsule in `_devicearray._DEVICEARRAY_API` does not.
When `import_devicearray()` is called, we're partway through importing `numba.cuda`. Therefore, the `PyCapsule_Import()` fails because it tries to access packages under `numba.cuda` during its initialization, which then fails due to this circularity. In order to work around this, we can get the `_DEVICEARRAY_API` attribute of the `_devicearray` module directly from its module dict. This was not a problem in upstream because `_devicearray` was not in the `numba.cuda` package. NOTE: This fix is not entirely complete - it is not good at cleaning up in the presence of errors occurring - this will be resolved in a later commit when the debug prints have been tidied up.
Contributor
|
/ok to test |
VijayKandiah
changed the title
Contributor
|
Heads up: I re-opened #373 which contains all of the changes in this PR except that it still has the import hacks in |
Contributor
|
Closing as #373 was merged. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Contains my changes to Vijay's #373 which allow tests to pass locally for me with and without an editable install, but I have not figured out the CI issues yet.
cc @VijayKandiah