Skip to content

feat[next-dace]: bump to dace v2.0 - #2937

Draft
edopao wants to merge 28 commits into
mainfrom
dace_noview_nested_sdfg
Draft

edopao wants to merge 28 commits into
mainfrom
dace_noview_nested_sdfg

Conversation

@edopao

@edopao edopao commented Oct 5, 2026 •

Copy link
Copy Markdown
Contributor

ICON4Py CI tests pass, see C2SM/icon4py#1471

@edopao
edopao added this pull request to stack #2935 October 5, 2026 07:23
@edopao
edopao force-pushed the dace_noview_nested_sdfg branch from d4c4686 to 3b0658d Compare October 5, 2026 13:01
Base automatically changed from bump_dace_version to main October 5, 2026 13:21
@edopao
edopao force-pushed the dace_noview_nested_sdfg branch from 3b0658d to 788525e Compare October 5, 2026 13:21
@edopao
edopao removed this pull request from stack #2935 October 5, 2026 13:22
@edopao

edopao commented Oct 9, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

1 similar comment
@edopao

edopao commented Oct 9, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao
edopao force-pushed the dace_noview_nested_sdfg branch from 1b31304 to d72f949 Compare October 9, 2026 14:54
@edopao

edopao commented Oct 9, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

1 similar comment
@edopao

edopao commented Oct 9, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao

edopao commented Oct 9, 2026

Copy link
Copy Markdown
Contributor Author

The failure in the RoCM CI (test_scan.py::test_solve_triag[dace.run_dace_gpu]) shows a problem in the lowering/optimization: some arrays are dynamically allocated inside the GPU kernel. This is supported in CUDA, not in HIP. However, it should never happen: it is a performance issue.

@edopao

edopao commented Oct 12, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

2 similar comments
@edopao

edopao commented Oct 12, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao

edopao commented Oct 12, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao

edopao commented Oct 12, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run dace-determinism

@edopao

edopao commented Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao

edopao commented Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run dace-determinism

@edopao

edopao commented Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

edopao added 16 commits October 13, 2026 16:50
GPUTransformLocalStorage (run with nested_seq=True) sets all
StorageType.Default transients inside a GPU map to StorageType.Register.
However, variable-length arrays (whose size is not known at compile
time, e.g. scan column temporaries) cannot be allocated in registers on
GPU — the DaCe code generator emits a call to the host-only
dace::aligned_new_array() inside the __global__ kernel, which nvcc
rejects.

Move such arrays to GPU_Global with the pool allocator instead.
…on in RemoveAccessNodeCopies

RemoveAccessNodeCopies repoints the edges of scan NestedSDFGs at the
full output field (e.g. next_w) by merging temporary access nodes.
The scan's inner output connector, however, has the domain shape,
which is windowed relative to the full field. This made
_gt_map_strides_into_nested_sdfg() raise:

  ValueError: The connector '__gtir_scan_output_x_kplus1' of
  NestedSDFG 'scan_1' is not equivalent to the data 'next_w'

The issue only surfaced with compile_time_domain/compile_time_vertical
variants, where the domain bounds are concrete and the pattern matcher
can prove the copy chain is a no-op; the default (symbolic) variant
does not trigger the pattern.

As the docstring of _gt_map_strides_into_nested_sdfg already states,
a transformation that connects different data to a connector must
restate it. Call restate_nested_sdfg_connector() to adopt the new
outer descriptor and shift the memlets inside by the offset of the
old domain-sized data within the full field, before propagating
strides.
@edopao
edopao force-pushed the dace_noview_nested_sdfg branch from 8c085c6 to 02259e7 Compare October 13, 2026 14:51
@edopao
edopao added this pull request to stack #2945 October 13, 2026 14:57
@edopao

edopao commented Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run default

@edopao

edopao commented Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

cscs-ci run dace-determinism

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant