Update to Cython 3.0.0 (#1313)

This PR contains the minimal set of changes to compile using Cython 3 without warnings. Future PRs can be made to take advantage of new or improved features.

The specific changes are:
- Ensuring `nogil` always comes after `except`. `except * nogil` is a compile-time error in Cython 3
- Adding `noexcept` or `except *` to any `cdef ` functions missing them. In Cython 0.29 these would default to `noexcept`, which meant that exceptions would not be properly propagated. In Cython 3.0.0, these default to `except *`, which incurs a performance penalty for reacquiring the GIL to check the exception value even for `nogil` functions. Being explicit here is important.

There are a large number of outstanding warnings due to NVIDIA/cuda-python#44. cuda-python for CUDA 12 has the necessary fix, but we will need a cuda-python 11.8.* bugfix with a backport to make those warnings go away.

Authors:
  - Vyas Ramasubramani (https://github.com/vyasr)

Approvers:
  - Ashwin Srinath (https://github.com/shwina)
  - Lawrence Mitchell (https://github.com/wence-)
  - Ray Douglass (https://github.com/raydouglass)

URL: #1313

conda/environments/all_cuda-118_arch-x86_64.yaml

@@ -8,7 +8,7 @@ dependencies:
 - cuda-python>=11.7.1,<12.0a0
 - cuda-version=11.8
 - cudatoolkit
-- cython>=0.29,<0.30
+- cython>=3.0.0
 - fmt>=9.1.0,<10
 - gcovr>=5.0
 - identify>=2.5.20

conda/environments/all_cuda-120_arch-x86_64.yaml

@@ -8,7 +8,7 @@ dependencies:
 - cuda-nvcc
 - cuda-python>=12.0,<13.0a0
 - cuda-version=12.0
-- cython>=0.29,<0.30
+- cython>=3.0.0
 - fmt>=9.1.0,<10
 - gcovr>=5.0
 - identify>=2.5.20

conda/recipes/rmm/meta.yaml

@@ -58,7 +58,7 @@ requirements:
     - cuda-cudart-dev
     - cuda-python ==12.0.0
     {% endif %}
-    - cython >=0.29,<0.30
+    - cython >=3.0.0
     - librmm ={{ version }}
     - python
     - scikit-build >=0.13.1

dependencies.yaml

@@ -62,7 +62,7 @@ dependencies:
       - output_types: [conda, requirements, pyproject]
         packages:
           - &cmake_ver cmake>=3.26.4
-          - cython>=0.29,<0.30
+          - cython>=3.0.0
           - ninja
           - scikit-build>=0.13.1
           - tomli

python/pyproject.toml

@@ -17,7 +17,7 @@ build-backend = "setuptools.build_meta"
 requires = [
     "cmake>=3.26.4",
     "cuda-python>=11.7.1,<12.0a0",
-    "cython>=0.29,<0.30",
+    "cython>=3.0.0",
     "ninja",
     "scikit-build>=0.13.1",
     "setuptools>=61.0.0",

python/rmm/_cuda/stream.pxd

@@ -26,8 +26,8 @@ cdef class Stream:
     @staticmethod
     cdef Stream _from_cudaStream_t(cudaStream_t s, object owner=*)
 
-    cdef cuda_stream_view view(self) nogil except *
-    cdef void c_synchronize(self) nogil except *
-    cdef bool c_is_default(self) nogil except *
+    cdef cuda_stream_view view(self) except * nogil
+    cdef void c_synchronize(self) except * nogil
+    cdef bool c_is_default(self) except * nogil
     cdef void _init_with_new_cuda_stream(self) except *
     cdef void _init_from_stream(self, Stream stream) except *

python/rmm/_cuda/stream.pyx

@@ -48,7 +48,7 @@ cdef class Stream:
                 self._init_from_cupy_stream(obj)
 
     @staticmethod
-    cdef Stream _from_cudaStream_t(cudaStream_t s, object owner=None):
+    cdef Stream _from_cudaStream_t(cudaStream_t s, object owner=None) except *:
         """
         Construct a Stream from a cudaStream_t.
         """
@@ -57,13 +57,13 @@ cdef class Stream:
         obj._owner = owner
         return obj
 
-    cdef cuda_stream_view view(self) nogil except *:
+    cdef cuda_stream_view view(self) except * nogil:
         """
         Generate a rmm::cuda_stream_view from this Stream instance
         """
         return cuda_stream_view(<cudaStream_t><uintptr_t>(self._cuda_stream))
 
-    cdef void c_synchronize(self) nogil except *:
+    cdef void c_synchronize(self) except * nogil:
         """
         Synchronize the CUDA stream.
         This function *must* be called in a `with nogil` block
@@ -77,7 +77,7 @@ cdef class Stream:
         with nogil:
             self.c_synchronize()
 
-    cdef bool c_is_default(self) nogil except *:
+    cdef bool c_is_default(self) except * nogil:
         """
         Check if we are the default CUDA stream
         """

python/rmm/_lib/cuda_stream.pxd

@@ -33,5 +33,5 @@ cdef extern from "rmm/cuda_stream.hpp" namespace "rmm" nogil:
 @cython.final
 cdef class CudaStream:
     cdef unique_ptr[cuda_stream] c_obj
-    cdef cudaStream_t value(self) nogil except *
-    cdef bool is_valid(self) nogil except *
+    cdef cudaStream_t value(self) except * nogil
+    cdef bool is_valid(self) except * nogil

python/rmm/_lib/cuda_stream.pyx

@@ -27,8 +27,8 @@ cdef class CudaStream:
     def __cinit__(self):
         self.c_obj.reset(new cuda_stream())
 
-    cdef cudaStream_t value(self) nogil except *:
+    cdef cudaStream_t value(self) except * nogil:
         return self.c_obj.get()[0].value()
 
-    cdef bool is_valid(self) nogil except *:
+    cdef bool is_valid(self) except * nogil:
         return self.c_obj.get()[0].is_valid()

python/rmm/_lib/device_buffer.pxd

@@ -56,7 +56,7 @@ cdef class DeviceBuffer:
 
     @staticmethod
     cdef DeviceBuffer c_to_device(const unsigned char[::1] b,
-                                  Stream stream=*)
+                                  Stream stream=*) except *
     cpdef copy_to_host(self, ary=*, Stream stream=*)
     cpdef copy_from_host(self, ary, Stream stream=*)
     cpdef copy_from_device(self, cuda_ary, Stream stream=*)

python/rmm/_lib/device_buffer.pyx

@@ -172,7 +172,7 @@ cdef class DeviceBuffer:
 
     @staticmethod
     cdef DeviceBuffer c_to_device(const unsigned char[::1] b,
-                                  Stream stream=DEFAULT_STREAM):
+                                  Stream stream=DEFAULT_STREAM) except *:
         """Calls ``to_device`` function on arguments provided"""
         return to_device(b, stream)
 
@@ -382,7 +382,7 @@ cdef void _copy_async(const void* src,
                       void* dst,
                       size_t count,
                       ccudart.cudaMemcpyKind kind,
-                      cuda_stream_view stream) nogil except *:
+                      cuda_stream_view stream) except * nogil:
     """
     Asynchronously copy data between host and/or device pointers.
 

python/rmm/_lib/memory_resource.pyx

@@ -14,10 +14,14 @@
 
 import os
 import warnings
+# This import is needed for Cython typing in translate_python_except_to_cpp
+# See https://github.com/cython/cython/issues/5589
+from builtins import BaseException
 from collections import defaultdict
 
 cimport cython
 from cython.operator cimport dereference as deref
+from libc.stddef cimport size_t
 from libc.stdint cimport int8_t, int64_t, uintptr_t
 from libcpp cimport bool
 from libcpp.memory cimport make_unique, unique_ptr
@@ -37,7 +41,7 @@ from rmm._lib.per_device_resource cimport (
 # Transparent handle of a C++ exception
 ctypedef pair[int, string] CppExcept
 
-cdef CppExcept translate_python_except_to_cpp(err: BaseException):
+cdef CppExcept translate_python_except_to_cpp(err: BaseException) noexcept:
     """Translate a Python exception into a C++ exception handle
 
     The returned exception handle can then be thrown by `throw_cpp_except()`,
@@ -526,7 +530,10 @@ cdef void* _allocate_callback_wrapper(
     size_t nbytes,
     cuda_stream_view stream,
     void* ctx
-) nogil:
+    # Note that this function is specifically designed to rethrow Python
+    # exceptions as C++ exceptions when called as a callback from C++, so it is
+    # noexcept from Cython's perspective.
+) noexcept nogil:
     cdef CppExcept err
     with gil:
         try:
@@ -540,7 +547,7 @@ cdef void _deallocate_callback_wrapper(
     size_t nbytes,
     cuda_stream_view stream,
     void* ctx
-) with gil:
+) except * with gil:
     (<object>ctx)(<uintptr_t>(ptr), nbytes)
 
 
@@ -796,7 +803,10 @@ cdef class TrackingResourceAdaptor(UpstreamResourceAdaptor):
             self.c_obj.get()))[0].log_outstanding_allocations()
 
 
-cdef bool _oom_callback_function(size_t bytes, void *callback_arg) nogil:
+# Note that this function is specifically designed to rethrow Python exceptions
+# as C++ exceptions when called as a callback from C++, so it is noexcept from
+# Cython's perspective.
+cdef bool _oom_callback_function(size_t bytes, void *callback_arg) noexcept nogil:
     cdef CppExcept err
     with gil:
         try: