Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
113 commits
Select commit Hold shift + click to select a range
cb33cf4
Add CopyLibraryNode / MemsetLibraryNode + InsertExplicitCopies pass
ThrudPrimrose May 13, 2026
9602b41
infer_types: default-schedule patch for CopyLibraryNode / MemsetLibra…
ThrudPrimrose May 13, 2026
6abc67b
Remove InsertExplicitCopies from SIMPLIFY_PASSES: lowering pass, not …
ThrudPrimrose May 13, 2026
8e807cf
tests/passes/iec: add xfail pins for view-lift bugs
ThrudPrimrose May 13, 2026
db43693
IEC: collapse AN->View->AN round-trips; CopyND fallback for rank-mism…
ThrudPrimrose May 13, 2026
f82811b
cleanup: dead code, docstrings, comments
ThrudPrimrose May 15, 2026
92f031c
docs: ASCII-only, repo backtick convention, concise docstrings
ThrudPrimrose May 15, 2026
9cf409e
docs: ASCII-only, repo backtick convention, concise docstrings
ThrudPrimrose May 15, 2026
26547d6
Unify libnode stream connector on __dace_current_stream (valid in bot…
ThrudPrimrose May 16, 2026
3812133
Remove libnode stream-input plumbing; rename STREAM_CONN -> CURRENT_S…
ThrudPrimrose May 16, 2026
dae3798
Reuse subset.num_elements(); unify memset tasklet builder
ThrudPrimrose May 17, 2026
6f26f73
Merge remote-tracking branch 'origin/explicit-copy-memset-nodes' into…
ThrudPrimrose May 17, 2026
18cc384
Drop stale stream-descriptor mention from CopyLibraryNode.validate do…
ThrudPrimrose May 18, 2026
a5540a2
Move no-cycle inline imports to module top in copy/memset libnodes
ThrudPrimrose May 18, 2026
fbc62ba
Move get_parent_map_and_loop_scopes to transformation.helpers; drop c…
ThrudPrimrose May 18, 2026
a40c343
Add length-1<->scalar conversion passes
ThrudPrimrose May 18, 2026
84a0079
Passes are the only API (drop standalone function)
ThrudPrimrose May 18, 2026
dad18c1
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose May 20, 2026
1cbc095
Move connector names to class constants; use dtypes.{CPU,GPU}_RESIDEN…
ThrudPrimrose May 20, 2026
8e39e71
Apply connector-contract + inner-literal consolidation to memset and …
ThrudPrimrose May 20, 2026
a9b8bea
helpers: use imported nodes module for MapEntry/Tasklet/LibraryNode i…
ThrudPrimrose May 20, 2026
ef777dd
Accept Fortran-packed layouts in CopyNDTemplate; route mixed C/F to M…
ThrudPrimrose May 20, 2026
192d6c9
MappedTasklet handles rank-mismatch via 1-D walker + int_floor/% deli…
ThrudPrimrose May 20, 2026
3074fd0
Revert MappedTasklet rank-mismatch; CopyNDTemplate is the only suppor…
ThrudPrimrose May 20, 2026
6a42deb
Refactor multi-dim copy tests onto shared helper + add unsupported-ca…
ThrudPrimrose May 20, 2026
31b963a
Extract auto_dispatch + merge memset test helpers + type hints on new…
ThrudPrimrose May 20, 2026
7b3e859
DRY/YAGNI: drop dead CopyExpansion fields, reuse collapse_shape_and_s…
ThrudPrimrose May 20, 2026
6c35792
Drop ExpandCopyNDTemplate; MappedTasklet handles rank-mismatch (CopyN…
ThrudPrimrose May 20, 2026
a8aceb7
Simplify _coarse_pick, _cuda2d_strides_are_supported, merge cudaMemcp…
ThrudPrimrose May 20, 2026
f42e134
Tailor _build_copynd_call to its sole shared-memory caller
ThrudPrimrose May 20, 2026
aef3558
Drop dynamic-input connectors from copy/memset libnodes; subset symbo…
ThrudPrimrose May 20, 2026
cb62493
Prune dead code: drop length-one<->scalar pass (unused), TODOs, requi…
ThrudPrimrose May 20, 2026
b55c5f5
Soften _is_cross_cpu_gpu docstring on Register handling assumption
ThrudPrimrose May 20, 2026
2e47679
Merge remote-tracking branch 'origin/main' into explicit-copy-memset-…
ThrudPrimrose May 20, 2026
567bd79
Inline _require_contiguous_subset as combined check at both call sites
ThrudPrimrose May 20, 2026
411b83a
Add type hints (StorageType, Range, SymExpr, LibraryNode forward refs…
ThrudPrimrose May 20, 2026
b29b36f
Improve _refine_cuda_impl_for_subsets docstring (accurate routing table)
ThrudPrimrose May 20, 2026
5e2ba18
Drop _memcpy_connector_typing: DaCe handles single-element pointer co…
ThrudPrimrose May 20, 2026
d99c6b0
Add 5 Auto-routed copy edge-case tests: 4D/1D flatten, 1D/4D inflate,…
ThrudPrimrose May 20, 2026
03f97ef
Trim inner-tasklet connector comments; unify memset inner connector t…
ThrudPrimrose May 20, 2026
3bf59c8
Rename memset tests: <expansion>_<rank>_<storage>; reject-test names …
ThrudPrimrose May 20, 2026
c8991f4
Drop new_gpu_codegen_only markers; remove stale connector-types test
ThrudPrimrose May 20, 2026
3a18bd6
Unify SDFG-construction helpers in copy_node_test.py
ThrudPrimrose May 20, 2026
4f723c6
Reject transpose pattern upfront in CopyLibraryNode + trim test docst…
ThrudPrimrose May 20, 2026
e891b81
Pin contract: same-rank copy needs matching per-dim subset sizes, not…
ThrudPrimrose May 20, 2026
ea15eca
Normalize storage-type references to the full dace.dtypes.StorageType…
ThrudPrimrose May 20, 2026
aa0adf7
MemsetLibraryNode: Auto falls back to 'pure' for non-contiguous subsets
ThrudPrimrose May 20, 2026
66f4624
Prune dead helpers, args, and stale docstring
ThrudPrimrose May 20, 2026
5c2f27c
InsertExplicitCopies: lift stage-in / stage-out copies to libnodes in…
ThrudPrimrose May 20, 2026
33e0fff
IEC cleanup: unify staging-lift methods + drop dead param
ThrudPrimrose May 20, 2026
8ceb7a4
copy_node_test: assert no dace::CopyND in generated C++ at every compile
ThrudPrimrose May 20, 2026
0686082
Inline single-callsite helpers, drop dead docs; simplify __main__ blo…
ThrudPrimrose May 20, 2026
22cb179
Trim verbose WHY comment on symbolic-< fallback in _is_consecutive_re…
ThrudPrimrose May 20, 2026
9451d5a
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose May 21, 2026
e0991d7
Derive owning SDFG from state in copy/memset libnodes and copy-insert…
ThrudPrimrose May 22, 2026
ee066ed
Unify contiguous-memcpy expansion across CPU and CUDA1D
ThrudPrimrose May 22, 2026
b716ee2
Audit InsertExplicitCopies and treat views as copy endpoints
ThrudPrimrose May 22, 2026
1a52336
Improve comments
ThrudPrimrose May 22, 2026
786eb82
Pre-commit
ThrudPrimrose May 22, 2026
7d02766
InsertExplicitCopies: memlet-path subset resolution, reuse is_in_scop…
ThrudPrimrose May 22, 2026
21d7ebb
Merge origin/main into explicit-copy-memset-nodes (prefer theirs)
ThrudPrimrose May 26, 2026
fd21b59
copy_node: scope-aware dispatcher + single-element invariant
ThrudPrimrose May 29, 2026
81f10cf
insert_explicit_copies: build outer-side Memlet from outer.data and r…
ThrudPrimrose May 29, 2026
7862281
fix(sdfg): resolve exit-write arg from the memlet tree root
ThrudPrimrose May 27, 2026
eddc2cc
argument_signature_test: split arglist assertion from GPU compile/run
ThrudPrimrose May 29, 2026
f696321
style: tighten outer-subset resolution comment in InsertExplicitCopies
ThrudPrimrose May 30, 2026
2e60b37
Improve register location detection
ThrudPrimrose Jun 4, 2026
aa0caba
Updated how `get_gpu_{backend, runtime}()` worked.
philip-paul-mueller Jun 8, 2026
334614a
The expansion of the copy library node always assumed CUDA and not Hi…
philip-paul-mueller Jun 8, 2026
f3f2a25
Removed a stray `print()`.
philip-paul-mueller Jun 4, 2026
a84c2d1
Fixed 1d memcopy (#2396)
philip-paul-mueller Jun 8, 2026
cca9637
copy_node: fix assumption-divergent shape equality + propagate stream…
ThrudPrimrose Jun 9, 2026
6fbb531
copy_node_test: sympify strides so 1D-single-element CUDA test parses
ThrudPrimrose Jun 9, 2026
02cfe08
update copy and memset node impl
ThrudPrimrose Jun 29, 2026
d88e2f7
Trim copy/memset PR to essentials and de-fluff
ThrudPrimrose Jun 29, 2026
0334e43
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Jul 9, 2026
f651312
copy_node: padded (1,N) row is a contiguous run under #2380 1D-slice …
ThrudPrimrose Jul 9, 2026
80d3815
test(copy_node): guard multi-row padded copy routes to pitched cudaMe…
ThrudPrimrose Jul 9, 2026
cbda96a
test(torch): drop unused ATen/NamedTensorUtils.h include
ThrudPrimrose Jul 9, 2026
dd68216
chore: gitignore local .mtoken secret
ThrudPrimrose Jul 9, 2026
df4ccdd
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Jul 13, 2026
c58df14
Fix MemsetLibraryNode tasklet expansion guard for GPU-resident storage
ThrudPrimrose Jul 13, 2026
0dc7bdb
Merge branch 'main' into explicit-copy-memset-nodes
tbennun Jul 14, 2026
9d8058b
Library-node codegen: single-source pitch/map params and guard dtype-…
ThrudPrimrose Jul 14, 2026
21c9c26
Merge remote-tracking branch 'origin/explicit-copy-memset-nodes' into…
ThrudPrimrose Jul 14, 2026
3e77df3
copy_node: route contiguous same-layout CPU copies to MemcpyCPU
ThrudPrimrose Jul 14, 2026
7e69a4b
insert_explicit_copies_test: exclude view-defining edges from other_s…
ThrudPrimrose Jul 14, 2026
38addf4
copy_node_test: fortran-packed same-rank copy now expects MemcpyCPU
ThrudPrimrose Jul 14, 2026
039f794
pre-commit: yapf-reflow copy_node picker and insert_explicit_copies v…
ThrudPrimrose Jul 14, 2026
ff6aa53
libnode: size-gate contiguous CPU copy/memset selection
ThrudPrimrose Jul 18, 2026
0af661b
libnode: shorten verbose comments on copy/memset expansions
ThrudPrimrose Jul 18, 2026
bd2d743
Update .gitignore
ThrudPrimrose Jul 18, 2026
b16b7c1
Merge remote-tracking branch 'origin/main' into explicit-copy-memset-…
ThrudPrimrose Jul 28, 2026
2f2ba28
Keep reference semantics and named subsets intact through explicit-co…
ThrudPrimrose Jul 28, 2026
dc241aa
Restore tests/polybench to its upstream form
ThrudPrimrose Jul 28, 2026
3375c16
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Jul 29, 2026
ba4598b
Merge remote-tracking branch 'origin/main' into upd_explicit-copy-mem…
Aug 3, 2026
6c0134e
Give each rank its own build cache root, under cache_distaware
ThrudPrimrose Aug 6, 2026
09d35d6
style: trim comments to density cap
ThrudPrimrose Aug 10, 2026
4e00e49
style: restore sphinx field lists
ThrudPrimrose Aug 10, 2026
978ab26
feat: parallel-for default for large/symbolic copies and memsets
ThrudPrimrose Aug 10, 2026
bcc8615
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Aug 10, 2026
16bd356
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Aug 11, 2026
aa69dc1
Merge branch 'feat/rank-aware-build-cache' into explicit-copy-memset-…
ThrudPrimrose Aug 11, 2026
ea49df7
ci: pin xdist workers and MPI ranks to GPUs
ThrudPrimrose Aug 11, 2026
fe1e51c
fix: unique cache mode survives pid recycling
ThrudPrimrose Aug 11, 2026
50ca24a
test: drop subsumed gpu pinning cases
ThrudPrimrose Aug 11, 2026
9be63eb
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Aug 11, 2026
51641e9
revert: drop pid cache token, stay close to main
ThrudPrimrose Aug 11, 2026
6d623d4
fix(copy): carry allow_oob through copy lifting and honour it
ThrudPrimrose Aug 11, 2026
c28869d
Merge branch 'main' into explicit-copy-memset-nodes
ThrudPrimrose Aug 11, 2026
e0dec12
fix(copy): never lift a WCR edge, and retarget the scope-boundary memlet
ThrudPrimrose Aug 13, 2026
af22ed9
Merge remote-tracking branch 'up/main' into explicit-copy-memset-nodes
ThrudPrimrose Aug 13, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions dace/config_schema.yml
Original file line number Diff line number Diff line change
Expand Up @@ -349,6 +349,14 @@ required:
generate "#pragma omp parallel sections" code around
them.

parallel_transfer_min_elements:
type: int
default: 262144
title: Parallel copy/memset element threshold
description: >
Minimum element count (symbolic size counts as large) for a
copy/memset to lower to an OpenMP element map instead of memcpy/memset.

#############################################
# GPU (CUDA/HIP) compiler
cuda:
Expand Down
1 change: 1 addition & 0 deletions dace/libraries/standard/environments/__init__.py
Original file line number Diff line number Diff line change
@@ -1,2 +1,3 @@
# Copyright 2019-2026 ETH Zurich and the DaCe authors. All rights reserved.
from .cpu import CPU
from .cuda import CUDA
23 changes: 23 additions & 0 deletions dace/libraries/standard/environments/cpu.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,23 @@
# Copyright 2019-2026 ETH Zurich and the DaCe authors. All rights reserved.
"""DaCe library environment exposing the C++ standard headers used by CPU-side libnode expansions."""
import dace.library


@dace.library.environment
class CPU:
"""Minimal library environment that pulls in ``<cstring>`` for plain CPU expansions."""

cmake_minimum_version = None
cmake_packages = []
cmake_variables = {}
cmake_includes = []
cmake_libraries = []
cmake_compile_flags = []
cmake_link_flags = []
cmake_files = []

headers = {'frame': ["cstring"]}
state_fields = []
init_code = ""
finalize_code = ""
dependencies = []
72 changes: 72 additions & 0 deletions dace/libraries/standard/helper.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,72 @@
# Copyright 2019-2026 ETH Zurich and the DaCe authors. All rights reserved.
"""Shared helpers for CopyLibraryNode and MemsetLibraryNode expansions."""
from typing import Callable, List, Tuple

import dace
from dace import dtypes
from dace.sdfg import nodes

# Both legacy and experimental codegens consume this exact name for stream wiring.
CURRENT_STREAM_NAME = "__dace_current_stream"

# Register is intentionally in neither set: resolves by scope (GPU register vs. host stack slot).
GPU_RESIDENT_STORAGES = frozenset({
dtypes.StorageType.GPU_Global,
dtypes.StorageType.GPU_Shared,
})
CPU_RESIDENT_STORAGES = frozenset({
dtypes.StorageType.CPU_Heap,
dtypes.StorageType.CPU_Pinned,
dtypes.StorageType.CPU_ThreadLocal,
})


def collapse_shape_and_strides(
subset: dace.subsets.Range,
strides: List[dace.symbolic.SymExpr]) -> Tuple[List[dace.symbolic.SymExpr], List[dace.symbolic.SymExpr]]:
"""Drop length-1 dims from a (subset, strides) pair; surviving strides scale by the subset step.

:param subset: The access range, one ``(begin, end, step)`` per dimension.
:param strides: The parent array strides, aligned with ``subset``.
:returns: ``(collapsed_shape, collapsed_strides)`` with singletons removed.
"""
collapsed_shape = []
collapsed_strides = []
for (b, e, s), stride in zip(subset, strides):
length = (e + 1 - b) // s
if length != 1:
collapsed_shape.append(length)
collapsed_strides.append(stride * s)
return collapsed_shape, collapsed_strides


def is_parallel_cpu_transfer_size(num_elements: dace.symbolic.SymbolicType) -> bool:
"""False only when ``num_elements`` is a compile-time constant below
``compiler.cpu.parallel_transfer_min_elements``; a symbolic (unknown-at-compile-time) size
is assumed large and takes the parallel path too.

:param num_elements: total contiguous element count (constant or symbolic).
:returns: ``True`` to route to the mapped expansion, ``False`` to keep the single libc call.
"""
threshold = int(dace.Config.get('compiler', 'cpu', 'parallel_transfer_min_elements'))
try:
return int(dace.symbolic.simplify(num_elements)) >= threshold
except (TypeError, ValueError):
return True


def auto_dispatch(node: nodes.LibraryNode, parent_state: dace.SDFGState,
select_fn: Callable[[nodes.LibraryNode, dace.SDFGState], str], library_cls: type):
"""Dispatch a library node's ``'Auto'`` implementation to the one ``select_fn`` picks, setting
``node.implementation`` so introspection reflects what was chosen.

:param node: the library node being expanded.
:param parent_state: state containing ``node`` (owning SDFG is ``parent_state.sdfg``).
:param select_fn: callable returning a concrete implementation name (not ``'Auto'``).
:param library_cls: the library node class with the ``implementations`` dict.
:returns: whatever the resolved expansion returns.
"""
impl_name = select_fn(node, parent_state)
assert impl_name != 'Auto', f"{select_fn.__name__} must not return 'Auto'."
node.implementation = impl_name
return library_cls.implementations[impl_name].expansion(node, parent_state, parent_state.sdfg)
2 changes: 2 additions & 0 deletions dace/libraries/standard/nodes/__init__.py
Original file line number Diff line number Diff line change
@@ -1,4 +1,6 @@
# Copyright 2019-2026 ETH Zurich and the DaCe authors. All rights reserved.
from .code import CodeLibraryNode
from .copy_node import CopyLibraryNode
from .memset_node import MemsetLibraryNode
from .gearbox import Gearbox
from .reduce import Reduce
Loading