Support python 3.14 (#599)

This PR vendors the relevant changes from upstream numba required to
support python 3.14.

Closes #595

---------

Co-authored-by: Phillip Cloud <417981+cpcloud@users.noreply.github.com>

Author: brandon-b-miller

.github/workflows/build-wheel.yml

@@ -32,7 +32,7 @@ jobs:
           - "3.11"
           - "3.12"
           - "3.13"
-          # - "3.14"
+          - "3.14"
           # - "3.14t"
     name: py${{ matrix.python-version }}
     runs-on: ${{ (inputs.host-platform == 'linux-64' && 'linux-amd64-cpu8') ||

ci/test-matrix.json

@@ -12,14 +12,18 @@
       { "ARCH": "amd64", "PY_VER": "3.12", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "l4", "DRIVER": "latest" },
       { "ARCH": "amd64", "PY_VER": "3.13", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "v100", "DRIVER": "latest" },
       { "ARCH": "amd64", "PY_VER": "3.13", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "rtxpro6000", "DRIVER": "latest" },
+      { "ARCH": "amd64", "PY_VER": "3.14", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "v100", "DRIVER": "latest" },
+      { "ARCH": "amd64", "PY_VER": "3.14", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "rtxpro6000", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.10", "CUDA_VER": "12.9.1", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.10", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.11", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.11", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.12", "CUDA_VER": "12.9.1", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.12", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "a100", "DRIVER": "latest" },
       { "ARCH": "arm64", "PY_VER": "3.13", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "a100", "DRIVER": "latest" },
-      { "ARCH": "arm64", "PY_VER": "3.13", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" }
+      { "ARCH": "arm64", "PY_VER": "3.13", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" },
+      { "ARCH": "arm64", "PY_VER": "3.14", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "a100", "DRIVER": "latest" },
+      { "ARCH": "arm64", "PY_VER": "3.14", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest" }
     ],
     "nightly": [],
     "special_runners": {
@@ -37,7 +41,9 @@
       { "ARCH": "amd64", "PY_VER": "3.12", "CUDA_VER": "12.9.1", "LOCAL_CTK": "0", "GPU": "l4", "DRIVER": "latest", "DRIVER_MODE": "MCDM" },
       { "ARCH": "amd64", "PY_VER": "3.12", "CUDA_VER": "13.0.2", "LOCAL_CTK": "1", "GPU": "a100", "DRIVER": "latest", "DRIVER_MODE": "TCC" },
       { "ARCH": "amd64", "PY_VER": "3.13", "CUDA_VER": "12.9.1", "LOCAL_CTK": "1", "GPU": "l4", "DRIVER": "latest", "DRIVER_MODE": "TCC" },
-      { "ARCH": "amd64", "PY_VER": "3.13", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "rtxpro6000", "DRIVER": "latest", "DRIVER_MODE": "MCDM" }
+      { "ARCH": "amd64", "PY_VER": "3.13", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "rtxpro6000", "DRIVER": "latest", "DRIVER_MODE": "MCDM" },
+      { "ARCH": "amd64", "PY_VER": "3.14", "CUDA_VER": "12.9.1", "LOCAL_CTK": "1", "GPU": "l4", "DRIVER": "latest", "DRIVER_MODE": "TCC" },
+      { "ARCH": "amd64", "PY_VER": "3.14", "CUDA_VER": "13.0.2", "LOCAL_CTK": "0", "GPU": "rtxpro6000", "DRIVER": "latest", "DRIVER_MODE": "MCDM" }
     ],
     "nightly": []
   }

numba_cuda/numba/cuda/cext/_dispatcher.cpp

@@ -29,7 +29,7 @@
  *
  */
 
-#if (PY_MAJOR_VERSION >= 3) && ((PY_MINOR_VERSION == 12) || (PY_MINOR_VERSION == 13))
+#if (PY_MAJOR_VERSION >= 3) && ((PY_MINOR_VERSION == 12) || (PY_MINOR_VERSION == 13) || (PY_MINOR_VERSION == 14))
 
 #ifndef Py_BUILD_CORE
     #define Py_BUILD_CORE 1
@@ -1004,12 +1004,18 @@ static PyTypeObject DispatcherType = {
     0,                                           /* tp_version_tag */
     0,                                           /* tp_finalize */
     0,                                           /* tp_vectorcall */
-#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION == 12)
+#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION >= 12)
 /* This was introduced first in 3.12
  * https://github.com/python/cpython/issues/91051
  */
     0,                                           /* tp_watched */
 #endif
+#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION >= 13)
+/* This was introduced in 3.13
+ * https://github.com/python/cpython/pull/114900
+ */
+    0,                                           /* tp_versions_used */
+#endif
 
 /* WARNING: Do not remove this, only modify it! It is a version guard to
  * act as a reminder to update this struct on Python version update! */

numba_cuda/numba/cuda/cext/_hashtable.cpp

@@ -110,7 +110,12 @@ _Numba_hashtable_hash_int(const void *key)
 extern "C" Py_uhash_t
 _Numba_hashtable_hash_ptr(const void *key)
 {
+/* Use public API on Python 3.13+; _Py_HashPointer is deprecated on 3.14+ */
+#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION >= 13)
+    return (Py_uhash_t)Py_HashPointer((void *)key);
+#else
     return (Py_uhash_t)_Py_HashPointer((void *)key);
+#endif
 }
 
 extern "C" int

numba_cuda/numba/cuda/cext/_pymodule.h

@@ -33,6 +33,6 @@
         Py_DECREF(tmp); } while (0)
 
 
-#define NB_SUPPORTED_PYTHON_MINOR ((PY_MINOR_VERSION == 10) || (PY_MINOR_VERSION == 11) || (PY_MINOR_VERSION == 12) || (PY_MINOR_VERSION == 13))
+#define NB_SUPPORTED_PYTHON_MINOR ((PY_MINOR_VERSION == 10) || (PY_MINOR_VERSION == 11) || (PY_MINOR_VERSION == 12) || (PY_MINOR_VERSION == 13) || (PY_MINOR_VERSION == 14))
 
 #endif /* NUMBA_PY_MODULE_H_ */

numba_cuda/numba/cuda/cext/_typeof.cpp

@@ -17,7 +17,7 @@
     #include <numpy/npy_2_compat.h>
 #endif
 
-#if (PY_MAJOR_VERSION >= 3) && (PY_MINOR_VERSION == 13)
+#if (PY_MAJOR_VERSION >= 3) && ((PY_MINOR_VERSION == 13) || (PY_MINOR_VERSION == 14))
     #ifndef Py_BUILD_CORE
         #define Py_BUILD_CORE 1
     #endif

numba_cuda/numba/cuda/cext/mviewbuf.c

@@ -337,12 +337,18 @@ static PyTypeObject MemAllocType = {
     0,                                          /* tp_version_tag */
     0,                                          /* tp_finalize */
     0,                                          /* tp_vectorcall */
-#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION == 12)
+#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION >= 12)
 /* This was introduced first in 3.12
  * https://github.com/python/cpython/issues/91051
  */
     0,                                           /* tp_watched */
 #endif
+#if (PY_MAJOR_VERSION == 3) && (PY_MINOR_VERSION >= 13)
+/* This was introduced in 3.13
+ * https://github.com/python/cpython/pull/114900
+ */
+    0,                                           /* tp_versions_used */
+#endif
 
 /* WARNING: Do not remove this, only modify it! It is a version guard to
  * act as a reminder to update this struct on Python version update! */

numba_cuda/numba/cuda/core/bytecode.py

@@ -14,7 +14,7 @@
 from numba.cuda.utils import PYVERSION
 
 
-if PYVERSION in ((3, 12), (3, 13)):
+if PYVERSION in ((3, 12), (3, 13), (3, 14)):
     from opcode import _inline_cache_entries
 
     # Instruction/opcode length in bytes
@@ -112,7 +112,7 @@ def get_jump_target(self):
         # https://bugs.python.org/issue27129
         # https://github.com/python/cpython/pull/25069
         assert self.is_jump
-        if PYVERSION in ((3, 13),):
+        if PYVERSION in ((3, 13), (3, 14)):
             if self.opcode in (
                 dis.opmap[k]
                 for k in ["JUMP_BACKWARD", "JUMP_BACKWARD_NO_INTERRUPT"]
@@ -141,7 +141,7 @@ def get_jump_target(self):
         else:
             raise NotImplementedError(PYVERSION)
 
-        if PYVERSION in ((3, 10), (3, 11), (3, 12), (3, 13)):
+        if PYVERSION in ((3, 10), (3, 11), (3, 12), (3, 13), (3, 14)):
             if self.opcode in JREL_OPS:
                 return self.next + self.arg * 2
             else:
@@ -179,7 +179,7 @@ def block_effect(self):
 OPCODE_NOP = dis.opname.index("NOP")
 
 
-if PYVERSION in ((3, 13),):
+if PYVERSION in ((3, 13), (3, 14)):
 
     def _unpack_opargs(code):
         buf = []
@@ -418,7 +418,7 @@ def get_used_globals(self):
 
 
 def _fix_LOAD_GLOBAL_arg(arg):
-    if PYVERSION in ((3, 11), (3, 12), (3, 13)):
+    if PYVERSION in ((3, 11), (3, 12), (3, 13), (3, 14)):
         return arg >> 1
     elif PYVERSION in (
         (3, 9),
@@ -592,7 +592,7 @@ def pop_and_merge_exceptions(
                 if not next_inst.opname == "FOR_ITER":
                     continue
 
-                if PYVERSION in ((3, 13),):
+                if PYVERSION in ((3, 13), (3, 14)):
                     # Check end of pattern, two instructions.
                     # Check for the corresponding END_FOR, exception table end
                     # is non-inclusive, so subtract one.
@@ -601,8 +601,14 @@ def pop_and_merge_exceptions(
                     if not curr_inst.opname == "END_FOR":
                         continue
                     next_inst = self.table[self.ordered_offsets[index - 1]]
-                    if not next_inst.opname == "POP_TOP":
-                        continue
+                    if PYVERSION in ((3, 13),):
+                        if not next_inst.opname == "POP_TOP":
+                            continue
+                    elif PYVERSION in ((3, 14),):
+                        if not next_inst.opname == "POP_ITER":
+                            continue
+                    else:
+                        raise NotImplementedError(PYVERSION)
                     # END_FOR must be followed by SWAP(2)
                     next_inst = self.table[self.ordered_offsets[index]]
                     if not next_inst.opname == "SWAP" and next_inst.arg == 2:
@@ -646,6 +652,7 @@ def pop_and_merge_exceptions(
 elif PYVERSION in (
     (3, 12),
     (3, 13),
+    (3, 14),
 ):
     ByteCode = ByteCodePy312
 elif PYVERSION < (3, 11):