[Draft, for CI] New GPU Codegen Complete PR - #2259
Draft
ThrudPrimrose wants to merge 776 commits into
Draft
ThrudPrimrose wants to merge 776 commits into
ThrudPrimrose wants to merge 776 commits into
Conversation
edopao
reviewed
May 11, 2026
| description: > | ||
| If enabled, inserts default __syncthreads() tasklets during preprocessing | ||
| in ExperimentalCUDACodeGen to ensure shared memory is ready before access. | ||
| This is a simple safeguard for correctness—it may not be complete, but it |
Collaborator
There was a problem hiding this comment.
Please check this hyphen, after "correctness". It is has UTF encoding and causes some issues in ICON4Py-related runs where the default encoding is ASCII.
Collaborator
There was a problem hiding this comment.
@philip-paul-mueller Maybe you can fix it in our fork repo? I mean on your branch (which I am using).
ThrudPrimrose
force-pushed
the
new-gpu-codegen-dev
branch
from
May 13, 2026 14:21
67d26eb to
eff356a
Compare
Cherry-pick of the ``tests/library/copy_node_test.py`` portion of 714dec0 (the PromoteGPUScalarsToArrays parts of that commit don't apply -- file doesn't exist on this branch). ``test_copy_cuda_1d_single_element`` passed raw strings as strides (``strides=["src_stride"]``) which ``Array.validate`` rejects ("Strides must be ... integer values or symbols"). The scaffold now sympifies stride entries -- which causes ``add_datadesc`` to auto-register the stride symbol -- so the test's own ``sdfg.add_symbol("src_stride", dace.int32)`` collided and is removed.
Follow-up to a16c2a7 (scheduling/wiring split). The Monolithic strategy returned a dict view of stream assignments but never wrote ``node.gpu_stream_id`` -- only ``NaiveGPUStreamScheduler`` did. ``GPUStreamWiring._collect_assignments`` reads from the per-node property, so the wiring pass saw zero consumers, allocated a 0-slot ``gpu_streams`` array, and the subsequent ``gpu_streams[0]`` memlet in the sync tasklet was out-of-bounds. ``test_monolithic_jacobi_2d_two_syncs_and_correctness`` and ``test_monolithic_heat_3d_two_syncs_and_correctness`` crashed at validation. Write the property alongside populating the dict view; keep the existing idempotency contract by only setting it when unset.
…eduler After the a16c2a7 scheduling/wiring split, the scheduler only writes ``Node.gpu_stream_id``; the gpu_streams array, connector hookup, and sync wiring are owned by ``GPUStreamWiring``. Five tests in this file construct ``Pipeline([NaiveGPUStreamScheduler()])`` directly without a wiring follow-up, so the assertions on a wired ``gpu_streams`` array and ``__stream`` connectors fired against a non-wired SDFG. Add the wiring pass after the scheduler in each affected test; production callers go through ``GPUStreamPipeline`` which already bundles both.
…ches ``sdfg.compile()`` calls ``get_folder_mode(build_folder, probe=True)`` to detect whether a previous cache exists; the contract is to return ``None`` when no usable cache is found so the caller falls through to regeneration. The old-style-cache sanity loop already honors that for the "no subfolder at all" branch, but raises ``NotADirectoryError`` unconditionally when one of ``[build, map, src, include, sample]`` is missing while a sibling exists. A partial cache (e.g. ``build/`` only, left over by an earlier failed compile or aborted run) therefore crashes every subsequent compile on the same SDFG name -- the actual failure mode reproduced across ``test_gpu_strided_2D_copy``, ``test_gesummv`` and ``test_copy_cuda_1d_single_element`` in cross-suite sweeps. Return ``None`` in that branch under ``probe=True`` so the caller regenerates instead.
New preprocess pass for ``AutoSingleStreamGPUScheduler``. For each state that mixes CPU and GPU work, classify the dataflow and rewrite via ``dace.transformation.helpers.state_fission`` so the resulting CFG holds class-pure states only. The classifier is shared with the strategy (``_classify_node`` / ``_Kind``) so MIXED detection is consistent. Rewrite rules: * multiple independent WCCs in one state, some pure CPU and some pure GPU -- lift all CPU WCCs into a new predecessor state. GPU WCCs stay. * one mixed WCC whose topological band structure is ``[CPU?, GPU, CPU?]`` -- lift the CPU prefix into the new predecessor state; lift the GPU middle too when there is a CPU suffix, so the original state is left holding only the suffix as the trailing CPU state. NEUTRAL nodes (AccessNodes, paired Exit nodes) attach to whichever band the topo order places them next to; state_fission duplicates the boundary AccessNodes. * GPU -> CPU -> GPU interleaving, cycles, or a ``MIXED`` interior node (e.g. a mixed NestedSDFG) -- refuse, leaving the strategy to fall back to ``NaiveGPUStreamScheduler`` for that state. The pass gates on ``is_stream_wiring_applied`` so it is a no-op when ``sdfg.compile()`` re-enters the codegen-preprocess pipeline on an SDFG that already had stream scheduling applied (e.g. the test driver ran the pipeline explicitly and then called ``.compile()``). Tests (``tests/gpu_specialization/split_state_by_gpu_class_test.py``) build the zoo of hand-written SDFGs covering each branch and assert the post-rewrite topological state-kind sequence.
Adds a third :class:`GPUStreamSchedulingStrategy` and switches the
:class:`GPUStreamPipeline` / :class:`GPUCodegenPreprocessPipeline` defaults
to it. Behaviour:
* every top-level dataflow node is classified ``CPU`` / ``GPU`` /
``NEUTRAL`` / ``MIXED`` (``_classify_node``). A ``Tasklet`` inside a
``GPU_Device`` scope is GPU; a free Tasklet is CPU. A ``MapEntry`` with
``GPU_Device`` schedule is GPU; other schedules fold over the scope body.
A ``NestedSDFG`` already inside a GPU kernel is GPU; otherwise recurse
via ``_classify_sdfg``.
* any node classified ``MIXED`` (typically a ``NestedSDFG`` that itself
mixes CPU and GPU work that this strategy can't single-stream) triggers
a one-line ``UserWarning`` and global delegation to
:class:`NaiveGPUStreamScheduler` for both ``assign_streams`` and
``insert_sync_tasklets``.
* otherwise every GPU consumer + every ``GPU_Global`` AccessNode is
stamped with ``gpu_stream_id = 0``. The AccessNode pass is necessary
because pool-backed transients route ``cudaMallocAsync`` /
``cudaFreeAsync`` through the AccessNode's stream
(see ``experimental_cuda.py``'s pool branch); the Naive strategy picked
this up implicitly via WCC membership.
* sync insertion creates dedicated sync states (containing one
``cudaStreamSynchronize`` Tasklet wired to ``gpu_streams[0]``) on:
- every interstate edge whose source is GPU and (a) the target is CPU
or (b) the iedge's condition / assignment reads a GPU-written array
(host-side iedge eval depending on kernel output);
- every region-level sink that is GPU and has no outgoing iedge.
CPU -> GPU iedges and CPU sinks get no sync (host is sequential).
``SplitStateByGPUClass`` is added to ``depends_on()`` so the pipeline
dependency-graph machinery pulls it in immediately before the strategy,
reducing the cases where the MIXED fallback fires.
Both ``assign_streams`` and ``insert_sync_tasklets`` gate on
``is_stream_wiring_applied`` so the pass is a no-op when ``sdfg.compile()``
re-enters the codegen-preprocess pipeline on an already-wired SDFG.
Test-fallout:
* ``gpu_stream_scheduler_registry_test.test_pipeline_default_strategy_is_auto``
pins the default-is-Auto contract.
* ``explicit_gpu_stream_management_test``'s module-level pipeline is now
constructed with ``NaiveGPUStreamScheduler()`` explicitly because those
tests pin Naive-specific behaviour (per-WCC streams, end-of-state fused
sync tasklets).
Six unused imports flagged by ruff on the just-pushed Auto + Split work; no behaviour change.
The tests pin behaviour specific to ``NaiveGPUStreamScheduler``: a single end-of-state fused sync tasklet with one ``gpu_streams[<i>]`` in-edge per stream id, asserted in-state. The pipeline default is now :class:`AutoSingleStreamGPUScheduler`, which places syncs in a dedicated sync state instead, so the assertions failed under the new default. Construct ``GPUStreamPipeline`` with ``NaiveGPUStreamScheduler()`` for the module-level pipeline, mirroring the fix already applied to ``explicit_gpu_stream_management_test``.
The previous fix for ``mempool_test`` tagged every ``GPU_Global`` AccessNode with ``gpu_stream_id = 0``. The predicate was too wide: it caught AccessNodes inside NestedSDFGs that the wiring pass propagates ``gpu_streams`` *into*, which then surfaced as ``Node validation failed: 'gpu_streams'`` on every NSDFG with an inner GPU array (``tests_cuda_block_test_nested2``, ``tests_codegen_gpu_memcpy_test_transpose_shared_to_global``, ``tests_codegen_nested_kernel_transient_test_transient``, ...). ``experimental_cuda.py``'s pool branch only consults the stream when ``desc.pool`` is true, so restrict the tagging to that case. The rest of the GPU_Global AccessNodes don't need a per-node stream id and were getting the tag for no reason. 47 local stream-suite tests still pass with the narrower predicate, including ``mempool_test_runs_correctly_and_emits_expected_calls``.
CI ``test-gpu-experimental`` surfaced ``Node validation failed: 'gpu_streams'`` errors across ~22 tests where the failing NSDFG sits inside a ``GPU_Device`` map (``cuda_block_test_nested2``, ``codegen_nested_kernel_transient_*``, ``codegen_gpu_memcpy_test_transpose_shared_to_global_*``, ``reduce_test`` GPUAuto, ...). ``gpu_streams`` is a host-side array; the wiring pass deliberately does not propagate it into kernel-internal NSDFGs (they execute on the kernel's own stream). Two places needed the same gate the wiring pass already uses: * ``AutoSingleStreamGPUScheduler.insert_sync_tasklets`` walked every CFR via ``all_control_flow_regions(recursive=True)`` and would splice or append a sync state into kernel-internal regions, planting a ``gpu_streams[0]`` memlet with no corresponding inner ``gpu_streams`` array. * ``SplitStateByGPUClass.apply_pass`` walked every nested SDFG and called ``state_fission`` on its states, generating new NSDFG-wrapped substates that again referenced ``gpu_streams`` without a connector. Both now skip via ``is_inside_gpu_device_kernel(sub_sdfg)``, mirroring ``wire_stream_connectors`` and ``_find_child_sdfgs_requiring_gpu_stream``. Local stream-suite still 47/47.
CI surfaced ``Node validation failed: 'gpu_streams'`` and ``Isolated node ..._cpu_before`` patterns across ~30 tests where the failing NSDFG sat at host level inside an outer ``GPU`` state. ``insert_sync_tasklets`` was walking ``all_control_flow_regions(recursive=True)``, splicing or appending a sync state inside that NSDFG's CFG. The inner NSDFG has no inner ``gpu_streams`` array (wiring only propagates it when there's a stream consumer to bind to, not for sync-state insertion), so validation broke. Per the user's design: * ``AutoSingleStreamGPUScheduler.assign_streams`` / ``insert_sync_tasklets`` now classify and operate over the root SDFG's ``.nodes()`` and ``.edges()`` only. Each root-level block (``SDFGState``, ``LoopRegion``, ``ConditionalBlock``) is classified by ``_classify_root_block``, which folds ``_classify_node`` for states and folds over child blocks for control-flow regions. NSDFGs are folded into their containing state's classification via ``_classify_node`` (already short-circuiting on ``is_inside_gpu_device_kernel`` -- no need to recurse). * Sync state placement: GPU -> non-GPU root iedges get a splice; root SDFG sinks that are GPU get a trailing sync state. No sync injection inside ``LoopRegion`` bodies (would have fired once per iteration) or inside NSDFGs. * ``_collect_gpu_written_arrays`` uses the existing ``read_and_write_sets()`` on each GPU root block instead of ad-hoc AccessNode walking. * ``SplitStateByGPUClass.apply_pass`` now only attempts to split root-level ``SDFGState`` blocks. Other top-level block kinds are opaque; the boundary between their GPU work and surrounding CPU work is already handled at root iedges. * Removed defensive ``try/except``, lazy imports, and ``getattr``/``hasattr`` patterns in ``_classify_node`` and the sync helpers. End-to-end test: a CPU -> GPU-NSDFG-state -> CPU 3-state SDFG (via ``/tmp/three_state_cpu_gpu_cpu.py``) gets the sync state spliced at the outer GPU -> CPU edge with the inner NSDFG untouched. 47 local stream-suite tests stay green.
Two hand-built SDFGs that exercise the failure shape from ``failed_validation.sdfg``: * ``test_sync_placed_at_root_not_inside_nsdfg`` -- simple ``cpu_pre -> gpu_mid (NSDFG wrapping a GPU map) -> cpu_post`` chain. * ``test_sync_at_root_with_nsdfg_inside_gpu_device_map`` -- the richer ICON-style shape: a ``GPU_Device`` map at the outer level with a nested ``GPU_ThreadBlock`` map and a ``NestedSDFG`` inside it. Both assert that after ``GPUStreamPipeline()``: * exactly one sync state exists at the root SDFG level, spliced between the GPU block and the trailing CPU block, named ``__gpu_sync_after_*``; * zero sync tasklets anywhere inside any inner SDFG; * the SDFG validates. The previous recursive walk in ``insert_sync_tasklets`` spliced a sync state inside the inner NSDFG, planting a ``gpu_streams[0]`` memlet in an SDFG with no inner ``gpu_streams`` array. Now restricted to root ``.nodes()`` / ``.edges()`` -- this test is the regression pin. Also verified that the 15 previously-failing CI tests (``cuda_block_test``, ``cuda_smem2d_test``, ``codegen/gpu_memcpy_test``, ``codegen/nested_kernel_transient_test``, ``transformations/gpu_transform_test``, ``transformations/gpu_grid_stride_tiling_test``, ``transformations/warp_tiling_test``, ``blockreduce_cudatest``) now all pass locally with the root-only walk fix.
* Add ``tests/gpu_specialization/failed_validation_e2e_test.py``: an exact structural reconstruction of the ICON ``failed_validation.sdfg`` ``__field_operator_pnabla`` shape -- ``metrics_entry`` ConditionalBlock, ``stmt_0`` compute state (free host deref + GPU_Device map with an inner NestedSDFG), ``metrics_exit`` ConditionalBlock -- built via the SDFG API. Assertion contract: after ``GPUStreamPipeline()`` runs, exactly one ``cudaStreamSynchronize`` Tasklet exists at the root level (in a ``__gpu_sync_after_*`` state), zero sync Tasklets inside any inner NestedSDFG, ``SplitStateByGPUClass`` lifted the host deref into a ``*_cpu_before`` state, and the SDFG validates clean. * Replace ``split_state_by_gpu_class._weakly_connected_components`` hand-roll with ``networkx.weakly_connected_components(state.nx)``. DaCe's ``OrderedDiGraph.nx`` property exposes the underlying ``nx.DiGraph``; DaCe's own ``sdfg.utils.weakly_connected_component`` is the singular form (one component for a seed node), so the multi-component case routes through networkx directly. Local stream suite + new e2e: 49/49 pass.
…mantics
New ``filter`` ``SetProperty`` on ``ConvertLengthOneArraysToScalars``.
Semantics:
* ``filter=None`` (the default) -- legacy behaviour: every length-1
``Array`` eligible under ``transient_only`` is scalarized.
* ``filter={names}`` -- the pass becomes an exclusive allow-list at the
root SDFG level: only arrays whose name appears in ``filter`` are
scalarized, and **being in the filter overrides the ``transient_only``
check** (a named non-transient array is still scalarized).
* The filter has no effect on the nested-SDFG recursion: inner
descriptors continue to follow the unchanged transient-only rule
(filter names refer to root-level data, not nested connectors).
``_rewrite`` now takes an ``apply_filter`` flag; the recursion calls it
with ``apply_filter=False`` so the outer filter is not erroneously
applied to inner names.
Five new unit tests cover: ``None`` == legacy; exclusive allow-list at
root; override of ``transient_only`` for listed names; unknown name
yields no rewrites; nested-SDFG recursion is unaffected by the outer
filter.
9/9 length-1 conversion unit tests pass; pre-commit clean.
…wrapping NSDFGs The pass on yakup/dev (commit 752cb69) rewrites the "GPU_Device nested inside GPU_Device" shape -- which the experimental CUDA codegen refuses ("Dynamic parallelism (nested GPU_Device schedules) is not supported") and the legacy codegen handles via a buggy ``__gbar`` grid-barrier -- into a single flat ``GPU_Device`` kernel whose outer map absorbs the union of the inner kernels' iteration ranges and each inner kernel body becomes an if-bound-checked ``NestedSDFG``. On the ICON ``native_functions_main`` reproducer (:repro:`https://github.com/spcl/dace/blob/new-gpu-codegen-dev/...`) two follow-up fixes are needed for the rewrite to validate clean: * When the outer ``GPU_Device`` map is range-expanded with the inner kernels' params (``__i``, ``__j`` in the reproducer), every ``NestedSDFG`` between the outer map and the inner kernel's enclosing state must also gain those symbols in its ``symbol_mapping`` and inner ``sdfg.symbols``. Without it, the validator fires ``Missing symbols on nested SDFG: ['__i', '__j']`` at the wrapping ``nested_sdfg``. * The inner map's own iteration params -- referenced by the if-bound-check Tasklet inside the newly created ``NestedSDFG`` -- must be threaded explicitly into that new ``NestedSDFG``'s ``symbol_mapping`` and ``inner_sdfg.symbols``. ``state.symbols_defined_at(map_entry)`` does not surface them reliably after the in-place ``map.params`` / ``map.range`` mutation. End-to-end: the reproducer now validates AND compiles through the experimental CUDA codegen (no ``__gbar``, no undefined-symbol errors). Unit test (``lower_nested_gpu_device_maps_test``) builds a miniature ``outer GPU_Device (0:K)`` with a wrapping NSDFG holding two sibling GPU_Device kernels ``(0:J+1, 0:I)`` and ``(0:J, 0:I)``; asserts post-pass that the outer map has absorbed ``__j``, ``__i`` and that no inner GPU_Device map remains. Validates clean.
Add ``dace/transformation/passes/unify_close_iteration_domains.py`` and its unit-test file. The pass has two surfaces: * ``UnifyCloseIterationDomains`` -- a ``SingleStateTransformation`` matching the vertical chain ``MapExit -> AccessNode -> MapEntry``. Picks two maps whose iteration ranges differ per-dimension only by a **constant** (a sympy ``Number``); the optional ``max_constant_diff`` knob (default ``0`` -- accept any constant difference) caps that diff. Extends both maps to their per-dim ``(min(start), max(end))`` union and wraps the body of whichever map was strictly smaller in a ``NestedSDFG`` whose top-level ``ConditionalBlock`` re-imposes the original range as a Python ``and``-joined bounds check on the map's own iteration params. * ``UnifyCloseIterationDomainsPass`` -- walks every state in the SDFG hierarchy, collects each vertical ``MapExit -> AccessNode -> MapEntry`` triple, and applies the transformation where the ranges are close. Independent siblings (no producer-consumer chain) are out of scope -- per the user's "vertical only" instruction. The body-wrap is built on top of the standard :func:`dace.transformation.helpers.nest_state_subgraph` helper: it encapsulates the body into a ``NestedSDFG`` with the boundary memlets wired through the surrounding ``MapEntry`` / ``MapExit``, then promotes the resulting NSDFG's start state into a ``ConditionalBlock`` branch guarded by the original-range bounds check. The map's own iteration params are threaded into ``inner_sdfg.symbols`` and ``nsdfg.symbol_mapping`` so the bounds-check expression resolves. After the pass the two maps share an identical iteration range -- the precondition vertical :class:`MapFusion` needs to fuse them. 11 unit tests covering: range-closeness predicates (constant accept, cap reject, symbolic diff reject, dim-count reject); per-dim union; the transformation's apply/refuse cases including the ``max_constant_diff`` knob; the pass form including a "no close pairs" no-op; and the post-unify precondition for vertical MapFusion. All 11 green; pre-commit clean.
Two follow-up fixes to ``UnifyCloseIterationDomains`` so the combined pipeline ``unify -> MapFusion`` actually fuses pointwise vertical chains end-to-end: 1. **No unnecessary ifs.** Replace ``_bounds_check_expr`` with a per-dimension minimal-guard builder. Per dim, emit ``p >= orig_start`` only when the new range's start moved below the original (i.e. an extension was introduced on that side) and ``p <= orig_end`` only when the new range's end moved above. Dimensions whose original range equals the new range contribute nothing to the guard and are dropped. When every dimension matches, ``_bounds_check_expr`` returns ``None`` and ``_wrap_map_body_in_bounds_check`` is a no-op for that map -- no wrapper NSDFG / ConditionalBlock is materialised at all. Removes the redundant ``i >= 0`` clauses the previous implementation always emitted when ``orig_start == new_start == 0``. 2. **``nest_state_subgraph(full_data=False)``.** The previous ``full_data=True`` promoted the wrapped map's boundary memlets to the full-array shape (e.g. ``arr[0:I, 0:J, 0:K]``), which broke vertical ``MapFusion.partition_first_outputs``' producer-covers-consumer check: the producer writes one element per iteration but the consumer would appear to read the whole array. Keeping the per-iteration subset (``arr[i, j]``) lets MapFusion fire. Combined-test upgrade: ``test_combined_unify_then_map_fusion_collapses_ to_single_map`` now actually runs ``MapFusionVertical`` after unify and asserts a single fused top-level map (instead of just asserting the precondition that ranges match). Plus two new minimal-guard unit tests (``test_transformation_only_guards_dimensions_actually_extended`` and ``test_transformation_no_guard_when_range_unchanged``). Also picks up uncommitted work from earlier in this session: * ``tests/gpu_specialization/failed_validation_e2e_test.py`` -- enrich the ICON ``failed_validation`` reproducer's compute state with a free pre-map deref, a short Sequential ``seq_neighbors_map[0:4]`` nested inside the GPU_Device map, and an in-map Tasklet pair before/after the NestedSDFG (mirrors the ICON ``tlet_6_V2E_neighbors_map[0:7]`` neighbour-iteration shape). * ``tests/passes/lower_nested_gpu_device_maps_test.py`` -- pure yapf formatting (CI's ``--all-files`` yapf wanted the ``state.add_map(...)`` call on one line). 14/14 pass unit tests (unify_close * 13 + lowering * 1) + 1 failed- validation e2e test green; ``pre-commit run --all-files`` clean.
…G boundaries
Five interlinked fixes to the sync-insertion pipeline plus a new GPU
correctness test file that pins five edge-case patterns end-to-end.
``dace/transformation/passes/gpu_specialization/gpu_stream_scheduling.py``
* **Philip's typo fix.** Two ``self._recursive(node.sdfg, ...)`` calls in
``_add_sync_state`` are renamed to ``self._add_sync_state(node.sdfg, ...)``
(the method that actually exists).
* **Program-end sync only at GPU sink states.** ``_add_sync_state`` used
to append a trailing ``cudaStreamSynchronize`` after *every*
``_Kind.GPU`` state -- including ``*_copyin`` / ``*_copyout`` scaffold
states that the classifier folds into ``_Kind.GPU`` because they
contain ``GPU_Global`` AccessNodes. Result: 2 syncs for a CPU -> GPU ->
CPU chain, 3 for a loop-body alternating pattern. Gate the append on
``state.parent_graph.out_degree(state) == 0`` -- non-sink GPU states
are already covered by the edge-splice loop in ``insert_sync_tasklets``.
* **Tuple-unpacking fix on the splice loop.** ``edges_to_splice`` carries
``(region, edge)`` tuples but the consume loop was ``for edge in
edges_to_splice: _splice_sync_state_on_edge(sdfg, edge, sdfg, ...)``,
passing the tuple where an ``Edge`` was expected. Failed silently
before because the loop's predecessors rejected every edge; once the
sink-only gate above started leaving edges for the splicer to handle,
the call fired and raised ``AttributeError: 'tuple' object has no
attribute 'src'``. Switch to ``for region, edge in edges_to_splice:
_splice_sync_state_on_edge(region, edge, sdfg, ...)``.
* **Accept any block kind on splice edges.** The splice loop was gated
on ``isinstance(src, SDFGState) and isinstance(dst, SDFGState)``,
silently skipping GPU-state -> ConditionalBlock and CFG-with-GPU-inside
-> state transitions. The classifier (``_classify_root_block``) already
returns the union of a CFG's descendant kinds, so just drop the kind
filter on src+dst.
``dace/transformation/passes/gpu_specialization/split_state_by_gpu_class.py``
* **Philip's ``state_fission(allow_isolated_nodes=False)`` fix.** Both
fission calls (the ``_cpu_before`` lift and the ``_gpu_middle`` lift)
now pass ``allow_isolated_nodes=False`` so isolated nodes move into
the new state rather than getting stranded.
``tests/gpu_specialization/sync_insertion_correctness_test.py`` (NEW)
Five ``@dace.program`` tests using ``range(...)`` for host loops and
``dace.map[...]`` for parallel maps, built through the standard path
``to_sdfg -> auto_optimize(GPU) -> GPUStreamPipeline``. Each asserts
both the expected ``cudaStreamSynchronize`` tasklet count (via the
``_count_sync_tasklets`` helper pattern from
``monolithic_single_stream_test.py``, using
``dace.codegen.common.get_gpu_backend()``) and ``np.testing.
assert_allclose`` against a numpy reference. Patterns covered:
* three-state CPU -> GPU -> CPU chain
* mixed-class single source state (CPU prefix + GPU kernel)
* two independent parallel writers feeding one host reader
* parallel output consumed by a host-side loop
* loop body alternating host init -> parallel kernel -> host accumulate
All five expect exactly one sync each; before the fixes they emitted
1, 1, 1, 2, and 3 respectively. Tests marked ``gpu + new_gpu_codegen
_only`` so they run only on GPU CI.
Targeted regression: 20/20 in
``tests/gpu_specialization/{three_state,failed_validation,split_state,
sync_insertion,auto_single_stream}*`` pass, including the previously
brittle ``test_sync_at_root_with_nsdfg_inside_gpu_device_map`` that
was hitting the tuple-unpack ``AttributeError`` on CI.
pre-commit run --all-files clean.
Three small refactors to swap hand-rolled code for existing DaCe utilities:
* ``tests/gpu_specialization/sync_insertion_correctness_test.py`` and
``tests/gpu_specialization/monolithic_single_stream_test.py``: replace
``_count_sync_tasklets``' ``f"{backend}StreamSynchronize(" in
node.code.as_string`` string-match with the canonical
:func:`dace.transformation.passes.gpu_specialization.helpers.gpu_helpers.
is_pipeline_sync_tasklet` predicate. This is the same helper the pipeline
uses internally to detect its own sync tasklets, so the test
``_count_sync_tasklets`` and the pipeline's own ``is_pipeline_sync_tasklet``
callsites stay in lockstep if the label set ever changes. Drops the
``from dace.codegen import common`` import.
* ``dace/transformation/passes/unify_close_iteration_domains.py``: replace
the local ``_same(r_a, r_b)`` + ``_flatten_range`` pair with a new
``_ranges_equal`` helper that delegates per-dim symbolic comparison to
:func:`dace.symbolic.equal` (the standard DaCe helper, which returns
``True`` / ``False`` / ``None`` and handles SymExpr unwrap, undefined
symbols, and ``is_length`` integer-positive assumptions). A conservative
``True`` requires every per-dim bound and stride to compare ``True``;
inconclusive ``None`` results yield ``False`` (no semantic change versus
the old ``sympy.simplify(a - b) == 0`` check, but gets DaCe's standard
symbol-equality protocol for free). ``_flatten_range`` is unused after
the refactor and is removed.
35/35 pass unit tests + 23/23 structural ``tests/gpu_specialization/``
tests pass after the refactor; pre-commit ``--all-files`` clean.
…workx Round-2 audit refactor on ``gpu_stream_scheduling.py``: swap hand-rolled patterns for established APIs or remove dead defensive code per the "NO ``getattr`` / ``hasattr`` / bare ``except Exception``" guideline (``feedback_no_core_dace_changes`` companion rule). * **``NaiveGPUStreamScheduler._weakly_connected``**: replace the hand-rolled DFS WCC implementation with ``networkx.weakly_connected_components(graph.nx)`` -- same refactor as ``split_state_by_gpu_class._weakly_connected_components``, same rationale (DaCe's ``OrderedDiGraph.nx`` exposes the underlying ``nx.DiGraph``, so the test stays in lockstep with DaCe's own graph internals). * **``_classify_node`` schedule lookup**: drop the ``getattr(node, 'schedule', None) or getattr(getattr(node, 'map', None), 'schedule', None)`` chain. ``MapEntry`` carries the schedule on ``.map``, ``ConsumeEntry`` on ``.consume`` -- both well-defined, both Property-backed. Dispatch by type and access directly. * **``_classify_node`` scope-subgraph fetch**: drop the broad ``try/except Exception`` around ``state.scope_subgraph(...)``. The branch already filters to ``MapEntry`` / ``ConsumeEntry``, both guaranteed scope entries; ``scope_subgraph`` won't raise. * **``MonolithicSingleStreamGPUScheduler._not_acceptable_reason``**: drop ``getattr(node, 'schedule', None)`` on the LibraryNode branch. ``LibraryNode.schedule`` is a Property with default ``ScheduleType.Default``; the attribute always exists. * **``_state_has_host_boundary_copy``**: drop ``hasattr(node.code, 'as_string')``. Tasklets' ``code`` is always a ``CodeBlock``; ``.as_string`` always exists. * **``_iedge_reads_gpu_array``**: drop the ``try/except Exception`` around ``InterstateEdge.read_symbols()``. ``read_symbols`` is a documented, non-raising method; the broad except hides real bugs without buying anything. * **Pool-narrow ``MonolithicSingleStreamGPUScheduler.assign_streams``**: swap ``getattr(desc, 'pool', False)`` for an explicit ``isinstance(desc, data.Array)`` guard. ``data.Scalar`` and ``data.Stream`` don't carry a ``pool`` property; the type guard is the honest predicate. 42/42 pass on the affected ``tests/gpu_specialization`` and ``tests/passes`` modules, ``pre-commit run --all-files`` clean.
The ``while enclosing_scope is not None`` loop in ``AutoSingleStreamGPUScheduler._add_sync_state`` never updated ``enclosing_scope`` inside the body. When the immediate parent map's schedule wasn't in ``GPU_SCHEDULES`` the loop spun forever, eventually hitting the CI test timeout on ``tests/transformations/gpu_grid_stride_tiling_test.py:: test_gpu_grid_stride_tiling_with_indirection``. Walk up the scope chain via ``scope_dict`` after each comparison so the loop terminates: either an ancestor map carries a GPU schedule (break), or we reach the top (``None``, then the ``while ... else`` clause runs the NSDFG-recursion). Reproduced locally with the failing test -- 24s pass after the fix; the full ``tests/gpu_specialization/`` + ``tests/passes/`` sweep is 599 passed, 2 skipped, 0 failed.
Audit round 3: continues the round-2 dead-defensive-code cleanup with three more found via deeper inspection. * **Dead edges-to-splice loop.** ``insert_sync_tasklets`` declared and populated an ``edges_to_splice`` list from ``sdfg.edges()`` (root-only), then **shadowed the variable** with a fresh ``List[Tuple['AbstractControlFlowRegion', any]] = []`` declaration and re-iterated ``all_control_flow_regions(recursive=True)``. The first loop's results were thrown away on every call. Drop the dead loop and fix the lingering comment that referenced "flat state machine". * **``getattr(node, 'label', node)``** (x2 sites). ``nodes.Node.label`` is a Property defined on every Node subclass; the fallback to ``node`` was masking either a real bug (some Node without a label?) or pointless defensiveness. Drop both. * **``hasattr(node.code, 'as_string')``** in ``helpers/gpu_helpers.is_gpu_stream_consumer``. ``Tasklet.code`` is always a ``CodeBlock``; ``.as_string`` always exists. Drop the guard and inline the lookup. 43/43 targeted tests pass (8 ``gpu_specialization/*`` + 2 ``passes/*`` + the previously-timing-out ``tests/transformations/gpu_grid_stride_tiling_test.py:: test_gpu_grid_stride_tiling_with_indirection``). pre-commit ``--all-files`` clean. Notable **latent concern** flagged but **not changed** here, for a separate decision: ``self._state_kinds`` is populated only for the root SDFG's blocks (line 645: ``for block in sdfg.nodes()``), but the edge-splice loop walks ``sdfg.all_control_flow_regions(recursive=True)`` and dereferences ``self._state_kinds.get(src)`` for inner-CFG-edge src states that are never classified. That branch falls through "not GPU" and is a no-op, which may mean a GPU-state-inside-LoopRegion -> host- state-inside-LoopRegion transition gets no explicit sync. Current ``sync_insertion_correctness_test.test_loop_body_alternating_host_ parallel_has_one_sync_and_matches_numpy`` passes (the runtime-side stream-0 memcpy semantics likely cover correctness), but the implicit reliance on runtime behaviour is worth a follow-up design pass.
…it syncs
AutoSingleStreamGPUScheduler appends a cudaStreamSynchronize at every GPU sink /
GPU->host boundary, making each GPU SDFG synchronous on return. For a host app
that shares one GPU stream across SDFG calls and synchronizes at its own
host-read boundaries (e.g. icon4py), that per-SDFG sync is pure overhead -- a
fixed ~0.66 ms/call host stall that dominates launch-bound stencils.
Add compiler.cuda.synchronize_on_exit (bool, default True = current behavior).
When False, the scheduler skips the exit sync only for GPU-resident outputs:
- program-end sink sync skipped when the sink writes no host-visible output
(_sink_writes_host_visible_output);
- GPU->host edge sync skipped when the host block reads no GPU-written array
(_block_reads_gpu_written) -- e.g. a trailing metrics state that only writes
the host-side gt_compute_time scalar.
Host-visible / copy-out outputs are always synchronized regardless, so numpy /
read-back SDFGs stay correct.
The flag is a scheduling-strategy constructor argument
(AutoSingleStreamGPUScheduler(synchronize_on_exit=...)); the experimental codegen
pipelines read the config and pass it down (None defers to the config).
Tests in tests/gpu_specialization/auto_single_stream_test.py cover default-keeps,
opt-out-drops (program-end and GPU->host edge), host-visible-always-synced, and
explicit-arg-overrides-config.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
ThrudPrimrose
force-pushed
the
new-gpu-codegen-dev
branch
from
June 13, 2026 16:33
097bd8a to
ced4111
Compare
…rcular import on Python 3.14) fparser 0.2.3 changed Fortran2008/__init__.py so that importing Fortran2008 first triggers Fortran2008 -> Fortran2003 -> (DynamicImport) label_do_stmt_r816 -> `from fparser.two.Fortran2008 import Loop_Control`, which on Python 3.14 raises ImportError from the partially-initialized Fortran2008 module. This broke every tests/fortran/* at collection on the general-ci Python 3.14 job (fparser 0.2.2 is unaffected; the unbounded `fparser >= 0.1.3` lets CI pull the broken 0.2.3). ast_components.py already imports both modules; loading Fortran2003 first breaks the cycle (Base is defined early in Fortran2003, so its partial init is fine) and is the isort-sorted order. No version pin needed -- the new fparser works. Verified on Python 3.14.0rc2 + fparser 0.2.3: the fortran frontend imports and tests/fortran/fortran_language_test.py passes (7/7, was ImportError at collection); Python 3.13 + fparser 0.2.2 still imports. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The ancient pin (fparser==0.1.4) is stale -- the codebase runs on fparser 0.2.x. With the Fortran2003-before-Fortran2008 import fix (commit 74110d6), the latest fparser (0.2.3) works on Python 3.14, so no version pin is needed. setup.py already declares an unbounded `fparser >= 0.1.3`. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…, and read persisted stream assignments from one helper
…enerate_Tasklet at main's complexity
# Conflicts: # dace/codegen/targets/cuda.py # dace/runtime/include/dace/cuda/dynmap.cuh # tests/codegen/gpu_index_type_test.py
The CUDA environment put cuda_runtime.h into the host frame whatever the backend, which a HIP build cannot include. dace.h already includes the runtime of the selected backend, so the environment declares no header again, as on main.
…GPUScheduler The stream lowering helpers move into gpu_stream_scheduling.py, and the monolithic strategy class becomes AutoSingleStreamGPUScheduler(monolithic=True), which keeps its on-device requirement and its per-state synchronization. NaiveGPUStreamScheduler reads compiler.cuda.max_concurrent_streams when it assigns streams instead of when it is built, so a pipeline built before a set_temporary honors it.
…een the GPU codegens The chiplet distribution of the legacy codegen (grid padding, block index remapping and guards) becomes module functions in cuda.py that the experimental kernel scope uses as well, and kernels without a thread-block map span threads. A pooled array released early is no longer freed a second time at deallocation in either codegen. Dynamic map range connectors reach the experimental kernel arguments through the helper the legacy launch uses, and a connector named after its container is not redefined at the launch site.
…dex types, warp size and shared memory planning The experimental codegen declares GPU map indices with common.gpu_map_index_types, the warp size with common.gpu_warp_size and warp ids with common.gpu_thread_id_type, which drops the gpu_index_type, cuda_warp_size and hip_warp_size config keys. It places every kernel's shared memory with gpu_shared_memory.plan_gpu_shared_memory through helpers shared with the legacy codegen, so symbolic and oversized shared arrays become dynamic shared memory, launched with their size and the opt-in beyond the static limit. LiftSharedOutOfNestedSDFG is removed: nested SDFGs at any depth and sibling nested SDFGs run correctly with the planner alone (tests/codegen/gpu_nested_shared_memory_test.py). InferGPUGridAndBlockSize enumerates kernels with gpu_helpers.gpu_kernels, and the kernel signature opens its body on the same line, as in the legacy codegen.
…y, give the dynamic-map test its own M32 symbol, and expect the experimental lowering of nested device maps
# Conflicts: # dace/codegen/targets/cuda.py # tests/codegen/gpu_index_type_test.py
ThrudPrimrose
added a commit
that referenced
this pull request
Sep 28, 2026
#2259's GPU codegen files win; extended-only GPU extras that depended on removed modules are dropped (chiplet module, AddThreadBlockMaps one-sweep tiling). Shared helpers keep extended's fixes: CUDA architecture resolution and host compiler in gpu_cmake_options, and the static+dynamic shared memory opt-in (now in dynamic_shared_memory_request, used by both codegens). Core files keep extended's side where it is the more general form; gpu_specialize_offloaded is kept for canonicalize finalize.
ThrudPrimrose
added a commit
that referenced
this pull request
Sep 28, 2026
…2259 CUDA codegen The #2259 merge took the new-gpu-codegen side of insert_explicit_copies.py, which left competing, converting and single-element WCR copies implicit, so CPF rendered dace::CopyND; and its experimental CUDA codegen dropped cpf_gpu_context and the one-unit output, so CPF device units included dace/dace.h. Existing tests in insert_explicit_copies_test, test_copy_lowering, sequential_wcr_accumulator_test and the CPF CUDA/HIP emission tests cover both (34 -> 10 failures in those files; the rest need cupy or predate the merge).
… stream through host map scopes MoveArrayOutOfKernel now lifts a symbolically-sized Register/Default transient inside a kernel, which would otherwise be a variable-length array nvcc rejects, and prefixes the array's subscripts in inlined tasklet bodies as well as memlets and control flow. The GPU stream array reaches a NestedSDFG inside a host map through the map's pass-through connectors instead of an edge crossing the scope, which scope_dict refused with 'Leftover nodes in queue'.
… whole identifier The per-stream connector __dace_current_stream_0 contains the legacy name as a substring, so every stream-synchronization tasklet declared an alias it never read.
ThrudPrimrose
added a commit
that referenced
this pull request
Sep 29, 2026
…DA folds map-exit WCR by BlockReduce - Cholesky: the merge dropped the _info access node. - ScalarToSymbolPromotion reads enclosing loop iterators' types, so a lossless widening of a scope symbol (dace.int64(ii) + W) promotes; tiled loops collapse again under #2605 iterator types. - Experimental CUDA thread-block and per-thread kernel scopes fold map-exit WCR accumulators with gpucub::BlockReduce and one atomic per block, through the helpers the legacy codegen uses. - Experimental kernel names carry the top-level SDFG name again; single-spaced __global__ signature. - An empty-memlet in-edge is not a write when CPF and scalar conversion decide a scalar is written. - CPF sorts with std::sort in C++ and an OpenMP task mergesort in C; ParallelSTL is provided. - s318 test pins the carried IV as the remaining lift blocker.
ThrudPrimrose
added a commit
that referenced
this pull request
Sep 29, 2026
…asses and codegen - MoveArrayOutOfKernel lifts a symbolically-sized Register/Default kernel transient (a VLA nvcc rejects) like a GPU_Global one, and prefixes the array's subscripts in inlined tasklet bodies too. - The stream array reaches a NestedSDFG inside a host map through the map's pass-through connectors. - A device-to-host copy is waited for on the stream it was issued on. - GPUCodegenPreprocessPipeline promotes pending warp tiles before AddThreadBlockMaps. - Restore extended's length-one scalar conversion pass and tests, copy_node and aligned_allocation tests, which the merge had reverted to older versions. - Tests: nested device map index types for both codegens; the lift-binding and typecast assertions no longer pin one codegen's declaration style; the chiplet thread-block-map requirement test goes.
…writes (#2614) ## Problem With `compiler.cuda.implementation = "experimental"`, GT4Py `concat_where` tests over dynamic domains segfault in `CompiledSDFG.fast_call`. The legacy codegen works. In the generated host code, a copy between two `GPU_Global` transients is emitted as a host `dace::CopyNDDynamic<...>::Copy(...)` over device pointers. Root cause: - `ExperimentalCUDACodeGen.copy_memory` assumes `InsertExplicitCopies` has already lifted every copy touching GPU memory to a `CopyLibraryNode`. It hands whatever is left to the CPU codegen. - `InsertExplicitCopies` leaves a direct copy implicit when `_competing_writer` reports that another edge may write an overlapping region of the destination. An *undecidable* overlap (`subsets.intersects(...) is None`) counts as a conflict. - In the GT4Py SDFG, the destination is also written by a GPU map. The two regions are disjoint in K, but their bounds are symbolic `Max`/`Min` expressions, so `intersects` returns `None`. The copy stays implicit and ends up as a host copy. ## Fix - **`InsertExplicitCopies._replace_direct_copies`:** a direct copy that involves GPU-resident storage is lifted even when a competing writer may overlap. - A plain copy edge is emitted when its source access node is visited, so it runs before every other consumer of that source. To keep that guarantee, the lifted copy gets ordering edges (empty memlets) to those consumers. This is done by the new helper `_consumers_to_precede`. - If that ordering would create a cycle with the edges carried by `_carry_write_ordering`, the copy is left implicit as before. - Copies that involve only CPU memory are unchanged, so the npbench `vadv` behaviour `_competing_writer` protects is not affected. - **`ExperimentalCUDACodeGen.copy_memory`:** a host-side copy between two access nodes that touches GPU memory now raises `CodegenError` instead of emitting a host copy that would segfault. ## Tests All new tests run without a GPU (they only generate code). - `tests/codegen/experimental_cuda_competing_copy_test.py` (new): - A reduced reproducer: a GPU map writes `C[0:Min(K, N)]` and a device-to-device copy writes `C[Max(K, 0):N]`. The test asserts the host code has no `CopyND` and does have `cudaMemcpyAsync`. - A test that a GPU copy left implicit raises `CodegenError`. - `tests/passes/insert_explicit_copies_test.py::test_device_copy_competing_with_another_write_is_lifted_ahead_of_the_source_consumers`: the copy is lifted and ordered before the source's other consumer. All three fail without the fix and pass with it. On the GT4Py SDFG, the host copy is replaced by a `cudaMemcpyDeviceToDevice` of the same size the legacy codegen emits. Locally (macOS, no GPU) the following pass: `insert_explicit_copies_test.py` including the CPU numeric runs, `intermediate_dead_store_test.py` (derived from vadv), the experimental CUDA codegen tests, and `tests/gpu_specialization`. The only failures are tests that need nvcc. ## Open points - Nothing has been run on a GPU yet. A GH200 run of the GT4Py `concat_where` tests would confirm the fix end to end. - The ordering edges only keep the guarantee the graph already gave: "before the source's other consumers". If the competing writer does not depend on the source, its order relative to the copy was never guaranteed, even for the implicit copy. - Copies staged in or out through map entries/exits (`_lift_staging_edge`) keep the old skip on a competing writer. A host-side one touching GPU memory now raises the new error instead of segfaulting. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…r once The lane partial was initialized by fissioning the state, which is only sound once the kernel body is nested; on a flat body it split the kernel scope (wrong result or a scope error). The partial is now set to the reduction identity inside the lane scope, ordered before the strided map, and the lane reduction folds the accumulator's value, read from the node that initialized it, exactly once.
…azard SplitStateByGPUClass now depends on InsertExplicitCopies: both were bare dependencies of the scheduler, so their order followed class hashes and a copy was classified GPU in some runs only (flaky sync counts). A GPU -> host edge now syncs only when the host block reads GPU output or writes an array the queued GPU work touches; exit visibility moves to a trailing sync after a host sink that GPU work still reaches.
…and a strided map Only the innermost nested SDFG got the symbol, so a strided map two levels down failed validation with "Missing symbols on nested SDFG: ['__tid']" (polybench trmm on extended).
ThrudPrimrose
added a commit
to philip-paul-mueller/dace
that referenced
this pull request
Sep 29, 2026
…flow is read before assignment (spcl#2530) Folds each loop block's own free symbols into LoopToMap's read-before-assigned set, so a wrap-around induction such as TSVC s291 (`im = N-1; a[i] = b[i] + b[im]; im = i`) is correctly refused instead of pinning `im` to `N-1`. Split out of the new GPU codegen PR (spcl#2259) as a prerequisite.
…ore changes, share one traversal and stream query, drop duplicate tests Fixes that are independent of the new target (legacy callback pinning, warp reduction identity, small-map scheduling, cuSolver selection, link flags, rocTX gating, map range symbols) are removed from the branch, and the prerequisite passes now equal their own PR branches. Symbol queries of GPU index inference go through the frame, the dual legacy and experimental hooks read the target in use instead of the configuration, and every test file calls its tests from its main block.
NaiveGPUStreamScheduler is PerComponentGPUStreamScheduler. AutoSingleStreamGPUScheduler splits into SingleStreamGPUScheduler, which raises on a node that mixes host and GPU work, and the default AutoGPUStreamScheduler, which falls back to per-component streams there as before.
…nch.*, not by file path
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
4 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.
The complete PR for the NEW GPU codegen. (To be split into individual branches until everything has been merged).