diff --git a/AbstractAPI.h b/AbstractAPI.h
index bde98ee..cef14d6 100644
--- a/AbstractAPI.h
+++ b/AbstractAPI.h
@@ -86,8 +86,46 @@ struct AbstractAPI {
 
   virtual bool isCapableOfGraphCapturing() = 0;
   virtual DeviceGraphHandle streamBeginCapture(std::vector<void*>& 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;
+
+  /**
+   * 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.
+   *
+   * 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 void graphBeginNode(const DeviceGraphHandle& graphHandle,
+                              const std::vector<DeviceGraphNodeHandle>& 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<DeviceGraphNodeHandle>& dependencies,
+                                     void* streamPtr,
+                                     const std::function<void(void*)>& recorder) {
+    graphBeginNode(graphHandle, dependencies, streamPtr);
+    recorder(streamPtr);
+    return graphEndNode(graphHandle, streamPtr);
+  }
 
   virtual void* createStream(double priority = NAN) = 0;
   virtual void destroyGenericStream(void* streamPtr) = 0;
diff --git a/DataTypes.h b/DataTypes.h
index 9f707cb..5aa723f 100644
--- a/DataTypes.h
+++ b/DataTypes.h
@@ -7,28 +7,64 @@
 
 #include <cstddef>
 #include <limits>
+#include <memory>
 
 namespace device {
-struct DeviceGraphHandle {
-  static const size_t invalidId{std::numeric_limits<size_t>::max()};
 
-  public:
-  explicit DeviceGraphHandle() : graphId(invalidId) {}
-  explicit DeviceGraphHandle(size_t id) : graphId(id) {}
+/**
+ * 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;
 
-  DeviceGraphHandle(const DeviceGraphHandle& other) = default;
-  DeviceGraphHandle& operator=(const DeviceGraphHandle& other) = default;
+/**
+ * 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:
+  DeviceGraphHandle() = default;
+  explicit DeviceGraphHandle(std::shared_ptr<DeviceGraph> graphPtr) : graph(std::move(graphPtr)) {}
 
-  bool isInitialized() const { return graphId != invalidId; }
+  [[nodiscard]] bool isInitialized() const { return static_cast<bool>(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:
+  std::shared_ptr<DeviceGraph> 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<size_t>::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 graphId{invalidId};
+  size_t nodeId{invalidId};
 };
 } // namespace device
 
diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h
index 6233474..a02f76f 100644
--- a/interfaces/cuda/CudaWrappedAPI.h
+++ b/interfaces/cuda/CudaWrappedAPI.h
@@ -83,8 +83,17 @@ class ConcreteAPI : public AbstractAPI {
 
   bool isCapableOfGraphCapturing() override;
   DeviceGraphHandle streamBeginCapture(std::vector<void*>& 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;
+
+  bool isCapableOfGraphNodes() override;
+  DeviceGraphHandle graphCreate() override;
+  void graphBeginNode(const DeviceGraphHandle& graphHandle,
+                      const std::vector<DeviceGraphNodeHandle>& dependencies,
+                      void* streamPtr) override;
+  DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle,
+                                     void* streamPtr) override;
+  void graphInstantiate(const DeviceGraphHandle& graphHandle) override;
 
   void* createStream(double priority) override;
   void destroyGenericStream(void* streamPtr) override;
@@ -127,14 +136,6 @@ class ConcreteAPI : public AbstractAPI {
 
   std::unordered_set<cudaStream_t> genericStreams{};
 
-  struct GraphDetails {
-    cudaGraph_t graph;
-    cudaGraphExec_t instance;
-    std::vector<void*> streamPtrs;
-    bool ready{false};
-  };
-  std::vector<GraphDetails> graphs;
-
   Statistics statistics{};
   std::unordered_map<void*, size_t> memToSizeMap{{nullptr, 0}};
 
diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu
index 263e3a7..c87575c 100644
--- a/interfaces/cuda/Graphs.cu
+++ b/interfaces/cuda/Graphs.cu
@@ -10,24 +10,58 @@
 #include <cassert>
 #include <cuda_runtime_api.h>
 #include <driver_types.h>
-#include <mutex>
+#include <functional>
+#include <memory>
+#include <vector>
 
 using namespace device;
 
-/* This is a wrapped graph capturing CUDA mechanism.
- * Call the following in order to capture a computational graph
- *    streamBeginCapture();              // 1
+/* Two ways of building a compute graph are offered.
  *
- *    // your GPU code here              // 2
+ * 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
  *
- *    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
+ * 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 {
+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<std::vector<cudaGraphNode_t>> nodes;
+
+  // only used by the whole-stream capture path
+  std::vector<void*> 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;
@@ -36,55 +70,146 @@ 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<void*>& 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<DeviceGraph>();
+  graphInstance->streamPtrs = streamPtrs;
 
   APIWRAP(cudaStreamBeginCapture(static_cast<cudaStream_t>(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<cudaStream_t>(graphInstance.streamPtrs[0]),
-                               &(graphInstance.graph)));
+  auto* graphInstance = handle.get();
+  assert(graphInstance != nullptr && "a capture must be started before it can be ended");
+
+  APIWRAP(cudaStreamEndCapture(static_cast<cudaStream_t>(graphInstance->streamPtrs[0]),
+                               &(graphInstance->graph)));
 
   APIWRAP(
-      cudaGraphInstantiate(&(graphInstance.instance), graphInstance.graph, nullptr, nullptr, 0));
+      cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0));
+
+  graphInstance->ready = true;
+#endif
+}
 
-  graphInstance.ready = true;
+DeviceGraphHandle ConcreteAPI::graphCreate() {
+#ifdef DEVICE_USE_GRAPH_CAPTURING
+  auto graphInstance = std::make_shared<DeviceGraph>();
+  APIWRAP(cudaGraphCreate(&(graphInstance->graph), 0));
+  return DeviceGraphHandle(std::move(graphInstance));
+#else
+  return DeviceGraphHandle();
+#endif
+}
 
-  {
-    std::lock_guard guard(apiMutex);
-    graphs[handle.getGraphId()] = graphInstance;
-  }
+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<cudaGraphNode_t> 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<cudaGraphNode_t>(frontier, frontier + frontierSize);
 }
+#endif
+} // namespace
 
-void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) {
+void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle,
+                                 const std::vector<DeviceGraphNodeHandle>& dependencies,
+                                 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()];
+  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<cudaGraphNode_t> 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());
   }
-  APIWRAP(cudaGraphLaunch(graphInstance.instance, reinterpret_cast<cudaStream_t>(streamPtr)));
+
+  APIWRAP(cudaStreamBeginCaptureToGraph(static_cast<cudaStream_t>(streamPtr),
+                                        graphInstance->graph,
+                                        nativeDependencies.data(),
+                                        nullptr,
+                                        nativeDependencies.size(),
+                                        cudaStreamCaptureModeThreadLocal));
+#endif
+}
+
+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");
+
+  auto stream = static_cast<cudaStream_t>(streamPtr);
+  auto produced = captureFrontier(stream);
+
+  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();
+  assert(graphInstance != nullptr && graphInstance->ready &&
+         "a graph must be captured before launching");
+
+  APIWRAP(cudaGraphLaunch(graphInstance->instance, static_cast<cudaStream_t>(streamPtr)));
 #endif
 }
diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp
index c7a5cf4..62c58ef 100644
--- a/interfaces/hip/Graphs.cpp
+++ b/interfaces/hip/Graphs.cpp
@@ -8,23 +8,60 @@
 #include "utils/logger.h"
 
 #include <cassert>
+#include <functional>
+#include <hip/hip_runtime.h>
+#include <hip/hip_runtime_api.h>
+#include <memory>
+#include <vector>
 
 using namespace device;
 
-/* This is a wrapped graph capturing CUDA mechanism.
- * Call the following in order to capture a computational graph
- *    streamBeginCapture();              // 1
+/* Two ways of building a compute graph are offered.
  *
- *    // your GPU code here              // 2
+ * 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
  *
- *    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
+ * 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 {
+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<std::vector<hipGraphNode_t>> nodes;
+
+  // only used by the whole-stream capture path
+  std::vector<void*> 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;
@@ -33,54 +70,137 @@ 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<void*>& 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<DeviceGraph>();
+  graphInstance->streamPtrs = streamPtrs;
 
   APIWRAP(hipStreamBeginCapture(static_cast<hipStream_t>(streamPtrs[0]),
                                 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<hipStream_t>(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<hipStream_t>(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) {
+DeviceGraphHandle ConcreteAPI::graphCreate() {
+#ifdef DEVICE_USE_GRAPH_CAPTURING
+  auto graphInstance = std::make_shared<DeviceGraph>();
+  APIWRAP(hipGraphCreate(&(graphInstance->graph), 0));
+  return DeviceGraphHandle(std::move(graphInstance));
+#else
+  return DeviceGraphHandle();
+#endif
+}
+
+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<hipGraphNode_t> 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<hipGraphNode_t>(frontier, frontier + frontierSize);
+}
+#endif
+} // namespace
+
+void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle,
+                                 const std::vector<DeviceGraphNodeHandle>& dependencies,
+                                 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()];
+  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<hipGraphNode_t> 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());
   }
-  APIWRAP(hipGraphLaunch(graphInstance.instance, reinterpret_cast<hipStream_t>(streamPtr)));
+
+  // the edge-data argument is not supported by HIP and has to stay a nullptr
+  APIWRAP(hipStreamBeginCaptureToGraph(static_cast<hipStream_t>(streamPtr),
+                                       graphInstance->graph,
+                                       nativeDependencies.data(),
+                                       nullptr,
+                                       nativeDependencies.size(),
+                                       hipStreamCaptureModeThreadLocal));
+#endif
+}
+
+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");
+
+  auto stream = static_cast<hipStream_t>(streamPtr);
+  auto produced = captureFrontier(stream);
+
+  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();
+  assert(graphInstance != nullptr && graphInstance->ready &&
+         "a graph must be captured before launching");
+
+  APIWRAP(hipGraphLaunch(graphInstance->instance, static_cast<hipStream_t>(streamPtr)));
 #endif
 }
diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h
index 121b55b..8519f7e 100644
--- a/interfaces/hip/HipWrappedAPI.h
+++ b/interfaces/hip/HipWrappedAPI.h
@@ -82,8 +82,17 @@ class ConcreteAPI : public AbstractAPI {
 
   bool isCapableOfGraphCapturing() override;
   DeviceGraphHandle streamBeginCapture(std::vector<void*>& 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;
+
+  bool isCapableOfGraphNodes() override;
+  DeviceGraphHandle graphCreate() override;
+  void graphBeginNode(const DeviceGraphHandle& graphHandle,
+                      const std::vector<DeviceGraphNodeHandle>& dependencies,
+                      void* streamPtr) override;
+  DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle,
+                                     void* streamPtr) override;
+  void graphInstantiate(const DeviceGraphHandle& graphHandle) override;
 
   void* createStream(double priority) override;
   void destroyGenericStream(void* streamPtr) override;
@@ -123,14 +132,6 @@ class ConcreteAPI : public AbstractAPI {
 
   std::unordered_set<hipStream_t> genericStreams{};
 
-  struct GraphDetails {
-    hipGraph_t graph;
-    hipGraphExec_t instance;
-    std::vector<void*> streamPtrs;
-    bool ready{false};
-  };
-  std::vector<GraphDetails> graphs;
-
   Statistics statistics{};
   std::unordered_map<void*, size_t> 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..9ae553f 100644
--- a/interfaces/sycl/Graphs.cpp
+++ b/interfaces/sycl/Graphs.cpp
@@ -8,24 +8,42 @@
 #include "utils/logger.h"
 
 #include <cassert>
+#include <functional>
+#include <memory>
 #include <vector>
 
 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<sycl::ext::oneapi::experimental::command_graph<
+      sycl::ext::oneapi::experimental::graph_state::executable>>
+      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;
@@ -34,52 +52,70 @@ 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<void*>& streamPtrs) {
-  auto handle = DeviceGraphHandle();
 #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT
   std::vector<sycl::queue> queues;
-
+  queues.reserve(streamPtrs.size());
   for (auto* streamPtr : streamPtrs) {
     queues.emplace_back(*static_cast<sycl::queue*>(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<DeviceGraph>(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<sycl::ext::oneapi::experimental::command_graph<
-      sycl::ext::oneapi::experimental::graph_state::executable>>(graphInstance.graph.finalize());
+  auto* graphInstance = handle.get();
+  assert(graphInstance != nullptr && "a capture must be started before it can be ended");
 
-  graphInstance.ready = true;
+  graphInstance->graph.end_recording();
+  graphInstance->instance = std::optional<sycl::ext::oneapi::experimental::command_graph<
+      sycl::ext::oneapi::experimental::graph_state::executable>>(graphInstance->graph.finalize());
+
+  graphInstance->ready = true;
 #endif
 }
 
-void ConcreteAPI::launchGraph(DeviceGraphHandle graphHandle, void* streamPtr) {
+DeviceGraphHandle ConcreteAPI::graphCreate() { return DeviceGraphHandle(); }
+
+void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle,
+                                 const std::vector<DeviceGraphNodeHandle>& 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();
+}
+
+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
-  assert(graphHandle.isInitialized() && "a graph must be captured before launching");
-  GraphDetails graphInstance = [&]() {
-    std::lock_guard guard(apiMutex);
-    return graphs[graphHandle.getGraphId()];
-  }();
+  auto* graphInstance = graphHandle.get();
+  assert(graphInstance != nullptr && graphInstance->ready &&
+         "a graph must be captured before launching");
+
   static_cast<sycl::queue*>(streamPtr)->submit(
-      [&](sycl::handler& handler) { handler.ext_oneapi_graph(graphInstance.instance.value()); });
+      [&](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..818a941 100644
--- a/interfaces/sycl/SyclWrappedAPI.h
+++ b/interfaces/sycl/SyclWrappedAPI.h
@@ -118,8 +118,17 @@ class ConcreteAPI : public AbstractAPI {
 
   bool isCapableOfGraphCapturing() override;
   DeviceGraphHandle streamBeginCapture(std::vector<void*>& 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;
+
+  bool isCapableOfGraphNodes() override;
+  DeviceGraphHandle graphCreate() override;
+  void graphBeginNode(const DeviceGraphHandle& graphHandle,
+                      const std::vector<DeviceGraphNodeHandle>& dependencies,
+                      void* streamPtr) override;
+  DeviceGraphNodeHandle graphEndNode(const DeviceGraphHandle& graphHandle,
+                                     void* streamPtr) override;
+  void graphInstantiate(const DeviceGraphHandle& graphHandle) override;
 
   void* createStream(double priority) override;
   void destroyGenericStream(void* streamPtr) override;
@@ -159,24 +168,6 @@ class ConcreteAPI : public AbstractAPI {
     return this->currentContext()->memoryToSizeMap;
   }
 
-#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT
-  struct GraphDetails {
-    std::optional<sycl::ext::oneapi::experimental::command_graph<
-        sycl::ext::oneapi::experimental::graph_state::executable>>
-        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<GraphDetails> graphs;
-
   void freeMem(void* devPtr);
 
   void initDevices();