From d0e33109c5fa31a30c6a2f86f866c18f331c2d3f Mon Sep 17 00:00:00 2001 From: David Schneller Date: Tue, 1 Sep 2026 11:53:13 +0200 Subject: [PATCH 1/4] refactor: make compute graphs own their handles DeviceGraphHandle was an index into a global std::vector 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 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 --- AbstractAPI.h | 4 +- DataTypes.h | 33 +++++++--- interfaces/cuda/CudaWrappedAPI.h | 12 +--- interfaces/cuda/Graphs.cu | 99 ++++++++++++++++-------------- interfaces/hip/Graphs.cpp | 101 +++++++++++++++++-------------- interfaces/hip/HipWrappedAPI.h | 12 +--- interfaces/sycl/Control.cpp | 2 - interfaces/sycl/Graphs.cpp | 89 +++++++++++++++------------ interfaces/sycl/SyclWrappedAPI.h | 22 +------ 9 files changed, 192 insertions(+), 182 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index bde98ee..c9d2150 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -86,8 +86,8 @@ struct AbstractAPI { virtual bool isCapableOfGraphCapturing() = 0; virtual DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) = 0; - virtual void streamEndCapture(DeviceGraphHandle handle) = 0; - virtual void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) = 0; + virtual void streamEndCapture(const DeviceGraphHandle& handle) = 0; + virtual void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) = 0; virtual void* createStream(double priority = NAN) = 0; virtual void destroyGenericStream(void* streamPtr) = 0; diff --git a/DataTypes.h b/DataTypes.h index 9f707cb..7ac23be 100644 --- a/DataTypes.h +++ b/DataTypes.h @@ -7,28 +7,41 @@ #include #include +#include namespace device { -struct DeviceGraphHandle { - static const size_t invalidId{std::numeric_limits::max()}; +/** + * Backend-specific payload of a compute graph. Only the active interface implementation defines + * this type; every other translation unit sees an incomplete type and reaches the graph through + * DeviceGraphHandle. + */ +struct DeviceGraph; + +/** + * Owning handle to a compute graph. + * + * The backend resources (the graph and its executable instance) are released once the last handle + * pointing to them goes out of scope. A graph that is dropped from a cache therefore also frees + * its device-side resources. + */ +class DeviceGraphHandle { public: - explicit DeviceGraphHandle() : graphId(invalidId) {} - explicit DeviceGraphHandle(size_t id) : graphId(id) {} + DeviceGraphHandle() = default; + explicit DeviceGraphHandle(std::shared_ptr graphPtr) : graph(std::move(graphPtr)) {} - DeviceGraphHandle(const DeviceGraphHandle& other) = default; - DeviceGraphHandle& operator=(const DeviceGraphHandle& other) = default; - - bool isInitialized() const { return graphId != invalidId; } + [[nodiscard]] bool isInitialized() const { return static_cast(graph); } operator bool() const { return isInitialized(); } bool operator!() const { return !isInitialized(); } - size_t getGraphId() { return graphId; } + [[nodiscard]] DeviceGraph* get() const { return graph.get(); } + + void reset() { graph.reset(); } private: - size_t graphId{invalidId}; + std::shared_ptr graph; }; } // namespace device diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index 6233474..c96eb1c 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -83,8 +83,8 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphCapturing() override; DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -127,14 +127,6 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; - struct GraphDetails { - cudaGraph_t graph; - cudaGraphExec_t instance; - std::vector streamPtrs; - bool ready{false}; - }; - std::vector graphs; - Statistics statistics{}; std::unordered_map memToSizeMap{{nullptr, 0}}; diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index 263e3a7..a9461d9 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -10,24 +10,47 @@ #include #include #include -#include +#include +#include using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. +/* This is a wrapped graph capturing mechanism. * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 + * auto graph = streamBeginCapture(streams); // 1 + * // your GPU code here // 2 + * streamEndCapture(graph); // 3 * - * // your GPU code here // 2 - * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 - * - * Once you have a coompute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * Once you have a compute-graph recorded you can invoke it as follows: + * launchGraph(graph, stream); // 1 * */ +namespace device { +struct DeviceGraph { + cudaGraph_t graph{nullptr}; + cudaGraphExec_t instance{nullptr}; + + std::vector streamPtrs; + + bool ready{false}; + + DeviceGraph() = default; + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; + + ~DeviceGraph() { + // deliberately unchecked: the graph may outlive the device context during teardown, and a + // failure here has nothing left to report to + if (instance != nullptr) { + cudaGraphExecDestroy(instance); + } + if (graph != nullptr) { + cudaGraphDestroy(graph); + } + } +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING return true; @@ -37,54 +60,40 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); #ifdef DEVICE_USE_GRAPH_CAPTURING - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{}); - handle = DeviceGraphHandle(graphs.size() - 1); - - GraphDetails& graphInstance = graphs[handle.getGraphId()]; - graphInstance.ready = false; - graphInstance.streamPtrs = streamPtrs; - } + auto graphInstance = std::make_shared(); + graphInstance->streamPtrs = streamPtrs; APIWRAP(cudaStreamBeginCapture(static_cast(streamPtrs[0]), cudaStreamCaptureModeThreadLocal)); + + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[handle.getGraphId()]; - } - APIWRAP(cudaStreamEndCapture(static_cast(graphInstance.streamPtrs[0]), - &(graphInstance.graph))); + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); - APIWRAP( - cudaGraphInstantiate(&(graphInstance.instance), graphInstance.graph, nullptr, nullptr, 0)); + APIWRAP(cudaStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), + &(graphInstance->graph))); - graphInstance.ready = true; + APIWRAP( + cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); - { - std::lock_guard guard(apiMutex); - graphs[handle.getGraphId()] = graphInstance; - } + graphInstance->ready = true; #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[graphHandle.getGraphId()]; - } - APIWRAP(cudaGraphLaunch(graphInstance.instance, reinterpret_cast(streamPtr))); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + + APIWRAP(cudaGraphLaunch(graphInstance->instance, static_cast(streamPtr))); #endif } diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index c7a5cf4..8f8eeb0 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -8,23 +8,49 @@ #include "utils/logger.h" #include +#include +#include +#include +#include using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. +/* This is a wrapped graph capturing mechanism. * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 + * auto graph = streamBeginCapture(streams); // 1 + * // your GPU code here // 2 + * streamEndCapture(graph); // 3 * - * // your GPU code here // 2 - * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 - * - * Once you have a coompute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * Once you have a compute-graph recorded you can invoke it as follows: + * launchGraph(graph, stream); // 1 * */ +namespace device { +struct DeviceGraph { + hipGraph_t graph{nullptr}; + hipGraphExec_t instance{nullptr}; + + std::vector streamPtrs; + + bool ready{false}; + + DeviceGraph() = default; + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; + + ~DeviceGraph() { + // deliberately unchecked: the graph may outlive the device context during teardown, and a + // failure here has nothing left to report to + if (instance != nullptr) { + hipGraphExecDestroy(instance); + } + if (graph != nullptr) { + hipGraphDestroy(graph); + } + } +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING return true; @@ -34,53 +60,40 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); #ifdef DEVICE_USE_GRAPH_CAPTURING - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{}); - handle = DeviceGraphHandle(graphs.size() - 1); - - GraphDetails& graphInstance = graphs[handle.getGraphId()]; - graphInstance.ready = false; - graphInstance.streamPtrs = streamPtrs; - } + auto graphInstance = std::make_shared(); + graphInstance->streamPtrs = streamPtrs; APIWRAP(hipStreamBeginCapture(static_cast(streamPtrs[0]), - hipStreamCaptureModeThreadLocal)); + hipStreamCaptureModeThreadLocal)); + + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[handle.getGraphId()]; - } - APIWRAP(hipStreamEndCapture(static_cast(graphInstance.streamPtrs[0]), - &(graphInstance.graph))); + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); - APIWRAP(hipGraphInstantiate(&(graphInstance.instance), graphInstance.graph, nullptr, nullptr, 0)); + APIWRAP(hipStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), + &(graphInstance->graph))); - graphInstance.ready = true; + APIWRAP( + hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); - { - std::lock_guard guard(apiMutex); - graphs[handle.getGraphId()] = graphInstance; - } + graphInstance->ready = true; #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance{}; - { - std::lock_guard guard(apiMutex); - graphInstance = graphs[graphHandle.getGraphId()]; - } - APIWRAP(hipGraphLaunch(graphInstance.instance, reinterpret_cast(streamPtr))); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + + APIWRAP(hipGraphLaunch(graphInstance->instance, static_cast(streamPtr))); #endif } diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index 121b55b..a0274e3 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -82,8 +82,8 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphCapturing() override; DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -123,14 +123,6 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; - struct GraphDetails { - hipGraph_t graph; - hipGraphExec_t instance; - std::vector streamPtrs; - bool ready{false}; - }; - std::vector graphs; - Statistics statistics{}; std::unordered_map memToSizeMap{{nullptr, 0}}; diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index aa1ada0..f3ec2b3 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -73,8 +73,6 @@ void ConcreteAPI::finalize() { this->availableDevices.clear(); this->availableDevices.shrink_to_fit(); - this->graphs.clear(); - this->m_isFinalized = true; this->deviceInitialized = false; } diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index 1905a6a..b11dbea 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -8,24 +8,42 @@ #include "utils/logger.h" #include +#include #include using namespace device; -/* This is a wrapped graph capturing CUDA mechanism. +/* This is a wrapped graph capturing mechanism. * Call the following in order to capture a computational graph - * streamBeginCapture(); // 1 - * - * // your GPU code here // 2 - * - * streamEndCapture(); // 3 - * auto graph = getGraphInstance(); // 4 + * auto graph = streamBeginCapture(streams); // 1 + * // your GPU code here // 2 + * streamEndCapture(graph); // 3 * * Once you have a compute-graph recorded you can invoke it as follows: - * launchGraph(graph) // 1 - * syncGraph(graph) // 2 + * launchGraph(graph, stream); // 1 * */ +namespace device { +struct DeviceGraph { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + std::optional> + instance; + sycl::ext::oneapi::experimental::command_graph< + sycl::ext::oneapi::experimental::graph_state::modifiable> + graph; + + DeviceGraph(const sycl::context& context, const sycl::device& device) + : graph(context, device) {} +#endif + + bool ready{false}; + + DeviceGraph(const DeviceGraph&) = delete; + DeviceGraph& operator=(const DeviceGraph&) = delete; +}; +} // namespace device + bool ConcreteAPI::isCapableOfGraphCapturing() { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT return true; @@ -35,51 +53,44 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { - auto handle = DeviceGraphHandle(); #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT std::vector queues; - + queues.reserve(streamPtrs.size()); for (auto* streamPtr : streamPtrs) { queues.emplace_back(*static_cast(streamPtr)); } - auto recordingGraph = sycl::ext::oneapi::experimental::command_graph< - sycl::ext::oneapi::experimental::graph_state::modifiable>(queues.at(0).get_context(), - queues.at(0).get_device()); - - { - std::lock_guard guard(apiMutex); - graphs.push_back(GraphDetails{std::nullopt, std::move(recordingGraph), false}); - handle = DeviceGraphHandle(graphs.size() - 1); + auto graphInstance = + std::make_shared(queues.at(0).get_context(), queues.at(0).get_device()); + graphInstance->graph.begin_recording(queues); - GraphDetails& graphInstance = graphs[handle.getGraphId()]; - - graphInstance.graph.begin_recording(queues); - } + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); #endif - return handle; } -void ConcreteAPI::streamEndCapture(DeviceGraphHandle handle) { +void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - std::lock_guard guard(apiMutex); - auto& graphInstance = graphs[handle.getGraphId()]; - graphInstance.graph.end_recording(); - graphInstance.instance = std::optional>(graphInstance.graph.finalize()); + auto* graphInstance = handle.get(); + assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + + graphInstance->graph.end_recording(); + graphInstance->instance = std::optional>(graphInstance->graph.finalize()); - graphInstance.ready = true; + graphInstance->ready = true; #endif } -void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) { +void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - assert(graphHandle.isInitialized() && "a graph must be captured before launching"); - GraphDetails graphInstance = [&]() { - std::lock_guard guard(apiMutex); - return graphs[graphHandle.getGraphId()]; - }(); - static_cast(streamPtr)->submit( - [&](sycl::handler& handler) { handler.ext_oneapi_graph(graphInstance.instance.value()); }); + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && graphInstance->ready && + "a graph must be captured before launching"); + + static_cast(streamPtr)->submit([&](sycl::handler& handler) { + handler.ext_oneapi_graph(graphInstance->instance.value()); + }); #endif } diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index 3f854b4..eaf69a2 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -118,8 +118,8 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphCapturing() override; DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; - void streamEndCapture(DeviceGraphHandle handle) override; - void launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) override; + void streamEndCapture(const DeviceGraphHandle& handle) override; + void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; @@ -159,24 +159,6 @@ class ConcreteAPI : public AbstractAPI { return this->currentContext()->memoryToSizeMap; } -#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT - struct GraphDetails { - std::optional> - instance; - sycl::ext::oneapi::experimental::command_graph< - sycl::ext::oneapi::experimental::graph_state::modifiable> - graph; - bool ready{false}; - }; -#else - struct GraphDetails { - bool ready{false}; - }; -#endif - - std::vector graphs; - void freeMem(void* devPtr); void initDevices(); From c6ba1ea118ea23f981ac0a78e980d9f1c5057864 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Tue, 1 Sep 2026 11:53:44 +0200 Subject: [PATCH 2/4] feat: add an explicit graph node API 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 --- AbstractAPI.h | 25 ++++++++ DataTypes.h | 23 ++++++++ interfaces/cuda/CudaWrappedAPI.h | 8 +++ interfaces/cuda/Graphs.cu | 98 +++++++++++++++++++++++++++++-- interfaces/hip/Graphs.cpp | 99 ++++++++++++++++++++++++++++++-- interfaces/hip/HipWrappedAPI.h | 8 +++ interfaces/sycl/Graphs.cpp | 24 ++++++++ interfaces/sycl/SyclWrappedAPI.h | 8 +++ 8 files changed, 285 insertions(+), 8 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index c9d2150..753451f 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -89,6 +89,31 @@ struct AbstractAPI { virtual void streamEndCapture(const DeviceGraphHandle& handle) = 0; virtual void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) = 0; + /** + * Explicit graph construction. + * + * Instead of recording a whole stream and letting the backend infer the dependency structure + * from events, the caller states the structure directly: every graphAddNode call contributes + * the work recorded by `recorder` and makes it depend on exactly `dependencies`. Fork/join is + * then a property of the graph rather than something that has to be expressed through streams + * and events. + * + * A single graph is built by one thread at a time. `recorder` receives a stream that is only a + * recording vehicle: the stream carries no ordering information beyond the extent of that one + * call, and the same stream may be reused for sibling nodes. + * + * If `recorder` enqueues nothing, the returned handle refers to `dependencies` themselves, so + * an empty recorder is a valid way to express a pure join node. + */ + virtual bool isCapableOfGraphNodes() = 0; + virtual DeviceGraphHandle graphCreate() = 0; + virtual DeviceGraphNodeHandle + graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) = 0; + virtual void graphInstantiate(const DeviceGraphHandle& graphHandle) = 0; + virtual void* createStream(double priority = NAN) = 0; virtual void destroyGenericStream(void* streamPtr) = 0; virtual void syncStreamWithHost(void* streamPtr) = 0; diff --git a/DataTypes.h b/DataTypes.h index 7ac23be..5aa723f 100644 --- a/DataTypes.h +++ b/DataTypes.h @@ -43,6 +43,29 @@ class DeviceGraphHandle { private: std::shared_ptr graph; }; + +/** + * Refers to the set of graph nodes produced by a single AbstractAPI::graphAddNode call. + * + * A node handle is an index into the graph that produced it and stays valid for that graph's + * lifetime. Passing it to a different graph is undefined. + */ +class DeviceGraphNodeHandle { + public: + static const size_t invalidId{std::numeric_limits::max()}; + + DeviceGraphNodeHandle() = default; + explicit DeviceGraphNodeHandle(size_t id) : nodeId(id) {} + + [[nodiscard]] bool isInitialized() const { return nodeId != invalidId; } + + operator bool() const { return isInitialized(); } + + [[nodiscard]] size_t getNodeId() const { return nodeId; } + + private: + size_t nodeId{invalidId}; +}; } // namespace device #endif // SEISSOLDEVICE_DATATYPES_H_ diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index c96eb1c..d69e87b 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -86,6 +86,14 @@ class ConcreteAPI : public AbstractAPI { void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; + void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; void syncStreamWithHost(void* streamPtr) override; diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index a9461d9..4de825d 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -10,19 +10,26 @@ #include #include #include +#include #include #include using namespace device; -/* This is a wrapped graph capturing mechanism. - * Call the following in order to capture a computational graph +/* Two ways of building a compute graph are offered. + * + * Whole-stream capture, for code that only wants to replay a fixed sequence: * auto graph = streamBeginCapture(streams); // 1 * // your GPU code here // 2 * streamEndCapture(graph); // 3 + * launchGraph(graph, stream); // 4 * - * Once you have a compute-graph recorded you can invoke it as follows: - * launchGraph(graph, stream); // 1 + * Explicit node construction, for code that knows its own dependency structure: + * auto graph = graphCreate(); // 1 + * auto a = graphAddNode(graph, {}, stream, recordA); // 2 + * auto b = graphAddNode(graph, {a}, stream, recordB); // 3 + * graphInstantiate(graph); // 4 + * launchGraph(graph, stream); // 5 * */ namespace device { @@ -30,6 +37,10 @@ struct DeviceGraph { cudaGraph_t graph{nullptr}; cudaGraphExec_t instance{nullptr}; + // one entry per graphAddNode call; an entry may hold zero, one or several native nodes + std::vector> nodes; + + // only used by the whole-stream capture path std::vector streamPtrs; bool ready{false}; @@ -59,6 +70,15 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } +bool ConcreteAPI::isCapableOfGraphNodes() { +#ifdef DEVICE_USE_GRAPH_CAPTURING + // requires cudaStreamBeginCaptureToGraph, i.e. CUDA >= 12.3 + return true; +#else + return false; +#endif +} + DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto graphInstance = std::make_shared(); @@ -88,6 +108,76 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #endif } +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto graphInstance = std::make_shared(); + APIWRAP(cudaGraphCreate(&(graphInstance->graph), 0)); + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); +#endif +} + +DeviceGraphNodeHandle + ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); + assert(!graphInstance->ready && "no nodes can be added to an instantiated graph"); + + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& nodes = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); + } + + auto stream = static_cast(streamPtr); + APIWRAP(cudaStreamBeginCaptureToGraph(stream, + graphInstance->graph, + nativeDependencies.data(), + nullptr, + nativeDependencies.size(), + cudaStreamCaptureModeThreadLocal)); + + recorder(streamPtr); + + // the capture frontier is what the next node has to depend on; it has to be read out before + // the capture is ended + cudaStreamCaptureStatus captureStatus{}; + unsigned long long captureId{}; + cudaGraph_t capturedGraph{nullptr}; + const cudaGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + APIWRAP(cudaStreamGetCaptureInfo_v2( + stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); + std::vector produced(frontier, frontier + frontierSize); + + cudaGraph_t endedGraph{nullptr}; + APIWRAP(cudaStreamEndCapture(stream, &endedGraph)); + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else + return DeviceGraphNodeHandle(); +#endif +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + + APIWRAP( + cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); + + graphInstance->ready = true; +#endif +} + void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = graphHandle.get(); diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index 8f8eeb0..8b0183a 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -8,6 +8,7 @@ #include "utils/logger.h" #include +#include #include #include #include @@ -15,14 +16,20 @@ using namespace device; -/* This is a wrapped graph capturing mechanism. - * Call the following in order to capture a computational graph +/* Two ways of building a compute graph are offered. + * + * Whole-stream capture, for code that only wants to replay a fixed sequence: * auto graph = streamBeginCapture(streams); // 1 * // your GPU code here // 2 * streamEndCapture(graph); // 3 + * launchGraph(graph, stream); // 4 * - * Once you have a compute-graph recorded you can invoke it as follows: - * launchGraph(graph, stream); // 1 + * Explicit node construction, for code that knows its own dependency structure: + * auto graph = graphCreate(); // 1 + * auto a = graphAddNode(graph, {}, stream, recordA); // 2 + * auto b = graphAddNode(graph, {a}, stream, recordB); // 3 + * graphInstantiate(graph); // 4 + * launchGraph(graph, stream); // 5 * */ namespace device { @@ -30,6 +37,10 @@ struct DeviceGraph { hipGraph_t graph{nullptr}; hipGraphExec_t instance{nullptr}; + // one entry per graphAddNode call; an entry may hold zero, one or several native nodes + std::vector> nodes; + + // only used by the whole-stream capture path std::vector streamPtrs; bool ready{false}; @@ -59,6 +70,15 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } +bool ConcreteAPI::isCapableOfGraphNodes() { +#ifdef DEVICE_USE_GRAPH_CAPTURING + // requires hipStreamBeginCaptureToGraph, i.e. ROCm >= 6.3 + return true; +#else + return false; +#endif +} + DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto graphInstance = std::make_shared(); @@ -88,6 +108,77 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #endif } +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto graphInstance = std::make_shared(); + APIWRAP(hipGraphCreate(&(graphInstance->graph), 0)); + return DeviceGraphHandle(std::move(graphInstance)); +#else + return DeviceGraphHandle(); +#endif +} + +DeviceGraphNodeHandle + ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); + assert(!graphInstance->ready && "no nodes can be added to an instantiated graph"); + + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& nodes = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); + } + + auto stream = static_cast(streamPtr); + // the edge-data argument is not supported by HIP and has to stay a nullptr + APIWRAP(hipStreamBeginCaptureToGraph(stream, + graphInstance->graph, + nativeDependencies.data(), + nullptr, + nativeDependencies.size(), + hipStreamCaptureModeThreadLocal)); + + recorder(streamPtr); + + // the capture frontier is what the next node has to depend on; it has to be read out before + // the capture is ended + hipStreamCaptureStatus captureStatus{}; + unsigned long long captureId{}; + hipGraph_t capturedGraph{nullptr}; + const hipGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + APIWRAP(hipStreamGetCaptureInfo_v2( + stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); + std::vector produced(frontier, frontier + frontierSize); + + hipGraph_t endedGraph{nullptr}; + APIWRAP(hipStreamEndCapture(stream, &endedGraph)); + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else + return DeviceGraphNodeHandle(); +#endif +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + + APIWRAP( + hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); + + graphInstance->ready = true; +#endif +} + void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = graphHandle.get(); diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index a0274e3..7df0244 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -85,6 +85,14 @@ class ConcreteAPI : public AbstractAPI { void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; + void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; void syncStreamWithHost(void* streamPtr) override; diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index b11dbea..b355641 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -8,6 +8,7 @@ #include "utils/logger.h" #include +#include #include #include @@ -52,6 +53,14 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { #endif } +bool ConcreteAPI::isCapableOfGraphNodes() { + // The oneAPI graph extension does expose an explicit node API, but it takes a sycl::handler + // rather than a queue, so it cannot record the queue-based kernel launches that the rest of + // SeisSol emits. Until those launches are expressed through a sink abstraction, this backend + // stays on whole-queue recording. + return false; +} + DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT std::vector queues; @@ -83,6 +92,21 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #endif } +DeviceGraphHandle ConcreteAPI::graphCreate() { return DeviceGraphHandle(); } + +DeviceGraphNodeHandle + ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) { + logError() << "Explicit graph nodes are not supported by the SYCL backend."; + return DeviceGraphNodeHandle(); +} + +void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { + logError() << "Explicit graph nodes are not supported by the SYCL backend."; +} + void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT auto* graphInstance = graphHandle.get(); diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index eaf69a2..ddbc229 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -121,6 +121,14 @@ class ConcreteAPI : public AbstractAPI { void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; + bool isCapableOfGraphNodes() override; + DeviceGraphHandle graphCreate() override; + DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) override; + void graphInstantiate(const DeviceGraphHandle& graphHandle) override; + void* createStream(double priority) override; void destroyGenericStream(void* streamPtr) override; void syncStreamWithHost(void* streamPtr) override; From bb149f591ce6a921548c192fe395791ad177b056 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Tue, 1 Sep 2026 12:38:57 +0200 Subject: [PATCH 3/4] fix: query the capture frontier through the unversioned CUDA entry point 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 --- AbstractAPI.h | 23 ++++++++--- interfaces/cuda/CudaWrappedAPI.h | 9 +++-- interfaces/cuda/Graphs.cu | 67 +++++++++++++++++++++++--------- interfaces/hip/Graphs.cpp | 62 ++++++++++++++++++----------- interfaces/hip/HipWrappedAPI.h | 9 +++-- interfaces/sycl/Graphs.cpp | 13 ++++--- interfaces/sycl/SyclWrappedAPI.h | 9 +++-- 7 files changed, 129 insertions(+), 63 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index 753451f..cef14d6 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -104,16 +104,29 @@ struct AbstractAPI { * * If `recorder` enqueues nothing, the returned handle refers to `dependencies` themselves, so * an empty recorder is a valid way to express a pure join node. + * + * graphBeginNode and graphEndNode are the same thing split in two, for callers that cannot + * wrap the recorded work in a callback and instead have to leave a node open across code they + * do not control. Only one node per stream may be open at a time. */ virtual bool isCapableOfGraphNodes() = 0; virtual DeviceGraphHandle graphCreate() = 0; - virtual DeviceGraphNodeHandle - graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) = 0; + virtual void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) = 0; + virtual DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) = 0; virtual void graphInstantiate(const DeviceGraphHandle& graphHandle) = 0; + DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr, + const std::function& recorder) { + graphBeginNode(graphHandle, dependencies, streamPtr); + recorder(streamPtr); + return graphEndNode(graphHandle, streamPtr); + } + virtual void* createStream(double priority = NAN) = 0; virtual void destroyGenericStream(void* streamPtr) = 0; virtual void syncStreamWithHost(void* streamPtr) = 0; diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index d69e87b..a02f76f 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -88,10 +88,11 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphNodes() override; DeviceGraphHandle graphCreate() override; - DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index 4de825d..41341a7 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -118,11 +118,45 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { #endif } -DeviceGraphNodeHandle - ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) { +namespace { +#ifdef DEVICE_USE_GRAPH_CAPTURING +/** + * Reads the capture frontier, i.e. the nodes a subsequently captured operation would depend on. + * Has to be called while the capture is still open. + * + * The unversioned name resolves to different signatures depending on the toolkit: up to CUDA + * 12.x it is the six-argument form, from CUDA 13 on it is the one that also reports edge data. + * cudaStreamGetCaptureInfo_v2 is not an option, as CUDA 13 no longer declares it. + */ +std::vector captureFrontier(cudaStream_t stream) { + cudaStreamCaptureStatus captureStatus{}; + unsigned long long captureId{}; + cudaGraph_t capturedGraph{nullptr}; + const cudaGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + +#if CUDART_VERSION >= 13000 + const cudaGraphEdgeData* edgeData{nullptr}; + APIWRAP(cudaStreamGetCaptureInfo(stream, + &captureStatus, + &captureId, + &capturedGraph, + &frontier, + &edgeData, + &frontierSize)); +#else + APIWRAP(cudaStreamGetCaptureInfo( + stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); +#endif + + return std::vector(frontier, frontier + frontierSize); +} +#endif +} // namespace + +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); @@ -135,26 +169,23 @@ DeviceGraphNodeHandle nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); } - auto stream = static_cast(streamPtr); - APIWRAP(cudaStreamBeginCaptureToGraph(stream, + APIWRAP(cudaStreamBeginCaptureToGraph(static_cast(streamPtr), graphInstance->graph, nativeDependencies.data(), nullptr, nativeDependencies.size(), cudaStreamCaptureModeThreadLocal)); +#endif +} - recorder(streamPtr); +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); - // the capture frontier is what the next node has to depend on; it has to be read out before - // the capture is ended - cudaStreamCaptureStatus captureStatus{}; - unsigned long long captureId{}; - cudaGraph_t capturedGraph{nullptr}; - const cudaGraphNode_t* frontier{nullptr}; - size_t frontierSize{0}; - APIWRAP(cudaStreamGetCaptureInfo_v2( - stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); - std::vector produced(frontier, frontier + frontierSize); + auto stream = static_cast(streamPtr); + auto produced = captureFrontier(stream); cudaGraph_t endedGraph{nullptr}; APIWRAP(cudaStreamEndCapture(stream, &endedGraph)); diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index 8b0183a..6c1af39 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -118,11 +118,30 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { #endif } -DeviceGraphNodeHandle - ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) { +namespace { +#ifdef DEVICE_USE_GRAPH_CAPTURING +/** + * Reads the capture frontier, i.e. the nodes a subsequently captured operation would depend on. + * Has to be called while the capture is still open. + */ +std::vector captureFrontier(hipStream_t stream) { + hipStreamCaptureStatus captureStatus{}; + unsigned long long captureId{}; + hipGraph_t capturedGraph{nullptr}; + const hipGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + + APIWRAP(hipStreamGetCaptureInfo_v2( + stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); + + return std::vector(frontier, frontier + frontierSize); +} +#endif +} // namespace + +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a graph must be created before nodes can be added"); @@ -135,27 +154,24 @@ DeviceGraphNodeHandle nativeDependencies.insert(nativeDependencies.end(), nodes.begin(), nodes.end()); } - auto stream = static_cast(streamPtr); // the edge-data argument is not supported by HIP and has to stay a nullptr - APIWRAP(hipStreamBeginCaptureToGraph(stream, - graphInstance->graph, - nativeDependencies.data(), - nullptr, - nativeDependencies.size(), - hipStreamCaptureModeThreadLocal)); + APIWRAP(hipStreamBeginCaptureToGraph(static_cast(streamPtr), + graphInstance->graph, + nativeDependencies.data(), + nullptr, + nativeDependencies.size(), + hipStreamCaptureModeThreadLocal)); +#endif +} - recorder(streamPtr); +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { +#ifdef DEVICE_USE_GRAPH_CAPTURING + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); - // the capture frontier is what the next node has to depend on; it has to be read out before - // the capture is ended - hipStreamCaptureStatus captureStatus{}; - unsigned long long captureId{}; - hipGraph_t capturedGraph{nullptr}; - const hipGraphNode_t* frontier{nullptr}; - size_t frontierSize{0}; - APIWRAP(hipStreamGetCaptureInfo_v2( - stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); - std::vector produced(frontier, frontier + frontierSize); + auto stream = static_cast(streamPtr); + auto produced = captureFrontier(stream); hipGraph_t endedGraph{nullptr}; APIWRAP(hipStreamEndCapture(stream, &endedGraph)); diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index 7df0244..8519f7e 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -87,10 +87,11 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphNodes() override; DeviceGraphHandle graphCreate() override; - DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index b355641..0087cf7 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -94,11 +94,14 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { DeviceGraphHandle ConcreteAPI::graphCreate() { return DeviceGraphHandle(); } -DeviceGraphNodeHandle - ConcreteAPI::graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) { +void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) { + logError() << "Explicit graph nodes are not supported by the SYCL backend."; +} + +DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) { logError() << "Explicit graph nodes are not supported by the SYCL backend."; return DeviceGraphNodeHandle(); } diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index ddbc229..818a941 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -123,10 +123,11 @@ class ConcreteAPI : public AbstractAPI { bool isCapableOfGraphNodes() override; DeviceGraphHandle graphCreate() override; - DeviceGraphNodeHandle graphAddNode(const DeviceGraphHandle& graphHandle, - const std::vector& dependencies, - void* streamPtr, - const std::function& recorder) override; + void graphBeginNode(const DeviceGraphHandle& graphHandle, + const std::vector& dependencies, + void* streamPtr) override; + DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle, + void* streamPtr) override; void graphInstantiate(const DeviceGraphHandle& graphHandle) override; void* createStream(double priority) override; From edd5b6a220430257df27d8925f52292387c4cebf Mon Sep 17 00:00:00 2001 From: David Schneller Date: Tue, 1 Sep 2026 15:15:38 +0200 Subject: [PATCH 4/4] refactor: apply pre-commit --- interfaces/cuda/Graphs.cu | 9 ++------- interfaces/hip/Graphs.cpp | 4 ++-- interfaces/sycl/Graphs.cpp | 8 +++----- 3 files changed, 7 insertions(+), 14 deletions(-) diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index 41341a7..c87575c 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -137,13 +137,8 @@ std::vector captureFrontier(cudaStream_t stream) { #if CUDART_VERSION >= 13000 const cudaGraphEdgeData* edgeData{nullptr}; - APIWRAP(cudaStreamGetCaptureInfo(stream, - &captureStatus, - &captureId, - &capturedGraph, - &frontier, - &edgeData, - &frontierSize)); + APIWRAP(cudaStreamGetCaptureInfo( + stream, &captureStatus, &captureId, &capturedGraph, &frontier, &edgeData, &frontierSize)); #else APIWRAP(cudaStreamGetCaptureInfo( stream, &captureStatus, &captureId, &capturedGraph, &frontier, &frontierSize)); diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index 6c1af39..62c58ef 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -85,7 +85,7 @@ DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs graphInstance->streamPtrs = streamPtrs; APIWRAP(hipStreamBeginCapture(static_cast(streamPtrs[0]), - hipStreamCaptureModeThreadLocal)); + hipStreamCaptureModeThreadLocal)); return DeviceGraphHandle(std::move(graphInstance)); #else @@ -99,7 +99,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { assert(graphInstance != nullptr && "a capture must be started before it can be ended"); APIWRAP(hipStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), - &(graphInstance->graph))); + &(graphInstance->graph))); APIWRAP( hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index 0087cf7..9ae553f 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -34,8 +34,7 @@ struct DeviceGraph { sycl::ext::oneapi::experimental::graph_state::modifiable> graph; - DeviceGraph(const sycl::context& context, const sycl::device& device) - : graph(context, device) {} + DeviceGraph(const sycl::context& context, const sycl::device& device) : graph(context, device) {} #endif bool ready{false}; @@ -116,8 +115,7 @@ void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* stream assert(graphInstance != nullptr && graphInstance->ready && "a graph must be captured before launching"); - static_cast(streamPtr)->submit([&](sycl::handler& handler) { - handler.ext_oneapi_graph(graphInstance->instance.value()); - }); + static_cast(streamPtr)->submit( + [&](sycl::handler& handler) { handler.ext_oneapi_graph(graphInstance->instance.value()); }); #endif }