From fbdb2ce1c9ee3e174aa86c16f46999b87084bd2d Mon Sep 17 00:00:00 2001 From: Atmn Patel Date: Wed, 27 Aug 2025 12:51:55 -0700 Subject: [PATCH 1/4] [Cleanup] remove dangling import to upstream cgutils --- numba_cuda/numba/cuda/lowering.py | 2 +- numba_cuda/numba/cuda/np/npyfuncs.py | 3 ++- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py index 8cdfe104e..dcf4066ca 100644 --- a/numba_cuda/numba/cuda/lowering.py +++ b/numba_cuda/numba/cuda/lowering.py @@ -17,10 +17,10 @@ funcdesc, generators, config, - cgutils, removerefctpass, targetconfig, ) +from numba.cuda import cgutils from numba.cuda.core import ir_utils from numba.core.errors import ( LoweringError, diff --git a/numba_cuda/numba/cuda/np/npyfuncs.py b/numba_cuda/numba/cuda/np/npyfuncs.py index 78cf49f0f..376ef7b17 100644 --- a/numba_cuda/numba/cuda/np/npyfuncs.py +++ b/numba_cuda/numba/cuda/np/npyfuncs.py @@ -14,7 +14,8 @@ from numba.core.extending import overload from numba.core.imputils import impl_ret_untracked -from numba.core import typing, types, errors, cgutils, config +from numba.core import typing, types, errors, config +from numba.cuda import cgutils from numba.core.extending import register_jitable from numba.np import npdatetime from numba.np.math import cmathimpl, mathimpl, numbers From 3ac146ec110e32ce06e3a0aa7187e86f9344df3c Mon Sep 17 00:00:00 2001 From: Atmn Patel Date: Wed, 27 Aug 2025 13:30:04 -0700 Subject: [PATCH 2/4] [Cleanup] Remove unsupported prange usage on CUDA target in ir_utils --- numba_cuda/numba/cuda/core/ir_utils.py | 3 --- 1 file changed, 3 deletions(-) diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py index e8425d91b..c5b079595 100644 --- a/numba_cuda/numba/cuda/core/ir_utils.py +++ b/numba_cuda/numba/cuda/core/ir_utils.py @@ -813,7 +813,6 @@ def has_no_side_effect(rhs, lives, call_table): """Returns True if this expression has no side effects that would prevent re-ordering. """ - from numba.misc.special import prange if isinstance(rhs, ir.Expr) and rhs.op == "call": func_name = rhs.func.name @@ -826,8 +825,6 @@ def has_no_side_effect(rhs, lives, call_table): or call_list == ["stencil", numba] or call_list == ["log", numpy] or call_list == ["dtype", numpy] - or call_list == [prange] - or call_list == ["prange", numba] or call_list == ["pndindex", numba] or call_list == ["ceil", math] or call_list == [max] From e1d43b558bf269c1c53f4777abdfc41dbadae779 Mon Sep 17 00:00:00 2001 From: Atmn Patel Date: Wed, 27 Aug 2025 13:33:00 -0700 Subject: [PATCH 3/4] [Cleanup] always assume legacy Numba type system --- numba_cuda/numba/cuda/core/ir_utils.py | 12 +---- numba_cuda/numba/cuda/np/npyfuncs.py | 7 +-- numba_cuda/numba/cuda/np/numpy_support.py | 55 ++++++++--------------- 3 files changed, 22 insertions(+), 52 deletions(-) diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py index c5b079595..69839c6ed 100644 --- a/numba_cuda/numba/cuda/core/ir_utils.py +++ b/numba_cuda/numba/cuda/core/ir_utils.py @@ -249,12 +249,7 @@ def mk_range_block(typemap, start, stop, step, calltypes, scope, loc): range_call_assign = ir.Assign(range_call, range_call_var, loc) # iter_var = getiter(range_call_var) iter_call = ir.Expr.getiter(range_call_var, loc) - if config.USE_LEGACY_TYPE_SYSTEM: - calltype_sig = signature( - types.range_iter64_type, types.range_state64_type - ) - else: - calltype_sig = signature(types.range_iter_type, types.range_state_type) + calltype_sig = signature(types.range_iter64_type, types.range_state64_type) calltypes[iter_call] = calltype_sig iter_var = ir.Var(scope, mk_unique_var("$iter_var"), loc) typemap[iter_var.name] = types.iterators.RangeIteratorType(types.intp) @@ -333,10 +328,7 @@ def mk_loop_header(typemap, phi_var, calltypes, scope, loc): types.intp, types.boolean ) iternext_call = ir.Expr.iternext(phi_var, loc) - if config.USE_LEGACY_TYPE_SYSTEM: - range_iter_type = types.range_iter64_type - else: - range_iter_type = types.range_iter_type + range_iter_type = types.range_iter64_type calltypes[iternext_call] = signature( types.containers.Pair(types.intp, types.boolean), range_iter_type ) diff --git a/numba_cuda/numba/cuda/np/npyfuncs.py b/numba_cuda/numba/cuda/np/npyfuncs.py index 376ef7b17..2eea1746d 100644 --- a/numba_cuda/numba/cuda/np/npyfuncs.py +++ b/numba_cuda/numba/cuda/np/npyfuncs.py @@ -14,7 +14,7 @@ from numba.core.extending import overload from numba.core.imputils import impl_ret_untracked -from numba.core import typing, types, errors, config +from numba.core import typing, types, errors from numba.cuda import cgutils from numba.core.extending import register_jitable from numba.np import npdatetime @@ -52,10 +52,7 @@ def _check_arity_and_homogeneity(sig, args, arity, return_type=None): assert False, msg -if getattr(config, "USE_LEGACY_TYPE_SYSTEM", True): - cast_arg_ty = types.float64 -else: - cast_arg_ty = types.np_float64 +cast_arg_ty = types.float64 def _call_func_by_name_with_cast( diff --git a/numba_cuda/numba/cuda/np/numpy_support.py b/numba_cuda/numba/cuda/np/numpy_support.py index 8f81e7cdc..3f037d037 100644 --- a/numba_cuda/numba/cuda/np/numpy_support.py +++ b/numba_cuda/numba/cuda/np/numpy_support.py @@ -3,48 +3,29 @@ import numpy as np import re -from numba.core import types, errors, config +from numba.core import types, errors numpy_version = tuple(map(int, np.__version__.split(".")[:2])) -if getattr(config, "USE_LEGACY_TYPE_SYSTEM", True): - FROM_DTYPE = { - np.dtype("bool"): types.boolean, - np.dtype("int8"): types.int8, - np.dtype("int16"): types.int16, - np.dtype("int32"): types.int32, - np.dtype("int64"): types.int64, - np.dtype("uint8"): types.uint8, - np.dtype("uint16"): types.uint16, - np.dtype("uint32"): types.uint32, - np.dtype("uint64"): types.uint64, - np.dtype("float32"): types.float32, - np.dtype("float64"): types.float64, - np.dtype("float16"): types.float16, - np.dtype("complex64"): types.complex64, - np.dtype("complex128"): types.complex128, - np.dtype(object): types.pyobject, - } -else: - FROM_DTYPE = { - np.dtype("bool"): types.np_bool_, - np.dtype("int8"): types.np_int8, - np.dtype("int16"): types.np_int16, - np.dtype("int32"): types.np_int32, - np.dtype("int64"): types.np_int64, - np.dtype("uint8"): types.np_uint8, - np.dtype("uint16"): types.np_uint16, - np.dtype("uint32"): types.np_uint32, - np.dtype("uint64"): types.np_uint64, - np.dtype("float32"): types.np_float32, - np.dtype("float64"): types.np_float64, - np.dtype("float16"): types.np_float16, - np.dtype("complex64"): types.np_complex64, - np.dtype("complex128"): types.np_complex128, - np.dtype(object): types.pyobject, - } +FROM_DTYPE = { + np.dtype("bool"): types.boolean, + np.dtype("int8"): types.int8, + np.dtype("int16"): types.int16, + np.dtype("int32"): types.int32, + np.dtype("int64"): types.int64, + np.dtype("uint8"): types.uint8, + np.dtype("uint16"): types.uint16, + np.dtype("uint32"): types.uint32, + np.dtype("uint64"): types.uint64, + np.dtype("float32"): types.float32, + np.dtype("float64"): types.float64, + np.dtype("float16"): types.float16, + np.dtype("complex64"): types.complex64, + np.dtype("complex128"): types.complex128, + np.dtype(object): types.pyobject, +} re_typestr = re.compile(r"[<>=\|]([a-z])(\d+)?$", re.I) From 6fe5f1e5562b2fab0fa4b56ac224efabdad39102 Mon Sep 17 00:00:00 2001 From: Atmn Patel Date: Wed, 27 Aug 2025 14:20:12 -0700 Subject: [PATCH 4/4] [Cleanup] Replace dangling import to upstream compile_result with vendored in cuda_compile_result --- numba_cuda/numba/cuda/core/typed_passes.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/numba_cuda/numba/cuda/core/typed_passes.py b/numba_cuda/numba/cuda/core/typed_passes.py index b56497f5a..9d17575d5 100644 --- a/numba_cuda/numba/cuda/core/typed_passes.py +++ b/numba_cuda/numba/cuda/core/typed_passes.py @@ -451,9 +451,9 @@ def run_pass(self, state): lowered = state["cr"] signature = typing.signature(state.return_type, *state.args) - from numba.core.compiler import compile_result + from numba.cuda.compiler import cuda_compile_result - state.cr = compile_result( + state.cr = cuda_compile_result( typing_context=state.typingctx, target_context=state.targetctx, entry_point=lowered.cfunc,