[Draft, for CI] FaCe - #2448
Draft
ThrudPrimrose wants to merge 522 commits into
Draft
Conversation
Exercises generated kernel_flat_dace binding on the ICON z_w_concorr_me stencil (flat explicit-shape arrays + dim/loop scalars), linked against the SDFG .so vs plain-gfortran reference. ICON-shaped counterpart to cloudsc_flux_bindings_e2e_test.
Pin the self-accumulate / loop-carried prefix-scan family the CLOUDSC Section-8 flux loop exercises: (1) flux(jk+1)=flux(jk)+ (qn-q0)*g with q0 saved EARLY and qn produced LATE across a big intervening LoopRegion, per inlined-callee block; (2) a long 8-operand + self accumulate mixing early-saved and late-produced operands. Both vs f2py/gfortran, both PASS -- documents the bridge's data-dependency ordering is correct for these shapes in isolation, so cloudsc_full flux divergence is full-kernel-context specific (carry+accumulate two-write fusion ruled out: unchanged).
…h legacy and experimental codegen)
…init The transient zero-init pass (gfortran -finit-local-zero parity) only descended one nested-SDFG level; `_all_sdfgs` now recurses to any depth so every transient Array/Scalar AccessNode in deeply nested SDFGs is stamped `setzero`. Robustness improvement (does NOT affect the cloudsc flux divergence -- proven independent of zero-init: zqxn2d/zqx0 are already memset and strengthening this pass leaves the 130-cell PFSQLF mismatch unchanged).
… explicit-copy-memset-nodes
…e-global auto-capture; py13 package imports cloudsc_full and the 5 bisection carve-outs now pass at strict rtol=atol=1e-15 against gfortran on physically-plausible, NaN-free inputs. The old uniform-random get_inputs drove microphysics intermediates (ZINEW) to NaN; Fortran MIN/MAX with a NaN operand is processor-dependent by the standard, so that comparison tested unspecified behaviour, not a bridge defect. get_inputs_physical (verified 0 NaN/Inf over 12 seeds incl. the ice-deposition chain, SDFG and gfortran) replaces it; the dead generator is deleted and the obsolete xfails removed. Bridge / bindings: - IntegerizePowerExponents post-cleanup pass: retype integer-valued float ** exponents to int so codegen uses repeated-multiply ipow (bit-matching gfortran) instead of libm pow. - F90-binding e2e coverage: SDFG-direct AND generated .f90 binding both vs gfortran -- LOGICAL kinds 1/2/4/8 (incl. KIND=2 copy-out), scalar/array logical in/out/inout, integer kinds, mixed-rank struct, full velocity_tendencies (hand iface + auto module-source). - Module-global provenance auto-detection (_QM<mod>E<entity>): extract_vars decode threaded through FrozenSignature -> block_builders; FlattenStructs peelMemberWrappers + per-leaf FlattenEntry for heterogeneous nested structs; scalar LOGICAL intent(out/inout) rank-0->length-1-array bridge fix. - FP flags reduced to the LLVM-flang-portable core -O0 -fno-fast-math -ffp-contract=off (gfortran-only -ffree-line-length-none isolated for the long-line source); bridge transient setzero stamping dropped (full sweep 0 regressions). Hygiene: - py13: bare build_bridge/hlfir_to_sdfg imports -> package-absolute, sys.path hacks removed across bridge + tests; no `import X as _X`. - Audit: cloudsc strict-FP fixture + flag string consolidated into a conftest + CLOUDSC_F90FLAGS constant; dead FrozenArg.module_origin dual-state removed; change-narration comments trimmed. - .gitignore: .dacecache_gw* (pytest-xdist caches) + stray root cmake build output. Full sweep: 784 passed / 1 xfailed / 1 xpassed / 0 failed.
…ackage (mirrors cloudsc_full/)
The full-ICON (velocity_tendencies) e2e set is moved into a dedicated
tests/hlfir/icon_full/ package, structured like tests/hlfir/cloudsc_full/:
colocated Fortran sources + test_*.py modules + __init__.py xdist marker.
velocity_full.f90, velocity_full_caller.f90 -> icon_full/
velocity_full_test.py -> icon_full/test_velocity_full.py
bindings/velocity_full_bindings_e2e_test.py
-> icon_full/test_velocity_full_bindings_e2e.py
bindings/velocity_full_auto_module_e2e_test.py
-> icon_full/test_velocity_full_auto_module_e2e.py
Path/import fixes after the move: the bindings e2e test now resolves
the driver/caller from its own dir (_HERE) instead of _HERE.parent;
the auto-module test points spec_from_file_location at the renamed
sibling. Same e2e workflow as before (build SDFG + emit_bindings +
gfortran reference, ctypes compare). icon_full: 3 passed.
Also: .gitignore now excludes the remaining Fortran/flang/object
build artifacts (*.smod, *.hlfir, *.o alongside the existing *.mod)
so test-run output cannot be accidentally staged; fix a stale test
docstring referencing the old auto-module module name.
…nto _util _compile_so and the flang-portable _FFLAGS list were copy-pasted (logically byte-identical) across the three new bindings e2e tests (logical_io, numeric_io, weird_struct). Consolidate into the shared tests/hlfir/_util.py those files already import: FLANG_PORTABLE_FFLAGS -- the -O0 -fno-fast-math -ffp-contract=off trio gfortran_compile_so(out_so, *sources, mod_dir=, link_so=) _build_binding_lib is intentionally left per-file: it genuinely differs between logical_io and numeric_io, so a shared abstraction would paper over a real divergence. Full sweep: 783 passed / 1 xfailed / 1 xpassed / 0 failed.
…ss + icon_full/_harness cloudsc_full: lower_keys / f2py_argnames / sdfg_call_args were logically identical across all six test_cloudsc_*.py (only local variable-name drift), and the build-SDFG + run-f2py-ref + scalar-ABI + call block was a ~20-line near-duplicate. Consolidate the three helpers + a run_cloudsc(src, name, f2py_ref, sdfg_dir) driver into cloudsc_full/_harness.py. The five standard tests use run_cloudsc; test_cloudsc_top_half keeps its bespoke body (it captures ZSOLQA / ZSOLQB through extra INTENT(OUT) dummies, which run_cloudsc cannot model) but still shares the three helpers -- forcing the driver onto that genuine outlier was reverted after the sweep caught it. icon_full: the flat-arg ABI order (_INIT_ARRAY_ORDER, correctness- critical), _OUTPUT_NAMES and the _allocate input builder were byte-identical in test_velocity_full and test_velocity_full_bindings_e2e. Move them to icon_full/_harness.py so the ABI order has one source of truth (verified byte-identical to the originals before the move); the auto-module test still reaches them via the bindings module's re-exported attributes. Trivial path 1-liners deliberately left per-file. Full sweep: 783 passed / 1 xfailed / 1 xpassed / 0 failed.
Port f2dace-windmill's single-translation-unit model: inline every externally-USE-d module's real source so flang sees one self-contained file. Reimplemented minimally as a text/comment-aware USE-graph walk (no fparser dependency) -- pass-through for self-contained input, so wiring it default-on in compile_to_hlfir is a verified no-op across the whole inline-source suite; only a real multi-file project under the scratch dir activates it. Also mkdir(out_dir) in compile_to_hlfir for symmetry with f2py_compile. module_merge_test: transitive physmod <- drivermod <- drv project, modelled on recursive_ast_improver_test. Asserts all-files-together f2py == merged-single-file f2py (bit-exact) and SDFG == reference, plus pass-through / idempotence / intrinsic-USE / dependency-order guards. Sweep 786P/1xF/1xP/0F (783 baseline + 3 new, 0 regressions).
Audit-driven, behaviour-preserving (full sweep 786P/1xF/1xP, baseline-identical): * delete dead freeze_signature (zero callers; built an incomplete object that contradicted the real freeze path) * delete dead _view_aliases/_view_shape_strs writes; correct the build() docstring that described copy-in/out staging which does not exist (view linking is lazy in acc()) * centralise the 3-site section-alias view_dim_map splice decode into iter_view_dim_map (intentional 0/1-based output divergence documented, not merged) * lift duplicated _fortran_strides to one module-level helper * lift _raw_hazard to a documented module-level _sibling_rw_hazard; cross-reference it from the realised-graph WAR/WAW guard * replace eval(reduce_identity) with _parse_reduce_identity (named-constant -> int -> float -> loud NotImplementedError); eval-free, covers every identity both producers emit plus any numeric literal, preserves the integer iinfo max/min path * drop redundant deferred acc import and double SimpleNamespace import; refresh stale docstrings
…ording, fix xdist collection race * derived_type_test: rename test_batched_csr_allocatable_xfail -> _jagged (it had no xfail marker and no assertion); add the f2py oracle compare every sibling test uses. Bridge is confirmed correct on the allocatable-array-of-records path (passes rtol=1e-12); it was untested, not buggy. Stale cross-ref fixed. * correct 13 comment sites that claimed 'bit-exact'/'bit-identical' where the check is assert_allclose / np.isclose (tolerance-based, never exact). Genuinely-exact checks (assert_array_equal, bool, algebraic/representational identity) left as-is. * remove the two redundant module-level sys.path.insert calls in bindings/multi_file_test.py and flatten_plan_roundtrip_test.py (conftest.py already puts tests/hlfir on the path); they ran at parallel-collection time and caused the xdist 'Different tests collected between gw0 and gwN' error under -n 8. Full sweep 786P/1xF/1xP/0F, baseline-identical.
… FaCe # Conflicts: # dace/libraries/standard/environments/cpu.py # dace/libraries/standard/helper.py # dace/libraries/standard/nodes/__init__.py # dace/libraries/standard/nodes/copy_node.py # dace/libraries/standard/nodes/memset_node.py # dace/sdfg/construction_utils.py # dace/transformation/passes/insert_explicit_copies.py # tests/library/copy_node_test.py # tests/library/memset_node_test.py # tests/passes/insert_explicit_copies_test.py
…onstruction_utils Use SDFG.parent (O(1)) instead of the recursive _get_parent_state scan to find the state containing a nested-SDFG node; verified equivalent. copy_node.py imports the helper from dace.transformation.helpers (top-level, no import cycle). dace/sdfg/construction_utils.py removed (it only held these two helpers on this branch).
… FaCe # Conflicts: # dace/transformation/helpers.py
…ent_map_and_loop_scopes to helpers
…nvertScalarsToLengthOneArrays; add unit tests
# Conflicts: # dace/libraries/standard/environments/__init__.py
The exact-leaf assertion assumed cache mode 'name'. A workflow whose DACE_cache resolves to anything else (env or a persisted config value the unlaunched fixture does not clear) flipped the leaf to a hash suffix and broke the path match. Pin it explicitly like the sibling tests pin their env, same env-wins-over-config precedence used to root-cause the distaware default flip.
Under srun/mpirun every build subprocess otherwise inherits (a) a blocked SIGCHLD mask, hanging cmake configure in select() forever, and (b) the launcher's PMI/PMIx rank identity, hanging any PMI-linked child in init. Both fixes ported from the extended branch at the single fork point so a new caller cannot reintroduce the hang. Required for warming and measuring dace-fortran kernels inside srun steps on daint.alps. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
All workers/ranks see every GPU and pile CUDA contexts onto device 0, which flakes as invalid device ordinal (101) under -n 32 on cscs CI.
# Conflicts: # dace/sdfg/utils.py
…t host copy-outs Steps 4 (map/library-node outputs) and 6 (top-level transients) flip containers to GPU_Global without recording them in data_already_on_gpu, so Step 8's interstate-edge machinery never saw them: a host-code interstate edge reading such a container (e.g. a loop condition indexing a GPU-written transient bool array) read device memory directly and failed validation with InvalidSDFGInterstateEdgeError during simplify. Both flip sites now route through a _move_to_gpu helper that flips storage once and registers the name, letting Step 8 clone a host copy and insert the copy-out state before each reading edge. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
… SDFG Nesting the map body replicates the LoopRegion inside a new nested SDFG, making its loop-bound symbols free there. The missing-symbol loop added them to symbol_mapping but never to the nested SDFG's symbols dict; validation only checks mapping coverage so this stayed green, and codegen first failed in arglist() with a KeyError when assembling GPU kernel arguments. Reuse the parent's symbol dtype when registering. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…gistration Two ways apply() broke the validator invariant (nested free symbols minus connectors must be covered by symbol_mapping): loop-assigned and shape symbols were registered into the nested SDFG's symbols after add_nested_sdfg froze the mapping, and a later application inside an existing nested SDFG rerouted the loop exit edge onto the new map state while tail bookkeeping only removed the symbol - leaving uses whose parent node mapping never gained the symbol. Both surfaced as "Missing symbols on nested SDFG" under specific transformation application orders only. Reconcile in both places: after symbol-type propagation, self-map any nested free symbol missing from the node mapping (mirroring dtype registration, sorted for determinism); after tail symbol removal, identity-map newly-free symbols into the parent node's mapping when one exists. Both blocks are no-ops on healthy graphs. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The rotation added here never ran in the GPU CI. That pipeline launches a single pytest process under slurm, which sets SLURM_PROCID=0, so the xdist CONTROLLER resolved a launcher rank and pinned itself to one device. Workers are spawned with the controller's environment, so all 32 inherited a one-device CUDA_VISIBLE_DEVICES, took the `len(device_pool) <= 1` early return, and piled onto GPU 0 -- precisely the storm this was written to remove. Measured on a 4-GPU GH200 node: 33 of 33 processes (controller plus 32 workers) on device 0. The controller runs no tests of its own, so it must never narrow the pool, and that has to win over any rank variable it happens to see. It now republishes the whole pool instead, which leaves the workers something to rotate over and saves each of them an nvidia-smi. The same node now reports 8 workers on each of the 4 GPUs. Paths without an xdist controller are unchanged: a plain session and the `mpirun -n 2 pytest --with-mpi` GPU workflow still pin by LAUNCHER_RANK_VARS rank. Also name the CUDA architecture for the GH200 runners. Left on `native`, cmake reports "No local CUDA-capable GPUs found" whenever the GPUs are busy and falls back to compute_60, which newer CUDA toolkits reject outright; that fallback, not the tests, produced most of the failures observed while the workers were stacked on one device. Plain 90 rather than 90a, whose SASS has no JIT fallback. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…match) The loop compared sympy symbols against the str-keyed sdfg.symbols dict, so it never executed; nested SDFGs relied on the add_nested_sdfg backfill, which registers enclosing-map parameters with a default int32 dtype instead of their defined type (e.g. int64), truncating in generated code. Restores the string-keyed logic lost in merge 9ac67d6 (original fix 03c71c3) and additionally includes the top scope entry's own parameters via new_symbols. WIP: regression test passes solo but fails under pytest -n 16 alongside other test files (suspected cross-test symbol dtype leakage), not yet resolved. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
ExpandPure declared the three input connectors of the inner SDFG with default packed-C strides while the output connector carried the parent descriptor's own strides. A connector is a view onto the caller's buffer, so an operand that is not packed C -- a Fortran-layout array, a sub-range of a larger array -- was read at the wrong addresses: the select returned other elements' values, and positions past the operand's real extent came back as zeros or uninitialized bytes. Give an array input the operand's strides, exactly as the output already does. A single-element (broadcast) input stays a length-1 view, where strides do not apply. Covered by two tests that fail before the fix (46/48 elements wrong): Fortran-strided operands, and the same with a broadcast scalar. Both assert the connector strides structurally, not just the numerics. Ported from dace extended e8d3f1a.
Frontends need somewhere to park JSON-serializable state that survives a save/load round-trip without dace interpreting it. dace-fortran uses it to carry a frozen argument signature through .sdfgz, so a later pipeline stage can check that the SDFG's arglist still matches what the Fortran binding was generated against. dace never reads the contents, so both directions are a plain deep copy -- enough to keep the stored dict and the emitted JSON from aliasing. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The read-set walk recursed only into LoopRegion / ConditionalBlock, so it
stopped at a ConditionalBlock's branch -- a branch is a plain
ControlFlowRegion, not either of those. Any data container read by a header
below that point never entered read_set, so it never became an input
connector on the loop-body NestedSDFG and stayed a free symbol of the body
instead. The parent's interstate edge was then left referencing a name that
exists only inside the nest.
Surfaced by enabling preserve_abi length-1 scalarization in the dace-fortran
post-cleanup: staging rewrites `lextra_diffu[0]` to the transient scalar
`scal_lextra_diffu`, which lands inside the nested SDFG, and the velocity
CPU build died in apply_transformations_repeated(LoopToMap) with
`Undefined symbols in edge: {'scal_lextra_diffu'}` at `block -> if_15`.
Walk all_control_flow_blocks() instead. Ported from the `extended` branch,
which carries this fix and consequently deletes two symbol self-mapping
workarounds in apply(); those depend on extended-only APIs
(symbolic.equalize_symbols, LoopRegion.pinned_sequential) and are left alone
here.
Verified: 34/34 loop_to_map tests; both velocity variants (loopexch,
noloopexch) complete phase A at 56 maps, having previously crashed ~2 min in.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Staging KEEPS the signature array by design -- that is what preserves the ABI -- so the array stays eligible on every subsequent run. Each re-application therefore created another transient scalar and another copy-in/copy-out pair, chaining a redundant hop onto the previous one. Measured on a velocity baseline: scalars 19 -> 41 -> 63 -> 85 and states 102 -> 104 -> 106 -> 108 over three runs. Every intermediate stayed valid with a byte-identical arglist, so this was silent growth rather than a failure. Skip a non-transient whose only references are a previous run's staging states, and skip one that is neither read nor written (staging it produced a dead scalar). Same baseline now holds flat at 20 scalars / 104 states across three runs. This inverts the contract of test_repeated_forward_finds_new_name_and_stays_correct, which asserted a second application MUST produce a fresh scalar. That assert tested the find_new_name mechanism; the test's stated intent is "must not collide", which idempotency satisfies more strongly. Renamed to test_repeated_forward_is_idempotent_and_stays_correct and now asserts the descriptor set and state count are unchanged, keeping the numerical check. Verified: 26/26 pass. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The header already aliased the types (gpuStream_t, gpuEvent_t, gpuError_t) and gpuGetLastError / gpuGetErrorString for both the HIP and CUDA arms, but shipped no function aliases -- DaCe's own codegen emits '%sStreamSynchronize' % self.backend and so never needed them. Backends that emit the backend-agnostic spelling directly (the f2dace fork's velocity passes) failed to compile with `identifier "gpuStreamSynchronize" is undefined`. Add gpuStreamSynchronize / gpuDeviceSynchronize / gpuEventSynchronize to both arms. Deliberately NOT adding gpuMalloc / gpuFree / gpuMemcpyAsync / gpuMemcpyD2H: consumers such as the velocity reductions kernel self-define those, and duplicating them here would collide. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…t access The 'i:i+stride' special case in AffineSMemlet.propagate returns the map range verbatim. It is meant for a range access whose length equals the map stride, but a single-element access has length one, which equals the stride of any unit-stride map range, and a*i at a zero map begin starts where the map range starts. Both of its tests therefore hold for an access it was never meant to cover, and returning the map range drops the multiplier: C[2*i] over i=0:N propagated to C[0:N] instead of C[0:2*N-2:2]. That under-approximates the write set, so it is unsound rather than merely imprecise. A non-zero offset takes the general path and was already correct, which is why this went unnoticed. Require a unit multiplier for the special case.
Innermost Sequential map gets #pragma omp simd on its innermost loop; innermost CPU_Multicore map folds simd into its parallel-for pragma. "Innermost" means the map body holds no inner Map and no NestedSDFG, checked non-recursively. Both gated by new default-on config keys (compiler.cpu.simd_sequential_maps / simd_innermost_multicore_maps) that serve only as an off-switch. min/max WCR and (Sequential-only) scatter WCR targets are excluded as unsafe. Adds tests/codegen/simd_pragma_test.py covering both rules, the inner-map/NestedSDFG disqualifiers, and the config off-switches.
The simd config keys were re-applied from a backup taken before the
merge that added cache_distaware, so committing that file removed the
key. Config.get_bool('cache_distaware') then raised KeyError on every
build_folder_root call.
…e fusion apply() re-verified can_be_applied in strict mode regardless of how the match was made, turning permissive-only matches into silent no-ops that apply_transformations_repeated kept re-matching forever. Record the mode can_be_applied ran in and replay it. Also route start-block updates through keep_start_block, which skips a redundant pin when the target is already the sole source, matching state_fusion.py's helper.
The innermost CPU_Multicore rule stamps simd onto the parallel-for directive, so the exact-string goldens now carry it. The invariant the test checks -- how schedule and num_threads compose on the directive -- is unchanged.
Port keep_start_block from extended: skip re-pinning start_block when the target is already the sole implicit source, instead of always writing an explicit (redundant) pin. Behavior-neutral for the getter (a unique source always wins over the explicit pin) and de-duplicates the three was_start call sites in apply(). Kept on FaCe's side: the networkx import (no dace.graphlib here, and none is planned) and the current is_start_block, which already carries a bounds check on a stale _start_block index that extended's simplified rewrite of the same fix dropped.
sympy_numeric_fix and two other call sites use math.isfinite, but the module never
imported math. Any SDFG whose interstate edge assignments reach symstr() ->
sympy_numeric_fix therefore failed to deserialize with
NameError: name 'math' is not defined
which serialize.from_json swallows into a SerializableObject, so the real failure
surfaced far away as
TypeError: Expected InterstateEdge, got <class 'dace.serialize.SerializableObject'>
Reproduced by loading the stored velocity_tendencies stage 3/4/5 SDFGs; they load
and validate again with this import, under DACE_testing_deserialize_exception=1.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
3 participants
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.
Draft PR to run CI on the FaCe branch. Base is
ci-fixfor now; retarget tomainlater.Not for review yet.