diff --git a/numba_cuda/numba/cuda/cgutils.py b/numba_cuda/numba/cuda/cgutils.py index 9fbadfff8..b07a83031 100644 --- a/numba_cuda/numba/cuda/cgutils.py +++ b/numba_cuda/numba/cuda/cgutils.py @@ -76,11 +76,12 @@ def create_struct_proxy(fe_type, kind="value"): return res -def copy_struct(dst, src, repl={}): +def copy_struct(dst, src, repl=None): """ Copy structure from *src* to *dst* with replacement from *repl*. """ - repl = repl.copy() + + repl = (repl or {}).copy() # copy data from src or use those in repl for k in src._datamodel._fields: v = repl.pop(k, getattr(src, k)) diff --git a/numba_cuda/numba/cuda/cloudpickle/cloudpickle.py b/numba_cuda/numba/cuda/cloudpickle/cloudpickle.py index dcf2722c3..1d7756d2c 100644 --- a/numba_cuda/numba/cuda/cloudpickle/cloudpickle.py +++ b/numba_cuda/numba/cuda/cloudpickle/cloudpickle.py @@ -969,7 +969,7 @@ def _code_reduce(obj): def _cell_reduce(obj): """Cell (containing values of a function's free variables) reducer.""" try: - obj.cell_contents + obj.cell_contents # noqa: B018 except ValueError: # cell is empty return _make_empty_cell, () else: diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py index 721672f81..4e84fe0f3 100644 --- a/numba_cuda/numba/cuda/compiler.py +++ b/numba_cuda/numba/cuda/compiler.py @@ -824,7 +824,7 @@ def kernel_fixup(kernel, debug): # ret void for block in kernel.blocks: - for i, inst in enumerate(block.instructions): + for inst in block.instructions: if isinstance(inst, ir.Ret): old_ret = block.instructions.pop() block.terminator = None diff --git a/numba_cuda/numba/cuda/core/base.py b/numba_cuda/numba/cuda/core/base.py index 3b7692066..617ced951 100644 --- a/numba_cuda/numba/cuda/core/base.py +++ b/numba_cuda/numba/cuda/core/base.py @@ -880,7 +880,7 @@ def get_dummy_type(self): return GENERIC_POINTER def _compile_subroutine_no_cache( - self, builder, impl, sig, locals={}, flags=None + self, builder, impl, sig, locals=None, flags=None ): """ Invoke the compiler to compile a function to be used inside a @@ -916,7 +916,7 @@ def _compile_subroutine_no_cache( sig.args, sig.return_type, flags, - locals=locals, + locals=locals or {}, ) # Allow inlining the function inside callers. @@ -924,7 +924,7 @@ def _compile_subroutine_no_cache( return cres def compile_subroutine( - self, builder, impl, sig, locals={}, flags=None, caching=True + self, builder, impl, sig, locals=None, flags=None, caching=True ): """ Compile the function *impl* for the given *sig* (in nopython mode). @@ -949,7 +949,7 @@ def compile_subroutine( cached = self.cached_internal_func.get(cache_key) if cached is None: cres = self._compile_subroutine_no_cache( - builder, impl, sig, locals=locals, flags=flags + builder, impl, sig, locals=locals or {}, flags=flags ) self.cached_internal_func[cache_key] = cres @@ -958,12 +958,12 @@ def compile_subroutine( self.active_code_library.add_linking_library(cres.library) return cres - def compile_internal(self, builder, impl, sig, args, locals={}): + def compile_internal(self, builder, impl, sig, args, locals=None): """ Like compile_subroutine(), but also call the function with the given *args*. """ - cres = self.compile_subroutine(builder, impl, sig, locals) + cres = self.compile_subroutine(builder, impl, sig, locals or {}) return self.call_internal(builder, cres.fndesc, sig, args) def call_internal(self, builder, fndesc, sig, args): diff --git a/numba_cuda/numba/cuda/core/boxing.py b/numba_cuda/numba/cuda/core/boxing.py index 5aca1428f..3269ba73b 100644 --- a/numba_cuda/numba/cuda/core/boxing.py +++ b/numba_cuda/numba/cuda/core/boxing.py @@ -1270,11 +1270,9 @@ def object_getattr_safely(obj, attr): ): handle_failure() - setattr( - struct_ptr, - "state_address", - c.unbox(types.uintp, interface_state_address).value, - ) + struct_ptr.state_address = c.unbox( + types.uintp, interface_state_address + ).value # Look up the "state" member and wire it into the struct interface_state = object_getattr_safely(ctypes_binding, "state") @@ -1286,11 +1284,7 @@ def object_getattr_safely(obj, attr): c.builder, stack, interface_state_value ): handle_failure() - setattr( - struct_ptr, - "state", - c.unbox(types.uintp, interface_state_value).value, - ) + struct_ptr.state = c.unbox(types.uintp, interface_state_value).value # Want to store callable function pointers to these CFunctionTypes, so # import ctypes and use it to cast the CFunctionTypes to c_void_p and diff --git a/numba_cuda/numba/cuda/core/bytecode.py b/numba_cuda/numba/cuda/core/bytecode.py index 5db62deff..13ad54e84 100644 --- a/numba_cuda/numba/cuda/core/bytecode.py +++ b/numba_cuda/numba/cuda/core/bytecode.py @@ -182,9 +182,10 @@ def block_effect(self): if PYVERSION in ((3, 13), (3, 14)): def _unpack_opargs(code): - buf = [] - for i, start_offset, op, arg in dis._unpack_opargs(code): - buf.append((start_offset, op, arg)) + buf = [ + (start_offset, op, arg) + for _, start_offset, op, arg in dis._unpack_opargs(code) + ] for i, (start_offset, op, arg) in enumerate(buf): if i + 1 < len(buf): next_offset = buf[i + 1][0] diff --git a/numba_cuda/numba/cuda/core/byteflow.py b/numba_cuda/numba/cuda/core/byteflow.py index 03350f0cd..8f8959caf 100644 --- a/numba_cuda/numba/cuda/core/byteflow.py +++ b/numba_cuda/numba/cuda/core/byteflow.py @@ -293,7 +293,7 @@ def propagate_phi_map(phismap): def apply_changes(used_phis, phismap): keep = {} - for state, used_set in used_phis.items(): + for used_set in used_phis.values(): for phi in used_set: keep[phi] = phismap[phi] _logger.debug("keep phismap: %s", _lazy_pformat(keep)) @@ -1139,7 +1139,7 @@ def op_RAISE_VARARGS(self, state, inst): def op_BEGIN_FINALLY(self, state, inst): temps = [] - for i in range(_EXCEPT_STACK_OFFSET): + for _ in range(_EXCEPT_STACK_OFFSET): tmp = state.make_temp() temps.append(tmp) state.push(tmp) @@ -1656,7 +1656,7 @@ def op_BUILD_MAP(self, state, inst): count = inst.arg items = [] # In 3.5+, BUILD_MAP takes pairs from the stack - for i in range(count): + for _ in range(count): v, k = state.pop(), state.pop() items.append((k, v)) state.append(inst, items=items[::-1], size=count, res=dct) @@ -2360,8 +2360,7 @@ def fork(self, pc, npop=0, npush=0, extra_block=None): stack = stack[:nstack] if npush: assert 0 <= npush - for i in range(npush): - stack.append(self.make_temp()) + stack.extend(self.make_temp() for _ in range(npush)) # Handle changes on the blockstack blockstack = list(self._blockstack) if PYVERSION in ((3, 11), (3, 12), (3, 13), (3, 14)): diff --git a/numba_cuda/numba/cuda/core/compiler_machinery.py b/numba_cuda/numba/cuda/core/compiler_machinery.py index c8d41308a..07b778c3c 100644 --- a/numba_cuda/numba/cuda/core/compiler_machinery.py +++ b/numba_cuda/numba/cuda/core/compiler_machinery.py @@ -101,9 +101,8 @@ def run_finalizer(self, *args, **kwargs): """ return False - def get_analysis_usage(self, AU): + def get_analysis_usage(self, AU): # noqa: B027 """Override to set analysis usage""" - pass def get_analysis(self, pass_name): """ @@ -230,7 +229,7 @@ def add_pass_after(self, pass_cls, location): assert self.passes self._validate_pass(pass_cls) self._validate_pass(location) - for idx, (x, _) in enumerate(self.passes): + for idx, (x, _) in enumerate(self.passes): # noqa: B007 if x == location: break else: diff --git a/numba_cuda/numba/cuda/core/cuda_errors.py b/numba_cuda/numba/cuda/core/cuda_errors.py index ebedec4e1..efca4e4ba 100644 --- a/numba_cuda/numba/cuda/core/cuda_errors.py +++ b/numba_cuda/numba/cuda/core/cuda_errors.py @@ -26,7 +26,7 @@ def _is_numba_core_config_loaded(): To detect if numba.core.config has been initialized due to circular imports. """ try: - numba.cuda.core.config + numba.cuda.core.config # noqa: B018 except AttributeError: return False else: @@ -261,7 +261,7 @@ def __enter__(self): reinit() def __exit__(self, *exc_detail): - Style.RESET_ALL + Style.RESET_ALL # noqa: B018 deinit() class reset_terminal(object): diff --git a/numba_cuda/numba/cuda/core/inline_closurecall.py b/numba_cuda/numba/cuda/core/inline_closurecall.py index 04876ae45..a6e952fe2 100644 --- a/numba_cuda/numba/cuda/core/inline_closurecall.py +++ b/numba_cuda/numba/cuda/core/inline_closurecall.py @@ -137,7 +137,7 @@ def run(self): sized_loops = [(k, len(loops[k].body)) for k in loops.keys()] visited = [] # We go over all loops, bigger loops first (outer first) - for k, s in sorted( + for k, _ in sorted( sized_loops, key=lambda tup: tup[1], reverse=True ): visited.append(k) @@ -1398,7 +1398,7 @@ def is_removed(val, removed): def _find_unsafe_empty_inferred(func_ir, expr): - unsafe_empty_inferred + unsafe_empty_inferred # noqa: B018 require(isinstance(expr, ir.expr_types) and expr.op == "call") callee = expr.func callee_def = get_definition(func_ir, callee) diff --git a/numba_cuda/numba/cuda/core/interpreter.py b/numba_cuda/numba/cuda/core/interpreter.py index 756434b94..0a7634715 100644 --- a/numba_cuda/numba/cuda/core/interpreter.py +++ b/numba_cuda/numba/cuda/core/interpreter.py @@ -779,7 +779,7 @@ def foo(a): while True: # first try and find a matching region # i.e. BUILD_LIST......LIST_TO_TUPLE - def find_postive_region(): + def find_postive_region(blk): found = False for idx in reversed(range(len(blk.body))): stmt = blk.body[idx] @@ -798,7 +798,7 @@ def find_postive_region(): region = (bt, (idx, stmt)) return region - region = find_postive_region() + region = find_postive_region(blk) # if there's a peep hole region then do something with it if region is not None: peep_hole = blk.body[region[1][0] : region[0][0]] @@ -836,7 +836,7 @@ def find_postive_region(): def append_and_fix(x): """Adds to the new_hole and fixes up definitions""" - new_hole.append(x) + new_hole.append(x) # noqa: B023 if x.target.name in func_ir._definitions: # if there's already a definition, drop it, should only # be 1 as the way cpython emits the sequence for @@ -1484,7 +1484,21 @@ def _end_try_blocks(self): See also: _insert_try_block_end """ assert PYVERSION in ((3, 11), (3, 12), (3, 13), (3, 14)) + + def do_change(remain, block): + while remain: + ent = remain.pop() + if ent["kind"] == BlockKind("TRY"): + # Extend block with marker for end of try + self.current_block = block + oldbody = list(block.body) + block.body.clear() + self._insert_try_block_end() + block.body.extend(oldbody) + return True + graph = self.cfa.graph + for offset, block in self.blocks.items(): # Get current blockstack cur_bs = self.dfa.infos[offset].blockstack @@ -1493,25 +1507,13 @@ def _end_try_blocks(self): inc_bs = self.dfa.infos[inc].blockstack # find first diff in the blockstack - for i, (x, y) in enumerate(zip(cur_bs, inc_bs)): + for i, (x, y) in enumerate(zip(cur_bs, inc_bs)): # noqa: B007 if x != y: break else: i = min(len(cur_bs), len(inc_bs)) - def do_change(remain): - while remain: - ent = remain.pop() - if ent["kind"] == BlockKind("TRY"): - # Extend block with marker for end of try - self.current_block = block - oldbody = list(block.body) - block.body.clear() - self._insert_try_block_end() - block.body.extend(oldbody) - return True - - if do_change(list(inc_bs[i:])): + if do_change(list(inc_bs[i:]), block): break def _legalize_exception_vars(self): diff --git a/numba_cuda/numba/cuda/core/ir.py b/numba_cuda/numba/cuda/core/ir.py index 2f66ee0c1..901473aba 100644 --- a/numba_cuda/numba/cuda/core/ir.py +++ b/numba_cuda/numba/cuda/core/ir.py @@ -159,10 +159,9 @@ def strformat(self, nlines_up=2): if lines and use_line > 0: def count_spaces(string): - spaces = 0 - for x in itertools.takewhile(str.isspace, str(string)): - spaces += 1 - return spaces + return sum( + 1 for _ in itertools.takewhile(str.isspace, str(string)) + ) # A few places in the code still use no `loc` or default to line 1 # this is often in places where exceptions are used for the purposes @@ -1547,15 +1546,15 @@ def diff_str(self, other): msg.append("Other block contains more statements") # find the indexes where they don't match - tmp = [] + tmp = set() for idx, stmts in enumerate( zip(block.body, other_blk.body) ): b_s, o_s = stmts if b_s != o_s: - tmp.append(idx) + tmp.add(idx) - def get_pad(ablock, l): + def get_pad(ablock, l, tmp=tmp): pointer = "-> " sp = len(pointer) * " " pad = [] diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py index d8ee41938..6408bfa7a 100644 --- a/numba_cuda/numba/cuda/core/ir_utils.py +++ b/numba_cuda/numba/cuda/core/ir_utils.py @@ -1487,7 +1487,7 @@ def find_single_branch(label): inst = blocks[label].body[0] predecessors = cfg.predecessors(label) delete_block = True - for p, q in predecessors: + for p, _ in predecessors: block = blocks[p] if isinstance(block.body[-1], ir.jump_types): block.body[-1] = copy.copy(inst) @@ -2617,7 +2617,7 @@ def fixup_var_define_in_scope(blocks): # Ensure the scope of each block defines all used variables. for blk in blocks.values(): scope = blk.scope - for var, inst in used_var.items(): + for var in used_var.keys(): # add this variable if it's not in scope if var.name not in scope.localvars: # Note: using a internal method to reuse the same diff --git a/numba_cuda/numba/cuda/core/options.py b/numba_cuda/numba/cuda/core/options.py index ef537e531..c3fd2d86c 100644 --- a/numba_cuda/numba/cuda/core/options.py +++ b/numba_cuda/numba/cuda/core/options.py @@ -151,7 +151,7 @@ def __init__(self, value): if value in ("always", "never"): ok = True else: - ok = hasattr(value, "__call__") + ok = callable(value) if ok: self._inline = value diff --git a/numba_cuda/numba/cuda/core/pythonapi.py b/numba_cuda/numba/cuda/core/pythonapi.py index 1e4fa5ed1..592e25780 100644 --- a/numba_cuda/numba/cuda/core/pythonapi.py +++ b/numba_cuda/numba/cuda/core/pythonapi.py @@ -133,7 +133,7 @@ def add_const(self, const): # All constants are frozen inside the environment if isinstance(const, str): const = sys.intern(const) - for index, val in enumerate(self.env.consts): + for val in self.env.consts: if val is const: break else: @@ -189,9 +189,7 @@ def __init__(self, context, builder): self.module = builder.basic_block.function.module # A unique mapping of serialized objects in this module - try: - self.module.__serialized - except AttributeError: + if not hasattr(self.module, "__serialized"): self.module.__serialized = {} # Initialize types diff --git a/numba_cuda/numba/cuda/core/tracing.py b/numba_cuda/numba/cuda/core/tracing.py index ea8059fc9..a3fc2ec28 100644 --- a/numba_cuda/numba/cuda/core/tracing.py +++ b/numba_cuda/numba/cuda/core/tracing.py @@ -170,7 +170,7 @@ def wrapper(*args, **kwds): if inspect.ismodule(arg0): for n, f in inspect.getmembers(arg0, inspect.isfunction): setattr(arg0, n, decorator(f)) - for n, c in inspect.getmembers(arg0, inspect.isclass): + for _, c in inspect.getmembers(arg0, inspect.isclass): dotrace(c, *args, recursive=recursive) elif inspect.isclass(arg0): for n, f in inspect.getmembers( diff --git a/numba_cuda/numba/cuda/core/transforms.py b/numba_cuda/numba/cuda/core/transforms.py index e8f05de28..ee46c3f19 100644 --- a/numba_cuda/numba/cuda/core/transforms.py +++ b/numba_cuda/numba/cuda/core/transforms.py @@ -892,7 +892,7 @@ def _fix_multi_exit_blocks(func_ir, exit_nodes, *, split_condition=None): # split the block if needed if split_condition is not None: - for pt, stmt in enumerate(blk.body): + for stmt in blk.body: if split_condition(stmt): break else: @@ -925,10 +925,7 @@ def _fix_multi_exit_blocks(func_ir, exit_nodes, *, split_condition=None): common_block.body.append(ir.Jump(post_label, loc=loc)) # Make if-else tree to jump to target - remain_blocks = [] - for remain in remainings: - remain_blocks.append(max_label) - max_label += 1 + remain_blocks = range(max_label, max_label + len(remainings)) switch_block = post_block loc = ir.unknown_loc diff --git a/numba_cuda/numba/cuda/core/typed_passes.py b/numba_cuda/numba/cuda/core/typed_passes.py index e4d046bf3..2d0f3ad9e 100644 --- a/numba_cuda/numba/cuda/core/typed_passes.py +++ b/numba_cuda/numba/cuda/core/typed_passes.py @@ -821,7 +821,7 @@ def _strip_phi_nodes(self, func_ir): exporters = defaultdict(list) phis = set() # Find all variables that needs to be exported - for label, block in func_ir.blocks.items(): + for block in func_ir.blocks.values(): for assign in block.find_insts(ir.assign_types): if isinstance(assign.value, ir.expr_types): if assign.value.op == "phi": diff --git a/numba_cuda/numba/cuda/core/unsafe/refcount.py b/numba_cuda/numba/cuda/core/unsafe/refcount.py index 8064e7264..c47ae7458 100644 --- a/numba_cuda/numba/cuda/core/unsafe/refcount.py +++ b/numba_cuda/numba/cuda/core/unsafe/refcount.py @@ -82,16 +82,14 @@ def codegen(context, builder, signature, args): # A sequence of (type, meminfo) meminfos = [] if context.enable_nrt: - tmp_mis = context.nrt.get_meminfos(builder, ty, obj) - meminfos.extend(tmp_mis) + meminfos.extend(context.nrt.get_meminfos(builder, ty, obj)) refcounts = [] - if meminfos: - for ty, mi in meminfos: - miptr = builder.bitcast(mi, _meminfo_struct_type.as_pointer()) - refctptr = cgutils.gep_inbounds(builder, miptr, 0, 0) - refct = builder.load(refctptr) - refct_32bit = builder.trunc(refct, ir.IntType(32)) - refcounts.append(refct_32bit) + for _, mi in meminfos: + miptr = builder.bitcast(mi, _meminfo_struct_type.as_pointer()) + refctptr = cgutils.gep_inbounds(builder, miptr, 0, 0) + refct = builder.load(refctptr) + refct_32bit = builder.trunc(refct, ir.IntType(32)) + refcounts.append(refct_32bit) return refcounts[0] sig = types.int32(obj) diff --git a/numba_cuda/numba/cuda/core/untyped_passes.py b/numba_cuda/numba/cuda/core/untyped_passes.py index 14a512a57..aab2db3f6 100644 --- a/numba_cuda/numba/cuda/core/untyped_passes.py +++ b/numba_cuda/numba/cuda/core/untyped_passes.py @@ -928,7 +928,7 @@ def inject_loop_body( sentinel_exits = set() sentinel_blocks = [] for lbl, blk in switch_ir.blocks.items(): - for i, stmt in enumerate(blk.body): + for stmt in blk.body: if isinstance(stmt, ir.assign_types): if "SENTINEL" in stmt.target.name: sentinel_blocks.append(lbl) @@ -941,8 +941,8 @@ def inject_loop_body( # find jumps that are non-local, we won't relabel these ignore_set = set() local_lbl = [x for x in loop_ir.blocks.keys()] - for lbl, blk in loop_ir.blocks.items(): - for i, stmt in enumerate(blk.body): + for blk in loop_ir.blocks.values(): + for stmt in blk.body: if isinstance(stmt, ir.jump_types): if stmt.target not in local_lbl: ignore_set.add(stmt.target) diff --git a/numba_cuda/numba/cuda/cudaimpl.py b/numba_cuda/numba/cuda/cudaimpl.py index 337d808b8..ea198010d 100644 --- a/numba_cuda/numba/cuda/cudaimpl.py +++ b/numba_cuda/numba/cuda/cudaimpl.py @@ -928,11 +928,11 @@ def _generic_array( # Compute strides laststride = itemsize - rstrides = [] - for i, lastsize in enumerate(reversed(shape)): - rstrides.append(laststride) + strides = [] + for lastsize in reversed(shape): + strides.append(laststride) laststride *= lastsize - strides = [s for s in reversed(rstrides)] + strides.reverse() kstrides = [context.get_constant(types.intp, s) for s in strides] # Compute shape diff --git a/numba_cuda/numba/cuda/debuginfo.py b/numba_cuda/numba/cuda/debuginfo.py index e4e0106c3..b7cd077e7 100644 --- a/numba_cuda/numba/cuda/debuginfo.py +++ b/numba_cuda/numba/cuda/debuginfo.py @@ -588,7 +588,7 @@ def _di_subroutine_type(self, line, function, argmap): llfunc = function md = [] - for idx, llarg in enumerate(llfunc.args): + for llarg in llfunc.args: if not llarg.name.startswith("arg."): name = llarg.name.replace(".", "$") # for gdb to work correctly lltype = llarg.type @@ -596,7 +596,7 @@ def _di_subroutine_type(self, line, function, argmap): mdtype = self._var_type(lltype, size, datamodel=None) md.append(mdtype) - for idx, (name, nbtype) in enumerate(argmap.items()): + for name, nbtype in argmap.items(): name = name.replace(".", "$") # for gdb to work correctly datamodel = self.cgctx.data_model_manager[nbtype] lltype = self.cgctx.get_value_type(nbtype) @@ -837,7 +837,7 @@ def _di_subroutine_type(self, line, function, argmap): md.append(mdtype) # Create metadata type for arguments - for idx, (name, nbtype) in enumerate(argmap.items()): + for nbtype in argmap.values(): datamodel = self.cgctx.data_model_manager[nbtype] lltype = self.cgctx.get_value_type(nbtype) size = self.cgctx.get_abi_sizeof(lltype) diff --git a/numba_cuda/numba/cuda/decorators.py b/numba_cuda/numba/cuda/decorators.py index 978fae4f6..b2ff5932e 100644 --- a/numba_cuda/numba/cuda/decorators.py +++ b/numba_cuda/numba/cuda/decorators.py @@ -23,7 +23,7 @@ def jit( device=False, inline="never", forceinline=False, - link=[], + link=(), debug=None, opt=None, lineinfo=False, diff --git a/numba_cuda/numba/cuda/deviceufunc.py b/numba_cuda/numba/cuda/deviceufunc.py index 27f711364..ca6be4bd4 100644 --- a/numba_cuda/numba/cuda/deviceufunc.py +++ b/numba_cuda/numba/cuda/deviceufunc.py @@ -128,10 +128,10 @@ def _multi_broadcast(*shapelist): result = shapelist[0] others = shapelist[1:] try: - for i, each in enumerate(others, start=1): + for i, each in enumerate(others, start=1): # noqa: B007 result = _pairwise_broadcast(result, each) except ValueError: - raise ValueError("failed to broadcast argument #{0}".format(i)) + raise ValueError(f"failed to broadcast argument #{i:d}") else: return result @@ -173,7 +173,7 @@ def _fill_argtypes(self): """ for i, ary in enumerate(self.arrays): if ary is not None: - dtype = getattr(ary, "dtype") + dtype = ary.dtype if dtype is None: dtype = np.asarray(ary).dtype self.argtypes[i] = dtype @@ -433,19 +433,21 @@ def to_dtype(ty): class DeviceVectorize(_BaseUFuncBuilder): - def __init__(self, func, identity=None, cache=False, targetoptions={}): + def __init__(self, func, identity=None, cache=False, targetoptions=None): if cache: raise TypeError("caching is not supported") - for opt in targetoptions: + for opt in targetoptions or {}: if opt == "nopython": warnings.warn( "nopython kwarg for cuda target is redundant", RuntimeWarning, ) else: - fmt = "Unrecognized options. " - fmt += "cuda vectorize target does not support option: '%s'" - raise KeyError(fmt % opt) + msg = ( + "Unrecognized options. " + f"cuda vectorize target does not support option: '{opt}'" + ) + raise KeyError(msg) self.py_func = func self.identity = parse_identity(identity) # { arg_dtype: (return_dtype), cudakernel } @@ -505,7 +507,7 @@ def __init__( sig, identity=None, cache=False, - targetoptions={}, + targetoptions=None, writable_args=(), ): if cache: @@ -514,6 +516,8 @@ def __init__( raise TypeError("writable_args are not supported") # Allow nopython flag to be set. + if not targetoptions: + targetoptions = {} if not targetoptions.pop("nopython", True): raise TypeError("nopython flag must be True") # Are there any more target options? diff --git a/numba_cuda/numba/cuda/extending.py b/numba_cuda/numba/cuda/extending.py index 11de8d096..0495cee7a 100644 --- a/numba_cuda/numba/cuda/extending.py +++ b/numba_cuda/numba/cuda/extending.py @@ -132,7 +132,7 @@ def generic(self): def overload( func, - jit_options={}, + jit_options=None, strict=True, inline="never", prefer_literal=False, @@ -201,7 +201,7 @@ def len_impl(seq): # set default options opts = _overload_default_jit_options.copy() - opts.update(jit_options) # let user options override + opts.update(jit_options or {}) # let user options override # TODO: abort now if the kwarg 'target' relates to an unregistered target, # this requires sorting out the circular imports first. diff --git a/numba_cuda/numba/cuda/np/arraymath.py b/numba_cuda/numba/cuda/np/arraymath.py index f5e5aecea..117e641f3 100644 --- a/numba_cuda/numba/cuda/np/arraymath.py +++ b/numba_cuda/numba/cuda/np/arraymath.py @@ -4037,18 +4037,18 @@ def _monotonicity(bins): if last_value < next_value: # Possibly monotonic increasing - for i in range(i + 1, len(bins)): + for j in range(i + 1, len(bins)): last_value = next_value - next_value = bins[i] + next_value = bins[j] if last_value > next_value: return 0 return 1 else: # last > next, possibly monotonic decreasing - for i in range(i + 1, len(bins)): + for j in range(i + 1, len(bins)): last_value = next_value - next_value = bins[i] + next_value = bins[j] if last_value < next_value: return 0 return -1 diff --git a/numba_cuda/numba/cuda/np/arrayobj.py b/numba_cuda/numba/cuda/np/arrayobj.py index 66c69920c..21793a1e1 100644 --- a/numba_cuda/numba/cuda/np/arrayobj.py +++ b/numba_cuda/numba/cuda/np/arrayobj.py @@ -452,7 +452,7 @@ def basic_indexing( if idxty is types.ellipsis: # Fill up missing dimensions at the middle n_missing = aryty.ndim - len(indices) + 1 + num_newaxes - for i in range(n_missing): + for _ in range(n_missing): output_indices.append(zero) output_shapes.append(shapes[ax]) output_strides.append(strides[ax]) @@ -1084,7 +1084,7 @@ def __init__(self, context, builder, aryty, ary, index_types, indices): if idxty is types.ellipsis: # Fill up missing dimensions at the middle n_missing = aryty.ndim - len(indices) + 1 + num_newaxes - for i in range(n_missing): + for _ in range(n_missing): indexer = EntireIndexer(context, builder, aryty, ary, ax) indexers.append(indexer) ax += 1 @@ -1564,12 +1564,8 @@ def codegen(context, builder, sig, args): ) # Hack to get np.broadcast_to to return a read-only array - setattr( - dest, - "parent", - Constant( - context.get_value_type(dest._datamodel.get_type("parent")), None - ), + dest.parent = Constant( + context.get_value_type(dest._datamodel.get_type("parent")), None ) res = dest._getvalue() @@ -1773,7 +1769,7 @@ def numpy_broadcast_arrays(*args): # number of dimensions m = 0 - for idx, arg in enumerate(args): + for arg in args: if isinstance(arg, types.ArrayCompatible): m = max(m, arg.ndim) elif isinstance(arg, (types.Number, types.Boolean, types.BaseTuple)): @@ -2506,7 +2502,7 @@ def impl(a, new_shape): repeats = -(-new_size // a.size) # ceil division res = a - for i in range(repeats - 1): + for _ in range(repeats - 1): res = np.concatenate((res, a)) res = res[:new_size] diff --git a/numba_cuda/numba/cuda/np/npyfuncs.py b/numba_cuda/numba/cuda/np/npyfuncs.py index 806dd301d..112aa342a 100644 --- a/numba_cuda/numba/cuda/np/npyfuncs.py +++ b/numba_cuda/numba/cuda/np/npyfuncs.py @@ -50,7 +50,7 @@ def _check_arity_and_homogeneity(sig, args, arity, return_type=None): fname = inspect.currentframe().f_back.f_code.co_name msg = "{0} called with invalid types: {1}".format(fname, sig) - assert False, msg + raise TypeError(msg) cast_arg_ty = types.float64 diff --git a/numba_cuda/numba/cuda/np/npyimpl.py b/numba_cuda/numba/cuda/np/npyimpl.py index 472f55645..5b1a5d275 100644 --- a/numba_cuda/numba/cuda/np/npyimpl.py +++ b/numba_cuda/numba/cuda/np/npyimpl.py @@ -156,7 +156,7 @@ def create_iter_indices(self): ZERO = ir.Constant(ir.IntType(intpty.width), 0) indices = [] - for i in range(self.ndim): + for _ in range(self.ndim): x = cgutils.alloca_once(self.builder, ir.IntType(intpty.width)) self.builder.store(ZERO, x) indices.append(x) @@ -215,7 +215,7 @@ def create_iter_indices(self): ZERO = ir.Constant(ir.IntType(intpty.width), 0) indices = [] - for i in range(self.ndim - self.inner_arr_ty.ndim): + for _ in range(self.ndim - self.inner_arr_ty.ndim): x = cgutils.alloca_once(self.builder, ir.IntType(intpty.width)) self.builder.store(ZERO, x) indices.append(x) diff --git a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py index cefc56fc1..2a0798aa5 100644 --- a/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py +++ b/numba_cuda/numba/cuda/np/polynomial/polynomial_functions.py @@ -328,7 +328,7 @@ def poly_polyint(c, m=1): def impl(c, m=1): c = np.asarray(c).astype(res_dtype) cdt = c.dtype - for i in range(m): + for _ in range(m): n = len(c) tmp = np.empty((n + 1,) + c.shape[1:], dtype=cdt) diff --git a/numba_cuda/numba/cuda/printimpl.py b/numba_cuda/numba/cuda/printimpl.py index 9141b7113..5f6f34461 100644 --- a/numba_cuda/numba/cuda/printimpl.py +++ b/numba_cuda/numba/cuda/printimpl.py @@ -122,7 +122,7 @@ def print_varargs(context, builder, sig, args): formats = [] values = [] - for i, (argtype, argval) in enumerate(zip(sig.args, args)): + for argtype, argval in zip(sig.args, args): argfmt, argvals = print_item(argtype, context, builder, argval) formats.append(argfmt) values.extend(argvals) diff --git a/numba_cuda/numba/cuda/simulator/kernel.py b/numba_cuda/numba/cuda/simulator/kernel.py index a654cb77b..5d859b684 100644 --- a/numba_cuda/numba/cuda/simulator/kernel.py +++ b/numba_cuda/numba/cuda/simulator/kernel.py @@ -66,7 +66,7 @@ class FakeCUDAKernel(object): Wraps a @cuda.jit-ed function. """ - def __init__(self, fn, device, fastmath=False, extensions=[], debug=False): + def __init__(self, fn, device, fastmath=False, extensions=(), debug=False): self.fn = fn self._device = device self._fastmath = fastmath diff --git a/numba_cuda/numba/cuda/tests/cudadrv/test_streams.py b/numba_cuda/numba/cuda/tests/cudadrv/test_streams.py index 526ac8674..019ab2f7f 100644 --- a/numba_cuda/numba/cuda/tests/cudadrv/test_streams.py +++ b/numba_cuda/numba/cuda/tests/cudadrv/test_streams.py @@ -94,34 +94,5 @@ async def test_cancelled_future(self): self.assertTrue(done2.done()) -@skip_on_cudasim("CUDA Driver API unsupported in the simulator") -class TestFailingStream(CUDATestCase): - # This test can only be run in isolation because it corrupts the CUDA - # context, which cannot be recovered from within the same process. It is - # left here so that it can be run manually for debugging / testing purposes - # - or may be re-enabled if in future there is infrastructure added for - # running tests in a separate process (a subprocess cannot be used because - # CUDA will have been initialized before the fork, so it cannot be used in - # the child process). - @unittest.skip - @with_asyncio_loop - async def test_failed_stream(self): - ctx = cuda.current_context() - module = ctx.create_module_ptx(""" - .version 6.5 - .target sm_30 - .address_size 64 - .visible .entry failing_kernel() { trap; } - """) - failing_kernel = module.get_function("failing_kernel") - - stream = cuda.stream() - failing_kernel.configure((1,), (1,), stream=stream).__call__() - done = stream.async_done() - with self.assertRaises(Exception): - await done - self.assertIsNotNone(done.exception()) - - if __name__ == "__main__": unittest.main() diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_analysis.py b/numba_cuda/numba/cuda/tests/cudapy/test_analysis.py index 70e2d1026..9ae545d48 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_analysis.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_analysis.py @@ -116,15 +116,15 @@ def assert_prune(self, func, args_tys, prune, *args, **kwargs): # what is expected to be pruned expect_removed = [] - for idx, prune in enumerate(prune): + for idx, p in enumerate(prune): branch = before_branches[idx] - if prune is True: + if p is True: expect_removed.append(branch.truebr) - elif prune is False: + elif p is False: expect_removed.append(branch.falsebr) - elif prune is None: + elif p is None: pass # nothing should be removed! - elif prune == "both": + elif p == "both": expect_removed.append(branch.falsebr) expect_removed.append(branch.truebr) else: diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_caching.py b/numba_cuda/numba/cuda/tests/cudapy/test_caching.py index 96f69f30b..3d3eadc32 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_caching.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_caching.py @@ -97,7 +97,7 @@ class DispatcherCacheUsecasesTest(BaseCacheTest): usecases_file = os.path.join(here, "cache_usecases.py") modname = "dispatcher_caching_test_fodder" - def run_in_separate_process(self, *, envvars={}): + def run_in_separate_process(self): # Cached functions can be run from a distinct process. # Also stresses issue #1603: uncached function calling cached function # shouldn't fail compiling. @@ -110,7 +110,6 @@ def run_in_separate_process(self, *, envvars={}): """ % dict(tempdir=self.tempdir, modname=self.modname) subp_env = os.environ.copy() - subp_env.update(envvars) popen = subprocess.Popen( [sys.executable, "-c", code], stdout=subprocess.PIPE, diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py index 68814f3fe..59fb6a0fe 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_extending.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_extending.py @@ -995,7 +995,7 @@ def checker(func, overload_func): def create_message(func, overload_func, func_sig, ol_sig): msg = [] s = ( - f"{func} from module '{getattr(func, '__module__')}' " + f"{func} from module '{func.__module__}' " "has mismatched sig." ) msg.append(s) diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py index 9d1e200bb..89ead77b5 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_ir.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_ir.py @@ -77,7 +77,7 @@ class CheckEquality(CUDATestCase): loc2 = ir.Loc("mock", 2, 0) loc3 = ir.Loc("mock", 3, 0) - def check(self, base, same=[], different=[]): + def check(self, base, same=(), different=()): for s in same: self.assertTrue(base == s) for d in different: @@ -435,7 +435,7 @@ def foo(a, b, c=12, d=1j, e=None): self.assertTrue(x_ir.equal_ir(y_ir)) - def check_diffstr(string, pointing_at=[]): + def check_diffstr(string, pointing_at=()): lines = string.splitlines() for item in pointing_at: for l in lines: diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py index 1f64b1b38..c3d85e8ac 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_ufuncs.py @@ -2,6 +2,7 @@ # SPDX-License-Identifier: BSD-2-Clause import functools +import itertools import warnings import numpy as np import unittest @@ -111,8 +112,8 @@ def _make_ufunc_usecase(self, ufunc): def basic_ufunc_test( self, ufunc, - skip_inputs=[], - additional_inputs=[], + skip_inputs=(), + additional_inputs=(), int_output_type=None, float_output_type=None, kinds="ifc", @@ -124,12 +125,9 @@ def basic_ufunc_test( pyfunc = self._make_ufunc_usecase(ufunc) - inputs = list(self.inputs) + additional_inputs - - for input_tuple in inputs: - input_operand = input_tuple[0] - input_type = input_tuple[1] - + for input_operand, input_type in itertools.chain( + self.inputs, additional_inputs + ): is_tuple = isinstance(input_operand, tuple) if is_tuple: args = input_operand diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py index 3a41b1234..222dee7d0 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py +++ b/numba_cuda/numba/cuda/tests/cudapy/test_vectorize.py @@ -274,8 +274,8 @@ def raising_transfer(*args, **kwargs): old_HtoD = getattr(driver, "cuMemcpyHtoD", None) old_DtoH = getattr(driver, "cuMemcpyDtoH", None) - setattr(driver, "cuMemcpyHtoD", raising_transfer) - setattr(driver, "cuMemcpyDtoH", raising_transfer) + driver.cuMemcpyHtoD = raising_transfer + driver.cuMemcpyDtoH = raising_transfer # Ensure that the mock functions are working as expected @@ -299,11 +299,11 @@ def func(noise): # no original implementation, simply remove ours. if old_HtoD is not None: - setattr(driver, "cuMemcpyHtoD", old_HtoD) + driver.cuMemcpyHtoD = old_HtoD else: del driver.cuMemcpyHtoD if old_DtoH is not None: - setattr(driver, "cuMemcpyDtoH", old_DtoH) + driver.cuMemcpyDtoH = old_DtoH else: del driver.cuMemcpyDtoH diff --git a/numba_cuda/numba/cuda/types/cuda_function_type.py b/numba_cuda/numba/cuda/types/cuda_function_type.py index a7ef4a48b..93b548545 100644 --- a/numba_cuda/numba/cuda/types/cuda_function_type.py +++ b/numba_cuda/numba/cuda/types/cuda_function_type.py @@ -212,7 +212,7 @@ def __init__(self, cres): attribute) """ self.cres = cres - name = getattr(cres.fndesc, "llvm_cfunc_wrapper_name") + name = cres.fndesc.llvm_cfunc_wrapper_name self.address = cres.library.get_pointer_to_function(name) def dump(self, tab=""): diff --git a/numba_cuda/numba/cuda/types/cuda_functions.py b/numba_cuda/numba/cuda/types/cuda_functions.py index 46c9bb2e5..6d6b21c1d 100644 --- a/numba_cuda/numba/cuda/types/cuda_functions.py +++ b/numba_cuda/numba/cuda/types/cuda_functions.py @@ -145,7 +145,7 @@ def template_info(tp): source_kind = src_info.get("kind", "Unknown template") return source_name, source_file, source_lines, source_kind - for i, (k, err_list) in enumerate(self._failures.items()): + for err_list in self._failures.values(): err = err_list[0] nduplicates = len(err_list) template, error = err.template, err.error diff --git a/numba_cuda/numba/cuda/typing/templates.py b/numba_cuda/numba/cuda/typing/templates.py index ac7ec1815..f9d83f95a 100644 --- a/numba_cuda/numba/cuda/typing/templates.py +++ b/numba_cuda/numba/cuda/typing/templates.py @@ -381,7 +381,7 @@ class AbstractTemplate(FunctionTemplate): """ def apply(self, args, kws): - generic = getattr(self, "generic") + generic = self.generic sig = generic(args, kws) # Enforce that *generic()* must return None or Signature if sig is not None: @@ -410,7 +410,7 @@ def unpack_opt(x): return sig def get_template_info(self): - impl = getattr(self, "generic") + impl = self.generic basepath = os.path.dirname( os.path.dirname(os.path.dirname(numba.cuda.__file__)) ) @@ -441,7 +441,7 @@ class CallableTemplate(FunctionTemplate): recvr = None def apply(self, args, kws): - generic = getattr(self, "generic") + generic = self.generic typer = generic() match_sig = inspect.signature(typer) try: @@ -499,7 +499,7 @@ def unpack_opt(x): return self._select(cases, bound.args, bound.kwargs) def get_template_info(self): - impl = getattr(self, "generic") + impl = self.generic basepath = os.path.dirname( os.path.dirname(os.path.dirname(numba.cuda.__file__)) ) @@ -527,7 +527,7 @@ class ConcreteTemplate(FunctionTemplate): """ def apply(self, args, kws): - cases = getattr(self, "cases") + cases = self.cases return self._select(cases, args, kws) def get_template_info(self): diff --git a/numba_cuda/numba/cuda/utils.py b/numba_cuda/numba/cuda/utils.py index ff84846fb..da20c661a 100644 --- a/numba_cuda/numba/cuda/utils.py +++ b/numba_cuda/numba/cuda/utils.py @@ -422,7 +422,7 @@ def __repr__(self): def format_time(tm): units = "s ms us ns ps".split() base = 1 - for unit in units[:-1]: + for _ in units[:-1]: if tm >= base: break base /= 1000 diff --git a/pixi.toml b/pixi.toml index 2974b1ed9..2b122b8b2 100644 --- a/pixi.toml +++ b/pixi.toml @@ -12,6 +12,9 @@ python = ["3.10.*", "3.11.*", "3.12.*", "3.13.*", "3.14.*"] # source numba-cuda = { path = "." } +[feature.ruff.dependencies] +ruff = ">=0.14.10" + [feature.cu.dependencies] cuda-nvcc = "*" cuda-nvcc-impl = "*" @@ -106,6 +109,7 @@ nvidia-sphinx-theme = "*" [environments] default = { features = ["cu-13-1", "test", "cu", "cu-13", "cu-rt", "nvvm", "py314"], solve-group = "default" } +dev = { features = ["ruff"], no-default-feature = true } bench-against = { features = ["test"], no-default-feature = true } # CUDA 12 cu-12-0-py310 = { features = [ @@ -252,9 +256,10 @@ benchcmp = { cmd = [ "--benchmark-group-by=name", "--benchmark-compare", ] } - -[target.linux.tasks.bench-against] -cmd = ["python", "$PIXI_PROJECT_ROOT/scripts/bench-against.py"] +bench-against = { cmd = [ + "python", + "$PIXI_PROJECT_ROOT/scripts/bench-against.py", +] } [target.linux.tasks.build-docs] cmd = ["make", "-C", "$PIXI_PROJECT_ROOT/docs", "html"] diff --git a/pyproject.toml b/pyproject.toml index 0f66fbb98..b1268dc7e 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -3,25 +3,23 @@ [build-system] build-backend = "setuptools.build_meta" -requires = [ - "setuptools", - "wheel", - "numpy", -] +requires = ["setuptools", "wheel", "numpy"] [project] name = "numba-cuda" dynamic = ["version"] description = "CUDA target for Numba" readme = { file = "README.md", content-type = "text/markdown" } -authors = [ - { name = "Anaconda Inc." }, - { name = "NVIDIA Corporation" } -] +authors = [{ name = "Anaconda Inc." }, { name = "NVIDIA Corporation" }] license = "BSD-2-Clause" license-files = ["LICENSE", "LICENSE.numba"] requires-python = ">=3.9" -dependencies = ["numba>=0.60.0", "cuda-bindings>=12.9.1,<14.0.0", "cuda-core>=0.3.2,<1.0.0", "packaging"] +dependencies = [ + "numba>=0.60.0", + "cuda-bindings>=12.9.1,<14.0.0", + "cuda-core>=0.3.2,<1.0.0", + "packaging", +] [project.optional-dependencies] cu12 = [ @@ -50,14 +48,8 @@ test = [ "ml_dtypes", "statistics", ] -test-cu12 = [ - "cuda-toolkit[curand]==12.*", - { include-group = "test" } -] -test-cu13 = [ - "cuda-toolkit[curand]==13.*", - { include-group = "test" } -] +test-cu12 = ["cuda-toolkit[curand]==12.*", { include-group = "test" }] +test-cu13 = ["cuda-toolkit[curand]==13.*", { include-group = "test" }] [project.urls] Homepage = "https://nvidia.github.io/numba-cuda/" @@ -67,7 +59,7 @@ License = "https://github.com/NVIDIA/numba-cuda/blob/main/LICENSE" Issues = "https://github.com/NVIDIA/numba-cuda/issues" [tool.setuptools.dynamic] -version = {attr = "numba_cuda.__version__"} +version = { attr = "numba_cuda.__version__" } [tool.setuptools.packages.find] include = ["numba_cuda*"] @@ -99,7 +91,12 @@ ignore = [ "E731", # Ambiguous variable names "E741", + "B028", + "B904", + "B905", + "B019", ] +select = ["B"] extend-select = ["PERF"] fixable = ["ALL"] @@ -119,7 +116,7 @@ exclude = [ [tool.ruff.lint.per-file-ignores] # Slightly long line in the standard version file -"**/test*.py" = ["PERF"] +"**/test*.py" = ["PERF", "B007", "B018", "B023"] "**/simulator/*" = ["PERF"] "numba_cuda/_version.py" = ["E501"] # "Unused" imports / potentially undefined names in init files