Update to numba-cuda >=0.21.2,<0.22.0a0

conda/environments/all_cuda-129_arch-aarch64.yaml

@@ -55,7 +55,7 @@ dependencies:
 - nbsphinx
 - ninja
 - notebook
-- numba-cuda>=0.19.1,<0.20.0a0
+- numba-cuda>=0.21.2,<0.22.0a0
 - numba>=0.60.0,<0.62.0a0
 - numpy>=1.23,<3.0a0
 - numpydoc

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

@@ -56,7 +56,7 @@ dependencies:
 - nbsphinx
 - ninja
 - notebook
-- numba-cuda>=0.19.1,<0.20.0a0
+- numba-cuda>=0.21.2,<0.22.0a0
 - numba>=0.60.0,<0.62.0a0
 - numpy>=1.23,<3.0a0
 - numpydoc

conda/environments/all_cuda-130_arch-aarch64.yaml

@@ -55,7 +55,7 @@ dependencies:
 - nbsphinx
 - ninja
 - notebook
-- numba-cuda>=0.19.1,<0.20.0a0
+- numba-cuda>=0.21.2,<0.22.0a0
 - numba>=0.60.0,<0.62.0a0
 - numpy>=1.23,<3.0a0
 - numpydoc

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

@@ -56,7 +56,7 @@ dependencies:
 - nbsphinx
 - ninja
 - notebook
-- numba-cuda>=0.19.1,<0.20.0a0
+- numba-cuda>=0.21.2,<0.22.0a0
 - numba>=0.60.0,<0.62.0a0
 - numpy>=1.23,<3.0a0
 - numpydoc

conda/recipes/cudf/recipe.yaml

@@ -71,7 +71,7 @@ requirements:
     - rapids-build-backend >=0.4.0,<0.5.0.dev0
     - scikit-build-core >=0.10.0
     - dlpack >=0.8,<1.0
-    - numba-cuda >=0.19.1,<0.20.0a0
+    - numba-cuda >=0.21.2,<0.22.0a0
     - libcudf =${{ version }}
     - pylibcudf =${{ version }}
     - rmm =${{ minor_version }}
@@ -86,7 +86,7 @@ requirements:
     - typing_extensions >=4.0.0
     - pandas >=2.0,<2.4.0dev0
     - cupy >=13.6.0
-    - numba-cuda >=0.19.1,<0.20.0a0
+    - numba-cuda >=0.21.2,<0.22.0a0
     - numba >=0.60.0,<0.62.0a0
     - numpy >=1.23,<3.0a0
     - pyarrow>=15.0.0,<22.0.0a0

dependencies.yaml

@@ -805,7 +805,7 @@ dependencies:
             packages:
               - numba==0.60.0
               - pandas==2.0.*
-              - numba-cuda==0.19.1
+              - numba-cuda==0.21.2
           - matrix: {dependencies: "latest"}
             packages:
               - pandas==2.3.3
@@ -948,18 +948,18 @@ dependencies:
     common:
       - output_types: [conda]
         packages:
-        - numba-cuda>=0.19.1,<0.20.0a0
+        - numba-cuda>=0.21.2,<0.22.0a0
     specific:
       - output_types: [requirements, pyproject]
         matrices:
           - matrix:
               cuda: "12.*"
             packages:
-              - numba-cuda[cu12]>=0.19.1,<0.20.0a0
+              - numba-cuda[cu12]>=0.21.2,<0.22.0a0
           # fallback to CUDA 13 versions if 'cuda' is '13.*' or not provided
           - matrix:
             packages:
-              - numba-cuda[cu13]>=0.19.1,<0.20.0a0
+              - numba-cuda[cu13]>=0.21.2,<0.22.0a0
   depends_on_pylibcudf:
     common:
       - output_types: conda

python/cudf/cudf/core/_internals/aggregation.py

@@ -5,7 +5,7 @@
 from typing import TYPE_CHECKING, Literal
 
 import numpy as np
-from numba.np import numpy_support
+from numba.cuda.np import numpy_support
 
 import pylibcudf as plc
 

python/cudf/cudf/core/udf/groupby_lowering.py

@@ -1,13 +1,13 @@
-# SPDX-FileCopyrightText: Copyright (c) 2022-2023, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 
 from functools import partial
 
 from numba import types
-from numba.core import cgutils
-from numba.core.extending import lower_builtin
-from numba.core.typing import signature as nb_signature
+from numba.cuda import cgutils
 from numba.cuda.cudaimpl import lower as cuda_lower
+from numba.cuda.extending import lower_builtin
+from numba.cuda.typing import signature as nb_signature
 
 from cudf.core.udf.groupby_typing import (
     SUPPORTED_GROUPBY_NUMBA_TYPES,

python/cudf/cudf/core/udf/groupby_typing.py

@@ -5,18 +5,19 @@
 from typing import Any
 
 import numba
-from numba import cuda, types
-from numba.core.extending import (
+from numba import cuda
+from numba.cuda import types
+from numba.cuda.cudadecl import registry as cuda_registry
+from numba.cuda.datamodel import cuda_models as models
+from numba.cuda.extending import (
     make_attribute_wrapper,
-    models,
     register_model,
     type_callable,
     typeof_impl,
 )
-from numba.core.typing import signature as nb_signature
-from numba.core.typing.templates import AbstractTemplate, AttributeTemplate
-from numba.cuda.cudadecl import registry as cuda_registry
-from numba.np import numpy_support
+from numba.cuda.np import numpy_support
+from numba.cuda.typing import signature as nb_signature
+from numba.cuda.typing.templates import AbstractTemplate, AttributeTemplate
 
 from cudf.core.udf._ops import arith_ops, comparison_ops, unary_ops
 from cudf.core.udf.utils import Row, UDFError

python/cudf/cudf/core/udf/groupby_utils.py

@@ -7,9 +7,9 @@
 import cupy as cp
 import numpy as np
 from numba import cuda, types
-from numba.core.errors import TypingError
+from numba.cuda.core.errors import TypingError
 from numba.cuda.cudadrv.devices import get_context
-from numba.np import numpy_support
+from numba.cuda.np import numpy_support
 
 from cudf.core.column import as_column, column_empty
 from cudf.core.udf.groupby_typing import (

python/cudf/cudf/core/udf/masked_lowering.py

@@ -4,13 +4,13 @@
 import operator
 
 from llvmlite import ir
-from numba.core import cgutils
-from numba.core.typing import signature as nb_signature
+from numba.cuda import cgutils
 from numba.cuda.cudaimpl import (
     lower as cuda_lower,
     registry as cuda_lowering_registry,
 )
-from numba.extending import lower_builtin, types
+from numba.cuda.extending import lower_builtin, types
+from numba.cuda.typing import signature as nb_signature
 
 from cudf.core.udf import api
 from cudf.core.udf._ops import (

python/cudf/cudf/core/udf/masked_typing.py

@@ -4,23 +4,23 @@
 import operator
 
 import numpy as np
-from numba import types
-from numba.core.datamodel import default_manager
-from numba.core.extending import (
+from numba.cuda import types
+from numba.cuda.cudadecl import registry as cuda_decl_registry
+from numba.cuda.datamodel import cuda_models as models
+from numba.cuda.descriptor import cuda_target
+from numba.cuda.extending import (
     make_attribute_wrapper,
-    models,
     register_model,
     typeof_impl,
 )
-from numba.core.typing import signature as nb_signature
-from numba.core.typing.templates import (
+from numba.cuda.np.numpy_support import from_dtype
+from numba.cuda.typing import signature as nb_signature
+from numba.cuda.typing.templates import (
     AbstractTemplate,
     AttributeTemplate,
     ConcreteTemplate,
 )
-from numba.core.typing.typeof import typeof
-from numba.cuda.cudadecl import registry as cuda_decl_registry
-from numba.np.numpy_support import from_dtype
+from numba.cuda.typing.typeof import typeof
 
 from cudf.core.missing import NA
 from cudf.core.udf import api
@@ -113,7 +113,9 @@ class MaskedType(types.Type):
     def __init__(self, value):
         # MaskedType in Numba shall be parameterized
         # with a value type
-        if default_manager[value].has_nrt_meminfo():
+        if cuda_target.target_context.data_model_manager[
+            value
+        ].has_nrt_meminfo():
             ctx = _current_nrt_context.get(None)
             if ctx is not None:
                 # we're in a compilation that is determining

python/cudf/cudf/core/udf/nrt_utils.py

@@ -4,7 +4,7 @@
 import contextvars
 from contextlib import contextmanager
 
-from numba import config as numba_config
+from numba.cuda import config as numba_config
 
 _current_nrt_context: contextvars.ContextVar = contextvars.ContextVar(
     "current_nrt_context"

python/cudf/cudf/core/udf/row_function.py

@@ -5,7 +5,7 @@
 
 import numpy as np
 from numba import cuda
-from numba.np import numpy_support
+from numba.cuda.np import numpy_support
 
 from cudf.core.udf.api import Masked, pack_return
 from cudf.core.udf.masked_typing import MaskedType
@@ -33,7 +33,7 @@ def _get_frame_row_type(dtype):
     containing these MaskedTypes. Large parts of this function are copied with
     comments from the Numba internals and slightly modified to account for
     validity bools to be present in the final struct. See
-    numba.np.numpy_support.from_struct_dtype for details.
+    numba.cuda.np.numpy_support.from_struct_dtype for details.
     """
 
     # Create the numpy structured type corresponding to the numpy dtype.

python/cudf/cudf/core/udf/scalar_function.py

@@ -4,7 +4,7 @@
 from functools import cache
 
 from numba import cuda
-from numba.np import numpy_support
+from numba.cuda.np import numpy_support
 
 from cudf.core.udf.api import Masked, pack_return
 from cudf.core.udf.masked_typing import MaskedType

python/cudf/cudf/core/udf/strings_lowering.py

@@ -6,12 +6,12 @@
 
 from llvmlite import ir
 from numba import cuda, types
-from numba.core import cgutils
-from numba.core.typing import signature as nb_signature
+from numba.cuda import cgutils
 from numba.cuda.cudaimpl import (
     lower as cuda_lower,
     registry as cuda_lowering_registry,
 )
+from numba.cuda.typing import signature as nb_signature
 
 from cudf._lib.strings_udf import (
     get_character_cases_table_ptr,

python/cudf/cudf/core/udf/strings_typing.py

@@ -5,12 +5,12 @@
 
 import numpy as np
 from numba import types
-from numba.core import cgutils
-from numba.core.extending import models, register_model
-from numba.core.typing import signature as nb_signature
-from numba.core.typing.templates import AbstractTemplate, AttributeTemplate
+from numba.cuda import cgutils
 from numba.cuda.cudadecl import registry as cuda_decl_registry
 from numba.cuda.descriptor import cuda_target
+from numba.cuda.extending import models, register_model
+from numba.cuda.typing import signature as nb_signature
+from numba.cuda.typing.templates import AbstractTemplate, AttributeTemplate
 
 import rmm
 

python/cudf/cudf/core/udf/udf_kernel_base.py

@@ -7,9 +7,9 @@
 
 import numpy as np
 from numba import cuda, typeof
-from numba.core.errors import TypingError
-from numba.np import numpy_support
-from numba.types import CPointer, Poison, Tuple, boolean, int64, void
+from numba.cuda.core.errors import TypingError
+from numba.cuda.np import numpy_support
+from numba.cuda.types import CPointer, Poison, Tuple, boolean, int64, void
 
 from cudf.api.types import is_scalar
 from cudf.core.udf.masked_typing import MaskedType

python/cudf/cudf/core/udf/utils.py

@@ -14,10 +14,11 @@
 import numpy as np
 from cuda.bindings import runtime
 from numba import cuda, typeof
-from numba.core.datamodel import default_manager, models
-from numba.core.extending import register_model
-from numba.np import numpy_support
-from numba.types import CPointer, Record, Tuple, int64, void
+from numba.cuda.datamodel import models
+from numba.cuda.descriptor import cuda_target
+from numba.cuda.extending import register_model
+from numba.cuda.np import numpy_support
+from numba.cuda.types import CPointer, Record, Tuple, int64, void
 
 import rmm
 
@@ -303,7 +304,7 @@ def _get_extensionty_size(ty):
     Return the size of an extension type in bytes
     """
     target_data = ll.create_target_data(_nvvm_data_layout)
-    llty = default_manager[ty].get_value_type()
+    llty = cuda_target.data_model_manager[ty].get_value_type()
     return llty.get_abi_size(target_data)
 
 

python/cudf/cudf/tests/dataframe/methods/test_apply.py

@@ -7,10 +7,10 @@
 import numpy as np
 import pytest
 from numba import cuda
-from numba.core.typing import signature as nb_signature
-from numba.core.typing.templates import AbstractTemplate
 from numba.cuda.cudadecl import registry as cuda_decl_registry
 from numba.cuda.cudaimpl import lower as cuda_lower
+from numba.cuda.typing import signature as nb_signature
+from numba.cuda.typing.templates import AbstractTemplate
 
 import cudf
 from cudf.core.missing import NA

python/cudf/cudf/tests/private_objects/test_compile_udf.py

@@ -2,7 +2,7 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import pytest
-from numba import types
+from numba.cuda import types
 
 from cudf.core.udf.utils import _udf_code_cache, compile_udf
 

python/cudf/cudf/tests/private_objects/test_extension_compilation.py

@@ -7,7 +7,7 @@
 import pytest
 from numba import cuda, types
 from numba.cuda import compile_ptx
-from numba.np.numpy_support import from_dtype
+from numba.cuda.np.numpy_support import from_dtype
 
 import cudf
 from cudf import NA

python/cudf/cudf/tests/private_objects/test_nrt_stats.py

@@ -1,7 +1,7 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 import pytest
-from numba import config
+from numba.cuda import config
 from numba.cuda.memory_management.nrt import rtsys
 
 import cudf

python/cudf/cudf/utils/_numba.py

@@ -3,7 +3,7 @@
 from __future__ import annotations
 
 import numba
-from numba import config as numba_config
+from numba.cuda import config as numba_config
 from packaging import version
 
 

python/cudf/cudf_pandas_tests/test_cudf_pandas.py

@@ -1764,7 +1764,7 @@ def add_one_ufunc(arr):
 
 
 @pytest.mark.filterwarnings(
-    "ignore:Grid size:numba.core.errors.NumbaPerformanceWarning"
+    "ignore:Grid size:numba.cuda.errors.NumbaPerformanceWarning"
 )
 def test_numpy_ndarray_numba_cuda_ufunc(array):
     arr1, arr2 = array

python/pylibcudf/tests/test_transform.py

@@ -86,7 +86,9 @@ def op(a, b, c):
         return (a + b) * c
 
     ptx, _ = cuda.compile_ptx_for_current_device(
-        op, (numba.float64, numba.float64, numba.float64), device=True
+        op,
+        (numba.cuda.float64, numba.cuda.float64, numba.cuda.float64),
+        device=True,
     )
 
     A = 5.0