From d0e33109c5fa31a30c6a2f86f866c18f331c2d3f Mon Sep 17 00:00:00 2001 From: David Schneller Date: Tue, 1 Sep 2026 11:53:13 +0200 Subject: [PATCH 01/31] 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 02/31] 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 03/31] 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 04/31] 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 } From 29a4211f8b0f8f09c95944aed4e783eed4135967 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 00:54:11 +0200 Subject: [PATCH 05/31] feat: support explicit graph nodes in the SYCL backend The oneAPI graph extension has an explicit node API, but its add() takes a sycl::handler, which does not fit code that submits to a queue - which is all of the kernel code. Record and replay does fit, and the extension is explicit that barriers are the way to express edges there: barrier commands are only allowed in nodes created through the record and replay API, and the explicit API rejects them outright. So a node becomes a recorded segment on a queue. graphBeginNode starts recording that queue if it is not recording yet and submits an empty node that pulls in the dependency events; graphEndNode closes the segment with a barrier whose event stands for the node. That is the same contract the CUDA and HIP backends provide, with events in place of node handles. The closing barrier goes through a command group rather than queue::ext_oneapi_submit_barrier: on an in-order queue the shortcut returns the last recorded event instead of the barrier's own, which is wrong whenever the preceding submission discarded its event. Two consequences of recording rather than capturing. The queues are in order, so submissions inside a node are chained for free - but so are two nodes recorded onto the same queue, which gains a redundant edge. Redundant edges cost concurrency, never correctness, and they cannot close a cycle because the recording order and the dependency order agree. Siblings that are meant to overlap therefore have to be recorded onto different queues. AI-generated. Model: Opus 5 --- interfaces/sycl/Graphs.cpp | 97 ++++++++++++++++++++++++++++++++------ 1 file changed, 82 insertions(+), 15 deletions(-) diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index 9ae553f..d14fc49 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -7,6 +7,7 @@ #include "SyclWrappedAPI.h" #include "utils/logger.h" +#include #include #include #include @@ -14,14 +15,16 @@ using namespace device; -/* This is a wrapped graph capturing mechanism. - * Call the following in order to capture a computational graph - * auto graph = streamBeginCapture(streams); // 1 - * // your GPU code here // 2 - * streamEndCapture(graph); // 3 +/* Two ways of building a compute graph are offered; see the CUDA backend for the shapes. * - * Once you have a compute-graph recorded you can invoke it as follows: - * launchGraph(graph, stream); // 1 + * Both rest on the same mechanism here: the oneAPI graph extension records queues. Whole-queue + * recording captures everything submitted between begin and end. Node construction records one + * segment per node and expresses the edges through barriers on the recorded events, which is + * what the extension offers in place of the node handles the CUDA and HIP backends hand out. + * + * The queues record in order, so submissions within one node are chained automatically, and two + * nodes recorded onto the same queue end up ordered even without an edge between them. Siblings + * that are meant to run concurrently therefore have to be recorded onto different queues. * */ namespace device { @@ -34,6 +37,13 @@ struct DeviceGraph { sycl::ext::oneapi::experimental::graph_state::modifiable> graph; + // one entry per graphEndNode call; SYCL expresses graph edges through the events of recorded + // submissions rather than through node objects + std::vector> nodes; + + // queues that recording has been started on, so that it is only started once per queue + std::vector recordedQueues; + DeviceGraph(const sycl::context& context, const sycl::device& device) : graph(context, device) {} #endif @@ -53,11 +63,11 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } 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. +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + return true; +#else return false; +#endif } DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { @@ -91,22 +101,79 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #endif } -DeviceGraphHandle ConcreteAPI::graphCreate() { return DeviceGraphHandle(); } +DeviceGraphHandle ConcreteAPI::graphCreate() { +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto& queue = this->currentDefaultQueue(); + return DeviceGraphHandle(std::make_shared(queue.get_context(), queue.get_device())); +#else + return DeviceGraphHandle(); +#endif +} void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, const std::vector& dependencies, void* streamPtr) { - logError() << "Explicit graph nodes are not supported by the SYCL backend."; +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + 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"); + + auto* queue = static_cast(streamPtr); + auto& recorded = graphInstance->recordedQueues; + if (std::find(recorded.begin(), recorded.end(), queue) == recorded.end()) { + graphInstance->graph.begin_recording(*queue); + recorded.push_back(queue); + } + + std::vector nativeDependencies; + for (const auto& dependency : dependencies) { + assert(dependency.isInitialized() && "an uninitialized node cannot be depended upon"); + const auto& events = graphInstance->nodes.at(dependency.getNodeId()); + nativeDependencies.insert(nativeDependencies.end(), events.begin(), events.end()); + } + + if (!nativeDependencies.empty()) { + // an empty node that pulls the dependencies onto this queue. Everything the caller records + // next follows it, because the queues are in order. + queue->submit([&nativeDependencies](sycl::handler& handler) { + handler.ext_oneapi_barrier(nativeDependencies); + }); + } +#endif } DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, void* streamPtr) { - logError() << "Explicit graph nodes are not supported by the SYCL backend."; +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a node must be opened before it can be closed"); + + auto* queue = static_cast(streamPtr); + + // Closing through a command group rather than through queue::ext_oneapi_submit_barrier: on an + // in-order queue the shortcut hands back the last recorded event, which is not the barrier's + // own event if the preceding submission discarded its event. + std::vector produced{ + queue->submit([](sycl::handler& handler) { handler.ext_oneapi_barrier(); })}; + + graphInstance->nodes.emplace_back(std::move(produced)); + return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); +#else return DeviceGraphNodeHandle(); +#endif } void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { - logError() << "Explicit graph nodes are not supported by the SYCL backend."; +#ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + auto* graphInstance = graphHandle.get(); + assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + + graphInstance->graph.end_recording(); + graphInstance->instance = std::optional>(graphInstance->graph.finalize()); + + graphInstance->ready = true; +#endif } void ConcreteAPI::launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) { From db6e6679b1c8a633197e9cb0681792b12951f9a9 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 00:54:46 +0200 Subject: [PATCH 06/31] test: cover streams, events, graph capture and graph nodes Three suites, all written so that a wrong answer means a wrong dependency rather than a wrong kernel: every case pairs operations that do not commute - a fill after a scale does not give the same array as a scale after a fill - so an edge that never made it into the graph shows up as a specific wrong value rather than as a flake. graphs.cpp checks that capturing records instead of executing, that a captured cross-stream event pair replays as edges, that node dependencies are honoured in a chain and across a fork/join, that a node recording nothing stands for its own dependencies, that a graph survives being launched repeatedly, and that handles own what they point at. The fork/join case puts each branch on its own stream, which is also what exercises multi-queue recording on SYCL. A loop that builds and drops a few hundred graphs guards the growth that unowned handles used to cause. streams.cpp covers event ordering between two streams, re-recording an event that has already been waited upon, completion after synchronizing, host callbacks running in stream order, and async allocations outliving the stream operations that use them. batch_transfer.cpp fills the gap around streamBatchedData, accumulateBatchedData and incrementalAdd, which had no coverage: batched strides and null entries are exactly the kind of thing that goes wrong quietly. Cases that need a capability the backend lacks skip rather than fail. AI-generated. Model: Opus 5 --- tests/CMakeLists.txt | 10 +- tests/batch_transfer.cpp | 168 ++++++++++++++++++++ tests/graphs.cpp | 331 +++++++++++++++++++++++++++++++++++++++ tests/streams.cpp | 161 +++++++++++++++++++ 4 files changed, 669 insertions(+), 1 deletion(-) create mode 100644 tests/batch_transfer.cpp create mode 100644 tests/graphs.cpp create mode 100644 tests/streams.cpp diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 7369f47..07afd5d 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -22,7 +22,15 @@ add_subdirectory(.. root) find_package(GTest REQUIRED) -add_executable(tests main.cpp reductions.cpp memory.cpp array_manip.cpp batch_manip.cpp) +add_executable(tests + main.cpp + reductions.cpp + memory.cpp + array_manip.cpp + batch_manip.cpp + batch_transfer.cpp + streams.cpp + graphs.cpp) target_link_libraries(tests PRIVATE device ${GTEST_BOTH_LIBRARIES}) target_include_directories(tests PRIVATE ${GTEST_INCLUDE_DIR}) diff --git a/tests/batch_transfer.cpp b/tests/batch_transfer.cpp new file mode 100644 index 0000000..4da1c3f --- /dev/null +++ b/tests/batch_transfer.cpp @@ -0,0 +1,168 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t BatchSize = 64; +constexpr std::size_t ElementSize = 48; +} // namespace + +/** + * Covers the batched transfers and the pointer arithmetic helper, which the other suites do not + * touch. These are the operations the solver uses to move data in and out of batched buffers, so + * a wrong stride or a skipped entry shows up as a wrong answer far away from here. + */ +class BatchTransfer : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + stream = device->api->createStream(); + src = static_cast(device->api->allocGlobMem(BatchSize * ElementSize * sizeof(float))); + dst = static_cast(device->api->allocGlobMem(BatchSize * ElementSize * sizeof(float))); + srcBatch = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + dstBatch = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + for (std::size_t i = 0; i < BatchSize; ++i) { + srcBatch[i] = src + i * ElementSize; + dstBatch[i] = dst + i * ElementSize; + } + } + + void TearDown() override { + device->api->freeUnifiedMem(dstBatch); + device->api->freeUnifiedMem(srcBatch); + device->api->freeGlobMem(dst); + device->api->freeGlobMem(src); + device->api->destroyGenericStream(stream); + } + + protected: + void upload(float* target, const std::vector& host) { + device->api->copyToAsync(target, host.data(), host.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + } + + std::vector download(const float* source) { + std::vector host(BatchSize * ElementSize, -1); + device->api->copyFromAsync(host.data(), source, host.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + return host; + } + + // a value that differs per batch entry and per element, so a mixed-up stride cannot pass + static float pattern(std::size_t entry, std::size_t element) { + return static_cast(entry * ElementSize + element); + } + + void* stream{nullptr}; + float* src{nullptr}; + float* dst{nullptr}; + float** srcBatch{nullptr}; + float** dstBatch{nullptr}; +}; + +TEST_F(BatchTransfer, streamBatchedDataCopiesEveryEntry) { + std::vector hostSrc(BatchSize * ElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + hostSrc[i * ElementSize + j] = pattern(i, j); + } + } + upload(src, hostSrc); + upload(dst, std::vector(BatchSize * ElementSize, 0.0F)); + + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(pattern(i, j), hostDst[i * ElementSize + j]) + << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, streamBatchedDataSkipsNullEntries) { + upload(src, std::vector(BatchSize * ElementSize, 1.0F)); + upload(dst, std::vector(BatchSize * ElementSize, -3.0F)); + + for (std::size_t i = 0; i < BatchSize; i += 2) { + srcBatch[i] = nullptr; + } + + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + const float expected = (i % 2 == 0) ? -3.0F : 1.0F; + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(expected, hostDst[i * ElementSize + j]) << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, accumulateBatchedDataAdds) { + std::vector hostSrc(BatchSize * ElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + hostSrc[i * ElementSize + j] = pattern(i, j); + } + } + upload(src, hostSrc); + upload(dst, std::vector(BatchSize * ElementSize, 5.0F)); + + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + const auto hostDst = download(dst); + for (std::size_t i = 0; i < BatchSize; ++i) { + for (std::size_t j = 0; j < ElementSize; ++j) { + ASSERT_EQ(5.0F + pattern(i, j), hostDst[i * ElementSize + j]) + << "at entry " << i << ", element " << j; + } + } +} + +TEST_F(BatchTransfer, accumulateBatchedDataIsRepeatable) { + upload(src, std::vector(BatchSize * ElementSize, 2.0F)); + upload(dst, std::vector(BatchSize * ElementSize, 0.0F)); + + for (int round = 0; round < 3; ++round) { + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, BatchSize, stream); + } + device->api->syncStreamWithHost(stream); + + for (const auto value : download(dst)) { + ASSERT_EQ(6.0F, value); + } +} + +TEST_F(BatchTransfer, incrementalAddBuildsAStridedPointerTable) { + auto** table = static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + device->algorithms.incrementalAdd(table, src, ElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + // the stride is given in elements, not bytes + for (std::size_t i = 0; i < BatchSize; ++i) { + ASSERT_EQ(src + i * ElementSize, table[i]) << "at entry " << i; + } + + device->api->freeUnifiedMem(table); +} diff --git a/tests/graphs.cpp b/tests/graphs.cpp new file mode 100644 index 0000000..18e2032 --- /dev/null +++ b/tests/graphs.cpp @@ -0,0 +1,331 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t ArraySize = 1 << 14; +constexpr std::size_t BranchCount = 4; +constexpr std::size_t ChunkSize = ArraySize / BranchCount; +} // namespace + +/** + * The tests below check what a graph guarantees, not how it is built: every one of them states a + * dependency structure and then asserts a result that only comes out right if that structure was + * honoured. Operations are picked so that swapping two of them changes the answer - a fill after a + * scale does not give the same array as a scale after a fill - because operations that commute + * would pass no matter how the edges came out. + */ +class Graphs : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + devArray = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + mainStream = device->api->createStream(); + for (auto& stream : branchStreams) { + stream = device->api->createStream(); + } + } + + void TearDown() override { + for (auto* stream : branchStreams) { + device->api->destroyGenericStream(stream); + } + device->api->destroyGenericStream(mainStream); + device->api->freeGlobMem(devArray); + } + + protected: + std::vector download() { + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), devArray, ArraySize * sizeof(float), mainStream); + device->api->syncStreamWithHost(mainStream); + return host; + } + + void fill(float value) { + device->algorithms.fillArray(devArray, value, ArraySize, mainStream); + device->api->syncStreamWithHost(mainStream); + } + + static void expectChunk(const std::vector& host, std::size_t chunk, float value) { + for (std::size_t i = chunk * ChunkSize; i < (chunk + 1) * ChunkSize; ++i) { + ASSERT_EQ(value, host[i]) << "at index " << i << " of chunk " << chunk; + } + } + + bool graphCapturingUnavailable() { return !device->api->isCapableOfGraphCapturing(); } + + bool graphNodesUnavailable() { return !device->api->isCapableOfGraphNodes(); } + + float* devArray{nullptr}; + void* mainStream{nullptr}; + std::array branchStreams{}; +}; + +TEST_F(Graphs, captureReplaysTheRecordedSequence) { + if (graphCapturingUnavailable()) { + GTEST_SKIP() << "the backend does not support graph capturing"; + } + + fill(0); + + std::vector streams{mainStream}; + auto graph = device->api->streamBeginCapture(streams); + device->algorithms.fillArray(devArray, 1.0F, ArraySize, mainStream); + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, mainStream); + device->api->streamEndCapture(graph); + + ASSERT_TRUE(graph.isInitialized()); + + // nothing has run yet: capturing records, it does not execute + for (const auto value : download()) { + ASSERT_EQ(0.0F, value); + } + + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(2.0F, value); + } + + // the fill is part of the graph, so a second replay lands on the same value rather than doubling + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(2.0F, value); + } +} + +TEST_F(Graphs, captureRecordsCrossStreamEvents) { + if (graphCapturingUnavailable()) { + GTEST_SKIP() << "the backend does not support graph capturing"; + } + + fill(0); + + // the fork/join shape that the stream path uses: an event hands work from the recorded stream + // to a side stream and back. Inside a capture these become graph edges rather than real waits. + auto* forkEvent = device->api->createEvent(); + auto* joinEvent = device->api->createEvent(); + + std::vector streams{mainStream, branchStreams[0]}; + auto graph = device->api->streamBeginCapture(streams); + + device->algorithms.fillArray(devArray, 3.0F, ArraySize, mainStream); + device->api->recordEventOnStream(forkEvent, mainStream); + device->api->syncStreamWithEvent(branchStreams[0], forkEvent); + device->algorithms.scaleArray(devArray, 4.0F, ArraySize, branchStreams[0]); + device->api->recordEventOnStream(joinEvent, branchStreams[0]); + device->api->syncStreamWithEvent(mainStream, joinEvent); + + device->api->streamEndCapture(graph); + ASSERT_TRUE(graph.isInitialized()); + + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(12.0F, value); + } + + device->api->destroyEvent(joinEvent); + device->api->destroyEvent(forkEvent); +} + +TEST_F(Graphs, nodesRunInDependencyOrder) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(0); + + auto graph = device->api->graphCreate(); + ASSERT_TRUE(graph.isInitialized()); + + const auto first = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 1.0F, ArraySize, stream); + }); + const auto second = device->api->graphAddNode(graph, {first}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 3.0F, ArraySize, stream); + }); + const auto third = device->api->graphAddNode(graph, {second}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 5.0F, ArraySize, stream); + }); + ASSERT_TRUE(third.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + // the fill has to come first: if it ran last, every entry would be 1 instead + for (const auto value : download()) { + ASSERT_EQ(15.0F, value); + } +} + +TEST_F(Graphs, nodesForkAndJoin) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + // each branch owns a disjoint chunk and runs on its own stream, so they may overlap + std::vector branches; + for (std::size_t i = 0; i < BranchCount; ++i) { + branches.push_back( + device->api->graphAddNode(graph, {root}, branchStreams[i], [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + const auto join = device->api->graphAddNode(graph, branches, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 10.0F, ArraySize, stream); + }); + ASSERT_TRUE(join.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + // a chunk holding i+1 means the join overtook its branch; a chunk holding 0 means the root + // overtook it + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 10.0F * static_cast(i + 1)); + } +} + +TEST_F(Graphs, anEmptyNodeJoinsItsDependencies) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + std::vector branches; + for (std::size_t i = 0; i < BranchCount; ++i) { + branches.push_back( + device->api->graphAddNode(graph, {root}, branchStreams[i], [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + // a node that records nothing stands for its own dependencies, which is what makes it usable + // as a join without costing a command + const auto join = device->api->graphAddNode(graph, branches, mainStream, [](void*) {}); + + const auto last = device->api->graphAddNode(graph, {join}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 100.0F, ArraySize, stream); + }); + ASSERT_TRUE(last.isInitialized()); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 100.0F * static_cast(i + 1)); + } +} + +TEST_F(Graphs, aNodeGraphCanBeLaunchedRepeatedly) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(1); + + // no fill inside the graph, so repeated launches accumulate and a graph that silently ran only + // once would be caught + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, stream); + }); + device->api->graphInstantiate(graph); + + for (int i = 0; i < 3; ++i) { + device->api->launchGraph(graph, mainStream); + } + device->api->syncStreamWithHost(mainStream); + + for (const auto value : download()) { + ASSERT_EQ(8.0F, value); + } +} + +TEST_F(Graphs, handlesOwnTheirGraph) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + DeviceGraphHandle empty; + EXPECT_FALSE(empty.isInitialized()); + EXPECT_TRUE(!empty); + + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 5.0F, ArraySize, stream); + }); + device->api->graphInstantiate(graph); + + auto copy = graph; + EXPECT_TRUE(copy.isInitialized()); + graph.reset(); + EXPECT_FALSE(graph.isInitialized()); + + // the graph is still alive through the second handle + device->api->launchGraph(copy, mainStream); + device->api->syncStreamWithHost(mainStream); + for (const auto value : download()) { + ASSERT_EQ(5.0F, value); + } +} + +TEST_F(Graphs, droppedGraphsReleaseTheirResources) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + // Graphs used to be held in a container that never gave anything back, so a workload that keys + // its graphs on something that varies - a time step width, say - grew without bound. Building + // and dropping many of them has to stay flat. + for (int i = 0; i < 256; ++i) { + auto graph = device->api->graphCreate(); + device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, static_cast(i), ArraySize, stream); + }); + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + } + device->api->syncStreamWithHost(mainStream); + + for (const auto value : download()) { + ASSERT_EQ(255.0F, value); + } +} diff --git a/tests/streams.cpp b/tests/streams.cpp new file mode 100644 index 0000000..17370a7 --- /dev/null +++ b/tests/streams.cpp @@ -0,0 +1,161 @@ +// SPDX-FileCopyrightText: 2026 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "BaseTestSuite.h" +#include "device.h" + +#include "gtest/gtest.h" +#include +#include +#include + +using namespace device; +using namespace ::testing; + +namespace { +constexpr std::size_t ArraySize = 1 << 14; +} // namespace + +class Streams : public BaseTestSuite { + public: + void SetUp() override { + BaseTestSuite::SetUp(); + devArray = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + streamA = device->api->createStream(); + streamB = device->api->createStream(); + } + + void TearDown() override { + device->api->destroyGenericStream(streamB); + device->api->destroyGenericStream(streamA); + device->api->freeGlobMem(devArray); + } + + protected: + std::vector download(void* stream) { + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), devArray, ArraySize * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + return host; + } + + float* devArray{nullptr}; + void* streamA{nullptr}; + void* streamB{nullptr}; +}; + +TEST_F(Streams, anEventOrdersTwoStreams) { + auto* event = device->api->createEvent(); + + device->algorithms.fillArray(devArray, 2.0F, ArraySize, streamA); + device->api->recordEventOnStream(event, streamA); + + device->api->syncStreamWithEvent(streamB, event); + device->algorithms.scaleArray(devArray, 7.0F, ArraySize, streamB); + + device->api->syncStreamWithHost(streamB); + + // without the event the scale could read the array before the fill wrote it + for (const auto value : download(streamB)) { + ASSERT_EQ(14.0F, value); + } + + device->api->destroyEvent(event); +} + +TEST_F(Streams, anEventCanBeRecordedAgain) { + // the stream runtime hands the same event out repeatedly, so re-recording one that has already + // been waited upon has to keep working + auto* event = device->api->createEvent(); + + for (int round = 1; round <= 4; ++round) { + device->algorithms.fillArray(devArray, static_cast(round), ArraySize, streamA); + device->api->recordEventOnStream(event, streamA); + device->api->syncStreamWithEvent(streamB, event); + device->algorithms.scaleArray(devArray, 2.0F, ArraySize, streamB); + device->api->syncStreamWithHost(streamB); + + for (const auto value : download(streamB)) { + ASSERT_EQ(2.0F * static_cast(round), value) << "in round " << round; + } + } + + device->api->destroyEvent(event); +} + +TEST_F(Streams, anEventIsCompleteOnceItsStreamIs) { + auto* event = device->api->createEvent(); + + device->algorithms.fillArray(devArray, 1.0F, ArraySize, streamA); + device->api->recordEventOnStream(event, streamA); + device->api->syncStreamWithHost(streamA); + + // only the direction after synchronizing is deterministic; whether the event is already + // complete beforehand depends on timing + EXPECT_TRUE(device->api->isEventCompleted(event)); + EXPECT_TRUE(device->api->isStreamWorkDone(streamA)); + + device->api->destroyEvent(event); +} + +TEST_F(Streams, aHostFunctionRunsInStreamOrder) { + std::vector staging(ArraySize, -1); + std::atomic sawFilledData{false}; + std::atomic ran{false}; + + device->algorithms.fillArray(devArray, 9.0F, ArraySize, streamA); + device->api->copyFromAsync(staging.data(), devArray, ArraySize * sizeof(float), streamA); + device->api->streamHostFunction(streamA, [&]() { + ran = true; + sawFilledData = (staging.front() == 9.0F) && (staging.back() == 9.0F); + }); + + device->api->syncStreamWithHost(streamA); + + EXPECT_TRUE(ran.load()); + // the callback must not observe the staging buffer before the copy that precedes it + EXPECT_TRUE(sawFilledData.load()); +} + +TEST_F(Streams, asyncAllocationsLiveOnTheStream) { + auto* scratch = + static_cast(device->api->allocMemAsync(ArraySize * sizeof(float), streamA)); + ASSERT_NE(nullptr, scratch); + + device->algorithms.fillArray(scratch, 4.0F, ArraySize, streamA); + + std::vector host(ArraySize, -1); + device->api->copyFromAsync(host.data(), scratch, ArraySize * sizeof(float), streamA); + + // the free is ordered behind the copy on the same stream + device->api->freeMemAsync(scratch, streamA); + device->api->syncStreamWithHost(streamA); + + for (const auto value : host) { + ASSERT_EQ(4.0F, value); + } +} + +TEST_F(Streams, workOnSeparateStreamsStaysSeparate) { + auto* other = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); + + device->algorithms.fillArray(devArray, 1.0F, ArraySize, streamA); + device->algorithms.fillArray(other, 2.0F, ArraySize, streamB); + + device->api->syncStreamWithHost(streamA); + device->api->syncStreamWithHost(streamB); + + std::vector hostOther(ArraySize, -1); + device->api->copyFromAsync(hostOther.data(), other, ArraySize * sizeof(float), streamB); + device->api->syncStreamWithHost(streamB); + + for (const auto value : download(streamA)) { + ASSERT_EQ(1.0F, value); + } + for (const auto value : hostOther) { + ASSERT_EQ(2.0F, value); + } + + device->api->freeGlobMem(other); +} From 16a3e09ef690e4fec3d796ce8082c6e2fd20dd13 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 01:19:36 +0200 Subject: [PATCH 07/31] fix: keep using enqueue barriers for SYCL events when graphs are available The barrier form of the empty operation was switched off whenever the oneAPI graph extension is present, which leaves a graph-capable build - every DPC++ build - on the fallback: an empty single_task for recording an event and another one, with depends_on, for waiting on it. An empty kernel is exactly the kind of command a runtime is free to drop, and a command that was dropped yields an event that is already complete, which makes the depends_on on it say nothing. Ordering across two streams then holds or does not hold depending on timing, which is what Streams.anEventCanBeRecordedAgain caught: the fourth round read the third round's data, scaled again, while the fill it should have waited for had not run yet. The exclusion is broader than the reason for it. The graph extension rejects barriers only in the explicit API; in record and replay - which is the only mode this backend uses, both for whole-queue capture and for node construction - barriers are the documented way to express an edge, precisely because they rest on events. AI-generated. Model: Opus 5 --- interfaces/sycl/SyclWrappedAPI.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index 818a941..9832376 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -42,7 +42,7 @@ #define DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(handle, event) \ handle.depends_on(event); \ handle.DEVICE_SYCL_DIRECT_OPERATION_NAME([=](...) {}); -#elif defined(SYCL_EXT_ONEAPI_ENQUEUE_BARRIER) && !defined(DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT) +#elif defined(SYCL_EXT_ONEAPI_ENQUEUE_BARRIER) #define DEVICE_SYCL_EMPTY_OPERATION(handle) handle.ext_oneapi_barrier(); #define DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(handle, event) handle.ext_oneapi_barrier({event}); #else From f6b3deaac24bb3d55f2d545cd5df8031599831a2 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 01:20:00 +0200 Subject: [PATCH 08/31] test: make the event ordering tests decide rather than guess Both tests recorded one fill and then immediately queued the dependent work, so the producing stream almost always won on its own and a missing dependency only surfaced now and then - the failure that started this was the fourth round of four. Queue enough work on the producing stream that the host is certain to get ahead of it. A dependency that holds still gives the same answer; one that does not now fails every run instead of one in a few, which is the difference between a test and a coin toss. The message spells out the value that a broken dependency produces, since "twice the previous round" is not obvious from the numbers alone. AI-generated. Model: Opus 5 --- tests/streams.cpp | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/tests/streams.cpp b/tests/streams.cpp index 17370a7..66fdb70 100644 --- a/tests/streams.cpp +++ b/tests/streams.cpp @@ -15,6 +15,11 @@ using namespace ::testing; namespace { constexpr std::size_t ArraySize = 1 << 14; + +// Enough queued work that the host is certain to run ahead of the producing stream. Without it, +// an ordering test passes whenever the producer happens to finish first, so a missing dependency +// shows up as an occasional failure rather than as a verdict. +constexpr int EnqueueDepth = 64; } // namespace class Streams : public BaseTestSuite { @@ -48,7 +53,9 @@ class Streams : public BaseTestSuite { TEST_F(Streams, anEventOrdersTwoStreams) { auto* event = device->api->createEvent(); - device->algorithms.fillArray(devArray, 2.0F, ArraySize, streamA); + for (int i = 0; i < EnqueueDepth; ++i) { + device->algorithms.fillArray(devArray, 2.0F, ArraySize, streamA); + } device->api->recordEventOnStream(event, streamA); device->api->syncStreamWithEvent(streamB, event); @@ -56,7 +63,7 @@ TEST_F(Streams, anEventOrdersTwoStreams) { device->api->syncStreamWithHost(streamB); - // without the event the scale could read the array before the fill wrote it + // a 7 here means the scale read the array before the fills wrote it for (const auto value : download(streamB)) { ASSERT_EQ(14.0F, value); } @@ -70,14 +77,19 @@ TEST_F(Streams, anEventCanBeRecordedAgain) { auto* event = device->api->createEvent(); for (int round = 1; round <= 4; ++round) { - device->algorithms.fillArray(devArray, static_cast(round), ArraySize, streamA); + for (int i = 0; i < EnqueueDepth; ++i) { + device->algorithms.fillArray(devArray, static_cast(round), ArraySize, streamA); + } device->api->recordEventOnStream(event, streamA); device->api->syncStreamWithEvent(streamB, event); device->algorithms.scaleArray(devArray, 2.0F, ArraySize, streamB); device->api->syncStreamWithHost(streamB); + // twice the previous round's value means the scale overtook this round's fills for (const auto value : download(streamB)) { - ASSERT_EQ(2.0F * static_cast(round), value) << "in round " << round; + ASSERT_EQ(2.0F * static_cast(round), value) + << "in round " << round << " (twice the previous round's value would be " + << 4.0F * static_cast(round - 1) << ")"; } } From 273811701e70d0989e6f64ce542a4c17851a3790 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 01:20:36 +0200 Subject: [PATCH 09/31] fix: take destroyed streams off the SYCL device's list newQueue records every stream it hands out in externalQueues, and syncAllQueuesWithHost walks that list - but deleteQueue freed the queue without removing the entry. Every destroyed stream therefore left a dangling pointer behind, and the next device-wide synchronization dereferenced it. The fault lands in syncDevice, arbitrarily far from the code that destroyed the stream; in the test binary that is after the last test has passed. Nothing hit this before because no test created a stream at all - the existing suites all work on the default one. Destroying a stream the device does not know about is now a warning and a no-op rather than a second free. exists() had the reserved-queue check inverted while it was open: it started from true and cleared on the first mismatch, so it only ever recognized the first reserved queue and reported every later one as foreign. copyToAsync throws on that, so a copy issued on anything but queues[0] from the round-robin buffer would have failed. AI-generated. Model: Opus 5 --- interfaces/sycl/DeviceCircularQueueBuffer.cpp | 20 ++++++++++++++++--- tests/streams.cpp | 16 +++++++++++++++ 2 files changed, 33 insertions(+), 3 deletions(-) diff --git a/interfaces/sycl/DeviceCircularQueueBuffer.cpp b/interfaces/sycl/DeviceCircularQueueBuffer.cpp index c83eac2..4025ce5 100644 --- a/interfaces/sycl/DeviceCircularQueueBuffer.cpp +++ b/interfaces/sycl/DeviceCircularQueueBuffer.cpp @@ -8,6 +8,7 @@ #include "SyclWrappedAPI.h" #include "utils/logger.h" +#include #include using namespace device::internals; @@ -101,6 +102,19 @@ sycl::queue* DeviceCircularQueueBuffer::newQueue(double priority) { void DeviceCircularQueueBuffer::deleteQueue(void* queue) { auto* queuePtr = static_cast(queue); + + // The queue has to leave the list before it is freed: syncAllQueuesWithHost walks that list, + // so a stale entry turns into a use-after-free at the next device-wide synchronization, far + // away from whoever destroyed the queue. + const auto entry = std::find(externalQueues.begin(), externalQueues.end(), queuePtr); + if (entry == externalQueues.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or belongs to a different device; not " + "freeing it again."; + return; + } + + externalQueues.erase(entry); delete queuePtr; } @@ -136,10 +150,10 @@ bool DeviceCircularQueueBuffer::exists(sycl::queue* queuePtr) { bool isDefaultQueue = queuePtr == (&defaultQueue.queue); bool isGenericQueue = queuePtr == (&genericQueue.queue); - bool isReservedQueue{true}; + bool isReservedQueue{false}; for (auto& reservedQueue : queues) { - if (queuePtr != (&reservedQueue.queue)) { - isReservedQueue = false; + if (queuePtr == (&reservedQueue.queue)) { + isReservedQueue = true; break; } } diff --git a/tests/streams.cpp b/tests/streams.cpp index 66fdb70..a756d7d 100644 --- a/tests/streams.cpp +++ b/tests/streams.cpp @@ -149,6 +149,22 @@ TEST_F(Streams, asyncAllocationsLiveOnTheStream) { } } +TEST_F(Streams, aDestroyedStreamIsForgotten) { + // A device-wide synchronization walks every stream the backend knows about. A stream that was + // destroyed therefore has to be off that list, or the walk runs into freed memory - long after + // the code that destroyed it, which is what makes this kind of fault hard to place. + auto* scratch = device->api->createStream(); + device->algorithms.fillArray(devArray, 1.0F, ArraySize, scratch); + device->api->syncStreamWithHost(scratch); + device->api->destroyGenericStream(scratch); + + device->api->syncDevice(); + + for (const auto value : download(streamA)) { + ASSERT_EQ(1.0F, value); + } +} + TEST_F(Streams, workOnSeparateStreamsStaysSeparate) { auto* other = static_cast(device->api->allocGlobMem(ArraySize * sizeof(float))); From fa228b20f454045c9997fcafc74272a5a81353e4 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 01:21:21 +0200 Subject: [PATCH 10/31] fix: ask the events, not the wait list, whether a SYCL stream is idle isStreamWorkDone reported an AdaptiveCpp queue as busy right after it had been synchronized. The wait list it consults holds the events a newly submitted operation would have to depend on, and those entries stay in place once they have been reached - so an empty list means nothing was ever submitted, not that nothing is outstanding. Checking the events themselves gives the answer the function promises, which is the one cudaStreamQuery gives on the other backends. The condition also gains the SYCL_EXT_ACPP_QUEUE_WAIT_LIST spelling that the rest of the interface already tests for. Without it, a build that only defines the newer name falls through to the branch that synchronizes and then reports true - a query that quietly blocks, which is the last thing the scheduling work wants. AI-generated. Model: Opus 5 --- interfaces/sycl/Streams.cpp | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/interfaces/sycl/Streams.cpp b/interfaces/sycl/Streams.cpp index 3e96c24..c09e1f7 100644 --- a/interfaces/sycl/Streams.cpp +++ b/interfaces/sycl/Streams.cpp @@ -43,8 +43,16 @@ bool ConcreteAPI::isStreamWorkDone(void* streamPtr) { // otherwise, synchronize #ifdef SYCL_EXT_ONEAPI_QUEUE_EMPTY return queuePtr->ext_oneapi_empty(); -#elif defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) - return queuePtr->get_wait_list().empty(); +#elif defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) || \ + defined(SYCL_EXT_ACPP_QUEUE_WAIT_LIST) + // The wait list holds the events a newly submitted operation would have to depend on. Those + // entries are not dropped once they have been reached, so an empty list means "nothing was + // ever submitted", not "nothing is outstanding". Ask the events themselves instead. + const auto waitList = queuePtr->get_wait_list(); + return std::all_of(waitList.begin(), waitList.end(), [](const sycl::event& event) { + return event.get_info() == + sycl::info::event_command_status::complete; + }); #else this->currentQueueBuffer().syncQueueWithHost(queuePtr); return true; From 6db232fc556ef5701ab5eae836ae09b931ce5ea9 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:41:54 +0200 Subject: [PATCH 11/31] fix: reduce over the whole value range For floating point types, numeric_limits::min() is the smallest positive normal value, so it is larger than every negative input. A max reduction over negative data therefore returned ~1.18e-38 instead of the maximum, and with overrideResult the same value was written as the starting point. lowest() is correct for integer and floating point types alike. AI-generated. Model: Opus 5 --- algorithms/cudahip/Reduction.cpp | 5 ++++- algorithms/sycl/Reduction.cpp | 4 +++- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/algorithms/cudahip/Reduction.cpp b/algorithms/cudahip/Reduction.cpp index 12f1bf7..d46ee06 100644 --- a/algorithms/cudahip/Reduction.cpp +++ b/algorithms/cudahip/Reduction.cpp @@ -7,6 +7,7 @@ #include #include +#include #include namespace device { @@ -21,7 +22,9 @@ struct Sum { template struct Max { - T defaultValue{std::numeric_limits::min()}; + // lowest(), not min(): for floating point types min() is the smallest positive normal value, + // which is larger than every negative input + T defaultValue{std::numeric_limits::lowest()}; __device__ __forceinline__ T operator()(T op1, T op2) { return op1 > op2 ? op1 : op2; } }; diff --git a/algorithms/sycl/Reduction.cpp b/algorithms/sycl/Reduction.cpp index e0d5d52..9ee0899 100644 --- a/algorithms/sycl/Reduction.cpp +++ b/algorithms/sycl/Reduction.cpp @@ -21,7 +21,9 @@ constexpr T neutral() { return T(0); } if constexpr (Type == ReductionType::Max) { - return std::numeric_limits::min(); + // lowest(), not min(): for floating point types min() is the smallest positive normal value, + // which is larger than every negative input + return std::numeric_limits::lowest(); } if constexpr (Type == ReductionType::Min) { return std::numeric_limits::max(); From 5f42f53eeed76b54992a4742655bb2d6db399a84 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:42:24 +0200 Subject: [PATCH 12/31] fix: match the atomic width to the accumulator type The generic atomicUpdate cast the result pointer to unsigned long long and ran an 8-byte compare-and-swap on it. Native specializations existed for Sum on int, float and double only, so max and min on int, unsigned and float took the generic path and read and wrote four bytes past the result, which also let the neighbouring memory decide the comparison. The compare-and-swap now runs on a word of exactly sizeof(T), and the native atomics are selected with if constexpr instead of explicit specializations, so the available set no longer differs between the host and the device pass. atomicMax and atomicMin now also cover the integer cases. AI-generated. Model. Opus 5 --- algorithms/cudahip/Reduction.cpp | 107 +++++++++++++++++++++---------- 1 file changed, 73 insertions(+), 34 deletions(-) diff --git a/algorithms/cudahip/Reduction.cpp b/algorithms/cudahip/Reduction.cpp index d46ee06..025a9b6 100644 --- a/algorithms/cudahip/Reduction.cpp +++ b/algorithms/cudahip/Reduction.cpp @@ -9,6 +9,8 @@ #include #include #include +#include +#include namespace device { @@ -60,38 +62,79 @@ __device__ __forceinline__ T warpReduce(T value, OperationT operation) { return value; } -// Helper function for Generic Atomic Update -// Fallback to atomicCAS-based implementation if atomic instruction is not available -// Picked from: https://docs.nvidia.com/cuda/cuda-c-programming-guide/#atomic-functions +template +struct AtomicWord { + static_assert(Size == 4 || Size == 8, "no atomic word of the size of the accumulator type"); +}; + +template <> +struct AtomicWord<4> { + using Type = unsigned int; +}; + +template <> +struct AtomicWord<8> { + using Type = unsigned long long; +}; + +template +__device__ __forceinline__ ToT reinterpretValue(const FromT& from) { + static_assert(sizeof(ToT) == sizeof(FromT), "reinterpretValue requires types of equal size"); + ToT to{}; + memcpy(&to, &from, sizeof(ToT)); + return to; +} + +// Fallback for the combinations without a native atomic. The compare-and-swap runs on a word of +// exactly the size of T: a wider word would read and write the memory next to the result. template -__device__ __forceinline__ void atomicUpdate(T* address, T val, OperationT operation) { - unsigned long long* address_as_ull = (unsigned long long*)address; - unsigned long long old = *address_as_ull, assumed; +__device__ __forceinline__ void atomicUpdateCas(T* address, T val, OperationT operation) { + using WordT = typename AtomicWord::Type; + auto* wordAddress = reinterpret_cast(address); + + WordT old = *wordAddress; + WordT assumed{}; do { assumed = old; - T calculatedRes = operation(*(T*)&assumed, val); - old = atomicCAS(address_as_ull, assumed, *(unsigned long long*)&calculatedRes); + const T updated = operation(reinterpretValue(assumed), val); + old = atomicCAS(wordAddress, assumed, reinterpretValue(updated)); } while (assumed != old); } -// Native atomics -template <> -__device__ __forceinline__ void - atomicUpdate>(int* address, int val, device::Sum operation) { - atomicAdd(address, val); -} -template <> -__device__ __forceinline__ void atomicUpdate>( - float* address, float val, device::Sum operation) { - atomicAdd(address, val); -} -#if __CUDA_ARCH__ >= 600 -template <> -__device__ __forceinline__ void atomicUpdate>( - double* address, double val, device::Sum operation) { - atomicAdd(address, val); -} +template +__device__ __forceinline__ void atomicUpdate(T* address, T val, OperationT operation) { + if constexpr (std::is_same_v>) { + if constexpr (std::is_same_v || std::is_same_v || + std::is_same_v || std::is_same_v) { + atomicAdd(address, val); + return; + } else if constexpr (std::is_integral_v && sizeof(T) == sizeof(unsigned long long)) { + // the unsigned addition wraps the same way, so it also gives the signed result + atomicAdd(reinterpret_cast(address), + static_cast(val)); + return; + } else if constexpr (std::is_same_v) { +// mirrors the guard the toolkit puts on the declaration itself +#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 600) + atomicAdd(address, val); + return; #endif + } + } + if constexpr (std::is_same_v> && + (std::is_same_v || std::is_same_v || + std::is_same_v)) { + atomicMax(address, val); + return; + } + if constexpr (std::is_same_v> && + (std::is_same_v || std::is_same_v || + std::is_same_v)) { + atomicMin(address, val); + return; + } + atomicUpdateCas(address, val, operation); +} // Block Reduce template @@ -125,8 +168,8 @@ __global__ void initKernel(T* result, OperationT operation) { } template -__launch_bounds__(BlockSize) void __global__ kernel_reduce( - AccT* result, const VecT* vector, size_t size, bool overrideResult, OperationT operation) { +__launch_bounds__(BlockSize) void __global__ + kernel_reduce(AccT* result, const VecT* vector, size_t size, OperationT operation) { // Maximum block size 1024, warp size 32 so 1024/32 = 32 chosen // For AMD, warp size 64, 1024/64 = 16, but 32 should work with a few idle memory addresses @@ -147,7 +190,6 @@ __launch_bounds__(BlockSize) void __global__ kernel_reduce( AccT blockAcc = blockReduce(threadAcc, shmem, operation); if (threadIdx.x == 0) { - (void)overrideResult; // to silence unused parameter warning for non-Add reductions atomicUpdate(result, blockAcc, operation); } } @@ -180,18 +222,15 @@ void Algorithms::reduceVector(AccT* result, switch (type) { case ReductionType::Add: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Sum()); + kernel_reduce<<>>(result, buffer, size, device::Sum()); break; } case ReductionType::Max: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Max()); + kernel_reduce<<>>(result, buffer, size, device::Max()); break; } case ReductionType::Min: { - kernel_reduce<<>>( - result, buffer, size, overrideResult, device::Min()); + kernel_reduce<<>>(result, buffer, size, device::Min()); break; } default: { From b46b9422f38cfb21e2e6272aa0a0854d81d53218 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:42:46 +0200 Subject: [PATCH 13/31] fix: map stream priorities onto the range the runtime reports cudaDeviceGetStreamPriorityRange and its HIP counterpart return the *lowest* priority first and the *highest* second, and the highest is the numerically smaller of the two. Passing them to mapPercentage as (minval, maxval) made its final clamp - max(min(x, maxval), minval) - collapse to minval for every input, so every stream ended up with the default priority. HIP additionally handed the raw double, defaulting to NAN, to hipStreamCreateWithPriority and left the mapped value unused. mapStreamPriority replaces mapPercentage, which had no other callers, states the direction it maps in, clamps the input rather than the output, and lives in namespace device like the rest of the header. AbstractAPI now writes the convention down: 0 lowest, 1 highest, NAN the runtime default. AI-generated. Model: Opus 5 --- AbstractAPI.h | 5 +++++ interfaces/common/Common.h | 25 ++++++++++++++++--------- interfaces/cuda/Control.cu | 2 +- interfaces/cuda/CudaWrappedAPI.h | 3 ++- interfaces/cuda/Streams.cu | 2 +- interfaces/hip/Control.cpp | 2 +- interfaces/hip/HipWrappedAPI.h | 3 ++- interfaces/hip/Streams.cpp | 4 ++-- 8 files changed, 30 insertions(+), 16 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index cef14d6..90706ca 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -127,6 +127,11 @@ struct AbstractAPI { return graphEndNode(graphHandle, streamPtr); } + /** + * Creates a stream. `priority` runs from 0 for the lowest to 1 for the highest priority the + * device offers; NAN asks for the runtime default. Backends that only know a few priority + * classes round to the nearest one, and a device without priority support ignores the value. + */ virtual void* createStream(double priority = NAN) = 0; virtual void destroyGenericStream(void* streamPtr) = 0; virtual void syncStreamWithHost(void* streamPtr) = 0; diff --git a/interfaces/common/Common.h b/interfaces/common/Common.h index 6fdc465..4415141 100644 --- a/interfaces/common/Common.h +++ b/interfaces/common/Common.h @@ -8,6 +8,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -38,18 +39,24 @@ U align(T number, U alignment) { size_t alignmentFactor = (number + alignment - 1) / alignment; return alignmentFactor * alignment; } -} // namespace device -constexpr auto mapPercentage(int minval, int maxval, double value) { - if (std::isnan(value)) { - return 0; +/** + * Maps a stream priority as the device API states it - 0 the lowest, 1 the highest, NaN the + * runtime default - onto the integer range the runtime reports. + * + * CUDA and HIP report that range as a pair in which the numerically *smaller* value is the + * *higher* priority, so the mapping runs downwards from leastPriority to greatestPriority. Both + * ends coincide on devices that do not support priorities, which then yields that single value. + */ +inline int mapStreamPriority(int leastPriority, int greatestPriority, double priority) { + if (std::isnan(priority)) { + return leastPriority; } - const auto convminval = static_cast(minval); - const auto convmaxval = static_cast(maxval); - - const auto transformed = value * (convmaxval - convminval + 1) + convminval; - return std::max(std::min(static_cast(std::floor(transformed)), maxval), minval); + const auto fraction = std::min(std::max(priority, 0.0), 1.0); + const auto span = static_cast(leastPriority) - static_cast(greatestPriority); + return static_cast(std::lround(static_cast(leastPriority) - fraction * span)); } +} // namespace device #endif // SEISSOLDEVICE_INTERFACES_COMMON_COMMON_H_ diff --git a/interfaces/cuda/Control.cu b/interfaces/cuda/Control.cu index 0fe1815..c1e5dcb 100644 --- a/interfaces/cuda/Control.cu +++ b/interfaces/cuda/Control.cu @@ -75,7 +75,7 @@ void ConcreteAPI::initialize() { usmDefault = properties[getDeviceId()].directManagedMemAccessFromHost != 0; - APIWRAP(cudaDeviceGetStreamPriorityRange(&priorityMin, &priorityMax)); + APIWRAP(cudaDeviceGetStreamPriorityRange(&priorityLeast, &priorityGreatest)); int canCompressProto = 0; DRVWRAP(cuDeviceGetAttribute( diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index a02f76f..51771a7 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -139,7 +139,8 @@ class ConcreteAPI : public AbstractAPI { Statistics statistics{}; std::unordered_map memToSizeMap{{nullptr, 0}}; - int priorityMin, priorityMax; + int priorityLeast{0}; + int priorityGreatest{0}; std::unordered_map allocationProperties; }; diff --git a/interfaces/cuda/Streams.cu b/interfaces/cuda/Streams.cu index b70156c..86cc162 100644 --- a/interfaces/cuda/Streams.cu +++ b/interfaces/cuda/Streams.cu @@ -29,7 +29,7 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); cudaStream_t stream; - const auto truePriority = mapPercentage(priorityMin, priorityMax, priority); + const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); APIWRAP(cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, truePriority)); genericStreams.insert(stream); return reinterpret_cast(stream); diff --git a/interfaces/hip/Control.cpp b/interfaces/hip/Control.cpp index 839c557..85c76a9 100644 --- a/interfaces/hip/Control.cpp +++ b/interfaces/hip/Control.cpp @@ -84,7 +84,7 @@ void ConcreteAPI::initialize() { properties[getDeviceId()].pageableMemoryAccessUsesHostPageTables != 0; } - APIWRAP(hipDeviceGetStreamPriorityRange(&priorityMin, &priorityMax)); + APIWRAP(hipDeviceGetStreamPriorityRange(&priorityLeast, &priorityGreatest)); } else { logWarning() << "Device Interface has already been initialized"; } diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index 8519f7e..ffcb52c 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -135,7 +135,8 @@ class ConcreteAPI : public AbstractAPI { Statistics statistics{}; std::unordered_map memToSizeMap{{nullptr, 0}}; - int priorityMin, priorityMax; + int priorityLeast{0}; + int priorityGreatest{0}; }; } // namespace device diff --git a/interfaces/hip/Streams.cpp b/interfaces/hip/Streams.cpp index 37a63e1..504d6f2 100644 --- a/interfaces/hip/Streams.cpp +++ b/interfaces/hip/Streams.cpp @@ -28,8 +28,8 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); hipStream_t stream; - const auto truePriority = mapPercentage(priorityMin, priorityMax, priority); - APIWRAP(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority)); + const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); + APIWRAP(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, truePriority)); genericStreams.insert(stream); return reinterpret_cast(stream); } From a46b717d6400d7defcb83e23a5ca2e282d05c814 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:43:10 +0200 Subject: [PATCH 14/31] fix: throw the exception, not a pointer to it throw new std::invalid_argument(...) throws an std::invalid_argument*, which no catch clause for const std::exception& matches, so the double initialization ends in std::terminate with the exception object leaked on the way out. AI-generated. Model: Opus 5 --- interfaces/sycl/Control.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index f3ec2b3..78f75ad 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -8,6 +8,7 @@ #include "utils/logger.h" #include +#include #include #include @@ -24,7 +25,7 @@ using namespace device; void ConcreteAPI::initDevices() { if (this->deviceInitialized) { - throw new std::invalid_argument("Cannot initialize the devices twice!"); + throw std::invalid_argument("Cannot initialize the devices twice!"); } for (const auto& platform : sycl::platform::get_platforms()) { From 045197c51eeae886b80b21bd695b27d512e6fab2 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:43:32 +0200 Subject: [PATCH 15/31] fix: order devices by a strict weak ordering compare() returned true for both argument orders once both devices matched PREFERRED_DEVICE_TYPE, and sorting on a comparator that does that is undefined. Both devices are now ranked by the same expression and the ranks are compared, so equal devices compare equal in either direction. The sort is stable now as well, which keeps the device ids the same from run to run when the comparator cannot tell two devices apart. AI-generated. Model: Opus 5 --- interfaces/sycl/Control.cpp | 15 ++++++++------ interfaces/sycl/DeviceType.cpp | 37 ++++++++++++++++------------------ interfaces/sycl/DeviceType.h | 2 +- 3 files changed, 27 insertions(+), 27 deletions(-) diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index 78f75ad..b6314f9 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -7,6 +7,7 @@ #include "SyclWrappedAPI.h" #include "utils/logger.h" +#include #include #include #include @@ -48,12 +49,14 @@ void ConcreteAPI::initDevices() { } } - sort(this->availableDevices.begin(), - this->availableDevices.end(), - [&](DeviceContext* c1, DeviceContext* c2) { - return compare(c1->queueBuffer.getDefaultQueue().get_device(), - c2->queueBuffer.getDefaultQueue().get_device()); - }); + // stable, so that devices the comparator sees as equal keep the order the platform reported + // them in and the device ids stay the same from run to run + std::stable_sort(this->availableDevices.begin(), + this->availableDevices.end(), + [](DeviceContext* c1, DeviceContext* c2) { + return compare(c1->queueBuffer.getDefaultQueue().get_device(), + c2->queueBuffer.getDefaultQueue().get_device()); + }); this->setDevice(0); this->deviceInitialized = true; diff --git a/interfaces/sycl/DeviceType.cpp b/interfaces/sycl/DeviceType.cpp index 6a99a8a..d1737e9 100644 --- a/interfaces/sycl/DeviceType.cpp +++ b/interfaces/sycl/DeviceType.cpp @@ -7,6 +7,8 @@ #include "utils/env.h" #include +#include +#include namespace device { @@ -26,26 +28,21 @@ DeviceType fromSyclType(sycl::info::device_type type) { return DeviceType::OTHERS; } -bool compare(sycl::device devA, sycl::device devB) { - std::string env; - env += utils::Env("").get("PREFERRED_DEVICE_TYPE", ""); - - auto typeA = devA.get_info(); - auto typeB = devB.get_info(); - - if (convertToString(typeA).compare(env) == 0) { - return true; - } - if (convertToString(typeB).compare(env) == 0) { - return false; - } - - // devices of same type are sorted by their cl::deviceid - if (typeA == typeB) { - // return devA.get() < devB.get(); - } - - return fromSyclType(typeA) < fromSyclType(typeB); +bool compare(const sycl::device& devA, const sycl::device& devB) { + std::string preferred; + preferred += utils::Env("").get("PREFERRED_DEVICE_TYPE", ""); + + // A device of the preferred type sorts ahead of every other device, and the rest follow the + // order of the DeviceType enum. Deciding the two directions independently - as in "A wins if it + // matches, B wins if it matches" - makes both compare(a, b) and compare(b, a) true for two + // devices of the preferred type, and sorting on such a comparator is undefined. + const auto rank = [&preferred](const sycl::device& device) { + const auto type = device.get_info(); + const auto matchesPreference = convertToString(type) == preferred ? 0 : 1; + return std::make_pair(matchesPreference, static_cast(fromSyclType(type))); + }; + + return rank(devA) < rank(devB); } std::string convertToString(sycl::info::device_type type) { diff --git a/interfaces/sycl/DeviceType.h b/interfaces/sycl/DeviceType.h index 4a0e288..bf3243a 100644 --- a/interfaces/sycl/DeviceType.h +++ b/interfaces/sycl/DeviceType.h @@ -13,7 +13,7 @@ enum class DeviceType { GPU = 0, CPU = 1, FPGA = 2, HOST = 3, OTHERS = 4 }; std::string convertToString(sycl::info::device_type type); DeviceType fromSyclType(sycl::info::device_type type); -bool compare(sycl::device devA, sycl::device devB); +bool compare(const sycl::device& devA, const sycl::device& devB); } // namespace device From 0df908f691f9d1e0c83101b945a8d2b636f0b5ed Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:43:56 +0200 Subject: [PATCH 16/31] fix: reduce the SYCL queue container to what it owns counter was never initialized and getNextQueue() incremented it, so the first call read an indeterminate value. It, getGenericQueue, allQueues, resetIndex, getCapacity and the two fork/join helpers had no callers, and with them gone the pool of six queues behind them has none either - the backend hands out queues through newQueue and shares one default queue. Building the pool also went through QueueWrapper's default constructor, which default-constructs a sycl::queue through the default selector before the real one overwrites it - once per pool entry per device. Queues from newQueue that the caller never destroys are now freed when the device goes away, with the same warning the CUDA backend prints. The class is called DeviceQueues now, since it no longer buffers anything circularly. AI-generated. Model: Opus 5 --- interfaces/sycl/Control.cpp | 2 +- interfaces/sycl/DeviceCircularQueueBuffer.cpp | 172 ------------------ interfaces/sycl/DeviceCircularQueueBuffer.h | 98 ---------- interfaces/sycl/DeviceContext.cpp | 5 +- interfaces/sycl/DeviceContext.h | 6 +- interfaces/sycl/DeviceQueues.cpp | 103 +++++++++++ interfaces/sycl/DeviceQueues.h | 66 +++++++ interfaces/sycl/SyclWrappedAPI.h | 2 +- sycl.cmake | 2 +- 9 files changed, 177 insertions(+), 279 deletions(-) delete mode 100644 interfaces/sycl/DeviceCircularQueueBuffer.cpp delete mode 100644 interfaces/sycl/DeviceCircularQueueBuffer.h create mode 100644 interfaces/sycl/DeviceQueues.cpp create mode 100644 interfaces/sycl/DeviceQueues.h diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index b6314f9..d3decec 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -44,7 +44,7 @@ void ConcreteAPI::initDevices() { } } - DeviceContext* context = new DeviceContext{device, 1}; + DeviceContext* context = new DeviceContext{device}; this->availableDevices.push_back(context); } } diff --git a/interfaces/sycl/DeviceCircularQueueBuffer.cpp b/interfaces/sycl/DeviceCircularQueueBuffer.cpp deleted file mode 100644 index 4025ce5..0000000 --- a/interfaces/sycl/DeviceCircularQueueBuffer.cpp +++ /dev/null @@ -1,172 +0,0 @@ -// SPDX-FileCopyrightText: 2022 SeisSol Group -// -// SPDX-License-Identifier: BSD-3-Clause - -#include "DeviceCircularQueueBuffer.h" - -#include "Internals.h" -#include "SyclWrappedAPI.h" -#include "utils/logger.h" - -#include -#include - -using namespace device::internals; - -namespace device { - -// very inconvenient, but AdaptiveCpp doesn't allow much freedom when constructing a property_list -#if defined(DEVICE_USE_GRAPH_CAPTURING) && defined(SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST) -#define BASE_QUEUE_PROPERTIES \ - sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::no_immediate_command_list {} -#else -#define BASE_QUEUE_PROPERTIES sycl::property::queue::in_order() -#endif - -QueueWrapper::QueueWrapper(const sycl::device& dev, - const std::function& handler) - : queue{dev, handler, sycl::property_list{BASE_QUEUE_PROPERTIES}} {} - -void QueueWrapper::synchronize() { waitCheck(queue); } -void QueueWrapper::dependency(QueueWrapper& other) { - // improvising... Adding an empty event here, mimicking a CUDA-like event dependency -#if defined(HIPSYCL_EXT_QUEUE_WAIT_LIST) || defined(ACPP_EXT_QUEUE_WAIT_LIST) || \ - defined(SYCL_EXT_ACPP_QUEUE_WAIT_LIST) - auto waitList1 = other.queue.get_wait_list(); - auto waitList2 = queue.get_wait_list(); - queue.submit([&](sycl::handler& h) { - h.depends_on(waitList1); - h.depends_on(waitList2); - DEVICE_SYCL_EMPTY_OPERATION(h); - }); -#else - auto queueEvent = other.queue.submit([&](sycl::handler& h) { DEVICE_SYCL_EMPTY_OPERATION(h); }); - queue.submit([&](sycl::handler& h) { DEVICE_SYCL_EMPTY_OPERATION_WITH_EVENT(h, queueEvent); }); -#endif -} - -DeviceCircularQueueBuffer::DeviceCircularQueueBuffer( - const sycl::device& dev, - const std::function& handler, - size_t capacity) - : queues{std::vector(capacity)}, deviceReference(dev), handlerReference(handler) { - if (capacity <= 0) - throw std::invalid_argument("Capacity must be at least 1!"); - - this->defaultQueue = QueueWrapper(dev, handler); - this->genericQueue = QueueWrapper(dev, handler); - for (size_t i = 0; i < capacity; i++) { - this->queues[i] = QueueWrapper(dev, handler); - } -} - -sycl::queue& DeviceCircularQueueBuffer::getDefaultQueue() { return defaultQueue.queue; } - -sycl::queue& DeviceCircularQueueBuffer::getGenericQueue() { return genericQueue.queue; } - -sycl::queue& DeviceCircularQueueBuffer::getNextQueue() { - (++this->counter) %= getCapacity(); - return (this->queues[this->counter].queue); -} - -std::vector DeviceCircularQueueBuffer::allQueues() { - std::vector queueCopy(queues.size() + 1); - queueCopy[0] = defaultQueue.queue; - for (size_t i = 0; i < queues.size(); ++i) { - queueCopy[i + 1] = queues[i].queue; - } - return queueCopy; -} - -sycl::queue* DeviceCircularQueueBuffer::newQueue(double priority) { - // missing for ACPP: how can we even find out the allowed priority range conveniently now? :/ - -#ifdef SYCL_EXT_ONEAPI_QUEUE_PRIORITY - const auto propertylist = [&]() -> sycl::property_list { - if (priority <= 0.33) { - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_low()}; - } - if (priority >= 0.67) { - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_high()}; - } - return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_normal()}; - }(); -#else - const auto propertylist{BASE_QUEUE_PROPERTIES}; -#endif - - auto* queue = new sycl::queue{deviceReference, handlerReference, propertylist}; - externalQueues.emplace_back(queue); - return queue; -} - -void DeviceCircularQueueBuffer::deleteQueue(void* queue) { - auto* queuePtr = static_cast(queue); - - // The queue has to leave the list before it is freed: syncAllQueuesWithHost walks that list, - // so a stale entry turns into a use-after-free at the next device-wide synchronization, far - // away from whoever destroyed the queue. - const auto entry = std::find(externalQueues.begin(), externalQueues.end(), queuePtr); - if (entry == externalQueues.end()) { - logWarning() << "Tried to destroy a stream that this device does not know about. It has " - "either been destroyed already or belongs to a different device; not " - "freeing it again."; - return; - } - - externalQueues.erase(entry); - delete queuePtr; -} - -void DeviceCircularQueueBuffer::resetIndex() { this->counter = 0; } - -size_t DeviceCircularQueueBuffer::getCapacity() { return queues.size(); } - -void DeviceCircularQueueBuffer::forkQueueDepencency() { - for (auto& queue : this->queues) { - queue.dependency(defaultQueue); - } -} - -void DeviceCircularQueueBuffer::joinQueueDepencency() { - for (auto& queue : this->queues) { - defaultQueue.dependency(queue); - } -} - -void DeviceCircularQueueBuffer::syncQueueWithHost(sycl::queue* queuePtr) { waitCheck(*queuePtr); } - -void DeviceCircularQueueBuffer::syncAllQueuesWithHost() { - defaultQueue.synchronize(); - for (auto& q : this->queues) { - q.synchronize(); - } - for (auto* q : this->externalQueues) { - waitCheck(*q); - } -} - -bool DeviceCircularQueueBuffer::exists(sycl::queue* queuePtr) { - bool isDefaultQueue = queuePtr == (&defaultQueue.queue); - bool isGenericQueue = queuePtr == (&genericQueue.queue); - - bool isReservedQueue{false}; - for (auto& reservedQueue : queues) { - if (queuePtr == (&reservedQueue.queue)) { - isReservedQueue = true; - break; - } - } - - bool isExternalQueue = false; - for (auto& queue : externalQueues) { - if (queuePtr == queue) { - isExternalQueue = true; - break; - } - } - - return isDefaultQueue || isGenericQueue || isReservedQueue || isExternalQueue; -} - -} // namespace device diff --git a/interfaces/sycl/DeviceCircularQueueBuffer.h b/interfaces/sycl/DeviceCircularQueueBuffer.h deleted file mode 100644 index b24aace..0000000 --- a/interfaces/sycl/DeviceCircularQueueBuffer.h +++ /dev/null @@ -1,98 +0,0 @@ -// SPDX-FileCopyrightText: 2021 SeisSol Group -// -// SPDX-License-Identifier: BSD-3-Clause - -#ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ -#define SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ - -#include -#include -#include -#include -#include - -namespace device { -struct QueueWrapper { - sycl::queue queue; - QueueWrapper() = default; - QueueWrapper(const sycl::device& dev, const std::function& f); - - void synchronize(); - void dependency(QueueWrapper& other); -}; - -class DeviceCircularQueueBuffer { - public: - /* - * Creates a new circular buffer containing sycl queues. The buffer needs - * an async exception handler and a capacity that is currently per default 8. - */ - DeviceCircularQueueBuffer(const sycl::device& dev, - const std::function& f, - size_t capacity = 6); - - /* - * Returns the default queue created by a device. - */ - sycl::queue& getDefaultQueue(); - - /* - * Returns the generic queue created by a device. - * Note, this queue can be used for asynchronous - * memory copies - */ - sycl::queue& getGenericQueue(); - - /* - * Returns the next queue within the capacity of this buffer. - */ - sycl::queue& getNextQueue(); - - sycl::queue* newQueue(double priority); - - void deleteQueue(void* queue); - - std::vector allQueues(); - - /* - * Resets the index to the current element. - */ - void resetIndex(); - - /* - * Returns the default queue created by a device. - */ - size_t getCapacity(); - - /* - * Synchronizes a queue from the buffer with the host device. - */ - void syncQueueWithHost(sycl::queue* queuePtr); - - /* - * Synchronizes all queues from the buffer with the host device. - */ - void syncAllQueuesWithHost(); - - /* - *Returns true if the queue pointer is available on the buffer. - */ - bool exists(sycl::queue* queuePtr); - - void forkQueueDepencency(); - - void joinQueueDepencency(); - - private: - QueueWrapper defaultQueue; - QueueWrapper genericQueue; - std::vector queues; - std::vector externalQueues; - size_t counter; - sycl::device deviceReference; - std::function handlerReference; -}; - -} // namespace device - -#endif // SEISSOLDEVICE_INTERFACES_SYCL_DEVICECIRCULARQUEUEBUFFER_H_ diff --git a/interfaces/sycl/DeviceContext.cpp b/interfaces/sycl/DeviceContext.cpp index 1638390..f1568f9 100644 --- a/interfaces/sycl/DeviceContext.cpp +++ b/interfaces/sycl/DeviceContext.cpp @@ -8,9 +8,8 @@ #include "utils/logger.h" namespace device { -DeviceContext::DeviceContext(const sycl::device& targetDevice, size_t concurrencyLevel) - : queueBuffer{DeviceCircularQueueBuffer{ - targetDevice, [&](sycl::exception_list l) { onExceptionOccurred(l); }, concurrencyLevel}}, +DeviceContext::DeviceContext(const sycl::device& targetDevice) + : queueBuffer{targetDevice, [&](sycl::exception_list l) { onExceptionOccurred(l); }}, statistics{Statistics{}} {} void DeviceContext::onExceptionOccurred(sycl::exception_list& exceptions) { diff --git a/interfaces/sycl/DeviceContext.h b/interfaces/sycl/DeviceContext.h index 174cd44..72e602c 100644 --- a/interfaces/sycl/DeviceContext.h +++ b/interfaces/sycl/DeviceContext.h @@ -5,7 +5,7 @@ #ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICECONTEXT_H_ #define SEISSOLDEVICE_INTERFACES_SYCL_DEVICECONTEXT_H_ -#include "DeviceCircularQueueBuffer.h" +#include "DeviceQueues.h" #include "Statistics.h" #include @@ -17,9 +17,9 @@ namespace device { */ class DeviceContext { public: - DeviceContext(const sycl::device& targetDevice, size_t concurrencyLevel); + explicit DeviceContext(const sycl::device& targetDevice); std::unordered_map memoryToSizeMap; - DeviceCircularQueueBuffer queueBuffer; + DeviceQueues queueBuffer; Statistics statistics; private: diff --git a/interfaces/sycl/DeviceQueues.cpp b/interfaces/sycl/DeviceQueues.cpp new file mode 100644 index 0000000..1050740 --- /dev/null +++ b/interfaces/sycl/DeviceQueues.cpp @@ -0,0 +1,103 @@ +// SPDX-FileCopyrightText: 2022 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#include "DeviceQueues.h" + +#include "Internals.h" +#include "SyclWrappedAPI.h" +#include "utils/logger.h" + +#include +#include +#include + +using namespace device::internals; + +namespace device { + +// very inconvenient, but AdaptiveCpp doesn't allow much freedom when constructing a property_list +#if defined(DEVICE_USE_GRAPH_CAPTURING) && defined(SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST) +#define BASE_QUEUE_PROPERTIES \ + sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::no_immediate_command_list {} +#else +#define BASE_QUEUE_PROPERTIES sycl::property::queue::in_order() +#endif + +DeviceQueues::DeviceQueues(const sycl::device& dev, + const std::function& handler) + : defaultQueue{dev, handler, sycl::property_list{BASE_QUEUE_PROPERTIES}}, deviceReference(dev), + handlerReference(handler) {} + +DeviceQueues::~DeviceQueues() { + if (!externalQueues.empty()) { + logInfo() << "DEVICE::WARNING:" << externalQueues.size() + << "device generic stream(s) were not deleted."; + } + for (auto* queue : externalQueues) { + delete queue; + } +} + +sycl::queue& DeviceQueues::getDefaultQueue() { return defaultQueue; } + +sycl::queue* DeviceQueues::newQueue(double priority) { + // missing for ACPP: how can we even find out the allowed priority range conveniently now? :/ + +#ifdef SYCL_EXT_ONEAPI_QUEUE_PRIORITY + const sycl::property_list propertylist = [&]() -> sycl::property_list { + if (std::isnan(priority)) { + return {BASE_QUEUE_PROPERTIES}; + } + if (priority <= 0.33) { + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_low()}; + } + if (priority >= 0.67) { + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_high()}; + } + return {BASE_QUEUE_PROPERTIES, sycl::ext::oneapi::property::queue::priority_normal()}; + }(); +#else + const sycl::property_list propertylist{BASE_QUEUE_PROPERTIES}; +#endif + + auto* queue = new sycl::queue{deviceReference, handlerReference, propertylist}; + externalQueues.emplace_back(queue); + return queue; +} + +void DeviceQueues::deleteQueue(void* queue) { + auto* queuePtr = static_cast(queue); + + // The queue has to leave the list before it is freed: syncAllQueuesWithHost walks that list, + // so a stale entry turns into a use-after-free at the next device-wide synchronization, far + // away from whoever destroyed the queue. + const auto entry = std::find(externalQueues.begin(), externalQueues.end(), queuePtr); + if (entry == externalQueues.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or belongs to a different device; not " + "freeing it again."; + return; + } + + externalQueues.erase(entry); + delete queuePtr; +} + +void DeviceQueues::syncQueueWithHost(sycl::queue* queuePtr) { waitCheck(*queuePtr); } + +void DeviceQueues::syncAllQueuesWithHost() { + waitCheck(defaultQueue); + for (auto* queue : this->externalQueues) { + waitCheck(*queue); + } +} + +bool DeviceQueues::exists(sycl::queue* queuePtr) { + if (queuePtr == &defaultQueue) { + return true; + } + return std::find(externalQueues.begin(), externalQueues.end(), queuePtr) != externalQueues.end(); +} + +} // namespace device diff --git a/interfaces/sycl/DeviceQueues.h b/interfaces/sycl/DeviceQueues.h new file mode 100644 index 0000000..ccf9a41 --- /dev/null +++ b/interfaces/sycl/DeviceQueues.h @@ -0,0 +1,66 @@ +// SPDX-FileCopyrightText: 2021 SeisSol Group +// +// SPDX-License-Identifier: BSD-3-Clause + +#ifndef SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ +#define SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ + +#include +#include +#include + +namespace device { + +/* + * Owns the queues of one device: the default queue every caller shares, and the queues handed out + * through newQueue. + */ +class DeviceQueues { + public: + DeviceQueues(const sycl::device& dev, const std::function& f); + ~DeviceQueues(); + + DeviceQueues(const DeviceQueues&) = delete; + DeviceQueues& operator=(const DeviceQueues&) = delete; + + /* + * Returns the default queue of this device. + */ + sycl::queue& getDefaultQueue(); + + /* + * Creates a queue owned by this device. `priority` follows the convention of + * AbstractAPI::createStream: 0 the lowest, 1 the highest, NaN the runtime default. + */ + sycl::queue* newQueue(double priority); + + /* + * Destroys a queue obtained from newQueue. + */ + void deleteQueue(void* queue); + + /* + * Synchronizes one queue with the host. + */ + void syncQueueWithHost(sycl::queue* queuePtr); + + /* + * Synchronizes every queue of this device with the host. + */ + void syncAllQueuesWithHost(); + + /* + * Returns true if the queue belongs to this device. + */ + bool exists(sycl::queue* queuePtr); + + private: + sycl::queue defaultQueue; + std::vector externalQueues; + sycl::device deviceReference; + std::function handlerReference; +}; + +} // namespace device + +#endif // SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index 9832376..0f41d8d 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -162,7 +162,7 @@ class ConcreteAPI : public AbstractAPI { DeviceContext* currentContext() { return this->availableDevices[getDeviceId()]; } sycl::queue& currentDefaultQueue() { return this->currentQueueBuffer().getDefaultQueue(); } - DeviceCircularQueueBuffer& currentQueueBuffer() { return this->currentContext()->queueBuffer; } + DeviceQueues& currentQueueBuffer() { return this->currentContext()->queueBuffer; } Statistics& currentStatistics() { return this->currentContext()->statistics; } std::unordered_map& currentMemoryToSizeMap() { return this->currentContext()->memoryToSizeMap; diff --git a/sycl.cmake b/sycl.cmake index 666e1e7..8b2ff28 100644 --- a/sycl.cmake +++ b/sycl.cmake @@ -10,7 +10,7 @@ set(DEVICE_SOURCE_FILES device.cpp interfaces/sycl/Memory.cpp interfaces/sycl/Streams.cpp interfaces/sycl/DeviceContext.cpp - interfaces/sycl/DeviceCircularQueueBuffer.cpp + interfaces/sycl/DeviceQueues.cpp interfaces/sycl/DeviceType.cpp algorithms/sycl/ArrayManip.cpp algorithms/sycl/BatchManip.cpp From c93882c2cbf54ea9ea860b97935f79dcef0d1406 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:44:29 +0200 Subject: [PATCH 17/31] fix: only advise a memory location that exists allocUnifiedMem and prefetchUnifiedMemTo left cudaMemLocation zeroed whenever the caller asked for the current device and the device had no concurrent managed access, and issued the call anyway. Up to CUDA 12 the zeroed id names device 0, so on a node with several GPUs the hint went to the wrong card; from CUDA 13 on the zeroed type is cudaMemLocationTypeInvalid and the call fails, which APIWRAP turns into an abort. The preferred location is now only set where there is one to set, and the prefetch returns early, since it needs concurrent managed access to begin with. The HIP backend passed a literal 1 as the device of the coarse-grain advice, which is not a device id on a single-GPU node. AI-generated. Model: Opus 5 --- interfaces/cuda/Copy.cu | 9 ++++++++- interfaces/cuda/Memory.cu | 35 +++++++++++++++++++++-------------- interfaces/hip/Memory.cpp | 2 +- 3 files changed, 30 insertions(+), 16 deletions(-) diff --git a/interfaces/cuda/Copy.cu b/interfaces/cuda/Copy.cu index 7776096..72adbaa 100644 --- a/interfaces/cuda/Copy.cu +++ b/interfaces/cuda/Copy.cu @@ -65,6 +65,13 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, size_t count, void* streamPtr) { isFlagSet(status); + + // Prefetching managed memory needs concurrent managed access, in either direction. A prefetch + // is a hint, so where the device cannot serve it, skipping is the whole handling. + if (!allowedConcurrentManagedAccess) { + return; + } + cudaStream_t stream = (streamPtr == nullptr) ? nullptr : (static_cast(streamPtr)); cudaMemLocation location{}; @@ -73,7 +80,7 @@ void ConcreteAPI::prefetchUnifiedMemTo(Destination type, #if CUDART_VERSION >= 13000 location.type = cudaMemLocationTypeHost; #endif - } else if (allowedConcurrentManagedAccess) { + } else { location.id = getDeviceId(); #if CUDART_VERSION >= 13000 location.type = cudaMemLocationTypeDevice; diff --git a/interfaces/cuda/Memory.cu b/interfaces/cuda/Memory.cu index df55a59..77431f0 100644 --- a/interfaces/cuda/Memory.cu +++ b/interfaces/cuda/Memory.cu @@ -90,28 +90,35 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) void* devPtr = nullptr; APIWRAP(cudaMallocManaged(&devPtr, size, cudaMemAttachGlobal)); - cudaMemLocation location{}; - if (hint == Destination::Host) { - location.id = cudaCpuDeviceId; + // Naming the device as the preferred location needs concurrent managed access. Where that is + // missing there is no location to name instead, so the allocation keeps the driver default - + // passing a zeroed location would advise for device 0 up to CUDA 12 and be rejected from CUDA + // 13 on. + const bool hasPreferredLocation = (hint == Destination::Host) || allowedConcurrentManagedAccess; + if (hasPreferredLocation) { + cudaMemLocation location{}; + if (hint == Destination::Host) { + location.id = cudaCpuDeviceId; #if CUDART_VERSION >= 13000 - location.type = cudaMemLocationTypeHost; + location.type = cudaMemLocationTypeHost; #endif - } else if (allowedConcurrentManagedAccess) { - location.id = getDeviceId(); + } else { + location.id = getDeviceId(); #if CUDART_VERSION >= 13000 - location.type = cudaMemLocationTypeDevice; + location.type = cudaMemLocationTypeDevice; #endif - } + } - APIWRAP(cudaMemAdvise(devPtr, - size, - cudaMemAdviseSetPreferredLocation, + APIWRAP(cudaMemAdvise(devPtr, + size, + cudaMemAdviseSetPreferredLocation, #if CUDART_VERSION >= 13000 - location + location #else - location.id + location.id #endif - )); + )); + } statistics.allocatedMemBytes += size; statistics.allocatedUnifiedMemBytes += size; diff --git a/interfaces/hip/Memory.cpp b/interfaces/hip/Memory.cpp index b4ef78a..96f67e7 100644 --- a/interfaces/hip/Memory.cpp +++ b/interfaces/hip/Memory.cpp @@ -28,7 +28,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) APIWRAP(hipMallocManaged(&devPtr, size, hipMemAttachGlobal)); // make coarse-grained memory access behavior the default (match with allocGlobMem) - APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetCoarseGrain, 1)); + APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetCoarseGrain, getDeviceId())); if (hint == Destination::Host) { APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, hipCpuDeviceId)); From 8aa2b2c067a6345e9e108cdc7d9928490a6bffa3 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:44:50 +0200 Subject: [PATCH 18/31] fix: release the allocation bookkeeping on free None of the free functions dropped the pointer from memToSizeMap, so the map grew for the lifetime of the process. freeGlobMem also kept the CUmemAllocationProp of a compressed allocation, which leaked it and, worse, left the address registered: the runtime is free to return the same address from a later cudaMalloc, and that allocation would then be released through cuMemUnmap. The properties are stored by value now and dropped with the allocation. freeUnifiedMem never subtracted from allocatedUnifiedMemBytes on any backend, so getCurrentlyOccupiedUnifiedMem only ever grew. Freeing a null pointer is handled up front, which is what the nullptr entry the maps were seeded with stood in for. UsmAllocator allocates through allocUnifiedMem and now frees through freeUnifiedMem rather than freeGlobMem. AI-generated. Model: Opus 5 --- UsmAllocator.h | 2 +- interfaces/cuda/CudaWrappedAPI.h | 7 +++-- interfaces/cuda/Memory.cu | 52 +++++++++++++++++++++++--------- interfaces/hip/HipWrappedAPI.h | 5 ++- interfaces/hip/Memory.cpp | 39 ++++++++++++++++++------ interfaces/sycl/Memory.cpp | 10 ++++-- interfaces/sycl/SyclWrappedAPI.h | 2 +- 7 files changed, 86 insertions(+), 31 deletions(-) diff --git a/UsmAllocator.h b/UsmAllocator.h index abb9efb..37d9ba0 100644 --- a/UsmAllocator.h +++ b/UsmAllocator.h @@ -44,7 +44,7 @@ class UsmAllocator { void deallocate(T* ptr, std::size_t) { if (ptr) { - api->freeGlobMem(ptr); + api->freeUnifiedMem(ptr); } } diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index 51771a7..3d616af 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -123,6 +123,9 @@ class ConcreteAPI : public AbstractAPI { void setupPrinting(int rank) override; private: + // Drops the allocation from the bookkeeping and returns its size. + size_t forgetAllocation(void* devPtr); + device::StatusT status{false}; std::vector properties; @@ -137,12 +140,12 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; Statistics statistics{}; - std::unordered_map memToSizeMap{{nullptr, 0}}; + std::unordered_map memToSizeMap; int priorityLeast{0}; int priorityGreatest{0}; - std::unordered_map allocationProperties; + std::unordered_map allocationProperties; }; } // namespace device diff --git a/interfaces/cuda/Memory.cu b/interfaces/cuda/Memory.cu index 77431f0..cc5d3aa 100644 --- a/interfaces/cuda/Memory.cu +++ b/interfaces/cuda/Memory.cu @@ -76,7 +76,7 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { prop.allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC; devPtr = driverAllocate(size, prop); - allocationProperties[devPtr] = reinterpret_cast(new CUmemAllocationProp(prop)); + allocationProperties[devPtr] = prop; } else { APIWRAP(cudaMalloc(&devPtr, size)); } @@ -136,15 +136,34 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) return devPtr; } +size_t ConcreteAPI::forgetAllocation(void* devPtr) { + const auto entry = memToSizeMap.find(devPtr); + if (entry == memToSizeMap.end()) { + assert(false && "DEVICE: an attempt to delete mem. which has not been allocated. unknown " + "pointer"); + return 0; + } + + const auto size = entry->second; + memToSizeMap.erase(entry); + statistics.deallocatedMemBytes += size; + return size; +} + void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; - if (allocationProperties.find(devPtr) != allocationProperties.end()) { - driverFree(devPtr, - memToSizeMap.at(devPtr), - *reinterpret_cast(allocationProperties.at(devPtr))); + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + + const auto properties = allocationProperties.find(devPtr); + if (properties != allocationProperties.end()) { + driverFree(devPtr, size, properties->second); + // the entry has to go with the allocation: the runtime is free to hand the same address out + // again, and a leftover entry would send that one down the driver path as well + allocationProperties.erase(properties); } else { APIWRAP(cudaFree(devPtr)); } @@ -152,17 +171,22 @@ void ConcreteAPI::freeGlobMem(void* devPtr) { void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + statistics.allocatedUnifiedMemBytes -= size; APIWRAP(cudaFree(devPtr)); } void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(cudaFreeHost(devPtr)); } diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index ffcb52c..075cead 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -122,6 +122,9 @@ class ConcreteAPI : public AbstractAPI { void setupPrinting(int rank) override; private: + // Drops the allocation from the bookkeeping and returns its size. + size_t forgetAllocation(void* devPtr); + device::StatusT status{false}; std::vector properties; @@ -133,7 +136,7 @@ class ConcreteAPI : public AbstractAPI { std::unordered_set genericStreams{}; Statistics statistics{}; - std::unordered_map memToSizeMap{{nullptr, 0}}; + std::unordered_map memToSizeMap; int priorityLeast{0}; int priorityGreatest{0}; diff --git a/interfaces/hip/Memory.cpp b/interfaces/hip/Memory.cpp index 96f67e7..e102f88 100644 --- a/interfaces/hip/Memory.cpp +++ b/interfaces/hip/Memory.cpp @@ -52,27 +52,48 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) return devPtr; } +size_t ConcreteAPI::forgetAllocation(void* devPtr) { + const auto entry = memToSizeMap.find(devPtr); + if (entry == memToSizeMap.end()) { + assert(false && "DEVICE: an attempt to delete mem. which has not been allocated. unknown " + "pointer"); + return 0; + } + + const auto size = entry->second; + memToSizeMap.erase(entry); + statistics.deallocatedMemBytes += size; + return size; +} + void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(hipFree(devPtr)); } void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + if (devPtr == nullptr) { + return; + } + + const auto size = forgetAllocation(devPtr); + statistics.allocatedUnifiedMemBytes -= size; APIWRAP(hipFree(devPtr)); } void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); - assert((memToSizeMap.find(devPtr) != memToSizeMap.end()) && - "DEVICE: an attempt to delete mem. which has not been allocated. unknown pointer"); - statistics.deallocatedMemBytes += memToSizeMap[devPtr]; + if (devPtr == nullptr) { + return; + } + + forgetAllocation(devPtr); APIWRAP(hipHostFree(devPtr)); } diff --git a/interfaces/sycl/Memory.cpp b/interfaces/sycl/Memory.cpp index 860365c..a379f16 100644 --- a/interfaces/sycl/Memory.cpp +++ b/interfaces/sycl/Memory.cpp @@ -35,7 +35,7 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) return ptr; } -void ConcreteAPI::freeMem(void* devPtr) { +void ConcreteAPI::freeMem(void* devPtr, bool unified) { // NOTE: Freeing nullptr results in segfault in oneAPI. It is an opposite behaviour // contrast to C++/CUDA/HIP if (devPtr == nullptr) { @@ -61,7 +61,11 @@ void ConcreteAPI::freeMem(void* devPtr) { return; // the std::throw is throwing some errors during the program finalization } - context->statistics.deallocatedMemBytes += map.at(devPtr); + const auto size = map.at(devPtr); + context->statistics.deallocatedMemBytes += size; + if (unified) { + context->statistics.allocatedUnifiedMemBytes -= size; + } map.erase(devPtr); auto& queue = context->queueBuffer.getDefaultQueue(); sycl::free(devPtr, queue.get_context()); @@ -80,7 +84,7 @@ void ConcreteAPI::freeUnifiedMem(void* devPtr) { // NOTE: Freeing nullptr results in segfault in oneAPI. It is an opposite behavior // contrast to C++/CUDA/HIP if (devPtr != nullptr) { - this->freeMem(devPtr); + this->freeMem(devPtr, true); } } diff --git a/interfaces/sycl/SyclWrappedAPI.h b/interfaces/sycl/SyclWrappedAPI.h index 0f41d8d..46c3c8c 100644 --- a/interfaces/sycl/SyclWrappedAPI.h +++ b/interfaces/sycl/SyclWrappedAPI.h @@ -168,7 +168,7 @@ class ConcreteAPI : public AbstractAPI { return this->currentContext()->memoryToSizeMap; } - void freeMem(void* devPtr); + void freeMem(void* devPtr, bool unified = false); void initDevices(); From e9301e25d80755b2a4a1790b445a0d3b74f28096 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:45:14 +0200 Subject: [PATCH 19/31] fix: return early on an empty range A reduction over zero elements computed a grid of zero blocks, which the runtime rejects as an invalid launch configuration, and the SYCL algorithms built an nd_range with a global size of zero. Where overrideResult is set, the result is still initialized before the launch is skipped. AI-generated. Model: Opus 5 --- algorithms/cudahip/Reduction.cpp | 7 +++++++ algorithms/sycl/ArrayManip.cpp | 20 ++++++++++++++++++++ algorithms/sycl/BatchManip.cpp | 30 ++++++++++++++++++++++++++++++ algorithms/sycl/Reduction.cpp | 5 +++++ 4 files changed, 62 insertions(+) diff --git a/algorithms/cudahip/Reduction.cpp b/algorithms/cudahip/Reduction.cpp index 025a9b6..40001af 100644 --- a/algorithms/cudahip/Reduction.cpp +++ b/algorithms/cudahip/Reduction.cpp @@ -220,6 +220,13 @@ void Algorithms::reduceVector(AccT* result, } } + // the result is set either way, but there is nothing to reduce into it, and a grid of zero + // blocks is not a valid launch configuration + if (size == 0) { + CHECK_ERR; + return; + } + switch (type) { case ReductionType::Add: { kernel_reduce<<>>(result, buffer, size, device::Sum()); diff --git a/algorithms/sycl/ArrayManip.cpp b/algorithms/sycl/ArrayManip.cpp index 4a72d23..5b68249 100644 --- a/algorithms/sycl/ArrayManip.cpp +++ b/algorithms/sycl/ArrayManip.cpp @@ -14,6 +14,11 @@ using namespace device::internals; namespace device { template void Algorithms::scaleArray(T* devArray, T scalar, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -44,6 +49,11 @@ template void template void Algorithms::fillArray(T* devArray, const T scalar, const size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -72,6 +82,11 @@ template void Algorithms::fillArray(char* devArray, char scalar, const size_t numElements, void* streamPtr); void Algorithms::touchMemoryI(void* ptr, size_t size, bool clean, void* streamPtr) { + // an empty range is not a valid launch configuration + if (size == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, size); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -94,6 +109,11 @@ void Algorithms::incrementalAddI( uintptr_t* oout = reinterpret_cast(out); uintptr_t obase = reinterpret_cast(base); + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = computeExecutionRange1D(device::internals::DefaultBlockDim, numElements); ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { diff --git a/algorithms/sycl/BatchManip.cpp b/algorithms/sycl/BatchManip.cpp index 0aab2b3..7531cb6 100644 --- a/algorithms/sycl/BatchManip.cpp +++ b/algorithms/sycl/BatchManip.cpp @@ -17,6 +17,11 @@ void Algorithms::streamBatchedDataI(const void** baseSrcPtr, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -39,6 +44,11 @@ void Algorithms::streamBatchedDataI(const void** baseSrcPtr, template void Algorithms::accumulateBatchedData( const T** baseSrcPtr, T** baseDstPtr, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -69,6 +79,11 @@ template void Algorithms::accumulateBatchedData(const double** baseSrcPtr, void Algorithms::touchBatchedMemoryI( void** basePtr, size_t elementSize, size_t numElements, bool clean, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -93,6 +108,11 @@ void Algorithms::touchBatchedMemoryI( template void Algorithms::setToValue( T** out, T value, size_t elementSize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { @@ -134,6 +154,11 @@ void Algorithms::copyUniformToScatterI(const void* src, size_t copySize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; @@ -158,6 +183,11 @@ void Algorithms::copyScatterToUniformI(const void** src, size_t copySize, size_t numElements, void* streamPtr) { + // an empty range is not a valid launch configuration + if (numElements == 0) { + return; + } + auto rng = sycl::nd_range<1>{numElements * device::internals::DefaultBlockDim, device::internals::DefaultBlockDim}; diff --git a/algorithms/sycl/Reduction.cpp b/algorithms/sycl/Reduction.cpp index 9ee0899..2951513 100644 --- a/algorithms/sycl/Reduction.cpp +++ b/algorithms/sycl/Reduction.cpp @@ -90,6 +90,11 @@ void launchReduction(AccT* result, }); } + // the result is set either way, but there is nothing to reduce into it + if (size == 0) { + return; + } + ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { const size_t numWorkGroups = (size + (workGroupSize * itemsPerWorkItem) - 1) / (workGroupSize * itemsPerWorkItem); From e5e8e8dd85e0fd6ee8ad71e1272bda7901e0d6c2 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:45:35 +0200 Subject: [PATCH 20/31] fix: pick the access width by the alignment that is there imemcpy and imemset opened with 16-byte vector accesses on the pointers they were handed. Batched buffers are addressed through a pointer table and an element stride, so an element only lands on a 16-byte boundary when the stride happens to be a multiple of 16, and a vector access below its own alignment faults. The width now steps down to what the addresses actually allow, and the narrower passes cover the rest as before. The two scatter kernels also dereferenced their table entries without checking them, unlike streamBatchedData next to them; accumulateBatchedData and setToValue did the same. AI-generated. Model: Opus 5 --- algorithms/Common.h | 43 ++++++++++++++++++++++++++----- algorithms/cudahip/BatchManip.cpp | 14 ++++++++-- algorithms/sycl/BatchManip.cpp | 24 ++++++++++------- 3 files changed, 63 insertions(+), 18 deletions(-) diff --git a/algorithms/Common.h b/algorithms/Common.h index 762b4f0..b7d9983 100644 --- a/algorithms/Common.h +++ b/algorithms/Common.h @@ -9,6 +9,7 @@ #include "Algorithms.h" #include "Internals.h" +#include #include #if defined(__ACPP__) @@ -140,12 +141,35 @@ DEVICE_DEVICEFUNC std::size_t iimemcpy(void* dst, return end * sizeof(T); } +/** + * Returns the address bits that keep the given pointers from being 16-byte aligned; zero means + * both of them are. A null pointer contributes nothing, which is how the single-pointer case is + * asked. + * + * The copy and fill routines below step down from 16-byte accesses to single bytes. Batched + * buffers are addressed through a pointer table and an element stride, so their elements are not + * guaranteed to sit on a 16-byte boundary, and a vector access to an address that is not aligned + * to its own width faults. + */ +DEVICE_DEVICEFUNC std::size_t unalignedBits(const void* first, const void* second) { + const auto bits = + reinterpret_cast(first) | reinterpret_cast(second); + return static_cast(bits & 15U); +} + DEVICE_DEVICEFUNC void imemcpy(void* dst, const void* src, std::size_t count, int local, std::size_t stride) { + const auto lowBits = unalignedBits(dst, src); std::size_t offset = 0; - offset += iimemcpy(dst, src, offset, count, local, stride); - offset += iimemcpy(dst, src, offset, count, local, stride); - offset += iimemcpy(dst, src, offset, count, local, stride); + if (lowBits % sizeof(LocalInt4) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } + if (lowBits % sizeof(LocalInt2) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } + if (lowBits % sizeof(int) == 0) { + offset += iimemcpy(dst, src, offset, count, local, stride); + } offset += iimemcpy(dst, src, offset, count, local, stride); } @@ -165,10 +189,17 @@ DEVICE_DEVICEFUNC std::size_t } DEVICE_DEVICEFUNC void imemset(void* dst, std::size_t count, int local, std::size_t stride) { + const auto lowBits = unalignedBits(dst, nullptr); std::size_t offset = 0; - offset += iimemset(dst, offset, count, local, stride); - offset += iimemset(dst, offset, count, local, stride); - offset += iimemset(dst, offset, count, local, stride); + if (lowBits % sizeof(LocalInt4) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } + if (lowBits % sizeof(LocalInt2) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } + if (lowBits % sizeof(int) == 0) { + offset += iimemset(dst, offset, count, local, stride); + } offset += iimemset(dst, offset, count, local, stride); } diff --git a/algorithms/cudahip/BatchManip.cpp b/algorithms/cudahip/BatchManip.cpp index 9ce8b04..ce060ed 100644 --- a/algorithms/cudahip/BatchManip.cpp +++ b/algorithms/cudahip/BatchManip.cpp @@ -45,6 +45,9 @@ __global__ void kernel_accumulateBatchedData(const T** baseSrcPtr, for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { const T* srcElement = baseSrcPtr[block]; T* dstElement = baseDstPtr[block]; + if (srcElement == nullptr || dstElement == nullptr) { + continue; + } #pragma unroll 4 for (int index = threadIdx.x; index < elementSize; index += device::internals::DefaultBlockDim) { @@ -105,6 +108,9 @@ template __global__ void kernel_setToValue(T** out, T value, size_t elementSize, size_t elementCount) { for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { T* element = out[block]; + if (element == nullptr) { + continue; + } #pragma unroll 4 for (int index = threadIdx.x; index < elementSize; index += device::internals::DefaultBlockDim) { @@ -148,7 +154,9 @@ __global__ void kernel_copyUniformToScatter( const void* srcElement = reinterpret_cast(&reinterpret_cast(src)[block * srcOffset]); void* dstElement = dst[block]; - imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + if (dstElement != nullptr) { + imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + } } } @@ -172,7 +180,9 @@ __global__ void kernel_copyScatterToUniform( for (size_t block = blockIdx.x; block < elementCount; block += gridDim.x) { const void* srcElement = src[block]; void* dstElement = reinterpret_cast(&reinterpret_cast(dst)[block * dstOffset]); - imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + if (srcElement != nullptr) { + imemcpy(dstElement, srcElement, copySize, threadIdx.x, device::internals::DefaultBlockDim); + } } } diff --git a/algorithms/sycl/BatchManip.cpp b/algorithms/sycl/BatchManip.cpp index 7531cb6..f129a47 100644 --- a/algorithms/sycl/BatchManip.cpp +++ b/algorithms/sycl/BatchManip.cpp @@ -168,11 +168,13 @@ void Algorithms::copyUniformToScatterI(const void* src, const void* srcElement = reinterpret_cast(&reinterpret_cast(src)[block * srcOffset]); void* dstElement = dst[block]; - imemcpy(dstElement, - srcElement, - copySize, - item.get_local_id(0), - device::internals::DefaultBlockDim); + if (dstElement != nullptr) { + imemcpy(dstElement, + srcElement, + copySize, + item.get_local_id(0), + device::internals::DefaultBlockDim); + } }); }); } @@ -196,11 +198,13 @@ void Algorithms::copyScatterToUniformI(const void** src, const auto block = item.get_group().get_group_id(0); const void* srcElement = src[block]; void* dstElement = reinterpret_cast(&reinterpret_cast(dst)[block * dstOffset]); - imemcpy(dstElement, - srcElement, - copySize, - item.get_local_id(0), - device::internals::DefaultBlockDim); + if (srcElement != nullptr) { + imemcpy(dstElement, + srcElement, + copySize, + item.get_local_id(0), + device::internals::DefaultBlockDim); + } }); }); } From 06f84c3969ef42d2f2ec681e6f0499b7206f0f37 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:46:00 +0200 Subject: [PATCH 21/31] fix: only promise SYCL event timings where they exist createEvent dropped its withTiming argument, and no queue carried sycl::property::queue::enable_profiling, so every timespanEvents call ended in an exception thrown from inside get_profiling_info. The queues take the profiling property when the module is built with ENABLE_PROFILING_MARKERS, which is the switch that already stands for "timings are wanted, the per-submission cost is acceptable". Without it, timespanEvents now says so instead of throwing, and the flag createEvent was given is kept and checked. AI-generated. Model: Opus 5 --- interfaces/sycl/DeviceQueues.cpp | 19 ++++++++++++++++--- interfaces/sycl/Events.cpp | 16 +++++++++++++++- 2 files changed, 31 insertions(+), 4 deletions(-) diff --git a/interfaces/sycl/DeviceQueues.cpp b/interfaces/sycl/DeviceQueues.cpp index 1050740..6849955 100644 --- a/interfaces/sycl/DeviceQueues.cpp +++ b/interfaces/sycl/DeviceQueues.cpp @@ -18,12 +18,25 @@ namespace device { // very inconvenient, but AdaptiveCpp doesn't allow much freedom when constructing a property_list #if defined(DEVICE_USE_GRAPH_CAPTURING) && defined(SYCL_EXT_INTEL_QUEUE_IMMEDIATE_COMMAND_LIST) -#define BASE_QUEUE_PROPERTIES \ - sycl::property::queue::in_order{}, sycl::ext::intel::property::queue::no_immediate_command_list {} +#define IMMEDIATE_COMMAND_LIST_PROPERTY \ + , sycl::ext::intel::property::queue::no_immediate_command_list {} +#else +#define IMMEDIATE_COMMAND_LIST_PROPERTY +#endif + +// profiling is what makes the timings of AbstractAPI::timespanEvents available, and it costs on +// every submission, so it follows the build option rather than being on by default +#ifdef PROFILING_ENABLED +#define PROFILING_PROPERTY \ + , sycl::property::queue::enable_profiling {} #else -#define BASE_QUEUE_PROPERTIES sycl::property::queue::in_order() +#define PROFILING_PROPERTY #endif +#define BASE_QUEUE_PROPERTIES \ + sycl::property::queue::in_order {} \ + IMMEDIATE_COMMAND_LIST_PROPERTY PROFILING_PROPERTY + DeviceQueues::DeviceQueues(const sycl::device& dev, const std::function& handler) : defaultQueue{dev, handler, sycl::property_list{BASE_QUEUE_PROPERTIES}}, deviceReference(dev), diff --git a/interfaces/sycl/Events.cpp b/interfaces/sycl/Events.cpp index 27dfb0d..02691c1 100644 --- a/interfaces/sycl/Events.cpp +++ b/interfaces/sycl/Events.cpp @@ -19,15 +19,28 @@ using namespace device::internals; namespace { struct Event { std::optional syclEvent; + bool withTiming{false}; }; } // namespace -void* ConcreteAPI::createEvent(bool withTiming) { return new Event(); } +void* ConcreteAPI::createEvent(bool withTiming) { return new Event{std::nullopt, withTiming}; } double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) { +#ifndef PROFILING_ENABLED + // Profiling information is only there when the queue carries the enable_profiling property, and + // asking an event without it throws from deep inside the runtime. + logError() << "The SYCL backend can only time events when it is built with " + "ENABLE_PROFILING_MARKERS=ON."; + return 0.0; +#else auto* start = static_cast(eventPtrStart); auto* end = static_cast(eventPtrEnd); + if (!(start->withTiming && end->withTiming)) { + logError() << "Timing was not requested for at least one of the events given for timing " + "calculation."; + } + if (!(start->syclEvent.has_value() && end->syclEvent.has_value())) { logError() << "Invalid events given for timing calculation."; } @@ -40,6 +53,7 @@ double ConcreteAPI::timespanEvents(void* eventPtrStart, void* eventPtrEnd) { // cf. https://oneapi-src.github.io/SYCLomatic/dev_guide/reference/diagnostic_ref/dpct1012.html return static_cast(endTime - startTime) / 1'000'000'000.0; +#endif } void ConcreteAPI::destroyEvent(void* eventPtr) { From 7397ffae74ee933d2be1169571a052754a356664 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:46:22 +0200 Subject: [PATCH 22/31] fix: report the node API as available only where it compiles isCapableOfGraphNodes answered for the graph capturing macro, while the node functions call cudaStreamBeginCaptureToGraph (CUDA 12.3) and hipStreamBeginCaptureToGraph (ROCm 6.3). On an older toolkit that is a compile error rather than a backend that says it cannot do it, so the node API gets its own macro and the capability follows that one. Three smaller things on the way through: capturing an empty list of streams indexed into it, a second instantiation of the same graph leaked the first executable graph, and the graph returned by end capture is now checked against the one that was captured into. The stream list is taken by const reference, since capturing does not change it. AI-generated. Model: Opus 5 --- AbstractAPI.h | 2 +- interfaces/cuda/CudaWrappedAPI.h | 2 +- interfaces/cuda/Graphs.cu | 29 +++++++++++++++++++++-------- interfaces/hip/Graphs.cpp | 30 ++++++++++++++++++++++-------- interfaces/hip/HipWrappedAPI.h | 2 +- interfaces/sycl/Graphs.cpp | 9 ++++++++- interfaces/sycl/SyclWrappedAPI.h | 2 +- 7 files changed, 55 insertions(+), 21 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index 90706ca..7ed505d 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -85,7 +85,7 @@ struct AbstractAPI { virtual void syncDefaultStreamWithHost() = 0; virtual bool isCapableOfGraphCapturing() = 0; - virtual DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) = 0; + virtual DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) = 0; virtual void streamEndCapture(const DeviceGraphHandle& handle) = 0; virtual void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) = 0; diff --git a/interfaces/cuda/CudaWrappedAPI.h b/interfaces/cuda/CudaWrappedAPI.h index 3d616af..451f342 100644 --- a/interfaces/cuda/CudaWrappedAPI.h +++ b/interfaces/cuda/CudaWrappedAPI.h @@ -82,7 +82,7 @@ class ConcreteAPI : public AbstractAPI { void syncDefaultStreamWithHost() override; bool isCapableOfGraphCapturing() override; - DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index c87575c..7724aaf 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -14,6 +14,12 @@ #include #include +// Explicit graph nodes rest on cudaStreamBeginCaptureToGraph, which the runtime gained in CUDA +// 12.3. Capturing whole streams works without it, so the two get their own macro. +#if defined(DEVICE_USE_GRAPH_CAPTURING) && (CUDART_VERSION >= 12030) +#define DEVICE_USE_GRAPH_NODES +#endif + using namespace device; /* Two ways of building a compute graph are offered. @@ -71,16 +77,20 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } bool ConcreteAPI::isCapableOfGraphNodes() { -#ifdef DEVICE_USE_GRAPH_CAPTURING - // requires cudaStreamBeginCaptureToGraph, i.e. CUDA >= 12.3 +#ifdef DEVICE_USE_GRAPH_NODES return true; #else return false; #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING + if (streamPtrs.empty()) { + logError() << "Graph capturing records streams, so it needs at least one."; + return DeviceGraphHandle(); + } + auto graphInstance = std::make_shared(); graphInstance->streamPtrs = streamPtrs; @@ -97,6 +107,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = handle.get(); assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); APIWRAP(cudaStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), &(graphInstance->graph))); @@ -109,7 +120,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { } DeviceGraphHandle ConcreteAPI::graphCreate() { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto graphInstance = std::make_shared(); APIWRAP(cudaGraphCreate(&(graphInstance->graph), 0)); return DeviceGraphHandle(std::move(graphInstance)); @@ -119,7 +130,7 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { } namespace { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES /** * 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. @@ -152,7 +163,7 @@ std::vector captureFrontier(cudaStream_t stream) { void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, const std::vector& dependencies, void* streamPtr) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES 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"); @@ -175,7 +186,7 @@ void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, void* streamPtr) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a node must be opened before it can be closed"); @@ -184,6 +195,7 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa cudaGraph_t endedGraph{nullptr}; APIWRAP(cudaStreamEndCapture(stream, &endedGraph)); + assert(endedGraph == graphInstance->graph && "capturing into a graph hands that same graph back"); graphInstance->nodes.emplace_back(std::move(produced)); return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); @@ -193,9 +205,10 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa } void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); APIWRAP( cudaGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index 62c58ef..6f3e798 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -14,6 +14,13 @@ #include #include +// Explicit graph nodes rest on hipStreamBeginCaptureToGraph, which HIP gained in ROCm 6.3. +// Capturing whole streams works without it, so the two get their own macro. +#if defined(DEVICE_USE_GRAPH_CAPTURING) && \ + (HIP_VERSION_MAJOR > 6 || (HIP_VERSION_MAJOR == 6 && HIP_VERSION_MINOR >= 3)) +#define DEVICE_USE_GRAPH_NODES +#endif + using namespace device; /* Two ways of building a compute graph are offered. @@ -71,16 +78,20 @@ bool ConcreteAPI::isCapableOfGraphCapturing() { } bool ConcreteAPI::isCapableOfGraphNodes() { -#ifdef DEVICE_USE_GRAPH_CAPTURING - // requires hipStreamBeginCaptureToGraph, i.e. ROCm >= 6.3 +#ifdef DEVICE_USE_GRAPH_NODES return true; #else return false; #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING + if (streamPtrs.empty()) { + logError() << "Graph capturing records streams, so it needs at least one."; + return DeviceGraphHandle(); + } + auto graphInstance = std::make_shared(); graphInstance->streamPtrs = streamPtrs; @@ -97,6 +108,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING auto* graphInstance = handle.get(); assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); APIWRAP(hipStreamEndCapture(static_cast(graphInstance->streamPtrs[0]), &(graphInstance->graph))); @@ -109,7 +121,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { } DeviceGraphHandle ConcreteAPI::graphCreate() { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto graphInstance = std::make_shared(); APIWRAP(hipGraphCreate(&(graphInstance->graph), 0)); return DeviceGraphHandle(std::move(graphInstance)); @@ -119,7 +131,7 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { } namespace { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES /** * 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. @@ -142,7 +154,7 @@ std::vector captureFrontier(hipStream_t stream) { void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, const std::vector& dependencies, void* streamPtr) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES 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"); @@ -166,7 +178,7 @@ void ConcreteAPI::graphBeginNode(const DeviceGraphHandle& graphHandle, DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHandle, void* streamPtr) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a node must be opened before it can be closed"); @@ -175,6 +187,7 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa hipGraph_t endedGraph{nullptr}; APIWRAP(hipStreamEndCapture(stream, &endedGraph)); + assert(endedGraph == graphInstance->graph && "capturing into a graph hands that same graph back"); graphInstance->nodes.emplace_back(std::move(produced)); return DeviceGraphNodeHandle(graphInstance->nodes.size() - 1); @@ -184,9 +197,10 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa } void ConcreteAPI::graphInstantiate(const DeviceGraphHandle& graphHandle) { -#ifdef DEVICE_USE_GRAPH_CAPTURING +#ifdef DEVICE_USE_GRAPH_NODES auto* graphInstance = graphHandle.get(); assert(graphInstance != nullptr && "a graph must be created before it is instantiated"); + assert(graphInstance->instance == nullptr && "a graph is instantiated once"); APIWRAP( hipGraphInstantiate(&(graphInstance->instance), graphInstance->graph, nullptr, nullptr, 0)); diff --git a/interfaces/hip/HipWrappedAPI.h b/interfaces/hip/HipWrappedAPI.h index 075cead..953b8b6 100644 --- a/interfaces/hip/HipWrappedAPI.h +++ b/interfaces/hip/HipWrappedAPI.h @@ -81,7 +81,7 @@ class ConcreteAPI : public AbstractAPI { void syncDefaultStreamWithHost() override; bool isCapableOfGraphCapturing() override; - DeviceGraphHandle streamBeginCapture(std::vector& streamPtrs) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; diff --git a/interfaces/sycl/Graphs.cpp b/interfaces/sycl/Graphs.cpp index d14fc49..8122501 100644 --- a/interfaces/sycl/Graphs.cpp +++ b/interfaces/sycl/Graphs.cpp @@ -70,8 +70,13 @@ bool ConcreteAPI::isCapableOfGraphNodes() { #endif } -DeviceGraphHandle ConcreteAPI::streamBeginCapture(std::vector& streamPtrs) { +DeviceGraphHandle ConcreteAPI::streamBeginCapture(const std::vector& streamPtrs) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT + if (streamPtrs.empty()) { + logError() << "Graph capturing records queues, so it needs at least one."; + return DeviceGraphHandle(); + } + std::vector queues; queues.reserve(streamPtrs.size()); for (auto* streamPtr : streamPtrs) { @@ -92,6 +97,7 @@ void ConcreteAPI::streamEndCapture(const DeviceGraphHandle& handle) { #ifdef DEVICE_USE_GRAPH_CAPTURING_ONEAPI_EXT auto* graphInstance = handle.get(); assert(graphInstance != nullptr && "a capture must be started before it can be ended"); + assert(!graphInstance->instance.has_value() && "a graph is instantiated once"); graphInstance->graph.end_recording(); graphInstance->instance = std::optionalinstance.has_value() && "a graph is instantiated once"); graphInstance->graph.end_recording(); graphInstance->instance = std::optional& streamPtrs) override; + DeviceGraphHandle streamBeginCapture(const std::vector& streamPtrs) override; void streamEndCapture(const DeviceGraphHandle& handle) override; void launchGraph(const DeviceGraphHandle& graphHandle, void* streamPtr) override; From 8bc46ab82b80697e00699bb2c7615bc4740f97ee Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:46:54 +0200 Subject: [PATCH 23/31] fix: state the same waiting and recording rules on every backend streamWaitMemory waited for the value to be equal on SYCL and for it to be at least as large on CUDA and HIP, which are two different things for a counter that a producer keeps incrementing. AbstractAPI now says which one callers get, and what kind of memory the location has to be. The note on graphAddNode promised that sibling nodes may share a stream. That holds where the backend can name node handles; the SYCL graph extension expresses edges through the recorded queue, so two nodes sharing a queue end up ordered there. The note now asks for what all three backends can keep. AI-generated. Model: Opus 5 --- AbstractAPI.h | 15 +++++++++++++-- interfaces/sycl/Streams.cpp | 4 +++- 2 files changed, 16 insertions(+), 3 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index 7ed505d..86a5077 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -99,8 +99,14 @@ struct AbstractAPI { * 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. + * recording vehicle: what it enqueues becomes the node, and the dependencies stated in the call + * are the ones the node is guaranteed to get. + * + * Nodes that are meant to run concurrently have to be recorded onto different streams. A stream + * can be reused for a later node, but a backend that expresses edges through the recorded + * stream rather than through node handles - the SYCL one does, since the graph extension has no + * node handles to hand out - adds an edge between two nodes that shared a stream, and those two + * then run one after the other. * * 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. @@ -139,6 +145,11 @@ struct AbstractAPI { virtual void syncStreamWithEvent(void* streamPtr, void* eventPtr) = 0; virtual void streamHostFunction(void* streamPtr, const std::function& function) = 0; + /** + * Blocks the stream until the value at `location` has reached at least `value`. `location` has + * to be host memory that the device can read, i.e. an allocation from allocPinnedMem with + * Destination::CurrentDevice. + */ virtual void streamWaitMemory(void* streamPtr, uint32_t* location, uint32_t value) = 0; virtual void* createEvent(bool withTiming = false) = 0; diff --git a/interfaces/sycl/Streams.cpp b/interfaces/sycl/Streams.cpp index c09e1f7..aaa6a32 100644 --- a/interfaces/sycl/Streams.cpp +++ b/interfaces/sycl/Streams.cpp @@ -75,7 +75,9 @@ void ConcreteAPI::streamWaitMemory(void* streamPtr, uint32_t* location, uint32_t volatile uint32_t* spinLocation = location; queuePtr->single_task([=]() { while (true) { - if (*spinLocation == value) { + // ">=", like the wait-value operations of the other backends: a counter that is written + // once per step is past the awaited value by the time this runs often enough + if (*spinLocation >= value) { return; } From 684edb25b833e99450d48df50735c991be6b4e66 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:47:22 +0200 Subject: [PATCH 24/31] fix: destroy only the streams the backend handed out destroyGenericStream removed the stream from the set when it found it there and destroyed it either way, so a pointer that came from somewhere else - the default stream, or a stream already destroyed - was passed to the runtime regardless. It now warns and returns, which is what the SYCL backend does. finalize sets m_isFinalized, so hasFinalized() answers on all three backends, and clears the set of streams it just destroyed. AI-generated. Model: Opus 5 --- interfaces/cuda/Control.cu | 4 ++++ interfaces/cuda/Streams.cu | 12 ++++++++++-- interfaces/hip/Control.cpp | 5 ++++- interfaces/hip/Streams.cpp | 12 ++++++++++-- 4 files changed, 28 insertions(+), 5 deletions(-) diff --git a/interfaces/cuda/Control.cu b/interfaces/cuda/Control.cu index c1e5dcb..91fb889 100644 --- a/interfaces/cuda/Control.cu +++ b/interfaces/cuda/Control.cu @@ -91,15 +91,19 @@ void ConcreteAPI::finalize() { CHECK_ERR; APIWRAP(cudaStreamDestroy(defaultStream)); + defaultStream = nullptr; + if (!genericStreams.empty()) { logInfo() << "DEVICE::WARNING:" << genericStreams.size() << "device generic stream(s) were not deleted."; for (auto stream : genericStreams) { APIWRAP(cudaStreamDestroy(stream)); } + genericStreams.clear(); } status[StatusID::InterfaceInitialized] = false; } + m_isFinalized = true; } int ConcreteAPI::getNumDevices() { return properties.size(); } diff --git a/interfaces/cuda/Streams.cu b/interfaces/cuda/Streams.cu index 86cc162..72bd0bd 100644 --- a/interfaces/cuda/Streams.cu +++ b/interfaces/cuda/Streams.cu @@ -38,10 +38,18 @@ void* ConcreteAPI::createStream(double priority) { void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); cudaStream_t stream = static_cast(streamPtr); + + // The stream has to leave the set before it is destroyed, and a stream that is not in it is not + // this backend's to destroy - the default stream, for one, would take the whole interface with + // it. auto it = genericStreams.find(stream); - if (it != genericStreams.end()) { - genericStreams.erase(it); + if (it == genericStreams.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or was not created here; not destroying it."; + return; } + + genericStreams.erase(it); APIWRAP(cudaStreamDestroy(stream)); } diff --git a/interfaces/hip/Control.cpp b/interfaces/hip/Control.cpp index 85c76a9..da87279 100644 --- a/interfaces/hip/Control.cpp +++ b/interfaces/hip/Control.cpp @@ -92,19 +92,22 @@ void ConcreteAPI::initialize() { void ConcreteAPI::finalize() { if (status[StatusID::InterfaceInitialized]) { - CHECK_ERR; APIWRAP(hipStreamDestroy(defaultStream)); + defaultStream = nullptr; + if (!genericStreams.empty()) { logInfo() << "DEVICE::WARNING:" << genericStreams.size() << "device generic stream(s) were not deleted."; for (auto stream : genericStreams) { APIWRAP(hipStreamDestroy(stream)); } + genericStreams.clear(); } status[StatusID::InterfaceInitialized] = false; } + m_isFinalized = true; } int ConcreteAPI::getNumDevices() { return properties.size(); } diff --git a/interfaces/hip/Streams.cpp b/interfaces/hip/Streams.cpp index 504d6f2..f06f62b 100644 --- a/interfaces/hip/Streams.cpp +++ b/interfaces/hip/Streams.cpp @@ -37,10 +37,18 @@ void* ConcreteAPI::createStream(double priority) { void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); hipStream_t stream = static_cast(streamPtr); + + // The stream has to leave the set before it is destroyed, and a stream that is not in it is not + // this backend's to destroy - the default stream, for one, would take the whole interface with + // it. auto it = genericStreams.find(stream); - if (it != genericStreams.end()) { - genericStreams.erase(it); + if (it == genericStreams.end()) { + logWarning() << "Tried to destroy a stream that this device does not know about. It has " + "either been destroyed already or was not created here; not destroying it."; + return; } + + genericStreams.erase(it); APIWRAP(hipStreamDestroy(stream)); } From 315b21ee0bc5b3b2e67611b378c122ccabda6a07 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:47:44 +0200 Subject: [PATCH 25/31] fix: guard the state that several threads reach genericStreams, memToSizeMap, allocationProperties, the statistics and the SYCL list of handed-out queues are all changed from whatever thread happens to call in, and syncDevice walks the last one while another thread may be adding to it. The mutex AbstractAPI already carries now guards them; the SYCL queues get their own, since they sit below the API object. syncAllQueuesWithHost copies the list under the lock and waits outside it, so waiting on one queue does not block creating another. The device id is thread-local, matching the runtime, which keeps its selected device per thread as well - so a worker thread that never called setDevice was working on device 0 whatever the process had selected. setDevice now also records the choice for the process, and a thread without one picks it up the first time it asks. AI-generated. Model: Opus 5 --- AbstractAPI.h | 5 +++++ interfaces/cuda/Control.cu | 17 ++++++++++++++++- interfaces/cuda/Memory.cu | 10 ++++++++++ interfaces/cuda/Streams.cu | 3 +++ interfaces/hip/Control.cpp | 11 +++++++++++ interfaces/hip/Memory.cpp | 10 ++++++++++ interfaces/hip/Streams.cpp | 3 +++ interfaces/sycl/DeviceQueues.cpp | 16 +++++++++++++++- interfaces/sycl/DeviceQueues.h | 3 +++ interfaces/sycl/Memory.cpp | 15 +++++++++++++++ 10 files changed, 91 insertions(+), 2 deletions(-) diff --git a/AbstractAPI.h b/AbstractAPI.h index 86a5077..f5b5178 100644 --- a/AbstractAPI.h +++ b/AbstractAPI.h @@ -33,6 +33,11 @@ enum class ProfilingColors : uint32_t { struct AbstractAPI { virtual ~AbstractAPI() = default; + /** + * Selects the device for the calling thread and makes it the choice of the process. Backends + * that keep the selected device per thread give a thread that has not called this the device + * the process selected, the first time that thread asks for the device id. + */ virtual void setDevice(int deviceId) = 0; virtual int getDeviceId() = 0; diff --git a/interfaces/cuda/Control.cu b/interfaces/cuda/Control.cu index 91fb889..408a228 100644 --- a/interfaces/cuda/Control.cu +++ b/interfaces/cuda/Control.cu @@ -5,6 +5,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -27,7 +28,11 @@ namespace { #ifdef DEVICE_CONTEXT_GLOBAL int currentDeviceId = 0; #else -thread_local int currentDeviceId = 0; +// The runtime keeps the selected device per thread, so a thread that has not selected one works +// on device 0 - which is the wrong card whenever the process picked another one. The device the +// process selected is kept here and picked up on the first request from a thread that has none. +std::atomic selectedDeviceId{0}; +thread_local int currentDeviceId = -1; #endif } // namespace @@ -52,6 +57,9 @@ ConcreteAPI::ConcreteAPI() { void ConcreteAPI::setDevice(int deviceId) { currentDeviceId = deviceId; +#ifndef DEVICE_CONTEXT_GLOBAL + selectedDeviceId.store(deviceId, std::memory_order_relaxed); +#endif APIWRAP(cudaSetDevice(deviceId)); @@ -87,6 +95,7 @@ void ConcreteAPI::initialize() { } void ConcreteAPI::finalize() { + const std::lock_guard lock(apiMutex); if (status[StatusID::InterfaceInitialized]) { CHECK_ERR; @@ -112,6 +121,12 @@ int ConcreteAPI::getDeviceId() { if (!status[StatusID::DeviceSelected]) { logError() << "Device has not been selected. Please, select device before requesting device Id"; } +#ifndef DEVICE_CONTEXT_GLOBAL + if (currentDeviceId < 0) { + currentDeviceId = selectedDeviceId.load(std::memory_order_relaxed); + APIWRAP(cudaSetDevice(currentDeviceId)); + } +#endif return currentDeviceId; } diff --git a/interfaces/cuda/Memory.cu b/interfaces/cuda/Memory.cu index cc5d3aa..6587525 100644 --- a/interfaces/cuda/Memory.cu +++ b/interfaces/cuda/Memory.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include using namespace device; @@ -66,6 +67,7 @@ void driverFree(void* ptr, std::size_t size, const CUmemAllocationProp& prop) { void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; if (compress && canCompress) { CUmemAllocationProp prop = {}; @@ -87,6 +89,7 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; APIWRAP(cudaMallocManaged(&devPtr, size, cudaMemAttachGlobal)); @@ -128,6 +131,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr = nullptr; const auto flag = hint == Destination::Host ? cudaHostAllocDefault : cudaHostAllocMapped; APIWRAP(cudaHostAlloc(&devPtr, size, flag)); @@ -152,6 +156,7 @@ size_t ConcreteAPI::forgetAllocation(void* devPtr) { void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -171,6 +176,7 @@ void ConcreteAPI::freeGlobMem(void* devPtr) { void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -182,6 +188,7 @@ void ConcreteAPI::freeUnifiedMem(void* devPtr) { void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -207,6 +214,7 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { std::string ConcreteAPI::getMemLeaksReport() { isFlagSet(status); + const std::lock_guard lock(apiMutex); std::ostringstream report{}; report << "Memory Leaks, bytes: " << (statistics.allocatedMemBytes - statistics.deallocatedMemBytes) << '\n'; @@ -217,11 +225,13 @@ size_t ConcreteAPI::getMaxAvailableMem() { return properties[getDeviceId()].tota size_t ConcreteAPI::getCurrentlyOccupiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedUnifiedMemBytes; } diff --git a/interfaces/cuda/Streams.cu b/interfaces/cuda/Streams.cu index 72bd0bd..6a00fbf 100644 --- a/interfaces/cuda/Streams.cu +++ b/interfaces/cuda/Streams.cu @@ -9,6 +9,7 @@ #include #include #include +#include #include using namespace device; @@ -28,6 +29,7 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); + const std::lock_guard lock(apiMutex); cudaStream_t stream; const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); APIWRAP(cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, truePriority)); @@ -37,6 +39,7 @@ void* ConcreteAPI::createStream(double priority) { void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); cudaStream_t stream = static_cast(streamPtr); // The stream has to leave the set before it is destroyed, and a stream that is not in it is not diff --git a/interfaces/hip/Control.cpp b/interfaces/hip/Control.cpp index da87279..bffd63e 100644 --- a/interfaces/hip/Control.cpp +++ b/interfaces/hip/Control.cpp @@ -6,6 +6,7 @@ #include "utils/env.h" #include "utils/logger.h" +#include #include #include #include @@ -49,6 +50,9 @@ ConcreteAPI::ConcreteAPI() { void ConcreteAPI::setDevice(int deviceId) { currentDeviceId = deviceId; +#ifndef DEVICE_CONTEXT_GLOBAL + selectedDeviceId.store(deviceId, std::memory_order_relaxed); +#endif APIWRAP(hipSetDevice(deviceId)); @@ -91,6 +95,7 @@ void ConcreteAPI::initialize() { } void ConcreteAPI::finalize() { + const std::lock_guard lock(apiMutex); if (status[StatusID::InterfaceInitialized]) { CHECK_ERR; @@ -116,6 +121,12 @@ int ConcreteAPI::getDeviceId() { if (!status[StatusID::DeviceSelected]) { logError() << "Device has not been selected. Please, select device before requesting device Id"; } +#ifndef DEVICE_CONTEXT_GLOBAL + if (currentDeviceId < 0) { + currentDeviceId = selectedDeviceId.load(std::memory_order_relaxed); + APIWRAP(hipSetDevice(currentDeviceId)); + } +#endif return currentDeviceId; } diff --git a/interfaces/hip/Memory.cpp b/interfaces/hip/Memory.cpp index e102f88..b5fa2e5 100644 --- a/interfaces/hip/Memory.cpp +++ b/interfaces/hip/Memory.cpp @@ -9,12 +9,14 @@ #include #include +#include #include using namespace device; void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; APIWRAP(hipMalloc(&devPtr, size)); statistics.allocatedMemBytes += size; @@ -24,6 +26,7 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; APIWRAP(hipMallocManaged(&devPtr, size, hipMemAttachGlobal)); @@ -44,6 +47,7 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { isFlagSet(status); + const std::lock_guard lock(apiMutex); void* devPtr; const auto flag = hint == Destination::Host ? hipHostMallocDefault : hipHostMallocMapped; APIWRAP(hipHostMalloc(&devPtr, size, flag)); @@ -68,6 +72,7 @@ size_t ConcreteAPI::forgetAllocation(void* devPtr) { void ConcreteAPI::freeGlobMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -78,6 +83,7 @@ void ConcreteAPI::freeGlobMem(void* devPtr) { void ConcreteAPI::freeUnifiedMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -89,6 +95,7 @@ void ConcreteAPI::freeUnifiedMem(void* devPtr) { void ConcreteAPI::freePinnedMem(void* devPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); if (devPtr == nullptr) { return; } @@ -114,6 +121,7 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { std::string ConcreteAPI::getMemLeaksReport() { isFlagSet(status); + const std::lock_guard lock(apiMutex); std::ostringstream report{}; report << "Memory Leaks, bytes: " << (statistics.allocatedMemBytes - statistics.deallocatedMemBytes) << '\n'; @@ -124,11 +132,13 @@ size_t ConcreteAPI::getMaxAvailableMem() { return properties[getDeviceId()].tota size_t ConcreteAPI::getCurrentlyOccupiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { isFlagSet(status); + const std::lock_guard lock(apiMutex); return statistics.allocatedUnifiedMemBytes; } diff --git a/interfaces/hip/Streams.cpp b/interfaces/hip/Streams.cpp index f06f62b..c0c8fa3 100644 --- a/interfaces/hip/Streams.cpp +++ b/interfaces/hip/Streams.cpp @@ -8,6 +8,7 @@ #include #include +#include #include using namespace device; @@ -27,6 +28,7 @@ void ConcreteAPI::syncDefaultStreamWithHost() { void* ConcreteAPI::createStream(double priority) { isFlagSet(status); + const std::lock_guard lock(apiMutex); hipStream_t stream; const auto truePriority = mapStreamPriority(priorityLeast, priorityGreatest, priority); APIWRAP(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, truePriority)); @@ -36,6 +38,7 @@ void* ConcreteAPI::createStream(double priority) { void ConcreteAPI::destroyGenericStream(void* streamPtr) { isFlagSet(status); + const std::lock_guard lock(apiMutex); hipStream_t stream = static_cast(streamPtr); // The stream has to leave the set before it is destroyed, and a stream that is not in it is not diff --git a/interfaces/sycl/DeviceQueues.cpp b/interfaces/sycl/DeviceQueues.cpp index 6849955..4046e0f 100644 --- a/interfaces/sycl/DeviceQueues.cpp +++ b/interfaces/sycl/DeviceQueues.cpp @@ -10,6 +10,7 @@ #include #include +#include #include using namespace device::internals; @@ -75,12 +76,15 @@ sycl::queue* DeviceQueues::newQueue(double priority) { #endif auto* queue = new sycl::queue{deviceReference, handlerReference, propertylist}; + + const std::lock_guard lock(queueMutex); externalQueues.emplace_back(queue); return queue; } void DeviceQueues::deleteQueue(void* queue) { auto* queuePtr = static_cast(queue); + const std::lock_guard lock(queueMutex); // The queue has to leave the list before it is freed: syncAllQueuesWithHost walks that list, // so a stale entry turns into a use-after-free at the next device-wide synchronization, far @@ -101,7 +105,15 @@ void DeviceQueues::syncQueueWithHost(sycl::queue* queuePtr) { waitCheck(*queuePt void DeviceQueues::syncAllQueuesWithHost() { waitCheck(defaultQueue); - for (auto* queue : this->externalQueues) { + + // copied under the lock: waiting on a queue takes as long as the work on it, and holding the + // lock for that would block every thread that wants to create or destroy one + const auto queues = [this]() { + const std::lock_guard lock(queueMutex); + return externalQueues; + }(); + + for (auto* queue : queues) { waitCheck(*queue); } } @@ -110,6 +122,8 @@ bool DeviceQueues::exists(sycl::queue* queuePtr) { if (queuePtr == &defaultQueue) { return true; } + + const std::lock_guard lock(queueMutex); return std::find(externalQueues.begin(), externalQueues.end(), queuePtr) != externalQueues.end(); } diff --git a/interfaces/sycl/DeviceQueues.h b/interfaces/sycl/DeviceQueues.h index ccf9a41..0566251 100644 --- a/interfaces/sycl/DeviceQueues.h +++ b/interfaces/sycl/DeviceQueues.h @@ -6,6 +6,7 @@ #define SEISSOLDEVICE_INTERFACES_SYCL_DEVICEQUEUES_H_ #include +#include #include #include @@ -56,6 +57,8 @@ class DeviceQueues { private: sycl::queue defaultQueue; + // guards externalQueues, which callers add to and remove from while other threads walk it + std::mutex queueMutex; std::vector externalQueues; sycl::device deviceReference; std::function handlerReference; diff --git a/interfaces/sycl/Memory.cpp b/interfaces/sycl/Memory.cpp index a379f16..48b7c02 100644 --- a/interfaces/sycl/Memory.cpp +++ b/interfaces/sycl/Memory.cpp @@ -6,11 +6,14 @@ #include "SyclWrappedAPI.h" #include +#include using namespace device; using namespace device::internals; void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_device(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); @@ -19,6 +22,8 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { } void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_shared(size, this->currentDefaultQueue()); this->currentStatistics().allocatedUnifiedMemBytes += size; this->currentStatistics().allocatedMemBytes += size; @@ -28,6 +33,8 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) } void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) { + const std::lock_guard lock(apiMutex); + auto* ptr = malloc_host(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); @@ -50,6 +57,8 @@ void ConcreteAPI::freeMem(void* devPtr, bool unified) { return; } + const std::lock_guard lock(apiMutex); + // Use the first device context to free memory DeviceContext* context = this->availableDevices[getDeviceId()]; if (!context) { @@ -110,6 +119,8 @@ void ConcreteAPI::freeMemAsync(void* devPtr, void* streamPtr) { } std::string ConcreteAPI::getMemLeaksReport() { + const std::lock_guard lock(apiMutex); + std::ostringstream report{}; report << "----MEMORY REPORT----\n"; @@ -128,10 +139,14 @@ size_t ConcreteAPI::getMaxAvailableMem() { } size_t ConcreteAPI::getCurrentlyOccupiedMem() { + const std::lock_guard lock(apiMutex); + return this->currentStatistics().allocatedMemBytes; } size_t ConcreteAPI::getCurrentlyOccupiedUnifiedMem() { + const std::lock_guard lock(apiMutex); + return this->currentStatistics().allocatedUnifiedMemBytes; } From 0b3381fbf3930a72e5b8930d17f8a4d4da238daf Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:48:09 +0200 Subject: [PATCH 26/31] perf: keep the API wrapper off the heap Every wrapped API call built a std::string from __FILE__ - a heap allocation for any path longer than the small-string buffer - and an empty std::unordered_set, before finding out that the call had succeeded. Both parameters are now a pointer to the string literal and a view of the caller's temporary array, so the success path allocates nothing. The kernel launch path runs through this a few thousand times per time step. AI-generated. Model: Opus 5 --- interfaces/cuda/Internals.cu | 21 +++++++++++---------- interfaces/cuda/Internals.h | 14 ++++++++------ interfaces/hip/Internals.cpp | 13 +++++++------ interfaces/hip/Internals.h | 10 ++++++---- 4 files changed, 32 insertions(+), 26 deletions(-) diff --git a/interfaces/cuda/Internals.cu b/interfaces/cuda/Internals.cu index 419b27f..44c54e3 100644 --- a/interfaces/cuda/Internals.cu +++ b/interfaces/cuda/Internals.cu @@ -4,26 +4,27 @@ #include "utils/logger.h" +#include #include +#include #include #include -#include namespace device::internals { -thread_local std::string prevFile{}; +thread_local const char* prevFile{nullptr}; thread_local int prevLine{-1}; cudaError_t checkResult(cudaError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != cudaSuccess && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != cudaSuccess && std::find(except.begin(), except.end(), error) == except.end()) { std::stringstream stream; stream << '\n' << file << ", line " << line << ": " << cudaGetErrorString(error) << " (" << error << ")\n"; - if (prevLine >= 0) { + if (prevFile != nullptr) { stream << "Previous CUDA API/Driver call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } @@ -35,10 +36,10 @@ cudaError_t checkResult(cudaError_t error, } CUresult checkResultDriver(CUresult error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != CUDA_SUCCESS && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != CUDA_SUCCESS && std::find(except.begin(), except.end(), error) == except.end()) { const char* errstr = nullptr; const auto errstrRes = cuGetErrorString(error, &errstr); @@ -50,7 +51,7 @@ CUresult checkResultDriver(CUresult error, stream << "[ERROR WHILE RETRIEVING ERROR STRING]"; } stream << " (" << error << ")\n"; - if (prevLine >= 0) { + if (prevFile != nullptr) { stream << "Previous CUDA API/Driver call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } diff --git a/interfaces/cuda/Internals.h b/interfaces/cuda/Internals.h index 497c739..a2c3652 100644 --- a/interfaces/cuda/Internals.h +++ b/interfaces/cuda/Internals.h @@ -6,8 +6,7 @@ #define SEISSOLDEVICE_INTERFACES_CUDA_INTERNALS_H_ #include -#include -#include +#include #include #define APIWRAP(call) (void)::device::internals::checkResult(call, __FILE__, __LINE__, {}) @@ -22,14 +21,17 @@ using DeviceStreamT = cudaStream_t; constexpr static int DefaultBlockDim = 1024; +// Every wrapped call goes through here, so the parameters stay free of anything that allocates: +// the file name is the string literal __FILE__ expands to, and the accepted errors are read from +// the caller's temporary array. cudaError_t checkResult(cudaError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); CUresult checkResultDriver(CUresult error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); inline dim3 computeGrid1D(const dim3& block, const size_t size) { int numBlocks = (size + block.x - 1) / block.x; diff --git a/interfaces/hip/Internals.cpp b/interfaces/hip/Internals.cpp index 7be2d80..a435e03 100644 --- a/interfaces/hip/Internals.cpp +++ b/interfaces/hip/Internals.cpp @@ -6,25 +6,26 @@ #include "utils/logger.h" +#include +#include #include -#include #include namespace device::internals { -thread_local std::string prevFile{}; +thread_local const char* prevFile{nullptr}; thread_local int prevLine{0}; hipError_t checkResult(hipError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except) { - if (error != hipSuccess && except.find(error) == except.end()) { + std::initializer_list except) { + if (error != hipSuccess && std::find(except.begin(), except.end(), error) == except.end()) { std::stringstream stream; stream << '\n' << file << ", line " << line << ": " << hipGetErrorString(error) << " (" << error << ")\n"; - if (prevLine > 0) { + if (prevFile != nullptr) { stream << "Previous HIP call:" << std::endl << prevFile << ", line " << prevLine << std::endl; } logError() << stream.str(); diff --git a/interfaces/hip/Internals.h b/interfaces/hip/Internals.h index a3ececd..2c0dbe3 100644 --- a/interfaces/hip/Internals.h +++ b/interfaces/hip/Internals.h @@ -7,8 +7,7 @@ #include "hip/hip_runtime.h" -#include -#include +#include #define APIWRAP(call) (void)::device::internals::checkResult(call, __FILE__, __LINE__, {}) #define APIWRAPX(call, except) ::device::internals::checkResult(call, __FILE__, __LINE__, except) @@ -19,10 +18,13 @@ namespace device::internals { constexpr static int DefaultBlockDim = 1024; using DeviceStreamT = hipStream_t; +// Every wrapped call goes through here, so the parameters stay free of anything that allocates: +// the file name is the string literal __FILE__ expands to, and the accepted errors are read from +// the caller's temporary array. hipError_t checkResult(hipError_t error, - const std::string& file, + const char* file, int line, - const std::unordered_set& except); + std::initializer_list except); inline dim3 computeGrid1D(const dim3& block, const size_t size) { int numBlocks = (size + block.x - 1) / block.x; return dim3(numBlocks, 1, 1); From 010ed9590379eba30e87c491f03e6ec197120d02 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:48:33 +0200 Subject: [PATCH 27/31] perf: ask for the occupancy and the allocation sync once blockcount ran cudaGetDevice, a device attribute query and an occupancy query in front of every algorithm launch, for a number that does not change while the program runs. It is now computed once per kernel, and clamped to at least one block, since a grid of zero is not a valid launch configuration. The SYCL allocations each ended with a device-wide wait. malloc_device and its siblings return once the allocation exists, so there is nothing to wait for. The wait in freeMem stays - the caller has no way to say that no queue is reading the memory any more - and now happens before the free rather than after it. AI-generated. Model: Opus 5 --- algorithms/Common.h | 48 +++++++++++++++++++++++++------------- interfaces/sycl/Memory.cpp | 7 +++--- 2 files changed, 35 insertions(+), 20 deletions(-) diff --git a/algorithms/Common.h b/algorithms/Common.h index b7d9983..b553cf6 100644 --- a/algorithms/Common.h +++ b/algorithms/Common.h @@ -9,8 +9,10 @@ #include "Algorithms.h" #include "Internals.h" +#include #include #include +#include #if defined(__ACPP__) #include @@ -58,14 +60,21 @@ DEVICE_DEVICEFUNC T ntload(const T* location) { template int blockcount(F&& func, int blocksize) { - int device = 0; - int smCount = 0; - int blocksPerSM = 0; - APIWRAP(cudaGetDevice(&device)); - APIWRAP(cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, device)); - APIWRAP(cudaOccupancyMaxActiveBlocksPerMultiprocessor( - &blocksPerSM, std::forward(func), blocksize, 0)); - return smCount * blocksPerSM; + // The occupancy of a kernel does not change while the program runs, and this sits in front of + // every algorithm launch, so it is asked once per kernel. Every caller uses the default block + // size, and the devices of a node are the same model. + static const int count = [&]() { + int device = 0; + int smCount = 0; + int blocksPerSM = 0; + APIWRAP(cudaGetDevice(&device)); + APIWRAP(cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, device)); + APIWRAP(cudaOccupancyMaxActiveBlocksPerMultiprocessor( + &blocksPerSM, std::forward(func), blocksize, 0)); + // a grid of zero blocks is not a valid launch configuration + return std::max(1, smCount * blocksPerSM); + }(); + return count; } using LocalInt4 = int4; @@ -84,14 +93,21 @@ DEVICE_DEVICEFUNC T ntload(const T* location) { template int blockcount(F&& func, int blocksize) { - int device = 0; - int smCount = 0; - int blocksPerSM = 0; - APIWRAP(hipGetDevice(&device)); - APIWRAP(hipDeviceGetAttribute(&smCount, hipDeviceAttributeMultiprocessorCount, device)); - APIWRAP(hipOccupancyMaxActiveBlocksPerMultiprocessor( - &blocksPerSM, std::forward(func), blocksize, 0)); - return smCount * blocksPerSM; + // The occupancy of a kernel does not change while the program runs, and this sits in front of + // every algorithm launch, so it is asked once per kernel. Every caller uses the default block + // size, and the devices of a node are the same model. + static const int count = [&]() { + int device = 0; + int smCount = 0; + int blocksPerSM = 0; + APIWRAP(hipGetDevice(&device)); + APIWRAP(hipDeviceGetAttribute(&smCount, hipDeviceAttributeMultiprocessorCount, device)); + APIWRAP(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &blocksPerSM, std::forward(func), blocksize, 0)); + // a grid of zero blocks is not a valid launch configuration + return std::max(1, smCount * blocksPerSM); + }(); + return count; } using LocalInt4 = __attribute__((vector_size(16))) int; diff --git a/interfaces/sycl/Memory.cpp b/interfaces/sycl/Memory.cpp index 48b7c02..cf6e0f5 100644 --- a/interfaces/sycl/Memory.cpp +++ b/interfaces/sycl/Memory.cpp @@ -17,7 +17,6 @@ void* ConcreteAPI::allocGlobMem(size_t size, bool compress) { auto* ptr = malloc_device(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } @@ -28,7 +27,6 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) this->currentStatistics().allocatedUnifiedMemBytes += size; this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } @@ -38,7 +36,6 @@ void* ConcreteAPI::allocPinnedMem(size_t size, bool compress, Destination hint) auto* ptr = malloc_host(size, this->currentDefaultQueue()); this->currentStatistics().allocatedMemBytes += size; this->currentMemoryToSizeMap().insert({ptr, size}); - waitCheck(this->currentDefaultQueue()); return ptr; } @@ -76,9 +73,11 @@ void ConcreteAPI::freeMem(void* devPtr, bool unified) { context->statistics.allocatedUnifiedMemBytes -= size; } map.erase(devPtr); + // freeing memory that a queue may still be reading from is undefined, and the caller has no + // way to state that it is done, so the wait stays auto& queue = context->queueBuffer.getDefaultQueue(); - sycl::free(devPtr, queue.get_context()); queue.wait(); + sycl::free(devPtr, queue.get_context()); } void ConcreteAPI::freeGlobMem(void* devPtr) { From f0df01e36f6c75e24db919578b601ab100e59dde Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:48:56 +0200 Subject: [PATCH 28/31] fix: let the graph own the host functions recorded into it A host function recorded into a graph is called on every replay, so its copy has to outlive the recording. It was allocated and never handed to anyone, which leaked one copy per recorded call for as long as the process ran. The graph that a stream records into now owns those copies and drops them when it is destroyed. Reading the capture status through cudaStreamGetCaptureInfo rather than cudaStreamIsCapturing gives the graph handle to key them on, and the same call already served the node API, so the two share it now. AI-generated. Model: Opus 5 --- interfaces/cuda/Graphs.cu | 86 +++++++++++++++++++++++-------------- interfaces/cuda/Internals.h | 22 ++++++++++ interfaces/cuda/Streams.cu | 26 +++++++---- interfaces/hip/Graphs.cpp | 67 +++++++++++++++++++---------- interfaces/hip/Internals.h | 22 ++++++++++ interfaces/hip/Streams.cpp | 24 +++++++---- 6 files changed, 175 insertions(+), 72 deletions(-) diff --git a/interfaces/cuda/Graphs.cu b/interfaces/cuda/Graphs.cu index 7724aaf..e1e6c1e 100644 --- a/interfaces/cuda/Graphs.cu +++ b/interfaces/cuda/Graphs.cu @@ -12,6 +12,9 @@ #include #include #include +#include +#include +#include #include // Explicit graph nodes rest on cudaStreamBeginCaptureToGraph, which the runtime gained in CUDA @@ -38,6 +41,54 @@ using namespace device; * launchGraph(graph, stream); // 5 * */ +namespace { +std::mutex hostFunctionMutex; +std::unordered_map>>> + capturedHostFunctions; +} // namespace + +namespace device::internals { +CaptureState captureState(cudaStream_t stream) { + CaptureState state{}; + unsigned long long captureId{}; + const cudaGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + + // 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. +#if CUDART_VERSION >= 13000 + const cudaGraphEdgeData* edgeData{nullptr}; + APIWRAP(cudaStreamGetCaptureInfo( + stream, &state.status, &captureId, &state.graph, &frontier, &edgeData, &frontierSize)); +#else + APIWRAP(cudaStreamGetCaptureInfo( + stream, &state.status, &captureId, &state.graph, &frontier, &frontierSize)); +#endif + + if (frontier != nullptr) { + state.frontier.assign(frontier, frontier + frontierSize); + } + return state; +} + +std::function* adoptHostFunction(cudaGraph_t graph, const std::function& function) { + if (graph == nullptr) { + return nullptr; + } + + const std::lock_guard lock(hostFunctionMutex); + auto& functions = capturedHostFunctions[graph]; + functions.emplace_back(std::make_unique>(function)); + return functions.back().get(); +} + +void forgetHostFunctions(cudaGraph_t graph) { + const std::lock_guard lock(hostFunctionMutex); + capturedHostFunctions.erase(graph); +} +} // namespace device::internals + namespace device { struct DeviceGraph { cudaGraph_t graph{nullptr}; @@ -56,6 +107,8 @@ struct DeviceGraph { DeviceGraph& operator=(const DeviceGraph&) = delete; ~DeviceGraph() { + internals::forgetHostFunctions(graph); + // deliberately unchecked: the graph may outlive the device context during teardown, and a // failure here has nothing left to report to if (instance != nullptr) { @@ -129,37 +182,6 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { #endif } -namespace { -#ifdef DEVICE_USE_GRAPH_NODES -/** - * 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) { @@ -191,7 +213,7 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa assert(graphInstance != nullptr && "a node must be opened before it can be closed"); auto stream = static_cast(streamPtr); - auto produced = captureFrontier(stream); + auto produced = internals::captureState(stream).frontier; cudaGraph_t endedGraph{nullptr}; APIWRAP(cudaStreamEndCapture(stream, &endedGraph)); diff --git a/interfaces/cuda/Internals.h b/interfaces/cuda/Internals.h index a2c3652..7918944 100644 --- a/interfaces/cuda/Internals.h +++ b/interfaces/cuda/Internals.h @@ -6,6 +6,8 @@ #define SEISSOLDEVICE_INTERFACES_CUDA_INTERNALS_H_ #include +#include +#include #include #include @@ -19,6 +21,26 @@ namespace device::internals { using DeviceStreamT = cudaStream_t; +/** + * What a stream is currently recording into: the capture status, the graph the operations end up + * in, and the nodes a subsequently recorded operation would depend on. + */ +struct CaptureState { + cudaStreamCaptureStatus status{}; + cudaGraph_t graph{nullptr}; + std::vector frontier; +}; + +CaptureState captureState(cudaStream_t stream); + +/** + * Hands a host function to the graph that is being recorded, which keeps it alive for as long as + * the graph can be replayed, and returns the copy to pass to the runtime. Returns nullptr if the + * graph is not one of ours. + */ +std::function* adoptHostFunction(cudaGraph_t graph, const std::function& function); +void forgetHostFunctions(cudaGraph_t graph); + constexpr static int DefaultBlockDim = 1024; // Every wrapped call goes through here, so the parameters stay free of anything that allocates: diff --git a/interfaces/cuda/Streams.cu b/interfaces/cuda/Streams.cu index 6a00fbf..6e6d69f 100644 --- a/interfaces/cuda/Streams.cu +++ b/interfaces/cuda/Streams.cu @@ -81,13 +81,15 @@ void ConcreteAPI::syncStreamWithEvent(void* streamPtr, void* eventPtr) { } namespace { +// Called once, so the copy goes away with the call. void streamCallbackEpheremal(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); delete function; } -void streamCallbackPermanent(void* data) { +// Called on every replay of the graph it was recorded into, which owns the copy. +void streamCallbackRecorded(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); } @@ -96,17 +98,23 @@ void streamCallbackPermanent(void* data) { void ConcreteAPI::streamHostFunction(void* streamPtr, const std::function& function) { cudaStream_t stream = static_cast(streamPtr); - cudaStreamCaptureStatus status{}; - APIWRAP(cudaStreamIsCapturing(stream, &status)); + const auto capture = internals::captureState(stream); + if (capture.status == cudaStreamCaptureStatusInvalidated) { + return; + } - if (status != cudaStreamCaptureStatusInvalidated) { - auto* functionData = new std::function(function); - if (status == cudaStreamCaptureStatusActive) { - APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackPermanent, functionData)); - } else { - APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackEpheremal, functionData)); + if (capture.status == cudaStreamCaptureStatusActive) { + auto* recorded = internals::adoptHostFunction(capture.graph, function); + if (recorded == nullptr) { + logError() << "A host function was recorded into a graph this backend does not know."; + return; } + APIWRAP(cudaLaunchHostFunc(stream, &streamCallbackRecorded, recorded)); + return; } + + APIWRAP( + cudaLaunchHostFunc(stream, &streamCallbackEpheremal, new std::function(function))); } namespace { diff --git a/interfaces/hip/Graphs.cpp b/interfaces/hip/Graphs.cpp index 6f3e798..3d2ec61 100644 --- a/interfaces/hip/Graphs.cpp +++ b/interfaces/hip/Graphs.cpp @@ -12,6 +12,9 @@ #include #include #include +#include +#include +#include #include // Explicit graph nodes rest on hipStreamBeginCaptureToGraph, which HIP gained in ROCm 6.3. @@ -39,6 +42,45 @@ using namespace device; * launchGraph(graph, stream); // 5 * */ +namespace { +std::mutex hostFunctionMutex; +std::unordered_map>>> + capturedHostFunctions; +} // namespace + +namespace device::internals { +CaptureState captureState(hipStream_t stream) { + CaptureState state{}; + unsigned long long captureId{}; + const hipGraphNode_t* frontier{nullptr}; + size_t frontierSize{0}; + + APIWRAP(hipStreamGetCaptureInfo_v2( + stream, &state.status, &captureId, &state.graph, &frontier, &frontierSize)); + + if (frontier != nullptr) { + state.frontier.assign(frontier, frontier + frontierSize); + } + return state; +} + +std::function* adoptHostFunction(hipGraph_t graph, const std::function& function) { + if (graph == nullptr) { + return nullptr; + } + + const std::lock_guard lock(hostFunctionMutex); + auto& functions = capturedHostFunctions[graph]; + functions.emplace_back(std::make_unique>(function)); + return functions.back().get(); +} + +void forgetHostFunctions(hipGraph_t graph) { + const std::lock_guard lock(hostFunctionMutex); + capturedHostFunctions.erase(graph); +} +} // namespace device::internals + namespace device { struct DeviceGraph { hipGraph_t graph{nullptr}; @@ -57,6 +99,8 @@ struct DeviceGraph { DeviceGraph& operator=(const DeviceGraph&) = delete; ~DeviceGraph() { + internals::forgetHostFunctions(graph); + // deliberately unchecked: the graph may outlive the device context during teardown, and a // failure here has nothing left to report to if (instance != nullptr) { @@ -130,27 +174,6 @@ DeviceGraphHandle ConcreteAPI::graphCreate() { #endif } -namespace { -#ifdef DEVICE_USE_GRAPH_NODES -/** - * 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) { @@ -183,7 +206,7 @@ DeviceGraphNodeHandle ConcreteAPI::graphEndNode(const DeviceGraphHandle& graphHa assert(graphInstance != nullptr && "a node must be opened before it can be closed"); auto stream = static_cast(streamPtr); - auto produced = captureFrontier(stream); + auto produced = internals::captureState(stream).frontier; hipGraph_t endedGraph{nullptr}; APIWRAP(hipStreamEndCapture(stream, &endedGraph)); diff --git a/interfaces/hip/Internals.h b/interfaces/hip/Internals.h index 2c0dbe3..7499e73 100644 --- a/interfaces/hip/Internals.h +++ b/interfaces/hip/Internals.h @@ -7,7 +7,9 @@ #include "hip/hip_runtime.h" +#include #include +#include #define APIWRAP(call) (void)::device::internals::checkResult(call, __FILE__, __LINE__, {}) #define APIWRAPX(call, except) ::device::internals::checkResult(call, __FILE__, __LINE__, except) @@ -18,6 +20,26 @@ namespace device::internals { constexpr static int DefaultBlockDim = 1024; using DeviceStreamT = hipStream_t; + +/** + * What a stream is currently recording into: the capture status, the graph the operations end up + * in, and the nodes a subsequently recorded operation would depend on. + */ +struct CaptureState { + hipStreamCaptureStatus status{}; + hipGraph_t graph{nullptr}; + std::vector frontier; +}; + +CaptureState captureState(hipStream_t stream); + +/** + * Hands a host function to the graph that is being recorded, which keeps it alive for as long as + * the graph can be replayed, and returns the copy to pass to the runtime. Returns nullptr if the + * graph is not one of ours. + */ +std::function* adoptHostFunction(hipGraph_t graph, const std::function& function); +void forgetHostFunctions(hipGraph_t graph); // Every wrapped call goes through here, so the parameters stay free of anything that allocates: // the file name is the string literal __FILE__ expands to, and the accepted errors are read from // the caller's temporary array. diff --git a/interfaces/hip/Streams.cpp b/interfaces/hip/Streams.cpp index c0c8fa3..5045691 100644 --- a/interfaces/hip/Streams.cpp +++ b/interfaces/hip/Streams.cpp @@ -80,13 +80,14 @@ void ConcreteAPI::syncStreamWithEvent(void* streamPtr, void* eventPtr) { } namespace { +// Called once, so the copy goes away with the call. void streamCallbackEpheremal(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); delete function; } -void streamCallbackPermanent(void* data) { +void streamCallbackRecorded(void* data) { auto* function = reinterpret_cast*>(data); (*function)(); } @@ -95,17 +96,22 @@ void streamCallbackPermanent(void* data) { void ConcreteAPI::streamHostFunction(void* streamPtr, const std::function& function) { hipStream_t stream = static_cast(streamPtr); - hipStreamCaptureStatus status{}; - APIWRAP(hipStreamIsCapturing(stream, &status)); + const auto capture = internals::captureState(stream); + if (capture.status == hipStreamCaptureStatusInvalidated) { + return; + } - if (status != hipStreamCaptureStatusInvalidated) { - auto* functionData = new std::function(function); - if (status == hipStreamCaptureStatusActive) { - APIWRAP(hipLaunchHostFunc(stream, &streamCallbackPermanent, functionData)); - } else { - APIWRAP(hipLaunchHostFunc(stream, &streamCallbackEpheremal, functionData)); + if (capture.status == hipStreamCaptureStatusActive) { + auto* recorded = internals::adoptHostFunction(capture.graph, function); + if (recorded == nullptr) { + logError() << "A host function was recorded into a graph this backend does not know."; + return; } + APIWRAP(hipLaunchHostFunc(stream, &streamCallbackRecorded, recorded)); + return; } + + APIWRAP(hipLaunchHostFunc(stream, &streamCallbackEpheremal, new std::function(function))); } namespace { From c013b4654299d1217d765c157a4ef9500db7968e Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:49:21 +0200 Subject: [PATCH 29/31] test: cover the cases the reductions and copies were wrong in The reductions ran over unsigned values only, where the neutral element of a maximum happens to coincide with the smallest representable value; signed and floating point maxima over negative data are the case that tells a wrong neutral element apart from a right one. The typed suite also puts max and min on 4-byte types under test, which is where an atomic wider than the type writes past the result - visible under compute-sanitizer. The batch tests used 48 floats per element, which is a multiple of 16 bytes, so every element landed on the boundary the copy routines assumed. 47 floats do not. Empty inputs, and streams created with a priority, had no coverage at all. Sibling nodes recorded onto one stream have to come out right whether or not the backend orders them, which is what the graph node contract asks for. AI-generated. Model: Opus 5 --- tests/array_manip.cpp | 13 +++++++++ tests/batch_transfer.cpp | 60 ++++++++++++++++++++++++++++++++++++++ tests/graphs.cpp | 44 ++++++++++++++++++++++++++-- tests/reductions.cpp | 62 ++++++++++++++++++++++++++++++++++++++++ tests/streams.cpp | 18 ++++++++++++ 5 files changed, 194 insertions(+), 3 deletions(-) diff --git a/tests/array_manip.cpp b/tests/array_manip.cpp index 71df856..18e30fd 100644 --- a/tests/array_manip.cpp +++ b/tests/array_manip.cpp @@ -131,3 +131,16 @@ TEST_F(ArrayManip, scale) { device->api->freeGlobMem(arr); } + +TEST_F(ArrayManip, anEmptyArrayIsNoWork) { + auto* devArray = static_cast(device->api->allocGlobMem(sizeof(float))); + + device->algorithms.fillArray(devArray, 1.0F, 0, device->api->getDefaultStream()); + device->algorithms.scaleArray(devArray, 2.0F, 0, device->api->getDefaultStream()); + device->algorithms.touchMemory(devArray, 0, true, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + + SUCCEED(); + + device->api->freeGlobMem(devArray); +} diff --git a/tests/batch_transfer.cpp b/tests/batch_transfer.cpp index 4da1c3f..b632f30 100644 --- a/tests/batch_transfer.cpp +++ b/tests/batch_transfer.cpp @@ -166,3 +166,63 @@ TEST_F(BatchTransfer, incrementalAddBuildsAStridedPointerTable) { device->api->freeUnifiedMem(table); } + +/** + * The copy routines step down from 16-byte accesses, and an element stride that is not a multiple + * of 16 bytes puts every second element off that boundary. 47 floats are 188 bytes, so entry 1 + * starts 4-byte aligned and a 16-byte access to it faults. + */ +TEST_F(BatchTransfer, unalignedElementsAreCopied) { + constexpr std::size_t OddElementSize = 47; + + auto* oddSrc = + static_cast(device->api->allocGlobMem(BatchSize * OddElementSize * sizeof(float))); + auto* oddDst = + static_cast(device->api->allocGlobMem(BatchSize * OddElementSize * sizeof(float))); + auto** oddSrcBatch = + static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + auto** oddDstBatch = + static_cast(device->api->allocUnifiedMem(BatchSize * sizeof(float*))); + + std::vector hostSrc(BatchSize * OddElementSize); + for (std::size_t i = 0; i < BatchSize; ++i) { + oddSrcBatch[i] = oddSrc + i * OddElementSize; + oddDstBatch[i] = oddDst + i * OddElementSize; + for (std::size_t j = 0; j < OddElementSize; ++j) { + hostSrc[i * OddElementSize + j] = static_cast(i * OddElementSize + j); + } + } + + const std::vector zeroes(hostSrc.size(), 0.0F); + device->api->copyToAsync(oddSrc, hostSrc.data(), hostSrc.size() * sizeof(float), stream); + device->api->copyToAsync(oddDst, zeroes.data(), zeroes.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + + device->algorithms.streamBatchedData( + const_cast(oddSrcBatch), oddDstBatch, OddElementSize, BatchSize, stream); + device->api->syncStreamWithHost(stream); + + std::vector hostDst(hostSrc.size(), -1); + device->api->copyFromAsync(hostDst.data(), oddDst, hostDst.size() * sizeof(float), stream); + device->api->syncStreamWithHost(stream); + + for (std::size_t i = 0; i < hostSrc.size(); ++i) { + ASSERT_EQ(hostSrc[i], hostDst[i]) << "at " << i; + } + + device->api->freeUnifiedMem(oddDstBatch); + device->api->freeUnifiedMem(oddSrcBatch); + device->api->freeGlobMem(oddDst); + device->api->freeGlobMem(oddSrc); +} + +TEST_F(BatchTransfer, anEmptyBatchIsNoWork) { + device->algorithms.streamBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, 0, stream); + device->algorithms.accumulateBatchedData( + const_cast(srcBatch), dstBatch, ElementSize, 0, stream); + device->algorithms.touchBatchedMemory(dstBatch, ElementSize, 0, true, stream); + device->api->syncStreamWithHost(stream); + + SUCCEED(); +} diff --git a/tests/graphs.cpp b/tests/graphs.cpp index 18e2032..cb161f5 100644 --- a/tests/graphs.cpp +++ b/tests/graphs.cpp @@ -312,9 +312,9 @@ TEST_F(Graphs, droppedGraphsReleaseTheirResources) { GTEST_SKIP() << "the backend does not support explicit graph nodes"; } - // Graphs used to be held in a container that never gave anything back, so a workload that keys - // its graphs on something that varies - a time step width, say - grew without bound. Building - // and dropping many of them has to stay flat. + // A workload that keys its graphs on something that varies - a time step width, say - builds + // and drops them all the time, and that has to stay flat rather than accumulate device-side + // resources. for (int i = 0; i < 256; ++i) { auto graph = device->api->graphCreate(); device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { @@ -329,3 +329,41 @@ TEST_F(Graphs, droppedGraphsReleaseTheirResources) { ASSERT_EQ(255.0F, value); } } + +TEST_F(Graphs, siblingNodesMayShareAStream) { + if (graphNodesUnavailable()) { + GTEST_SKIP() << "the backend does not support explicit graph nodes"; + } + + fill(-1); + + // Two nodes with the same dependency and no edge between them, recorded onto one stream. They + // may end up ordered - a backend that expresses edges through the recorded stream orders them - + // but they write disjoint chunks, so the result is the same either way and neither may be lost. + auto graph = device->api->graphCreate(); + + const auto root = device->api->graphAddNode(graph, {}, mainStream, [&](void* stream) { + device->algorithms.fillArray(devArray, 0.0F, ArraySize, stream); + }); + + std::vector siblings; + for (std::size_t i = 0; i < BranchCount; ++i) { + siblings.push_back(device->api->graphAddNode(graph, {root}, mainStream, [&, i](void* stream) { + device->algorithms.fillArray( + devArray + i * ChunkSize, static_cast(i + 1), ChunkSize, stream); + })); + } + + device->api->graphAddNode(graph, siblings, mainStream, [&](void* stream) { + device->algorithms.scaleArray(devArray, 10.0F, ArraySize, stream); + }); + + device->api->graphInstantiate(graph); + device->api->launchGraph(graph, mainStream); + device->api->syncStreamWithHost(mainStream); + + const auto host = download(); + for (std::size_t i = 0; i < BranchCount; ++i) { + expectChunk(host, i, 10.0F * static_cast(i + 1)); + } +} diff --git a/tests/reductions.cpp b/tests/reductions.cpp index 0aaac48..00b9243 100644 --- a/tests/reductions.cpp +++ b/tests/reductions.cpp @@ -96,3 +96,65 @@ TEST_F(Reductions, Min) { device->api->freePinnedMem(testResult); device->api->freeGlobMem(devVector); } + +/** + * The reductions above run over unsigned values only, where the neutral element of a maximum and + * the smallest representable value are the same thing. They are not for signed integers, and for + * floating point types numeric_limits::min() is the smallest positive normal value, so a + * maximum over negative data has to start below all of them to come out right. + */ +template +class SignedReductions : public BaseTestSuite { + protected: + void run(ReductionType type, const std::vector& host, T expected) { + auto* devVector = static_cast(device->api->allocGlobMem(sizeof(T) * host.size())); + device->api->copyTo(devVector, host.data(), sizeof(T) * host.size()); + + auto* result = static_cast(device->api->allocPinnedMem(sizeof(T))); + *result = T{0}; + + device->algorithms.reduceVector( + result, devVector, true, host.size(), type, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + + EXPECT_EQ(expected, *result); + + device->api->freePinnedMem(result); + device->api->freeGlobMem(devVector); + } +}; + +using SignedTypes = ::testing::Types; +TYPED_TEST_SUITE(SignedReductions, SignedTypes); + +TYPED_TEST(SignedReductions, maxOverNegativeValues) { + std::vector host(100000, TypeParam{-7}); + host[host.size() / 3] = TypeParam{-2}; + this->run(ReductionType::Max, host, TypeParam{-2}); +} + +TYPED_TEST(SignedReductions, minOverNegativeValues) { + std::vector host(100000, TypeParam{-7}); + host[host.size() / 3] = TypeParam{-11}; + this->run(ReductionType::Min, host, TypeParam{-11}); +} + +TYPED_TEST(SignedReductions, addOverNegativeValues) { + std::vector host(1000, TypeParam{-3}); + this->run(ReductionType::Add, host, TypeParam{-3000}); +} + +TEST_F(Reductions, emptyInput) { + auto* devVector = static_cast(device->api->allocGlobMem(sizeof(float))); + auto* result = static_cast(device->api->allocPinnedMem(sizeof(float))); + *result = 123.0F; + + // nothing to reduce, but the result is still initialized to the neutral element + device->algorithms.reduceVector( + result, devVector, true, 0, ReductionType::Add, device->api->getDefaultStream()); + device->api->syncDefaultStreamWithHost(); + EXPECT_EQ(0.0F, *result); + + device->api->freePinnedMem(result); + device->api->freeGlobMem(devVector); +} diff --git a/tests/streams.cpp b/tests/streams.cpp index a756d7d..459c6c7 100644 --- a/tests/streams.cpp +++ b/tests/streams.cpp @@ -187,3 +187,21 @@ TEST_F(Streams, workOnSeparateStreamsStaysSeparate) { device->api->freeGlobMem(other); } + +TEST_F(Streams, streamsCanBeGivenAPriority) { + // 0 is the lowest priority the device offers, 1 the highest, and the default is whatever the + // runtime picks; all three have to give a stream that works + for (const double priority : {0.0, 0.5, 1.0}) { + auto* stream = device->api->createStream(priority); + ASSERT_NE(nullptr, stream) << "at priority " << priority; + + device->algorithms.fillArray(devArray, 6.0F, ArraySize, stream); + device->api->syncStreamWithHost(stream); + + for (const auto value : download(stream)) { + ASSERT_EQ(6.0F, value) << "at priority " << priority; + } + + device->api->destroyGenericStream(stream); + } +} From 991e7751ce11915a9ece718288645434d55f37e5 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:49:45 +0200 Subject: [PATCH 30/31] ci: treat warnings as errors and run the sanitizer The workflow carried a note to move to -Wall -Werror. Doing so catches, among other things, a mapped stream priority that is computed and then not passed on. A memcheck run on the NVIDIA runners covers what the assertions cannot see: an atomic that reaches past its result, or a vector access below its own alignment. It runs on the debug build only, and skips the three large reductions, which take minutes under the sanitizer and exercise no addressing the smaller suites do not. The GitLab pipeline still built the hipsycl backend against an image from three years ago, next to a GitHub workflow that covers the same three backends on current ones. Dropping it as its own commit, so it can be left out of the series if it is still wanted somewhere. The architecture list in the test CMakeLists had a stray space in gfx908 and stopped before the current generations. AI-generated. Model: Opus 5 --- .github/workflows/test.yml | 17 +++++++- .gitlab-ci.yml | 89 -------------------------------------- tests/CMakeLists.txt | 4 +- 3 files changed, 17 insertions(+), 93 deletions(-) delete mode 100644 .gitlab-ci.yml diff --git a/.github/workflows/test.yml b/.github/workflows/test.yml index bdd29f5..2bafc70 100644 --- a/.github/workflows/test.yml +++ b/.github/workflows/test.yml @@ -35,6 +35,7 @@ jobs: container-options: --runtime=nvidia --gpus=all pythonbreak: true test: true + sanitizer: compute-sanitizer - arch: sm_86 backend: acpp cc: gcc-13 @@ -55,6 +56,7 @@ jobs: container-options: --runtime=nvidia --gpus=all pythonbreak: true test: true + sanitizer: compute-sanitizer - arch: sm_86 backend: cuda cc: nvc @@ -122,8 +124,7 @@ jobs: cd tests mkdir build && cd build - # TODO: change to `-Wall -Werror` at some point - EXTRA_FLAGS="-Wall" + EXTRA_FLAGS="-Wall -Werror" export CFLAGS="${EXTRA_FLAGS} ${CFLAGS}" export CXXFLAGS="${EXTRA_FLAGS} ${CXXFLAGS}" @@ -148,3 +149,15 @@ jobs: cd build ./tests + + - id: sanitize + name: sanitize-device + if: ${{matrix.setup.test && matrix.setup.sanitizer && matrix.build_type == 'Debug'}} + run: | + cd tests + cd build + + # The three large reductions take minutes under the sanitizer and cover no pointer + # arithmetic the smaller suites do not, so they sit this one out. + ${{matrix.setup.sanitizer}} --tool memcheck --error-exitcode 1 ./tests \ + --gtest_filter=-Reductions.Add:Reductions.Max:Reductions.Min diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml deleted file mode 100644 index feb69cd..0000000 --- a/.gitlab-ci.yml +++ /dev/null @@ -1,89 +0,0 @@ -# SPDX-FileCopyrightText: 2019 SeisSol Group -# -# SPDX-License-Identifier: BSD-3-Clause - -stages: - - adjust - - test - - miscellaneous - - -image: - name: "seissol/gpu-ci-image:v0.0.4" - entrypoint: [""] - - -submodules: - stage: adjust - tags: - - atsccs68-docker-executor - script: - - cat ./.gitmodules - - git submodule update --init --recursive - artifacts: - paths: - - submodules - - -.common_test_steps: &common_test_script - - export CTEST_OUTPUT_ON_FAILURE=1 - - if [[ "${BACKEND}" == "hip" ]]; then - . /etc/profile.d/rocm.sh ; - export HIP_PLATFORM=nvidia ; - fi ; - - -tests: - stage: test - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - script: - - *common_test_script - - mkdir -p tests/build - - cd tests/build - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 - - make -j4 - - ./tests - - -basic_example: - stage: miscellaneous - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - script: - - *common_test_script - - export WORKDIR=./examples/basic/build - - mkdir -p ${WORKDIR} && cd ${WORKDIR} - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 - - make -j4 - - ./basic - - -jacobi_example: - stage: miscellaneous - tags: - - atsccs68-docker-executor - parallel: - matrix: - - BACKEND: [cuda, hip, hipsycl] - before_script: - - git clone --depth 1 --branch release-1.10.0 https://github.com/google/googletest - - mkdir -p googletest/build && cd googletest/build - - cmake .. -DBUILD_GTEST=ON -DBUILD_GMOCK=ON -Dgtest_disable_pthreads=ON -DBUILD_SHARED_LIBS=ON - - make -j $(nproc) - - make install - - cd ../.. - script: - - *common_test_script - - export WORKDIR=./examples/jacobi/build - - mkdir -p ${WORKDIR} && cd ${WORKDIR} - - cmake .. -DDEVICE_BACKEND=${BACKEND} -DSM=${GPU_MODEL} -DREAL_SIZE_IN_BYTES=4 -DWITH_MPI=OFF -DWITH_TESTS=ON - - make -j4 - - ./tests - - ./solver ./config.yaml diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 07afd5d..d4607d3 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -8,8 +8,8 @@ project(examples) #set(CMAKE_CXX_CLANG_TIDY clang-tidy) set(SM "sm_60" CACHE STRING "model of Nvidia Streaming Multiprocessor") -set(SM_OPTIONS "sm_50" "sm_60" "sm_70" "sm_71" "sm_75" "sm_80" "sm_86" - "gfx906" "gfx908 " +set(SM_OPTIONS "sm_50" "sm_60" "sm_70" "sm_71" "sm_75" "sm_80" "sm_86" "sm_89" "sm_90" + "gfx906" "gfx908" "gfx90a" "gfx942" "dg1" "bdw" "skl" "Gen8" "Gen9" "Gen11" "Gen12LP") set_property(CACHE SM PROPERTY STRINGS ${SM_OPTIONS}) From 2cbbfb3385e16773ba61b44cac24a6632345dec5 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Sat, 5 Sep 2026 08:50:10 +0200 Subject: [PATCH 31/31] chore: drop what nothing reads The SYCL reduction carried a second implementation behind #if 1 / #else. The CUDA and HIP cmake files defined DEVICE_CUDA_LANG and DEVICE_HIP_LANG, while the code reads DEVICE_LANG_CUDA and DEVICE_LANG_HIP, which the top-level file sets; cuda.cmake also repeated the graph capturing definition and the C++ standard. C++17 becomes a public requirement, since the installed headers use it. The rest is a stray semicolon after a function body, an unused device query, and a test fixture flag that was written and never read. AI-generated. Model: Opus 5 --- CMakeLists.txt | 4 +-- algorithms/sycl/Reduction.cpp | 54 ----------------------------------- cuda.cmake | 10 ------- hip.cmake | 3 +- interfaces/common/Common.h | 2 +- interfaces/sycl/Control.cpp | 4 +-- tests/BaseTestSuite.h | 9 ++---- 7 files changed, 8 insertions(+), 78 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e58157a..24c1f69 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -53,8 +53,8 @@ elseif((${DEVICE_BACKEND} STREQUAL "oneapi") OR (${DEVICE_BACKEND} STREQUAL "hip include(sycl.cmake) endif() -# common options -target_compile_features(device PRIVATE cxx_std_17) +# common options; the public headers use C++17 as well, so consumers get it too +target_compile_features(device PUBLIC cxx_std_17) if (USE_GRAPH_CAPTURING) target_compile_definitions(device PRIVATE DEVICE_USE_GRAPH_CAPTURING) diff --git a/algorithms/sycl/Reduction.cpp b/algorithms/sycl/Reduction.cpp index 2951513..91421cd 100644 --- a/algorithms/sycl/Reduction.cpp +++ b/algorithms/sycl/Reduction.cpp @@ -10,8 +10,6 @@ #include #include -#if 1 - namespace { using namespace device; @@ -155,58 +153,6 @@ void Algorithms::reduceVector(AccT* result, } } -#else - -namespace { -template -void launchReduction(AccT* result, const VecT* buffer, size_t size, S reducer, void* streamPtr) { - ((sycl::queue*)streamPtr)->submit([&](sycl::handler& cgh) { - cgh.parallel_for(sycl::range<1>{size}, reducer, [=](sycl::id<1> idx, auto& redval) { - redval.combine(static_cast(buffer[idx])); - }); - }); -} -} // namespace - -namespace device { -template -void Algorithms::reduceVector(AccT* result, - const VecT* buffer, - bool overrideResult, - size_t size, - ReductionType type, - void* streamPtr) { - auto properties = [&]() -> sycl::property_list { - if (overrideResult) { - return sycl::property_list{sycl::property::reduction::initialize_to_identity()}; - } else { - return sycl::property_list{}; - } - }(); - switch (type) { - case ReductionType::Add: { - return launchReduction( - result, buffer, size, sycl::reduction(result, sycl::plus(), properties), streamPtr); - } - case ReductionType::Max: { - return launchReduction(result, - buffer, - size, - sycl::reduction(result, sycl::maximum(), properties), - streamPtr); - } - case ReductionType::Min: { - return launchReduction(result, - buffer, - size, - sycl::reduction(result, sycl::minimum(), properties), - streamPtr); - } - } -} - -#endif - template void Algorithms::reduceVector(int* result, const int* buffer, bool overrideResult, diff --git a/cuda.cmake b/cuda.cmake index 3c5e0d6..d4eabb7 100644 --- a/cuda.cmake +++ b/cuda.cmake @@ -30,16 +30,6 @@ set_source_files_properties(device.cpp string(REPLACE "sm_" "" CUDA_DEVICE_ARCH "${DEVICE_ARCH}") set_target_properties(device PROPERTIES CUDA_ARCHITECTURES "${CUDA_DEVICE_ARCH}") -target_compile_features(device PRIVATE cxx_std_17) - -target_compile_definitions(device PRIVATE $<$: - -DDEVICE_${BACKEND_UPPER_CASE}_LANG; - >) - -if (USE_GRAPH_CAPTURING) - target_compile_definitions(device PRIVATE DEVICE_USE_GRAPH_CAPTURING) -endif() - target_compile_options(device PRIVATE $<$: --expt-relaxed-constexpr; >) diff --git a/hip.cmake b/hip.cmake index f8b3d88..2171f81 100644 --- a/hip.cmake +++ b/hip.cmake @@ -45,8 +45,7 @@ if (IS_NVCC_PLATFORM) else() set(DEVICE_HIPCC -std=c++17; -O3; - --offload-arch=${DEVICE_ARCH}; - -DDEVICE_${BACKEND_UPPER_CASE}_LANG) + --offload-arch=${DEVICE_ARCH}) endif() if (DEVICE_KERNEL_INFOPRINT) diff --git a/interfaces/common/Common.h b/interfaces/common/Common.h index 4415141..a789209 100644 --- a/interfaces/common/Common.h +++ b/interfaces/common/Common.h @@ -32,7 +32,7 @@ using StatusT = std::array; template void isFlagSet(const StatusT& status) { assert(status[ID]); -}; +} template U align(T number, U alignment) { diff --git a/interfaces/sycl/Control.cpp b/interfaces/sycl/Control.cpp index d3decec..2402a42 100644 --- a/interfaces/sycl/Control.cpp +++ b/interfaces/sycl/Control.cpp @@ -91,8 +91,8 @@ int ConcreteAPI::getDeviceId() { } unsigned int ConcreteAPI::getGlobMemAlignment() { - auto device = this->currentDefaultQueue().get_device(); - return 128; // ToDo: find attribute; not: device.get_info(); + // ToDo: find attribute; not: device.get_info(); + return 128; } void ConcreteAPI::syncDevice() { this->currentQueueBuffer().syncAllQueuesWithHost(); } diff --git a/tests/BaseTestSuite.h b/tests/BaseTestSuite.h index 4d2c2ae..7583a94 100644 --- a/tests/BaseTestSuite.h +++ b/tests/BaseTestSuite.h @@ -14,18 +14,13 @@ using namespace device; using namespace ::testing; -static bool setUp = false; - class BaseTestSuite : public ::testing::Test { public: - DeviceInstance* device; + DeviceInstance* device{nullptr}; BaseTestSuite() { randomEngine.seed(randomDevice()); } - void SetUp() { - device = &DeviceInstance::getInstance(); - setUp = true; - } + void SetUp() override { device = &DeviceInstance::getInstance(); } protected: std::random_device randomDevice;