Fix three GPU codegen bugs that silently produce a wrong or uncompilable program - #2480
Merged
Conversation
…return zeros CUB reads a null d_temp_storage as "only report the size and do nothing". The device reduction left three ways to reach that state unchecked: the size query, whose failure leaves the size at zero; the allocation itself; and a zero-byte allocation, which cudaMalloc answers with a null pointer and cudaSuccess. In each case the reduction that follows performs no work and leaves the output as it found it -- on freshly allocated device memory, a clean array of zeros with no error raised anywhere. Check all three, never request zero bytes, and return CUB's status from the reduce helper so the call site can check that too. DACE_GPU_CHECK named its argument a second time to format the message, which would have performed the failing call twice; it now formats from the saved status.
…argument Everything DaCe emits ends up in a dlopen-ed shared library, so PIC is not optional -- but it was carried as an -fPIC in the configurable compiler arguments, which a user is free to replace, and which nvcc needs spelled -Xcompiler=-fPIC. CMAKE_POSITION_INDEPENDENT_CODE covers every language at once and gets each spelling right, so the flag comes out of the defaults.
A GPU-touching callback that is not stream-aware is pinned to the legacy default stream, because that is the only stream its own device work (cupy, a vendor library) can be assumed to use. The state footer skipped that stream when collecting the streams to synchronize, and so did presynchronize_streams, so a state whose sink nodes all sit on it emitted no synchronization at all: the sink node was excluded, fell through to the in-edge scan, and those sources are in the same connected component and so were excluded too. Nothing else covers it. The streams are created StreamNonBlocking, which opts them out of the legacy stream's implicit serialization, and DeviceSynchronize runs only in __dace_exit_cuda at teardown, not when __program_* returns. So neither the following state nor the host reading the result was ordered against the callback's data movement. It has been surviving on cudaFree synchronizing and on async copies to pageable host memory staging; pooled memory or pinned host memory removes both. Rather than special-case the default stream at each site, gpu_stream_expr() renders a _cuda_stream annotation as the C expression naming that stream, and every site that turns the annotation into code goes through it. That also removes three places where a node on the default stream would have crashed codegen formatting 'nullptr' with %d, one that would have emitted streams[nullptr], and the unused __dace_current_stream_id local that forced the integer format. streams_to_sync becomes an insertion-ordered dict: it now holds both stream indices and a string, and a set would order the emitted synchronizations by hash.
…ays an argument unordered_arglist collects the data a subgraph uses but does not contain. For a write leaving through an exit node it looked one hop past the matching outgoing connector and took that memlet's data name. That is enough only while the write leaves through a single exit. The GPU code generator tiles a device map, and the edge one hop out then still names the transient the tasklet wrote, which is already a known descriptor -- so nothing was added. The array, and the strides reachable only through it, were then missing from the kernel signature while the kernel body still referenced them, and the generated CUDA did not compile. Each outgoing edge is now followed to the end of its memlet path, mirroring the incoming case just above it, which has always used memlet_path to reach the outermost source. Only compiling catches this: the argument list of the untiled map is already correct, which is why the existing assertions passed throughout. So the compile moves into a GPU-marked test that runs the result, instead of returning early whenever cupy happens to be absent.
ThrudPrimrose
marked this pull request as ready for review
August 5, 2026 12:06
With optimizer.automatic_simplification off the callback gets a state to itself and the copies around it sit on a created stream, so there is no null-stream copy to anchor on and max() was handed an empty sequence. The callback is the one thing both settings emit alike, and fencing it is what the synchronization is for; the check also now refuses any null-stream work queued after the last synchronization of that stream. Verified against all three simplification settings the matrix runs, on the fix and on main.
…roke A failed call leaves state that makes later ones fail too, and those later errors describe consequences. cscs/GPU showed both: a CUB size query failed with a masked device error, leaving its workspace unsized, and the reduction that then read that workspace reported "invalid argument" -- which is the one the caller was handed, about the workspace rather than the query.
acalotoiu
enabled auto-merge
August 5, 2026 15:39
ThrudPrimrose
added a commit
that referenced
this pull request
Aug 5, 2026
Resolved per hunk, preferring main's design and keeping our additions on top. Taken from main: the CMake configure/command cache and precompiled header (#2453) throughout compiler.py, CMakeLists and the schema; the GPU codegen fixes (#2480), whose state.py hunk is the reviewed evolution of work that originated here; the optimized LiftTrivialIf (#2460); dropping -fPIC from the cpu/hip default args, since CMake supplies it via CMAKE_POSITION_INDEPENDENT_CODE; and the CI install blocks, which are supersets of ours. Kept ours, where we add on top of that design: - cudacommon.cuh -- our DACE_GPU_CHECK_RETURN/_RETURN_VAL need a bool, and report_error() is already main's macro body factored into a function. - reduce.py -- the per-stream CUB scratch pool. Main's single per-node __cub_storage is shared by every stream that runs the node, so concurrent streams would race on one workspace. - CMakeLists static_archive target, config_schema build_mode, setup.py's ordered-set/pygments/numpy pin, gpu-ci's cutensor include/lib wiring, cpu.py's CodeObject import (still used four times here). Two things the merge itself surfaced: - config_schema had build_jobs, precompiled_header and configure_cache defined on BOTH sides; keeping both would have been a duplicate YAML key. Main's definitions are kept, ours dropped. - get_scratch() returned a null pointer for a zero-byte request, because 0 > 0 skips the allocation and cudaMalloc(0) yields null anyway. CUB reads a null workspace as "only report the size", so the reduction became a silent no-op leaving the output untouched -- the same failure #2480 guards on its own path. It now allocates a one-byte floor and reallocates whenever the entry is null. The default/null stream needs no special case: pool_map keys on the stream, so 0 is simply another key. Gates: import clean, config_schema parses, ruff/yapf/pre-commit clean; 124 passed across the conflicted test files, 47 across config/frontend.
philip-paul-mueller
pushed a commit
to philip-paul-mueller/dace
that referenced
this pull request
Aug 7, 2026
…g xdist workers position_independent_code_test (new in spcl#2480) reads build/compile_commands.json, which the production folder mode CI runs under deletes after compiling - the same class as the build-cache test fixes, same fix: pin the development mode with the env var dropped first, since environment variables beat set_temporary. The recurring 'worker crashed / node down' CI failures on operational_intensity_test's tiled_mmm cases are not crashes at all: they are minutes of pure sympy on a slow runner with coverage enabled, and CI's --timeout=300 with --timeout_method=thread cannot cancel a test - on expiry it kills the whole xdist worker. A 900s timeout mark gives the analysis the time it legitimately needs.
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
2 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.
Fixes three GPU code generation bugs that each yield a silently wrong or non-compiling program: unchecked CUB calls in the device reduction let a null workspace return all zeros, the legacy default stream that a stream-unaware callback is pinned to was never synchronized at the end of a state, and an array reached only indirectly through a tiled map's exits was dropped from the kernel signature while its body still referenced it. Position-independent code is also enforced through
CMAKE_POSITION_INDEPENDENT_CODErather than a configurable compiler argument, and every change carries tests that fail on current main.