[Refactor] Vendor in _dispatcher, _devicearray, mviewbuf C extension for CUDA-specific customization

[VijayKandiah]

This PR vendors in the _dispatcher, _devicearray, mviewbuf C extensions for CUDA-specific changes.

[copy-pr-bot]

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.

[VijayKandiah]

Currently, pip install -e . does not seem to build the extension, so we have to use python setup.py build_ext --inplace to build the extension. This will probably fail CI testing until we either update CI or make pip install -e . automatically build the extension.

[gmarkall]

Currently, pip install -e . does not seem to build the extension, so we have to use python setup.py build_ext --inplace to build the extension. This will probably fail CI testing until we either update CI or make pip install -e . automatically build the extension.

I just had a vague recollection about this - do you have setuptools > 78 installed? If so, does going back to setuptools 78 or earlier make pip install -e . build the extension?

[VijayKandiah]

Currently, pip install -e . does not seem to build the extension, so we have to use python setup.py build_ext --inplace to build the extension. This will probably fail CI testing until we either update CI or make pip install -e . automatically build the extension.

I just had a vague recollection about this - do you have setuptools > 78 installed? If so, does going back to setuptools 78 or earlier make pip install -e . build the extension?

The issue was not a newer setuptools version. It was simply that is_building method in setup.py was not looking for editable_wheel as one of the arguments to enable building the C extension for. Fixed it with my recent commit.

[VijayKandiah]

/ok to test

[VijayKandiah]

I had to include numpy as a build requirement for setuptools.build_meta in pyproject.toml. However, I am seeing build errors in our CI with numpy not found.

ModuleNotFoundError: No module named 'numpy'
  error: subprocess-exited-with-error
  
  × Building wheel for numba-cuda (pyproject.toml) did not run successfully.
  │ exit code: 1

Is there a way to update CI to add this build requirement?

[VijayKandiah]

/ok to test

[VijayKandiah]

My latest commit also vendors in _devicearray and mviewbuf C extension. This should encompass all the C extensions we use in numba-cuda.

I noticed a circular import issue if _devicearray cext was placed at numba_cuda.numba.cuda, and moving it to the top level numba_cuda fixed this. For uniformity, I ended up placing all cexts at numba_cuda.cext. @gmarkall please let me know your thoughts on the updated folder structure here.

[VijayKandiah]

/ok to test

[VijayKandiah]

/ok to test

[VijayKandiah]

/ok to test

[VijayKandiah]

This PR currently contains some import hacks in cext/__init__.py to help with our CI testing. The attempt to remove import hacks is ongoing at #451. As long as CI passes here, and everything else looks good to reviewers, I am in favor of getting this PR in and create a follow up PR for removing import hacks.

[VijayKandiah]

/ok to test

[VijayKandiah]

/ok to test

[VijayKandiah]

/ok to test

[VijayKandiah]

/ok to test

[gmarkall]

I'm going to merge this, but there are two required fixups to remove the cext hacks, that need to be done in a subsequent PR.

• Change the CI setup so we don't install the package then run tests from the repo. For a combination of reasons involving the redirector, pytest's package discovery and import logic, and the fact that it's weird to try and test an installed package when the source repo is the current directory, the cext hacks are required. Changing the setup so that tests run from the installed package without the source tree in the path should remove the need for the cext hacks.
• Fix the capsule name for the device array, which is presently numba_cuda._devicearray, instead of numba.cuda.cext._devicearray. Simply changing the name doesn't work, because it then creates a circular import - a change to the import_devicearray() function is needed to avoid this.

A start of a fix for the circular import issue looks like:

diff --git a/numba_cuda/numba/cuda/cext/_devicearray.h b/numba_cuda/numba/cuda/cext/_devicearray.h
index e1672698..2e9df6cb 100644
--- a/numba_cuda/numba/cuda/cext/_devicearray.h
+++ b/numba_cuda/numba/cuda/cext/_devicearray.h
@@ -8,7 +8,7 @@
     extern "C" {
 #endif
 
-#define NUMBA_DEVICEARRAY_IMPORT_NAME "numba_cuda._devicearray"
+#define NUMBA_DEVICEARRAY_IMPORT_NAME "numba.cuda.cext._devicearray"
 /* These definitions should only be used by consumers of the Device Array API.
  * Consumers access the API through the opaque pointer stored in
  * _devicearray._DEVICEARRAY_API.  We don't want these definitions in
diff --git a/numba_cuda/numba/cuda/cext/_dispatcher.cpp b/numba_cuda/numba/cuda/cext/_dispatcher.cpp
index bfd3c651..653421f9 100644
--- a/numba_cuda/numba/cuda/cext/_dispatcher.cpp
+++ b/numba_cuda/numba/cuda/cext/_dispatcher.cpp
@@ -947,14 +947,22 @@ import_devicearray(void)
     if (devicearray == NULL) {
         return -1;
     }
-    Py_DECREF(devicearray);
 
-    DeviceArray_API = (void**)PyCapsule_Import(NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API", 0);
-    if (DeviceArray_API == NULL) {
-        return -1;
+    PyObject *d = PyModule_GetDict(devicearray);
+    if (d == NULL) {
+      Py_DECREF(devicearray);
+      return -1;
     }
 
-    return 0;
+    PyObject *c_api = PyDict_GetItemString(d, "_DEVICEARRAY_API");
+    if (PyCapsule_IsValid(c_api, NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API")) {
+      DeviceArray_API = (void**)PyCapsule_GetPointer(c_api, NUMBA_DEVICEARRAY_IMPORT_NAME "._DEVICEARRAY_API");
+      Py_DECREF(devicearray);
+      return 0;
+    } else {
+      Py_DECREF(devicearray);
+      return -1;
+    }
 }
 
 static PyMethodDef Dispatcher_methods[] = {