Add an explicit graph node API - #64
Open
davschneller wants to merge 4 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
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.
No description provided.