[Only for CI] Extensions - #2475
Draft
ThrudPrimrose wants to merge 4087 commits into
Draft
Conversation
ThrudPrimrose
added
no-ci
'ik,ik->i' and 'xyzk,xyzk->xyz' share every index between the two operands:
the batch indices reach the output, the rest is contracted, and neither
operand has a private index. That class is C[batch] = sum_k A[batch,k] *
B[batch,k] -- no BLAS-2/3 form expresses it. is_bmm() called it a batched
matmul anyway, so the contraction path minted a degenerate M=N=1 MatMul whose
operand views collapse under simplify into two equal 2-D shapes the MatMul
dispatch rejects (NotImplementedError). Classify it as not-a-BMM, next to the
existing 'ij,i->i' carve-out; the unbatched dot 'i,i->' keeps its BLAS path.
Three pure-path defects the reroute walks into, fixed here as well:
- a repeated operand (a, a) added a read node the dict comprehension then
dropped, leaving an isolated node that fails validation;
- the index letters named the map parameters, so a letter that matches an SDFG
symbol ('k' contracted over a [..., k] array) shadowed it and emitted the
self-referential bound k = 0:k -- zero iterations, silently wrong results;
- the accumulator init state was prepended to the SDFG rather than to the
state's own control-flow region, so an einsum inside a loop body raised
KeyError.
Fixes issues/einsum_rowdot_matmul_dispatch.md.
… gate floor Complex-typed chains>1 failed to compile: the UDRs in dace::scan::detail are invisible to unqualified lookup from the spliced multi-chain pragmas. Emit the declare-reduction at the tasklet's own block scope instead (GCC and clang both accept it with inscan). The corpus gate's constant fp64 atol=1e-11 is smaller than the scalar oracle's own rounding error at the paper preset (2.7e-11 vs float128), so it rejected a candidate 10x MORE accurate than the reference (scan_multi_5carry, canon arms only -- the arms that reproduce the oracle bit-for-bit passed). The absolute term becomes max(constant, 1e-12 * max|ref|), still elementwise; explicit atol overrides are never floored; integer gates stay exact; a non-finite reference keeps the constant (overflowing wavefront kernels would make the floor vacuous). Exact-int64 sweep over K x n x threads x seeded proves the lowering itself correct.
…ed benchmarks
Two recovery gaps: a rung tiling a parent window emitted T3*int_ceil(T2,T3) as
its bound, which the next sweep no longer recognized as T2, stalling 3-level
symbolic cascades half-collapsed (fix records Mod(span,K)==0 under the same
contract as the cascade-stride assumption); and dace.map-tiled nests were never
seen at all (pipeline runs map_roundtrip=False, matcher reads LoopRegions only)
-- now detected in-pass and round-tripped behind a probe-on-deepcopy gate that
declines if any map would come back as a loop (fires on 0/76 corpus kernels).
New corpus kernels jacobi2d_triple_tiled_{const,sym} and
heat3d_double_tiled_{const,sym}; untile/unroll tests now assert static
structure (collapse counts, one perfect chain, no leftover tile index or
remainder guard, generated-C++ loop counts, exact unroll body-copy counts),
not just values.
…join cost model moves to the specialization bands The canonical representation is the maximally parallel one; making a scope sequential again is a target decision. Copy/memset expansions now emit a parallel element map unconditionally -- a symbolic count is assumed big (only a provably-small constant keeps the libc call), and the expansion no longer reasons about its enclosing scope at all. The sequentialization that lived in canonicalize/finalize moves to cpu_specialization/SequentializeParallelScopes, the single home of the CPU fork/join cost model: it pins maps whose work per region cannot pay for a parallel region and scopes re-entered by a parallel map or a long loop (s115/s119 forked 768 teams per call; stockham_fft 349,525). The re-entry and short-loop rules live once, in libraries/standard/helper, consumed by the band. SpecializeCpuTransfers then hands sequentialized contiguous transfers their single memcpy/memset back, and runs again in codegen for the late-born explicit copies. The GPU counterpart moves to gpu_specialization/SequentializeNestedDeviceScopes, so canonical output keeps nested parallelism intact for both targets. Placed at the end of the pipeline: the verdict must read the final map shapes, and early Sequential pins would block map fusion.
…oMap guards NestedCall.add_state chained new states onto the visitor's cfg_target, which can disagree with the region owning last_state inside a loop or branch -- the interstate edge then joined a node the graph does not own. Use last_state.parent_graph. Einsum lifting and LoopToMap carry the matching guards for the shapes the corpus gate exposed.
…c GPU offload tests
SymPy folds assumptions into symbol identity, so one name spelled two ways is two distinct objects: index arithmetic does not cancel and dependence predicates silently answer wrong. Report it at validation and when a symbol is registered, behind an opt-in experimental config flag, default off.
Frontend minted loop iterators with explicit-None assumption kwargs, which sympy treats as distinct from omitted kwargs: two same-named symbols coexist and free_symbols/match silently give wrong answers. Omit None-valued kwargs at minting. Add opt-in validation walk and add_symbol check that raise on same-name different-assumption symbols (flag experimental.check_symbol_assumption_collisions; stays opt-in here until canonicalize assumption stamping moves to a registry).
Fusion removed the entry block and pinned start_block even when the region could derive it from its single source; BlockFusion pinned before remove_node, which clears the cached value. New helper pins only when underivable, and after removal.
LICM Python-parsed C++ tasklet code, got an empty symbol set, and hoisted tasklets out of the map defining their symbols; the resulting degenerate staging shape was then corrupted by whole-tree renames in ScalarFission and replicate_scope. Identifier-scan fallback for non-Python code, value-carrying write counts, and rename trees that stop at a different access node of the same container.
Clean-pattern passes counted scalar reuse only in other states; LoopToScan accepted carrier reads at foreign offsets; wavefront skew window widened with ISL containment check.
Bare a/b/o connectors in the sequential reduce expansion shadowed same-named parent arrays after codegen inlining; dotted structure member names reached copy/memset labels and produced invalid C++ function names.
C++ tasklet bodies lose their free symbols when rebuilt as Python AST, so such maps stay scalar; symbol definition map refuses names with WCR or multi-writer producers instead of resolving them to a stale seed.
Canonicalization stamps nonnegative on integer symbols while the rebuilt loop symbol is unassumed; match then binds the wildcard and reads per-iteration writes as invariant. Normalize same-named symbols to one instance before matching; only ever recognizes more a*i+b subsets.
Scope generation allocates arrays outside the tasklet/nested-SDFG calls that set calling_codegen, so idx helpers flushed there under the host key were re-emitted in the .cu as C++ redefinitions. Test regex updated to the 3-dim BlockReduce form with structural asserts on the register-partial protocol.
Entry canonicalize no longer short-unrolls constant-trip loops the tiler was called to widen; producer-only remainder bodies get ordering edges so topological emission drains; copy writes no longer count as supersedes in stage-global.
Assert parallel reduce entry points and reduction clauses instead of absence markers, pin the cost-model sequential default with the zero-threshold arm, wire the serialize path through explicit copies, and drop stale xfails.
Changed default value of cache_distaware from false to true.
Bare 'out' connector collides with caller array named 'out' after inlining.
SymPy folds assumptions into symbol identity, so one name spelled two ways is two distinct objects: index arithmetic does not cancel and dependence predicates silently answer wrong. The frontend minted loop iterators with bound-derived nonnegative/positive assumptions, and with unestablished ones passed as an explicit None, both of which differ from the plain symbol every reparse yields. Mint the canonical spelling instead; subset covering re-derives nonnegativity in subsets.nng. Validation now rejects a collision unconditionally, walking the SDFG once and deriving assumptions only for distinct objects that share a name.
Ranked-vs-shared assertions assumed distaware defaulted off, so clearing the env override fell through to the new true default and compared a rank-suffixed path against itself. Wrap the old off-path assertions in an explicit distaware=False context and add structural per-rank-root assertions for the new on-by-default behavior, for every cache mode.
insert_scatter_guard hard-required 1-D index arrays, crashing on ICON 2-D connectivity (edge_blk[jb,jc]). Classify a single contiguous varying dim per loop and guard that 1-D window instead; anything else stays un-lifted.
… the score function minted a fresh DaCe symbol for each parameter, so coeff extraction failed when the memlet range used a plain sympy Symbol with the same name. Now the actual symbol instance from the expression is used, and strides are stringified before pystr_to_symbolic.
cudacommon defined the gpu* typedefs and the two error macros but not the synchronization ones, so any generated code that wants a backend-agnostic stream/device/event sync had to name the CUDA symbol directly and would not build under HIP. Context also allocated internal_streams and never freed them.
… a nested tasklet/library connector had the same name as an outer array or symbol, inlining produced a validation error or wrong code because the connector name was not renamed. The pass now renames clashing connectors and rewrites the assignment targets in tasklet code.
…, explanation: exit code freed communicators and grids unconditionally, but a host process (mpi4py) or an earlier program in the same process may have already finalized MPI, which aborts. Guard every teardown call with MPI_Finalized and only call MPI_Finalize when DaCe itself initialized MPI.
…erted a cross-NSDFG hoist that v1 never promised. v1 hoists within the owning SDFG and stops at the NSDFG boundary; the test now pins that behaviour structurally instead of an xfail.
…planation: a fully split element gets a length-1 array descriptor, but _get_corresponding_array dropped every split dimension from the subset and returned an empty Range. Validation rejected the resulting edge (expected 1 dim, got 0) whenever automatic simplification was off and the slice state survived. Keep one dimension so the memlet matches the (1,) descriptor.
…uce library node is lowered to a nested SDFG by expand_library_nodes before LiftEinsum runs, so LiftEinsum cannot match it. Capture the Reduce axes before lowering and assert zero matches plus numeric correctness.
…anation: the parametrize marker only set gpu/mkl labels, so the suite failed on runners without the vendor libraries. Add skipif guards so CI lanes without those backends pass cleanly.
… connector, explanation: the recursion passed the enclosing SDFG while the names it collected are the NestedSDFG node's connectors, which name arrays of the nested SDFG, so the lookup raised KeyError. Out edges also read dst_conn, which belongs to the destination node rather than the NestedSDFG. Simplification hid both by inlining the nested SDFG before the walk. Also drop a leftover debug print.
…y pulled in the HIP runtime, explanation: the header typedefs hipStream_t/hipEvent_t/hipError_t under WITH_HIP but never included hip/hip_runtime.h. The precompiled runtime header reaches it through a forced -include, ahead of the generated file's own includes, so the types were undefined. Include the backend runtime header here instead.
…, explanation: DACE_USE_GPU_ATOMICS was keyed off __HIPCC__, which hip_common.h defines whenever __CUDACC__ is set -- including the host pass. The host pass then called __int_as_float and __longlong_as_double, which are __device__ only under nvcc. Use __HIP_DEVICE_COMPILE__, HIP's __CUDA_ARCH__ equivalent, so the host pass takes the omp critical path as it does under CUDA.
…form, explanation: atomicAggInc/atomicAggDec were gated on ifndef __HIPCC__, and hip_common.h defines __HIPCC__ under nvcc too, so the 64-bit-warp AMD path (__shfl, __ballot) was selected where those intrinsics do not exist. Gate on __CUDACC__, matching warpReduce in reduction.h.
…d architecture, explanation: cmake_options emitted -DEXTRA_HIP_FLAGS, which no CMake code reads, so compiler.cuda.hip_args and hip_arch were silently discarded on any branch defaulting to this backend. The architecture was also hardcoded to the AMD gfx spelling. Emit DACE_HIP_ARCHITECTURES_DEFAULT and CMAKE_HIP_FLAGS, the variables CMakeLists actually consumes, and forward the host args like the CUDA path does.
…enerate_constants built array elements with str(), so a numpy bool printed as True and a complex as (1+2j), neither of which compiles. Route them through sym2cpp like the scalar branch already did, and give sym2cpp the two literal rules symstr cannot carry: a bool becomes true/false, and a complex keeps its own width instead of widening to complex128. cppunparse._Num gains the same rules for the AST path, where numpy bools reach it because _Constant only catches Python True/False by identity.
…GState.to_json emits scope_dict for the viewer and from_json never reads it back, but keyword_remover did not strip it, so a cached scope tree was part of the content hash and of every SDFG cache key. It is also ordered before nodes, so it masked the real first divergence in any digest diff.
…the typedefs sat behind an inner ifdef __CUDACC__ inside the GPU branch, and the thrust include that backs them was CUDA-only, so neither the thrust nor the std fallback reached a HIP build. Every translation unit that includes reduction.h then failed on is_reducible<complex64>. Include rocThrust when it is present (__has_include, as reduction.h already does for cub) and fall back to std::complex otherwise; verified compiling complex device arithmetic on gfx1103 both with and without rocThrust, and unchanged for CUDA.
The special case in unsqueeze_memlet that keeps one dimension when the internal memlet selects a single element and the outer memlet uses all of its size-1 dimensions tested for the literal index 0. A squeezed view read inside a reduce expansion is indexed by the map parameter instead (tmp_max[_o0] over a one-element range), which selects the same element but fell through to the general path and raised NotImplementedError. Test the element count symbolically.
…licate memlets, explanation: keep the fast per-subset edge path when there is no explicit out= argument, and use per-array tasklets with propagation only when out= is supplied. This prevents the simplifier from dropping partial writes to the output array and avoids shared Memlet objects that fail validation.
…nectors inlined vector connectors as scalar array accesses; reject vector and pointer connectors so the classic vector lowering is kept.
…FG did not scale scalar indices for vector element width, so inlining a vector-typed nested SDFG into a scalar outer array wrote overlapping pairs; skip inlining when inner/outer connector dtypes disagree on vector-ness.
…braryNode ignored tile sizes in strided ranges and rejected non-contiguous rank-mismatch copies; split tiled dims in collapse_shape_and_strides and let the 1-D walker handle any rank mismatch using the collapsed strides.
…tion: experimental readable codegen validated the SDFG unconditionally even when compile(validate=False) was requested; thread the validate flag through inline_host_nested_sdfgs and gate the final validate call.
…the pure-reduce expansion only writes the scalar result, so the cache-miss denominator should add one cache line (64 bytes) for it, not two.
…on: the finite OCP fp8 formats had no numeric_limits specializations; add DACE_LP_LIMITS entries for float8_e5m2 (has infinity) and float8_e4m3fn (no infinity) with correct bounds and bit-pattern static_asserts.
…anation: broadening the 1-D walker to all rank mismatches broke existing tests for mixed layouts, padded strides, and strided subsets; only use the permissive walker when tile sizes create the extra dimensions, otherwise keep the original packed/contiguous guards.
The inline_after path migrated the whole state's dataflow into a fresh successor before checking whether InlineMultistateSDFG would apply, to keep external references to the original state valid. When the inline was then refused -- which is every StencilTiling use -- the original state stayed behind empty, the next simplify deleted it, and the reference was lost anyway, minus its dataflow. Gate the migration on inline.can_be_applied and run the widening step against the state itself. MapToForLoop now also reports target_state, the state holding the dataflow that stayed outside the loop, since the caller cannot know whether the migration happened. StencilTiling follows it instead of holding one state across its map loop, and addresses StripMining and MapCollapse with graph.block_id rather than the state id it matched on, which the migration invalidates. Fixes the 8 unroll=True cases in tiling_pool_test and tiling_stencil_test.
…tion: can_be_applied used nested_sdfg.sdfg but named the variable nsdfg; add the missing local binding so the vector dtype check can run.
Innermost Sequential maps get "#pragma omp simd" on the innermost loop; innermost CPU_Multicore maps get the simd clause stamped onto the same "parallel for" directive. "Innermost" is a non-recursive leaf-body check (no inner Map, no NestedSDFG); an uncovered (scatter) WCR withholds the pragma on Sequential but not on CPU_Multicore, which already lowers it through an atomic. Both gated on by compiler.cpu.simd_sequential_maps / simd_innermost_multicore_maps (default true, off-switch only). experimental_readable (ExperimentalCPUCodeGen) needs no separate implementation: it subclasses CPUCodeGen and its _generate_MapEntry override calls super(), so the shared emitter runs unchanged; verified with the same structural tests against both compiler.cpu.implementation settings. Updates tests/openmp_test.py's exact-pragma-text assertions for the arrayop leaf map, which now also qualifies for the default-on simd clause.
_start_block holds a node ID, and removing a node shifts every later ID, so a stale one can index a different block or fall off the end. Bounds-check it before resolving through node(). Matches the version on the FaCe fork, which kept the guard.
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
2 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.
No description provided.