diff --git a/numba_cuda/numba/cuda/compiler.py b/numba_cuda/numba/cuda/compiler.py index a40ae2e45..033f7105c 100644 --- a/numba_cuda/numba/cuda/compiler.py +++ b/numba_cuda/numba/cuda/compiler.py @@ -57,15 +57,9 @@ IRLegalization, NopythonTypeInference, NopythonRewrites, - PreParforPass, - ParforPass, - DumpParforDiagnostics, InlineOverloads, PreLowerStripPhis, - NativeParforLowering, NoPythonSupportedFeatureValidation, - ParforFusionPass, - ParforPreLoweringPass, ) @@ -163,12 +157,8 @@ def define_nopython_lowering_pipeline(state, name="nopython_lowering"): # Annotate only once legalized pm.add_pass(AnnotateTypes, "annotate types") # lower - if state.flags.auto_parallel.enabled: - pm.add_pass(NativeParforLowering, "native parfor lowering") - else: - pm.add_pass(NativeLowering, "native lowering") + pm.add_pass(NativeLowering, "native lowering") pm.add_pass(CUDABackend, "nopython mode backend") - pm.add_pass(DumpParforDiagnostics, "dump parfor diagnostics") pm.finalize() return pm @@ -186,10 +176,7 @@ def define_parfor_gufunc_nopython_lowering_pipeline( # Annotate only once legalized pm.add_pass(AnnotateTypes, "annotate types") # lower - if state.flags.auto_parallel.enabled: - pm.add_pass(NativeParforLowering, "native parfor lowering") - else: - pm.add_pass(NativeLowering, "native lowering") + pm.add_pass(NativeLowering, "native lowering") pm.add_pass(CUDABackend, "nopython mode backend") pm.finalize() return pm @@ -206,26 +193,8 @@ def define_typed_pipeline(state, name="typed"): # optimisation pm.add_pass(InlineOverloads, "inline overloaded functions") - if state.flags.auto_parallel.enabled: - pm.add_pass(PreParforPass, "Preprocessing for parfors") if not state.flags.no_rewrites: pm.add_pass(NopythonRewrites, "nopython rewrites") - if state.flags.auto_parallel.enabled: - pm.add_pass(ParforPass, "convert to parfors") - pm.add_pass(ParforFusionPass, "fuse parfors") - pm.add_pass(ParforPreLoweringPass, "parfor prelowering") - - pm.finalize() - return pm - - @staticmethod - def define_parfor_gufunc_pipeline(state, name="parfor_gufunc_typed"): - """Returns the typed part of the nopython pipeline""" - pm = PassManager(name) - assert state.func_ir - pm.add_pass(IRProcessing, "processing IR") - pm.add_pass(NopythonTypeInference, "nopython frontend") - pm.add_pass(ParforPreLoweringPass, "parfor prelowering") pm.finalize() return pm diff --git a/numba_cuda/numba/cuda/core/compiler.py b/numba_cuda/numba/cuda/core/compiler.py index 68f2556bc..0700cf97c 100644 --- a/numba_cuda/numba/cuda/core/compiler.py +++ b/numba_cuda/numba/cuda/core/compiler.py @@ -6,7 +6,6 @@ from numba.cuda.core import bytecode from numba.core import callconv, config, errors from numba.core.errors import CompilerError -from numba.parfors.parfor import ParforDiagnostics from numba.cuda.core.untyped_passes import ExtractByteCode, FixupArgs from numba.core.targetconfig import ConfigStack @@ -66,8 +65,6 @@ def _make_subtarget(targetctx, flags): subtargetoptions["enable_boundscheck"] = True if flags.nrt: subtargetoptions["enable_nrt"] = True - if flags.auto_parallel: - subtargetoptions["auto_parallel"] = flags.auto_parallel if flags.fastmath: subtargetoptions["fastmath"] = flags.fastmath error_model = callconv.create_error_model(flags.error_model, targetctx) @@ -114,13 +111,6 @@ def __init__( # hold this for e.g. with_lifting, null out on exit self.state.pipeline = self - # parfor diagnostics info, add to metadata - self.state.parfor_diagnostics = ParforDiagnostics() - self.state.metadata["parfor_diagnostics"] = ( - self.state.parfor_diagnostics - ) - self.state.metadata["parfors"] = {} - self.state.status = _CompileStatus( can_fallback=self.state.flags.enable_pyobject ) diff --git a/numba_cuda/numba/cuda/core/ir_utils.py b/numba_cuda/numba/cuda/core/ir_utils.py index 09b810527..e8425d91b 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.parfors import array_analysis, parfor from numba.misc.special import prange if isinstance(rhs, ir.Expr) and rhs.op == "call": @@ -827,11 +826,9 @@ 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 == [array_analysis.wrap_index] or call_list == [prange] or call_list == ["prange", numba] or call_list == ["pndindex", numba] - or call_list == [parfor.internal_prange] or call_list == ["ceil", math] or call_list == [max] or call_list == [int] diff --git a/numba_cuda/numba/cuda/core/typed_passes.py b/numba_cuda/numba/cuda/core/typed_passes.py index f2b9f4d1b..b56497f5a 100644 --- a/numba_cuda/numba/cuda/core/typed_passes.py +++ b/numba_cuda/numba/cuda/core/typed_passes.py @@ -4,7 +4,6 @@ import abc from contextlib import contextmanager from collections import defaultdict, namedtuple -from functools import partial from copy import copy import warnings @@ -20,15 +19,6 @@ lowering, ) -from numba.parfors.parfor import PreParforPass as _parfor_PreParforPass -from numba.parfors.parfor import ParforPass as _parfor_ParforPass -from numba.parfors.parfor import ParforFusionPass as _parfor_ParforFusionPass -from numba.parfors.parfor import ( - ParforPreLoweringPass as _parfor_ParforPreLoweringPass, -) -from numba.parfors.parfor import Parfor -from numba.parfors.parfor_lowering import ParforLower - from numba.cuda.core.compiler_machinery import ( FunctionPass, LoweringPass, @@ -44,10 +34,8 @@ dead_code_elimination, simplify_CFG, get_definition, - build_definitions, compute_cfg_from_blocks, is_operator_or_getitem, - replace_vars, ) from numba.cuda.core import postproc from llvmlite import binding as llvm @@ -304,182 +292,6 @@ def run_pass(self, state): return True -@register_pass(mutates_CFG=True, analysis_only=False) -class PreParforPass(FunctionPass): - _name = "pre_parfor_pass" - - def __init__(self): - FunctionPass.__init__(self) - - def run_pass(self, state): - """ - Preprocessing for data-parallel computations. - """ - # Ensure we have an IR and type information. - assert state.func_ir - preparfor_pass = _parfor_PreParforPass( - state.func_ir, - state.typemap, - state.calltypes, - state.typingctx, - state.targetctx, - state.flags.auto_parallel, - state.parfor_diagnostics.replaced_fns, - ) - - preparfor_pass.run() - return True - - -# this is here so it pickles and for no other reason -def _reload_parfors(): - """Reloader for cached parfors""" - # Re-initialize the parallel backend when load from cache. - from numba.np.ufunc.parallel import _launch_threads - - _launch_threads() - - -@register_pass(mutates_CFG=True, analysis_only=False) -class ParforPass(FunctionPass): - _name = "parfor_pass" - - def __init__(self): - FunctionPass.__init__(self) - - def run_pass(self, state): - """ - Convert data-parallel computations into Parfor nodes - """ - # Ensure we have an IR and type information. - assert state.func_ir - parfor_pass = _parfor_ParforPass( - state.func_ir, - state.typemap, - state.calltypes, - state.return_type, - state.typingctx, - state.targetctx, - state.flags.auto_parallel, - state.flags, - state.metadata, - state.parfor_diagnostics, - ) - parfor_pass.run() - - # check the parfor pass worked and warn if it didn't - has_parfor = False - for blk in state.func_ir.blocks.values(): - for stmnt in blk.body: - if isinstance(stmnt, Parfor): - has_parfor = True - break - else: - continue - break - - if not has_parfor: - # parfor calls the compiler chain again with a string - if not ( - config.DISABLE_PERFORMANCE_WARNINGS - or state.func_ir.loc.filename == "" - ): - url = ( - "https://numba.readthedocs.io/en/stable/user/" - "parallel.html#diagnostics" - ) - msg = ( - "\nThe keyword argument 'parallel=True' was specified " - "but no transformation for parallel execution was " - "possible.\n\nTo find out why, try turning on parallel " - "diagnostics, see %s for help." % url - ) - warnings.warn( - errors.NumbaPerformanceWarning(msg, state.func_ir.loc) - ) - - # Add reload function to initialize the parallel backend. - state.reload_init.append(_reload_parfors) - return True - - -@register_pass(mutates_CFG=True, analysis_only=False) -class ParforFusionPass(FunctionPass): - _name = "parfor_fusion_pass" - - def __init__(self): - FunctionPass.__init__(self) - - def run_pass(self, state): - """ - Do fusion of parfor nodes. - """ - # Ensure we have an IR and type information. - assert state.func_ir - parfor_pass = _parfor_ParforFusionPass( - state.func_ir, - state.typemap, - state.calltypes, - state.return_type, - state.typingctx, - state.targetctx, - state.flags.auto_parallel, - state.flags, - state.metadata, - state.parfor_diagnostics, - ) - parfor_pass.run() - - return True - - -@register_pass(mutates_CFG=True, analysis_only=False) -class ParforPreLoweringPass(FunctionPass): - _name = "parfor_prelowering_pass" - - def __init__(self): - FunctionPass.__init__(self) - - def run_pass(self, state): - """ - Prepare parfors for lowering. - """ - # Ensure we have an IR and type information. - assert state.func_ir - parfor_pass = _parfor_ParforPreLoweringPass( - state.func_ir, - state.typemap, - state.calltypes, - state.return_type, - state.typingctx, - state.targetctx, - state.flags.auto_parallel, - state.flags, - state.metadata, - state.parfor_diagnostics, - ) - parfor_pass.run() - - return True - - -@register_pass(mutates_CFG=False, analysis_only=True) -class DumpParforDiagnostics(AnalysisPass): - _name = "dump_parfor_diagnostics" - - def __init__(self): - AnalysisPass.__init__(self) - - def run_pass(self, state): - if state.flags.auto_parallel.enabled: - if config.PARALLEL_DIAGNOSTICS: - if state.parfor_diagnostics is not None: - state.parfor_diagnostics.dump(config.PARALLEL_DIAGNOSTICS) - else: - raise RuntimeError("Diagnostics failed.") - return True - - class BaseNativeLowering(abc.ABC, LoweringPass): """The base class for a lowering pass. The lowering functionality must be specified in inheriting classes by providing an appropriate lowering class @@ -596,18 +408,6 @@ def lowering_class(self): return lowering.Lower -@register_pass(mutates_CFG=True, analysis_only=False) -class NativeParforLowering(BaseNativeLowering): - """Lowering pass for a native function IR described using Numba's standard - `numba.core.ir` nodes and also parfor.Parfor nodes.""" - - _name = "native_parfor_lowering" - - @property - def lowering_class(self): - return ParforLower - - @register_pass(mutates_CFG=False, analysis_only=True) class NoPythonSupportedFeatureValidation(AnalysisPass): """NoPython Mode check: Validates the IR to ensure that features in use are @@ -989,10 +789,6 @@ def __init__(self): def run_pass(self, state): state.func_ir = self._strip_phi_nodes(state.func_ir) - state.func_ir._definitions = build_definitions(state.func_ir.blocks) - if "flags" in state and state.flags.auto_parallel.enabled: - self._simplify_conditionally_defined_variable(state.func_ir) - state.func_ir._definitions = build_definitions(state.func_ir.blocks) # Rerun postprocessor to update metadata post_proc = postproc.PostProcessor(state.func_ir) @@ -1071,78 +867,3 @@ def _strip_phi_nodes(self, func_ir): func_ir.blocks = newblocks return func_ir - - def _simplify_conditionally_defined_variable(self, func_ir): - """ - Rewrite assignments like: - - ver1 = null() - ... - ver1 = ver - ... - uses(ver1) - - into: - # delete all assignments to ver1 - uses(ver) - - This is only needed for parfors because the SSA pass will create extra - variable assignments that the parfor code does not expect. - This pass helps avoid problems by reverting the effect of SSA. - """ - any_block = next(iter(func_ir.blocks.values())) - scope = any_block.scope - defs = func_ir._definitions - - def unver_or_undef(unver, defn): - # Is the definition undefined or pointing to the unversioned name? - if isinstance(defn, ir.Var): - if defn.unversioned_name == unver: - return True - elif isinstance(defn, ir.Expr): - if defn.op == "null": - return True - return False - - def legalize_all_versioned_names(var): - # Are all versioned names undefined or defined to the same - # variable chain? - if not var.versioned_names: - return False - for versioned in var.versioned_names: - vs = defs.get(versioned, ()) - if not all(map(partial(unver_or_undef, k), vs)): - return False - return True - - # Find unversioned variables that met the conditions - suspects = set() - for k in defs: - try: - # This may fail? - var = scope.get_exact(k) - except errors.NotDefinedError: - continue - # is the var name unversioned? - if var.unversioned_name == k: - if legalize_all_versioned_names(var): - suspects.add(var) - - delete_set = set() - replace_map = {} - for var in suspects: - # rewrite Var uses to the unversioned name - for versioned in var.versioned_names: - ver_var = scope.get_exact(versioned) - # delete assignment to the versioned name - delete_set.add(ver_var) - # replace references to versioned name with the unversioned - replace_map[versioned] = var - - # remove assignments to the versioned names - for _label, blk in func_ir.blocks.items(): - for assign in blk.find_insts(ir.Assign): - if assign.target in delete_set: - blk.remove(assign) - # do variable replacement - replace_vars(func_ir.blocks, replace_map) diff --git a/numba_cuda/numba/cuda/core/untyped_passes.py b/numba_cuda/numba/cuda/core/untyped_passes.py index 7d0f739d1..d21996d27 100644 --- a/numba_cuda/numba/cuda/core/untyped_passes.py +++ b/numba_cuda/numba/cuda/core/untyped_passes.py @@ -240,7 +240,7 @@ def run_pass(self, state): inline_pass = InlineClosureCallPass( state.func_ir, state.flags.auto_parallel, - state.parfor_diagnostics.replaced_fns, + None, typed_pass, ) inline_pass.run() diff --git a/numba_cuda/numba/cuda/dispatcher.py b/numba_cuda/numba/cuda/dispatcher.py index 5958206fe..0fb04d685 100644 --- a/numba_cuda/numba/cuda/dispatcher.py +++ b/numba_cuda/numba/cuda/dispatcher.py @@ -213,13 +213,11 @@ def link_to_library_functions( # The following are referred to by the cache implementation. Note: # - There are no referenced environments in CUDA. # - Kernels don't have lifted code. - # - reload_init is only for parfors. self.target_context = tgt_ctx self.fndesc = cres.fndesc self.environment = cres.environment self._referenced_environments = [] self.lifted = [] - self.reload_init = [] def maybe_link_nrt(self, link, tgt_ctx, asm): """ @@ -1347,27 +1345,6 @@ def stats(self): cache_misses=self._cache_misses, ) - def parallel_diagnostics(self, signature=None, level=1): - """ - Print parallel diagnostic information for the given signature. If no - signature is present it is printed for all known signatures. level is - used to adjust the verbosity, level=1 (default) is minimal verbosity, - and 2, 3, and 4 provide increasing levels of verbosity. - """ - - def dump(sig): - ol = self.overloads[sig] - pfdiag = ol.metadata.get("parfor_diagnostics", None) - if pfdiag is None: - msg = "No parfors diagnostic available, is 'parallel=True' set?" - raise ValueError(msg) - pfdiag.dump(level) - - if signature is not None: - dump(signature) - else: - [dump(sig) for sig in self.signatures] - def get_metadata(self, signature=None): """ Obtain the compilation metadata for a given signature.