Repository navigation
Conversation
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.
…PU_Global" This reverts commit d72f949.
…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
force-pushed
the
dace_external_stream
branch
from
October 13, 2026 16:19
f383e83 to
ab24527
Compare
Contributor
Author
|
cscs-ci run default |
edopao
force-pushed
the
dace_noview_nested_sdfg
branch
2 times, most recently
from
October 28, 2026 16:20
b11af05 to
ae9a3cb
Compare
This reverts commit 5bb48ef.
Contributor
Author
|
cscs-ci run default |
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
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.
Adds support for running the DaCe-generated SDFG on a GPU stream provided by the caller, instead of always using DaCe's default stream. This lets the SDFG's GPU work be interleaved/synchronized with an external (non-GT4Py) workload and share that workload's stream memory pool.
make_dace_backend()gains a newexternal_gpu_stream: int | Noneparameter. The stream handle is a pointer, which in general changes from run to run of the application (except for the default stream), therefore it is a runtime argument of the SDFG, not a compile-time value:gt_external_gpu_stream(int64) to every GPU SDFG, viaadd_external_gpu_stream_arg(). The SDFG init code casts it withreinterpret_cast<gpuStream_t>and passes it to DaCe's__dace_gpu_set_all_streams(), which writes it into every entry of__state->gpu_context->streams; the value0is the default stream (nullptr). It is a symbol, not a scalar data container, because DaCe only passes the symbols to the SDFG init functions.DaCeBackend.external_gpu_streamand injected into the loaded program byload_artifact()throughDaCeDecoratedProgram.set_external_gpu_stream(), in the same way as the external workspace. On the first call, it is passed to the SDFG argument vector (0when not set), and applied by the SDFG initialization.DaCeTranslator.async_sdfg_call/make_dace_backend(async_sdfg_call=...)tosync_sdfg_call, inverting the flag so its name matches its effect (sync_sdfg_call=Trueforces a synchronous SDFG call; the defaultFalsepreserves the previous async-by-default behavior).make_sdfg_call_async(): setting all DaCe streams is no longer tied to the async/sync choice — it is now handled unconditionally (for both sync and async calls) by the SDFG init code added inadd_external_gpu_stream_arg().make_sdfg_call_sync()'s synchronization tasklet now loops over every stream tracked in DaCe's GPU context (__state->gpu_context->streams[0..num_streams)) and callsStreamSynchronizeon each one, rather than only the default stream, since an external stream may not be DaCe's stream 0.use_metrics) now synchronize only the SDFG GPU streams, instead of callingDeviceSynchronize(), so that the measured time does not include unrelated work on other streams, e.g. external workload queued by the caller.Switch to the experimental DaCe GPU codegen
With
compiler.cuda.max_concurrent_streams = -1, the legacy DaCe CUDA codegen hard-codesnullptr(the default stream) for kernel launches and copies, and never readsgpu_context->streams. Overriding the stream array therefore had no effect. To make the external stream actually used:set_dace_config()selectscompiler.cuda.implementation = "experimental". The experimental codegen launches kernels and copies on__state->gpu_context->streams[...]also withmax_concurrent_streams = -1, and the SDFG init code runs after the stream array is initialized, so the override takes effect.daceis temporarily pinned to thefix-experimental-cuda-gpu-copy-fallbackbranch ofedopao/dace, which is based on thenew-gpu-codegen-devbranch ofspcl/daceand fixes a bug in the experimental codegen: a GPU-to-GPU copy whose write overlap with another writer could not be decided symbolically (e.g.concat_wherewith dynamic domains) was emitted as a host-side copy between device buffers, which segfaulted.set_dace_config()setscompiler.cuda.synchronize_on_exit = False. Otherwise the experimental codegen synchronizes the GPU stream at the end of every SDFG call, making the call always synchronous. Whensync_sdfg_call=True, the synchronization is still added bymake_sdfg_call_sync(). Note that DaCe still synchronizes the stream before the host reads data computed on GPU.Tests
test_generate_sdfg_external_gpu_streamis parametrized over a program lowered to a GPU kernel and a program lowered to a GPU copy. It checks that the SDFG declares thegt_external_gpu_streamsymbol, that it is in the init signature and applied by the init code, and it generates the code with the GT4Py DaCe configuration to check that kernel launches andMemcpy*Asynccalls take their stream from__state->gpu_context->streams(i.e. nonullptr). This test fails with the legacy codegen.test_decorated_program_passes_external_gpu_stream_argchecks that the stream set on the decorated program is passed as SDFG argument on the first call (0when not set), and that it cannot be changed after the first call.test_decorated_program_skips_external_gpu_stream_arg_on_cpuchecks that the argument is not passed to CPU SDFGs._are_streams_synchronized(): it usedre.match, which only matches at the beginning of the generated file and thus always returnedFalse, and it generated code outside of the GT4Py DaCe configuration context. It now usesre.searchwithindace_context().test_generate_sdfg_async_call_multi_statenow checks that the GPU stream is synchronized after the device-to-host copy of a value that is used on the InterState edge or in a CPU tasklet (this was previously documented as a known wrong behavior, see No Synchronization When Using The "DefaultStream" spcl/dace#2120). The case where the value is only used in the InterState edge condition is marked asxfail(strict=True), because the experimental codegen does not emit the synchronization there.Open points
set_external_gpu_stream()raises an error if called after the first call.test_dace_translation.pyonly generate code, thus therequires_gpumarks were removed from the CUDA/ROCM parametrization so that they also run in CPU-only CI.test_nexton CSCS) is the real validation.dacepin on theedopao/dacefork is temporary, until the copy fix is merged upstream innew-gpu-codegen-dev.xfailabove) — to be reported upstream.nullptr) for kernels in library nodes expanded to nested SDFGs, which have no GPU stream connected; such kernels would not run on the external stream.