codegen: decide the OpenMP simd clause in a preprocessing pass - #2503
Draft
ThrudPrimrose wants to merge 7 commits into
Draft
codegen: decide the OpenMP simd clause in a preprocessing pass#2503ThrudPrimrose wants to merge 7 commits into
ThrudPrimrose wants to merge 7 commits into
Conversation
MarkSIMDMaps sets Map.omp_simd on the innermost Sequential and CPU_Multicore maps whose body holds no further loop or map, splitting a multidimensional map's innermost dimension off with MapExpansion first, since the clause vectorizes the loop it precedes. min/max reductions and Sequential scatters keep the clause off. Code generation only renders what the pass decided.
… leaf map now carries it The Fortran frontend outlines every map body into a NestedSDFG, so the pass has to open one rather than refuse it on sight; the parametrized case pins that, including the loop and map payloads that withhold the clause. openmp_test asserts the emitted pragma text, which now reads 'parallel for simd' on the leaf map it builds.
A Sequential map lowers WCR to wcr_fixed::reduce, a read-modify-write in the loop body: an accumulation into a fixed location carries across iterations and a scatter can alias across them, neither of which simd permits. CPU_Multicore keeps the clause, since its conflicted WCR goes through reduce_atomic - pinned by running the sum.
…x what it uncovered An unexpanded library node lowers to a Map after this pass runs, and the OpenMP Reduce expansion writes "#pragma omp parallel for" into a tasklet body; either lands inside the simd region, which gcc rejects. MapExpansion left a stale collapse count on the map it reuses, and omp_simd did not serialize for Sequential maps - the schedule this pass marks most often.
PatternTransformation.apply_to propagates memlets across the entire SDFG after every application unless told otherwise, so one whole-SDFG propagation rode along with each candidate map: quadratic in the map count, and code generation for a LLaMA-sized model never finished. MapExpansion already propagates the scope it rewrote, and hands back the new nest, so the pass turns the annotation off and takes the innermost entry from the return value instead of walking the scope tree again.
…tion guard for what it is Nothing checked the nest MapExpansion leaves behind when the pass splits a marked multidimensional map, so a permuted or mis-nested expansion would have gone through with the clause on it. The new parametrized test pins that structure for a 2-D and a 3-D map: one single-dimension scope per dimension in the original order and range, entries nesting and exits re-joining in mirror order, each scope handing the next its slice, one vectorized loop in the nest, and the same numbers as numpy. The whole-SDFG propagation guard says in its name that it guards cost, and swaps a hand-rolled global swap for the monkeypatch fixture.
ThrudPrimrose
added a commit
that referenced
this pull request
Aug 18, 2026
The merge of main into extended took upstream's version of several files whole, dropping extended-side definitions their own callers still use, and left two pragma emitters both firing. Every hunk here is one of those, resolved toward upstream's implementation: * ProgramVisitor.defined went back to a dict-building property while DefinedNames -- the view every caller reaches through .materialize() -- stayed. Restores the view, and finishes the PR #2501 port inside it (SDFG arrays already carry process grids). * cpu.py defined _use_aligned_operator_new but calls it, and experimental_cpu imports it, under the public name. * codegen.py calls RegionBoundaryStates without importing it. * SDFG grew a second UsedNames class while find_new_name_avoiding_connectors, which sdfg_nesting calls, disappeared. Keeps upstream's _UsedNames, teaches it the connector walk, restores the method. * Autodiff lost get_symbol_upper_bound_from_loop, init_grad and a typing import, and gained a duplicate connector_symbol and TYPE_CHECKING. connector_dict goes rather than comes back: no caller, and the call sites already spell sorted() over an OrderedSet. * MarkSIMDMaps (PR #2503) and this target's own simd_sequential_maps both wrote "#pragma omp simd" before the same loop, so the second one preceded no loop at all and every map_expansion case failed to compile ("loop nest expected"). One emission site now, fed by either decision. Two GPU-specialization defects the spmv sample hits under the transformation tester, both edges wired across a scope boundary: state_fission asserted a boundary node has NO in-edges when an empty ORDERING edge between two boundary nodes legitimately stays, and stream propagation wired gpu_streams straight into a NestedSDFG sitting inside a host map, leaving scope_dict with an undequeued MapExit. samples/simple/spmv.py goes 27/29 -> 29/29. The SMT oracle (z3, optional -- absent it returns None and every caller keeps its existing refusal) answers the two questions LoopToMap cannot: is a non-affine write injective over the iteration domain, and does a read/write pair carry a RAW. Only a proven 'none' admits a pair: BreakAntiDependence runs earlier and has no solver, so a WAR still standing here was never snapshot-renamed. tests/canonicalize/smt_required_parallel_test.py's quadratic scatter XPASSes and becomes a plain assertion; the other two cases stay xfail. CI: the disabled whole-suite rustworkx job becomes a graphlib-only one on 3.14 with simplify=1, so General Tests keeps the networkx backend and nothing runs twice.
ThrudPrimrose
added a commit
that referenced
this pull request
Aug 18, 2026
The Code Quality lane was red on four files ``pre-commit run --all-files`` reformats; running it over the tree is the fix (two of them also lose an import ruff proves unused). The same merge that dropped those had dropped the block emitting ``reduction(op:target)`` on a CPU_Multicore ``parallel for`` -- the clause targets are what ``_omp_reduction_scope_stack`` feeds to ``write_and_resolve_expr``, so with nothing pushing to it every reduction fell back to ``wcr_fixed::reduce_atomic``. Restored, and ``normalize_wcr_test``'s indirect-read reduction goes back to a clause instead of an atomic. The simd emission then had two owners. ``MarkSIMDMaps`` (PR #2503) decides whether a map vectorizes and records it in ``Map.omp_simd``; this target carried its own leaf-body / WCR-coverage rules behind two config knobs and stamped the clause independently, so the two disagreed on min/max WCR maps and on tasklets carrying their own OpenMP directive -- and on any leaf map they both accepted, the clause was written twice, leaving the second ``#pragma omp simd`` facing no loop ("loop nest expected", every map_expansion case). The pass owns the analysis now: both CPU targets read ``omp_simd`` and render, ``simd_sequential_maps`` / ``simd_innermost_multicore_maps`` are gone with the rules they gated (``simd_maps`` is the switch), and the tests that pinned the old rules pin the pass's contract instead -- a multidimensional map reaches the clause through expansion rather than ``collapse``, and a Sequential map with any WCR keeps none.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
MarkSIMDMaps marks the innermost Sequential and CPU_Multicore maps whose body holds no further loop or map, splitting a multidimensional map with MapExpansion first since the clause vectorizes the loop it precedes, and withholding it from min/max reductions and Sequential scatters. Code generation only renders what the pass decided.