Repository navigation
[Only for CI] Extensions - #2475
Draft
ThrudPrimrose wants to merge 7832 commits into
Draft
ThrudPrimrose wants to merge 7832 commits into
ThrudPrimrose wants to merge 7832 commits into
Conversation
ThrudPrimrose
added
no-ci
ThrudPrimrose
force-pushed
the
extended
branch
from
August 12, 2026 11:16
269d9cf to
f9c1368
Compare
A strided operand of np.outer is flattened by a copy map into a transient, and the product map read that transient from a second, unconnected access node in the same state. Nothing ordered the read after the copy: the frontend's emission order happened to, and RemoveViews reordered it, so the npbench LU spelling A[k+1:, k+1:] -= np.outer(A[k+1:, k], A[k, k+1:]) read the previous column for every intermediate graph from the first cleanup stage on. The product now goes into a step state, the pattern the other multi-map lowerings use; simplification fuses the states once the access nodes are linked. The augmented LU spelling now lifts to GetrfNoPivot. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
at_most_one_trip read only the count's numeric bounds, which a registry symbol never gives (no SymPy sign), so a map whose extent the SDFG's facts bound by one stayed a parallel region. It now falls back to asking count <= 1 under the facts at the state. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… the config, and prove a relation whose extremum sat in a division ChunkAntiDependence matched a constant read-ahead only. A symbolic offset that is nonnegative under the canonical symbol assumption now gets a seam row of K cells per chunk, the trailing row copied by a map so K = 0 runs it zero times. Its snapshot copy shares a state with BreakAntiDependence's K >= 0 trap, so the matcher also finds the copy in the one state before the map. Bit-exact for K = 0..300 including N <= K. A D >= 2 chunk count is int_ceil over a Max; prove() generalized the division away and then indexed the Max it had just removed (IndexError in the fork/join threshold test). It now asks the generalized relation directly. canonicalize's CPU and GPU presets pinned peel_limit to 4 while the config default is 8; the presets no longer carry it, so every entry point reads optimizer.canonicalization.peel_limit, like the unroll limits.
Fold the Python spelling of each tile op into the existing op tables as a template field, rename an in-out read connector in place instead of in a postprocessing hook, split Emitter.scope below the complexity limit, and route the three tile-kernel expansion filters through one helper. Add the design document of the Python, tile DSL and CUDA Tile C++ backends.
The binary-operator result type classed every float outside float16/32/64 as an unsigned integer, so a bfloat16 array divided by an integer size was a "true division of integers" and came out float64, and ml_dtypes' floats (numpy kind V, not inexact) kept an integer symbol a strong int64. HB's bf16 kernels then fed a double block into a bf16 tensordot (conv_transpose3d_avg_pool_clamp_softmax_multiply). dtypes.FLOAT_TYPES is the float set now, for both rules. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
``b:b:s`` held ceiling(1/s) elements unless s was 1 and b a plain sum of symbols, so a strided window's per-iteration element (``x[..., i*p + t :: p]``, ``max(0, lo) * s + k``) read as a possibly-larger range and its tasklet connector became a pointer the body used as a value (HB conv_transpose3d_avg_pool_clamp_softmax_multiply through CPF hip, at any precision). ``b + c - b`` cancels to ``c`` for any finite expression without floats; that is the condition now. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The loop expansion passed the operands' numpy scalar types to add_array, which maps only numpy's own: a bf16 GEMM raised ``Invalid type "type" for property dtype`` (HB gemm_add_relu at bf16). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
HB grades its ML kernels in bf16, and CPF refused the type outright. Device dialects spell dace::bfloat16 as the toolkit's __hip_bfloat16 / __nv_bfloat16 -- what dace::bfloat16 IS in a device build -- and include hip_bf16.h / cuda_bf16.h. An integer reaches it through cpf_narrow<T> (``__hip_bfloat16(n)`` is ambiguous between its converting constructors), and cpf_gpu_atomic swaps the 32-bit word holding a 2-byte accumulator. The C++ dialect spells bfloat16 / float16 as the compiler's __bf16 / _Float16 (GCC 13+, Clang 18+: arithmetic, conversions and OpenMP reductions). Its math runs in float through cpf_lowp_* wrappers, emitted only for a unit that holds a 16-bit float, because std::exp(__bf16) is ambiguous (Clang in every standard, GCC before C++23). The C dialect, whose helpers are typed by name, refuses them up front with the way out. All four types are two bytes in the bit layout of HB's __npb_bf16 (uint16_t) storage, so an entry taking them takes the caller's storage pointers as they are. A vendor library is not offered a 16-bit float operand: the BLAS / FFT expansions take single and double precision. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
# Conflicts: # dace/subsets.py
Define py_mod, c_mod and CMod in the Python prelude, so an interstate condition that calls them runs. Refuse a C++ tasklet inside a CUDA Tile C++ kernel with a CodegenError instead of a Python SyntaxError. Type a one-element GPU operand of a CuPy expansion as a device pointer, as the cuBLAS expansions do, so a GPU scalar result validates. Add CuPy expansions of TensorTranspose and ScatterConflictCheck.
…uch one element of an argument array AnnotateLoopKinds called a loop sequential only on a constant-distance access pair or a carried symbol, so pivot rows and triangular solves (lu, cholesky, trisolv, ludcmp, gramschmidt, durbin) stayed unsure. It now also asks Z3 for two iterations of the loop, each side with its own copy of every inner iterator and its bounds, at which a write and another access of the same argument array touch one element. The witness may choose only array sizes: any other symbol is a fixed unknown (a[i] = a[i - k] stays unsure, since k = 0 carries nothing). Only accesses every iteration makes count: no branch bodies, no loops with an early exit or without a constant-stride counter, no transients. An array copy contributes its whole ranges, since it moves every element; tasklet and library ranges do not. The hint names the access and the witness values. test_writes_the_test_cannot_separate_are_unsure_not_a_proof asserted a[5i] / a[3i] stays unsure; i = 3 and i = 5 write the same element once N >= 6, so it is now test_writes_the_test_cannot_separate_are_proven_by_a_witness and checks the witness.
…4-bit strides CopyLibraryNode gets aligned_words: rows contiguous on both sides move through dace::BlockCollective3DVec in aligned 16-byte words, peeling each row's unaligned head and tail. A region of rank above 3 is a loop nest of rank-3 collectives with one barrier after it (was a ValueError). dace::BlockCollective3D/3DVec take 64-bit strides: a tile of a 144-level x 16M-column field spans more than 2**31 elements. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Code generation reorganizes states and host nodes, so their events timed nothing reliable. The provider now times GPU_Device maps only, recording on the stream the kernel is launched on; validation rejects GPU_Events on SDFGs, states and other nodes; dace.instrument annotates kernels only; the tuners instrument a state's kernels and measure() sums them. AddThreadBlockMap keeps the timing on the kernel map. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
StageThreadBlockData stages the bounded regions a kernel's thread block reads (and, on request, writes) through shared memory with block-collective copies and one barrier per direction. CoarsenKernelThreads gives every thread a serial tile of one or two map parameters. DefaultSharedMemorySync adds no barrier where one already follows; block_tile_kernels does not stride SMALL-only axes across the lanes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
# Conflicts: # dace/symbolic_facts.py
A `tile` marker on the cuTile/TileLang/Triton and CUDA Tile C++ test modules, deselected in the GPU CI; the CPU jobs already deselect them as `gpu`. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… passes device buffers itself The stage applied CloudSC's hand-written mirror, which leaves an array written by both a host loop the pipeline cannot parallelize (nssopt left symbolic) and a kernel unmirrored, so kernels read host memory (illegal access in the parallelize leg). The pipeline's own offload places copies both ways. run_candidate reads the device-resident arguments off the descriptors (moved from the target test).
…() lookups LoopToMap declares lift_context, symbol_resolver and contract_arrays (unset on a plain match) instead of reading them with vars(self).get; ParallelizeLoopsHoldingSiblingChains copies them directly. BreakAntiDependence declares its per-loop memos in __init__. LoopStridePermutation's revert restores the loop metadata tuple its snapshot now holds.
…r-element scalars instead of refusing A fused map's per-iteration buffer (CloudSC's __map_fusion_zpfplsx[2, 1, 1]: tasklets write and read single elements, one copy moves a range out) is lane-dependent with more than one element, and the GPU vectorizer aborted the whole pipeline on it. When step 3 meets such a buffer it is split into one scalar per element, its range copies into one element copy each, and the propagation reruns; each element is then a per-lane temporary. Only a buffer the widening would refuse is split. The pipeline offload stage test reads the device-resident arguments the real offload leaves.
# Conflicts: # dace/dtypes.py # tests/codegen/gpu_schedule_groups_test.py
…e SDFG declares A symbol is its name, so the parse and SymPy caches can hand back an object of the same name another program declared at another dtype: after a program that (correctly) refused a float scalar promoted to __sym_n_0 as an extent, an integer program's __sym_n_0 extent was refused too (an xdist-order failure of size_scalar_shape_promotion_test). The shape is now validated inside the SDFG's symbol-dtype authority, as symbol_dtype is designed to be read.
… candidate array LoopToMap records the container a dependence refusal names (last_refusal_container); a refusal on any other container, or for structure, bounds or symbols, survives the privatization, so the pass no longer promotes, re-probes and undoes those loops.
… once per map split MapExpansion propagates its own scope; the whole-SDFG propagation apply_to ran after every split now runs once at the end. Same generated code on the corpus; pb_correlation's inner-scope volume becomes the exact j instead of a stale N.
…d of on copies The split search isolated the loop into a mini SDFG (JSON round trip), deep-copied it per candidate, split the copy, ran the LiftTrivialIf/EmptyStateElimination/modulo cleanup and PrivatizeScalars/SymbolPropagation/ConstantPropagation on it, and only then probed LoopToMap. On the CloudSC monolith that cost 18.7 s of thread time and split nothing. The search now analyzes first and transforms only the winner. Each candidate's segments come from _split_plan, the header computation _split_loop_at now builds from, so the scored segments are the ones the chosen split emits. Each non-singleton segment is probed on the real loop: its init statement and condition are swapped to the segment's, the guards the range-aware LiftTrivialIf would resolve over that range are resolved in the branch lists, the modulo subsets the range makes affine are folded, LoopToMap.can_be_applied_to decides, and every edit is put back in a finally block (original CodeBlocks, branch entries, subset ranges). Selection is unchanged: the unsplit loop maps nothing, the first candidate to reach a count keeps it, and the bare-bound (guarded) form wins only by mapping strictly more. LiftTrivialIf exposes its per-conditional decision as dead_branch/taken_branch, used by the pass itself and by the virtual resolution, so the decision exists once. _mappable_loop_count, _loop_to_map_prober and mappable_count_upper_bound had no other user and are removed. The modulo peel/split stages keep their copy-based probes (_isolate_loop). Corpus (tsvc, tsvc25, poly; --peel 4): all 257 kernels identical. canonicalize hashes on the 27 canon-speed2 inputs identical. Replay of the peel unit: mono 18.7 s -> 1.7 s, cloudsc 8.5 s -> 1.2 s, same output.
…nnector's own access node
…uceExpansion does not validate, codegen resolves scope symbols and nested-SDFG arguments once
…tself The modulo stages isolated the loop into a mini SDFG and, per (count, direction) peel or per split point, deep-copied it, peeled or split the copy, ran the cleanup and asked whether a wrapping modulo survived. They now decide by range analysis on the real loop and transform only the winner with the existing _peel_one_loop / _specialize_modulo_split. A peel's segments come from LoopPeeling.peel_plan, the header and iterate values LoopPeeling.apply now builds from: the remaining loop's header and, when the body holds a loop, a single-iteration header per peeled iteration (only loops can keep a wrapping modulo, and a peeled iteration is no loop). A split's segments come from _split_plan. Each segment is probed under _as_segment, which swaps the loop header, resolves the guards and folds the modulo subsets the range decides (the cleanup the real rewrite runs), and restores everything when the block exits; the candidate is accepted when no loop of any segment keeps a wrapping modulo, and the folds' offset < modulus relations are the split's branch condition as before. Selection order is unchanged: smallest count first, front/back/both, then the first split point that folds every wrap. _peel_one_loop's gates move to _peel_sides, shared with the search. _isolate_loop has no user left and is removed with its test.
…ument symbols through the frame's free_symbols
…r handled wherever sympy expressions are inspected
# Conflicts: # dace/transformation/helpers.py # tests/nest_subgraph_test.py # tests/sdfg/state_test.py
…o-extended # Conflicts: # dace/codegen/targets/cpu.py # dace/codegen/targets/framecode.py # dace/sdfg/state.py
…stop repeating the nest per region, the descriptor per access node and the parse per assignment; nested-SDFG argument symbols through the frame's free_symbols
# Conflicts: # dace/codegen/targets/framecode.py # dace/sdfg/sdfg.py # dace/sdfg/state.py # dace/transformation/passes/prune_symbols.py
… FtnModulo as frontend builtins
…read before assignment Extended already carries #2706 as a port (41640a3, f1747c3) and its own free_symbols walk, so each conflict keeps extended's side; from main it takes the program-callable modulo spellings (PyMod, FtnMod, FtnModulo through operators.MODULO_CALLS and binop_tasklet_code) and main's tests. Two of main's modulo tests are adapted to extended's documented semantics: a CMod of operands not proven integer prints as c_mod, and sign facts of a dace.symbol live in the symbol registry, so the CMod->Mod fold is checked on SymPy symbols with assumptions. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ed dace uv counts the pin's equal version as satisfied, so --no-deps -e alone left the pinned commit installed. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…v sources From the bench directory uv resolved the reinstall through the bench's dace pin again. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The bench's [tool.uv] override-dependencies re-pin dace for any install run from the bench's directory. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…elper warnings A one-element GPU-resident container crosses an ExternCall as its device pointer instead of the address of a host copy, and a GPU-scheduled call waits for the SDFG's stream before it and the device after it, so no device-wide fence (and warning) is added. Standalone CPF CUDA/HIP units mark their optional helpers [[maybe_unused]] / inline constexpr, so nvcc reports no unreferenced function.
…d SymPy's Mod prints as PyMod float(i * 7 % N) handed the cast the call's one-element result list, so its tasklet kept a dangling __inp (tsvc s315 failed validation). A call argument now unwraps a single symbolic or numeric result; a list literal stays a list. astutils.unparse printed SymPy's floored Mod as %, which an SDFG reads as C's truncating modulo; it prints PyMod. The Python backend knows every modulo spelling. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…GEMM shares out its batch examinimd's np.sum(w[:, :, None] * d, axis=1) lifts to an Einsum whose BatchedMatMul ran 10976 3x1xK cblas_dgemm calls in a row: 452 ms, 82 ms now. - ExpandTransformation: the top-level maps and library nodes of an expansion, through top-level nested SDFGs, take the expanded node's schedule; inferred as nested, the Einsum's BatchedMatMul was Sequential. - BatchedMatMul OpenBLAS: a top-level node over a batch that is not small runs the batch loop under omp parallel for, one single-core GEMM per element. - libnode_is_sequential follows the map rules: sequential under a parallel parent map, over a small domain, or over a data segment under a sequential loop; a sequential loop alone no longer serializes it. - symbolic.is_small_extent: a SMALL symbol or a literal below 64, for schedules only (the canonical form keeps literals parallel). SequentializeSmallAxes and expansions isolate literal small axes without a declared SMALL symbol. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This branch has not been deployed
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.
No description provided.