diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml
index f668c4a4c..0240b46e5 100644
--- a/conda/environments/all_cuda-118_arch-x86_64.yaml
+++ b/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
diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml
index 40291ba33..96b8c8065 100644
--- a/conda/environments/all_cuda-120_arch-x86_64.yaml
+++ b/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
diff --git a/conda/recipes/rmm/meta.yaml b/conda/recipes/rmm/meta.yaml
index 5d905b853..e13254d4c 100644
--- a/conda/recipes/rmm/meta.yaml
+++ b/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
diff --git a/dependencies.yaml b/dependencies.yaml
index 91eed8280..17fb1d66d 100644
--- a/dependencies.yaml
+++ b/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
diff --git a/python/pyproject.toml b/python/pyproject.toml
index 3cac6c043..f83e4fec4 100644
--- a/python/pyproject.toml
+++ b/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",
diff --git a/python/rmm/_cuda/stream.pxd b/python/rmm/_cuda/stream.pxd
index 6aa4e0b24..3c3d3aa6f 100644
--- a/python/rmm/_cuda/stream.pxd
+++ b/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 *
diff --git a/python/rmm/_cuda/stream.pyx b/python/rmm/_cuda/stream.pyx
index d60dde4e1..4795cbb9f 100644
--- a/python/rmm/_cuda/stream.pyx
+++ b/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
         """
diff --git a/python/rmm/_lib/cuda_stream.pxd b/python/rmm/_lib/cuda_stream.pxd
index 1eed1cefb..e224cf9af 100644
--- a/python/rmm/_lib/cuda_stream.pxd
+++ b/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
diff --git a/python/rmm/_lib/cuda_stream.pyx b/python/rmm/_lib/cuda_stream.pyx
index fb35ec11f..0861f0663 100644
--- a/python/rmm/_lib/cuda_stream.pyx
+++ b/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()
diff --git a/python/rmm/_lib/device_buffer.pxd b/python/rmm/_lib/device_buffer.pxd
index 364dbb2c0..3d5f29f9a 100644
--- a/python/rmm/_lib/device_buffer.pxd
+++ b/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=*)
diff --git a/python/rmm/_lib/device_buffer.pyx b/python/rmm/_lib/device_buffer.pyx
index c2bfd1459..d248d01ab 100644
--- a/python/rmm/_lib/device_buffer.pyx
+++ b/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.
 
diff --git a/python/rmm/_lib/memory_resource.pyx b/python/rmm/_lib/memory_resource.pyx
index 774db374a..aa0d21d1f 100644
--- a/python/rmm/_lib/memory_resource.pyx
+++ b/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: