Add an explicit graph node API - #64
Open
davschneller wants to merge 31 commits into
Open
Conversation
DeviceGraphHandle was an index into a global std::vector<GraphDetails> held by the API
object. Three consequences:
* the vector reallocates on push_back, so a GraphDetails& taken from it is only valid
while the lock is held,
* launchGraph copied the whole GraphDetails (including its std::vector<void*> of
streams) under a global mutex on every single launch, which is one allocation and one
global lock per graph launch, and
* there was no way to release a graph, so anything that dropped a handle leaked both the
graph and its executable instance for the rest of the run.
Turn the handle into a shared_ptr to a backend-defined DeviceGraph instead. Ownership now
follows the handle, so dropping a handle frees the backend resources, and the global
vector and its mutex disappear together with the per-launch copy. The payload type stays
incomplete outside the active backend, which keeps the public header free of CUDA, HIP and
SYCL types.
streamEndCapture and launchGraph take the handle by const reference to avoid refcount
traffic on the hot path.
AI-generated. Model: Opus 5
Fork/join is currently expressed by recording a whole stream and letting the backend infer the structure from events. Outside a capture those events are real work, and inside one the topology has to be rediscovered on every rebuild. Both go away if the caller states the dependency structure directly. graphAddNode records the work of one callback into an existing graph with an explicit dependency set and returns a handle to the nodes it produced. On CUDA and HIP this uses cudaStreamBeginCaptureToGraph / hipStreamBeginCaptureToGraph, so the callback keeps taking a stream and no kernel launch has to change; the capture frontier read back through StreamGetCaptureInfo_v2 becomes the node handle. The stream passed to the callback is only a recording vehicle - it carries no ordering, so sibling nodes may share one. An empty callback is meaningful: the resulting handle refers to its own dependencies, which makes a pure join node a one-liner. The SYCL backend reports isCapableOfGraphNodes() == false for now. Its node API takes a sycl::handler rather than a queue and therefore cannot record queue-based launches; it keeps using whole-queue recording until the kernel launches go through a sink abstraction. AI-generated. Model: Opus 5
cudaStreamGetCaptureInfo_v2 is no longer declared by CUDA 13. The unversioned name is what survives, but its signature moves: up to CUDA 12.x it is the six-argument form, from CUDA 13 on it resolves to the variant that also reports edge data. Pick between the two on CUDART_VERSION. HIP keeps declaring hipStreamGetCaptureInfo_v2 and stays as it is. While the query moves into a helper anyway, split graphAddNode into graphBeginNode and graphEndNode, with graphAddNode as a non-virtual convenience on top. A caller that cannot wrap its work in a callback - because the work is spread over code it does not control - can then leave a node open across that code. AI-generated. Model: Opus 5
The oneAPI graph extension has an explicit node API, but its add() takes a sycl::handler, which does not fit code that submits to a queue - which is all of the kernel code. Record and replay does fit, and the extension is explicit that barriers are the way to express edges there: barrier commands are only allowed in nodes created through the record and replay API, and the explicit API rejects them outright. So a node becomes a recorded segment on a queue. graphBeginNode starts recording that queue if it is not recording yet and submits an empty node that pulls in the dependency events; graphEndNode closes the segment with a barrier whose event stands for the node. That is the same contract the CUDA and HIP backends provide, with events in place of node handles. The closing barrier goes through a command group rather than queue::ext_oneapi_submit_barrier: on an in-order queue the shortcut returns the last recorded event instead of the barrier's own, which is wrong whenever the preceding submission discarded its event. Two consequences of recording rather than capturing. The queues are in order, so submissions inside a node are chained for free - but so are two nodes recorded onto the same queue, which gains a redundant edge. Redundant edges cost concurrency, never correctness, and they cannot close a cycle because the recording order and the dependency order agree. Siblings that are meant to overlap therefore have to be recorded onto different queues. AI-generated. Model: Opus 5
Three suites, all written so that a wrong answer means a wrong dependency rather than a wrong kernel: every case pairs operations that do not commute - a fill after a scale does not give the same array as a scale after a fill - so an edge that never made it into the graph shows up as a specific wrong value rather than as a flake. graphs.cpp checks that capturing records instead of executing, that a captured cross-stream event pair replays as edges, that node dependencies are honoured in a chain and across a fork/join, that a node recording nothing stands for its own dependencies, that a graph survives being launched repeatedly, and that handles own what they point at. The fork/join case puts each branch on its own stream, which is also what exercises multi-queue recording on SYCL. A loop that builds and drops a few hundred graphs guards the growth that unowned handles used to cause. streams.cpp covers event ordering between two streams, re-recording an event that has already been waited upon, completion after synchronizing, host callbacks running in stream order, and async allocations outliving the stream operations that use them. batch_transfer.cpp fills the gap around streamBatchedData, accumulateBatchedData and incrementalAdd, which had no coverage: batched strides and null entries are exactly the kind of thing that goes wrong quietly. Cases that need a capability the backend lacks skip rather than fail. AI-generated. Model: Opus 5
…lable The barrier form of the empty operation was switched off whenever the oneAPI graph extension is present, which leaves a graph-capable build - every DPC++ build - on the fallback: an empty single_task for recording an event and another one, with depends_on, for waiting on it. An empty kernel is exactly the kind of command a runtime is free to drop, and a command that was dropped yields an event that is already complete, which makes the depends_on on it say nothing. Ordering across two streams then holds or does not hold depending on timing, which is what Streams.anEventCanBeRecordedAgain caught: the fourth round read the third round's data, scaled again, while the fill it should have waited for had not run yet. The exclusion is broader than the reason for it. The graph extension rejects barriers only in the explicit API; in record and replay - which is the only mode this backend uses, both for whole-queue capture and for node construction - barriers are the documented way to express an edge, precisely because they rest on events. AI-generated. Model: Opus 5
Both tests recorded one fill and then immediately queued the dependent work, so the producing stream almost always won on its own and a missing dependency only surfaced now and then - the failure that started this was the fourth round of four. Queue enough work on the producing stream that the host is certain to get ahead of it. A dependency that holds still gives the same answer; one that does not now fails every run instead of one in a few, which is the difference between a test and a coin toss. The message spells out the value that a broken dependency produces, since "twice the previous round" is not obvious from the numbers alone. AI-generated. Model: Opus 5
newQueue records every stream it hands out in externalQueues, and syncAllQueuesWithHost walks that list - but deleteQueue freed the queue without removing the entry. Every destroyed stream therefore left a dangling pointer behind, and the next device-wide synchronization dereferenced it. The fault lands in syncDevice, arbitrarily far from the code that destroyed the stream; in the test binary that is after the last test has passed. Nothing hit this before because no test created a stream at all - the existing suites all work on the default one. Destroying a stream the device does not know about is now a warning and a no-op rather than a second free. exists() had the reserved-queue check inverted while it was open: it started from true and cleared on the first mismatch, so it only ever recognized the first reserved queue and reported every later one as foreign. copyToAsync throws on that, so a copy issued on anything but queues[0] from the round-robin buffer would have failed. AI-generated. Model: Opus 5
isStreamWorkDone reported an AdaptiveCpp queue as busy right after it had been synchronized. The wait list it consults holds the events a newly submitted operation would have to depend on, and those entries stay in place once they have been reached - so an empty list means nothing was ever submitted, not that nothing is outstanding. Checking the events themselves gives the answer the function promises, which is the one cudaStreamQuery gives on the other backends. The condition also gains the SYCL_EXT_ACPP_QUEUE_WAIT_LIST spelling that the rest of the interface already tests for. Without it, a build that only defines the newer name falls through to the branch that synchronizes and then reports true - a query that quietly blocks, which is the last thing the scheduling work wants. AI-generated. Model: Opus 5
davschneller
force-pushed
the
davschneller/graph-rework
branch
from
September 4, 2026 23:21
f62ce3c to
fa228b2
Compare
For floating point types, numeric_limits<T>::min() is the smallest positive normal value, so it is larger than every negative input. A max reduction over negative data therefore returned ~1.18e-38 instead of the maximum, and with overrideResult the same value was written as the starting point. lowest() is correct for integer and floating point types alike. AI-generated. Model: Opus 5
The generic atomicUpdate cast the result pointer to unsigned long long and ran an 8-byte compare-and-swap on it. Native specializations existed for Sum on int, float and double only, so max and min on int, unsigned and float took the generic path and read and wrote four bytes past the result, which also let the neighbouring memory decide the comparison. The compare-and-swap now runs on a word of exactly sizeof(T), and the native atomics are selected with if constexpr instead of explicit specializations, so the available set no longer differs between the host and the device pass. atomicMax and atomicMin now also cover the integer cases. AI-generated. Model. Opus 5
cudaDeviceGetStreamPriorityRange and its HIP counterpart return the *lowest* priority first and the *highest* second, and the highest is the numerically smaller of the two. Passing them to mapPercentage as (minval, maxval) made its final clamp - max(min(x, maxval), minval) - collapse to minval for every input, so every stream ended up with the default priority. HIP additionally handed the raw double, defaulting to NAN, to hipStreamCreateWithPriority and left the mapped value unused. mapStreamPriority replaces mapPercentage, which had no other callers, states the direction it maps in, clamps the input rather than the output, and lives in namespace device like the rest of the header. AbstractAPI now writes the convention down: 0 lowest, 1 highest, NAN the runtime default. AI-generated. Model: Opus 5
throw new std::invalid_argument(...) throws an std::invalid_argument*, which no catch clause for const std::exception& matches, so the double initialization ends in std::terminate with the exception object leaked on the way out. AI-generated. Model: Opus 5
compare() returned true for both argument orders once both devices matched PREFERRED_DEVICE_TYPE, and sorting on a comparator that does that is undefined. Both devices are now ranked by the same expression and the ranks are compared, so equal devices compare equal in either direction. The sort is stable now as well, which keeps the device ids the same from run to run when the comparator cannot tell two devices apart. AI-generated. Model: Opus 5
counter was never initialized and getNextQueue() incremented it, so the first call read an indeterminate value. It, getGenericQueue, allQueues, resetIndex, getCapacity and the two fork/join helpers had no callers, and with them gone the pool of six queues behind them has none either - the backend hands out queues through newQueue and shares one default queue. Building the pool also went through QueueWrapper's default constructor, which default-constructs a sycl::queue through the default selector before the real one overwrites it - once per pool entry per device. Queues from newQueue that the caller never destroys are now freed when the device goes away, with the same warning the CUDA backend prints. The class is called DeviceQueues now, since it no longer buffers anything circularly. AI-generated. Model: Opus 5
allocUnifiedMem and prefetchUnifiedMemTo left cudaMemLocation zeroed whenever the caller asked for the current device and the device had no concurrent managed access, and issued the call anyway. Up to CUDA 12 the zeroed id names device 0, so on a node with several GPUs the hint went to the wrong card; from CUDA 13 on the zeroed type is cudaMemLocationTypeInvalid and the call fails, which APIWRAP turns into an abort. The preferred location is now only set where there is one to set, and the prefetch returns early, since it needs concurrent managed access to begin with. The HIP backend passed a literal 1 as the device of the coarse-grain advice, which is not a device id on a single-GPU node. AI-generated. Model: Opus 5
None of the free functions dropped the pointer from memToSizeMap, so the map grew for the lifetime of the process. freeGlobMem also kept the CUmemAllocationProp of a compressed allocation, which leaked it and, worse, left the address registered: the runtime is free to return the same address from a later cudaMalloc, and that allocation would then be released through cuMemUnmap. The properties are stored by value now and dropped with the allocation. freeUnifiedMem never subtracted from allocatedUnifiedMemBytes on any backend, so getCurrentlyOccupiedUnifiedMem only ever grew. Freeing a null pointer is handled up front, which is what the nullptr entry the maps were seeded with stood in for. UsmAllocator allocates through allocUnifiedMem and now frees through freeUnifiedMem rather than freeGlobMem. AI-generated. Model: Opus 5
A reduction over zero elements computed a grid of zero blocks, which the runtime rejects as an invalid launch configuration, and the SYCL algorithms built an nd_range with a global size of zero. Where overrideResult is set, the result is still initialized before the launch is skipped. AI-generated. Model: Opus 5
imemcpy and imemset opened with 16-byte vector accesses on the pointers they were handed. Batched buffers are addressed through a pointer table and an element stride, so an element only lands on a 16-byte boundary when the stride happens to be a multiple of 16, and a vector access below its own alignment faults. The width now steps down to what the addresses actually allow, and the narrower passes cover the rest as before. The two scatter kernels also dereferenced their table entries without checking them, unlike streamBatchedData next to them; accumulateBatchedData and setToValue did the same. AI-generated. Model: Opus 5
createEvent dropped its withTiming argument, and no queue carried sycl::property::queue::enable_profiling, so every timespanEvents call ended in an exception thrown from inside get_profiling_info. The queues take the profiling property when the module is built with ENABLE_PROFILING_MARKERS, which is the switch that already stands for "timings are wanted, the per-submission cost is acceptable". Without it, timespanEvents now says so instead of throwing, and the flag createEvent was given is kept and checked. AI-generated. Model: Opus 5
isCapableOfGraphNodes answered for the graph capturing macro, while the node functions call cudaStreamBeginCaptureToGraph (CUDA 12.3) and hipStreamBeginCaptureToGraph (ROCm 6.3). On an older toolkit that is a compile error rather than a backend that says it cannot do it, so the node API gets its own macro and the capability follows that one. Three smaller things on the way through: capturing an empty list of streams indexed into it, a second instantiation of the same graph leaked the first executable graph, and the graph returned by end capture is now checked against the one that was captured into. The stream list is taken by const reference, since capturing does not change it. AI-generated. Model: Opus 5
streamWaitMemory waited for the value to be equal on SYCL and for it to be at least as large on CUDA and HIP, which are two different things for a counter that a producer keeps incrementing. AbstractAPI now says which one callers get, and what kind of memory the location has to be. The note on graphAddNode promised that sibling nodes may share a stream. That holds where the backend can name node handles; the SYCL graph extension expresses edges through the recorded queue, so two nodes sharing a queue end up ordered there. The note now asks for what all three backends can keep. AI-generated. Model: Opus 5
destroyGenericStream removed the stream from the set when it found it there and destroyed it either way, so a pointer that came from somewhere else - the default stream, or a stream already destroyed - was passed to the runtime regardless. It now warns and returns, which is what the SYCL backend does. finalize sets m_isFinalized, so hasFinalized() answers on all three backends, and clears the set of streams it just destroyed. AI-generated. Model: Opus 5
genericStreams, memToSizeMap, allocationProperties, the statistics and the SYCL list of handed-out queues are all changed from whatever thread happens to call in, and syncDevice walks the last one while another thread may be adding to it. The mutex AbstractAPI already carries now guards them; the SYCL queues get their own, since they sit below the API object. syncAllQueuesWithHost copies the list under the lock and waits outside it, so waiting on one queue does not block creating another. The device id is thread-local, matching the runtime, which keeps its selected device per thread as well - so a worker thread that never called setDevice was working on device 0 whatever the process had selected. setDevice now also records the choice for the process, and a thread without one picks it up the first time it asks. AI-generated. Model: Opus 5
Every wrapped API call built a std::string from __FILE__ - a heap allocation for any path longer than the small-string buffer - and an empty std::unordered_set, before finding out that the call had succeeded. Both parameters are now a pointer to the string literal and a view of the caller's temporary array, so the success path allocates nothing. The kernel launch path runs through this a few thousand times per time step. AI-generated. Model: Opus 5
blockcount ran cudaGetDevice, a device attribute query and an occupancy query in front of every algorithm launch, for a number that does not change while the program runs. It is now computed once per kernel, and clamped to at least one block, since a grid of zero is not a valid launch configuration. The SYCL allocations each ended with a device-wide wait. malloc_device and its siblings return once the allocation exists, so there is nothing to wait for. The wait in freeMem stays - the caller has no way to say that no queue is reading the memory any more - and now happens before the free rather than after it. AI-generated. Model: Opus 5
A host function recorded into a graph is called on every replay, so its copy has to outlive the recording. It was allocated and never handed to anyone, which leaked one copy per recorded call for as long as the process ran. The graph that a stream records into now owns those copies and drops them when it is destroyed. Reading the capture status through cudaStreamGetCaptureInfo rather than cudaStreamIsCapturing gives the graph handle to key them on, and the same call already served the node API, so the two share it now. AI-generated. Model: Opus 5
The reductions ran over unsigned values only, where the neutral element of a maximum happens to coincide with the smallest representable value; signed and floating point maxima over negative data are the case that tells a wrong neutral element apart from a right one. The typed suite also puts max and min on 4-byte types under test, which is where an atomic wider than the type writes past the result - visible under compute-sanitizer. The batch tests used 48 floats per element, which is a multiple of 16 bytes, so every element landed on the boundary the copy routines assumed. 47 floats do not. Empty inputs, and streams created with a priority, had no coverage at all. Sibling nodes recorded onto one stream have to come out right whether or not the backend orders them, which is what the graph node contract asks for. AI-generated. Model: Opus 5
The workflow carried a note to move to -Wall -Werror. Doing so catches, among other things, a mapped stream priority that is computed and then not passed on. A memcheck run on the NVIDIA runners covers what the assertions cannot see: an atomic that reaches past its result, or a vector access below its own alignment. It runs on the debug build only, and skips the three large reductions, which take minutes under the sanitizer and exercise no addressing the smaller suites do not. The GitLab pipeline still built the hipsycl backend against an image from three years ago, next to a GitHub workflow that covers the same three backends on current ones. Dropping it as its own commit, so it can be left out of the series if it is still wanted somewhere. The architecture list in the test CMakeLists had a stray space in gfx908 and stopped before the current generations. AI-generated. Model: Opus 5
The SYCL reduction carried a second implementation behind #if 1 / #else. The CUDA and HIP cmake files defined DEVICE_CUDA_LANG and DEVICE_HIP_LANG, while the code reads DEVICE_LANG_CUDA and DEVICE_LANG_HIP, which the top-level file sets; cuda.cmake also repeated the graph capturing definition and the C++ standard. C++17 becomes a public requirement, since the installed headers use it. The rest is a stray semicolon after a function body, an unused device query, and a test fixture flag that was written and never read. AI-generated. Model: Opus 5
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
1 participant
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.
Also fix some bugs.