diff --git a/CMakeLists.txt b/CMakeLists.txt index 77f106a3f2..132840d4ce 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -59,6 +59,7 @@ option(CAF_ENABLE_PROTOBUF_EXAMPLES "Build examples with Google Protobuf" OFF) option(CAF_ENABLE_QT6_EXAMPLES "Build examples with the Qt6 framework" OFF) option(CAF_ENABLE_ROBOT_TESTS "Add the Robot tests to CTest " OFF) option(CAF_ENABLE_RUNTIME_CHECKS "Build CAF with extra runtime assertions" OFF) +option(CAF_ENABLE_SC26_BENCHMARKS "Build experimental SC26 CUDA benchmarks" OFF) option(CAF_USE_STD_FORMAT "Enable std::format support" OFF) option(CAF_ENABLE_CUDA "Build caf with cuda support" ON) @@ -110,6 +111,10 @@ if(MSVC AND CAF_SANITIZERS) message(FATAL_ERROR "Sanitizer builds are currently not supported on MSVC") endif() +if(CAF_ENABLE_SC26_BENCHMARKS AND NOT CAF_ENABLE_CUDA) + message(FATAL_ERROR "SC26 benchmarks require CAF_ENABLE_CUDA=ON") +endif() + # -- doxygen setup ------------------------------------------------------------- configure_file("${CMAKE_CURRENT_SOURCE_DIR}/cmake/Doxyfile.in" @@ -388,6 +393,10 @@ if(CAF_ENABLE_CUDA) add_subdirectory(libcaf_cuda) endif() +if(CAF_ENABLE_SC26_BENCHMARKS) + add_subdirectory(libcaf_cuda/sc26) +endif() + # -- optionally add the Robot tests to CTest ----------------------------------- diff --git a/libcaf_cuda/CMakeLists.txt b/libcaf_cuda/CMakeLists.txt index d1282578e2..b49cc84aeb 100644 --- a/libcaf_cuda/CMakeLists.txt +++ b/libcaf_cuda/CMakeLists.txt @@ -60,13 +60,15 @@ caf_add_component( cuda DEPENDENCIES PUBLIC - CAF::core + CAF::core $<$:ws2_32> + CUDA::cuda_driver + CUDA::cublas + CUDA::cusparse PRIVATE CAF::internal #CAF::CORE CUDA::nvrtc - CUDA::cuda_driver HEADERS ${CAF_CUDA_HEADERS} SOURCES diff --git a/libcaf_cuda/caf/cuda/command_runner.hpp b/libcaf_cuda/caf/cuda/command_runner.hpp index 1def06d3f8..87cb8db49d 100644 --- a/libcaf_cuda/caf/cuda/command_runner.hpp +++ b/libcaf_cuda/caf/cuda/command_runner.hpp @@ -11,6 +11,11 @@ #include "caf/cuda/control-layer/launch_response_token.hpp" #include "caf/cuda/control-layer/memory_response_token.hpp" #include "caf/cuda/event.hpp" +#include "caf/cuda/detail/context_guard.hpp" + +#include +#include +#include namespace caf::cuda { @@ -20,6 +25,9 @@ namespace caf::cuda { // gpu actors if they wish // Manages synchronous and asynchronous command execution with overloads // for actor_id, shared_memory, and device_number. +// An actor_id is a submission serialization domain: submissions and stream +// release for the same ID must not overlap. Different IDs may submit in +// parallel. // =========================================================================== template class command_runner { @@ -37,11 +45,9 @@ class command_runner { int actor_id, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - std::forward(xs)...); - return cmd->enqueue(); + auto result = submit_impl(program, dims, actor_id, 0, -1, + std::forward(xs)...); + return result.device->collect_output_buffers(result.memory); } // ------------------------------- @@ -54,12 +60,9 @@ class command_runner { int shared_memory, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - shared_memory, - std::forward(xs)...); - return cmd->enqueue(); + auto result = submit_impl(program, dims, actor_id, shared_memory, -1, + std::forward(xs)...); + return result.device->collect_output_buffers(result.memory); } // ------------------------------- @@ -73,13 +76,9 @@ class command_runner { int device_number, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - shared_memory, - device_number, - std::forward(xs)...); - return cmd->enqueue(); + auto result = submit_impl(program, dims, actor_id, shared_memory, + device_number, std::forward(xs)...); + return result.device->collect_output_buffers(result.memory); } @@ -118,11 +117,9 @@ class command_runner { int actor_id, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - std::forward(xs)...); - return cmd->base_enqueue(); + auto result = submit_impl(program, dims, actor_id, 0, -1, + std::forward(xs)...); + return std::move(result.memory); } // ------------------------------- @@ -135,12 +132,9 @@ class command_runner { int shared_memory, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - shared_memory, - std::forward(xs)...); - return cmd->base_enqueue(); + auto result = submit_impl(program, dims, actor_id, shared_memory, -1, + std::forward(xs)...); + return std::move(result.memory); } // ------------------------------- @@ -154,13 +148,9 @@ class command_runner { int device_number, Us&&... xs) { - auto cmd = caf::make_counted(std::move(program), - std::move(dims), - actor_id, - shared_memory, - device_number, - std::forward(xs)...); - return cmd->base_enqueue(); + auto result = submit_impl(program, dims, actor_id, shared_memory, + device_number, std::forward(xs)...); + return std::move(result.memory); } @@ -212,9 +202,9 @@ class command_runner { int stream_id, Us&&... args) { - auto cmd = caf::make_counted...>>( - device_number, stream_id, std::forward(args)...); - return cmd->enqueue(); + bulk_memory_command...> cmd{ + device_number, stream_id, std::forward(args)...}; + return cmd.enqueue(); } // ------------------------------------------------------------------------- @@ -257,25 +247,28 @@ class command_runner { cmd.enqueue(dst, count); } - // Asynchronous copy back (default) + // Asynchronous copy back (default). CUDA host callbacks must not call CUDA + // APIs or release the final reference to an object whose destructor does. template void copy_to_host_async(mem_ptr ptr, F callback) { - auto cmd = caf::make_counted>(std::move(ptr)); - cmd->run_async(std::move(callback)); + copy_back_command cmd{std::move(ptr)}; + cmd.run_async(std::move(callback)); } - // Asynchronous copy back with explicit stream/actor ID + // Asynchronous copy back with explicit stream/actor ID. CUDA host callbacks + // must not call CUDA APIs or release the final CUDA-owning reference. template void copy_to_host_async(mem_ptr ptr, int stream_id, F callback) { - auto cmd = caf::make_counted>(std::move(ptr), stream_id); - cmd->run_async(std::move(callback)); + copy_back_command cmd{std::move(ptr), stream_id}; + cmd.run_async(std::move(callback)); } - // Asynchronous copy back to user-provided buffer + // Asynchronous copy back to user-provided buffer. CUDA host callbacks must + // not call CUDA APIs or release the final CUDA-owning reference. template void copy_to_host_async(mem_ptr ptr, T* dst, size_t count, F callback) { - auto cmd = caf::make_counted>(std::move(ptr)); - cmd->run_async(dst, count, std::move(callback)); + copy_back_command cmd{std::move(ptr)}; + cmd.run_async(dst, count, std::move(callback)); } // Enqueue an asynchronous free operation on the given stream @@ -289,31 +282,47 @@ class command_runner { // ------------------------------- // Destroy streams for a given actor ID // ------------------------------- + // The caller must first ensure that no submission for this ID is in flight. void release_stream_for_actor(int actor_id) { auto plat = platform::create(); plat->release_streams_for_actor(actor_id); } // ------------------------------- - // Register a callback on the actor's stream + // Register a callback on the actor's stream. The callback must not call CUDA + // APIs or release the final reference to an object whose destructor does. // ------------------------------- template void add_callback(int stream_id, int device_number, F callback) { auto plat = platform::create(); auto dev = plat->schedule(stream_id, device_number); auto stream = dev->get_stream_for_actor(stream_id); - auto* f_ptr = new F(std::move(callback)); - auto res = cuLaunchHostFunc(stream, [](void* data) { - auto* f = static_cast(data); - (*f)(); - delete f; - }, f_ptr); - if (res != CUDA_SUCCESS) { delete f_ptr; check(res, "cuLaunchHostFunc"); } + auto state = std::unique_ptr{new F(std::move(callback))}; + detail::context_guard context_scope{dev->getContext()}; + auto res = cuLaunchHostFunc(stream, [](void* data) noexcept { + auto fn = std::unique_ptr{static_cast(data)}; + try { + (*fn)(); + } catch (const std::exception& error) { + std::fprintf(stderr, "CUDA host callback failed: %s\n", + error.what()); + } catch (...) { + std::fprintf(stderr, + "CUDA host callback failed: unknown error\n"); + } + }, state.get()); + if (res != CUDA_SUCCESS) + throw detail::cuda_error(res, "cuLaunchHostFunc"); + state.release(); } // ------------------------------- // Resets the CUDA context for a given device number. - // This will force the device to flush its stream pool and create a new context. + // This will force the device to flush its context-bound resources and create + // a new context. The caller must first quiesce submissions, callbacks, + // program/memory/event lifecycle, and raw handle use on this device. + // Existing non-scalar memory, events, and raw handles become stale; retained + // programs reload lazily on their next launch. // ------------------------------- void reset_context(int device_number) { auto plat = platform::create(); @@ -377,6 +386,34 @@ class command_runner { return dev->get_stream_for_actor(stream_number); } +private: + using result_type = std::tuple>...>; + + struct submission_result { + device_ptr device; + result_type memory; + }; + + template + submission_result submit_impl(const program_ptr& program, + const nd_range& dims, + int actor_id, + int shared_memory, + int device_number, + Us&&... xs) { + static_assert(sizeof...(Us) == sizeof...(Ts), "Argument count mismatch"); + + // Match base_command's conversion behavior while keeping only one tuple. + std::tuple arguments{ + std::make_tuple(std::forward(xs)...)}; + + auto dev = program->device_for_launch(actor_id, device_number); + auto kernel = program->get_kernel(dev->getId()); + auto memory = dev->launch_kernel_mem_ref_impl( + kernel, dims, arguments, actor_id, shared_memory); + return {std::move(dev), std::move(memory)}; + } + }; } // namespace caf::cuda diff --git a/libcaf_cuda/caf/cuda/control-layer/request_token.hpp b/libcaf_cuda/caf/cuda/control-layer/request_token.hpp index af4444f807..5ad620081e 100644 --- a/libcaf_cuda/caf/cuda/control-layer/request_token.hpp +++ b/libcaf_cuda/caf/cuda/control-layer/request_token.hpp @@ -20,13 +20,6 @@ class CAF_CUDA_EXPORT request_token : public token { public: request_token(int dependency = INDEPENDENT) : token(dependency) {} - - // Required for CAF message passing - request_token() = default; - - -private: - }; using request_token_ptr = caf::intrusive_ptr; diff --git a/libcaf_cuda/caf/cuda/detail/context_guard.hpp b/libcaf_cuda/caf/cuda/detail/context_guard.hpp new file mode 100644 index 0000000000..e5ef54d1a5 --- /dev/null +++ b/libcaf_cuda/caf/cuda/detail/context_guard.hpp @@ -0,0 +1,57 @@ +#pragma once + +#include + +#include +#include +#include + +namespace caf::cuda::detail { + +inline std::runtime_error cuda_error(CUresult result, const char* operation) { + const char* description = nullptr; + cuGetErrorString(result, &description); + return std::runtime_error(std::string{operation} + " failed: " + + (description ? description : "unknown CUDA error")); +} + +/// Temporarily makes a CUDA context current and restores the caller's context. +class context_guard { +public: + explicit context_guard(CUcontext context) { + if (!context) + throw std::invalid_argument("context_guard requires a valid context"); + CUcontext current = nullptr; + auto result = cuCtxGetCurrent(¤t); + if (result != CUDA_SUCCESS) + throw cuda_error(result, "cuCtxGetCurrent"); + if (current != context) { + result = cuCtxPushCurrent(context); + if (result != CUDA_SUCCESS) + throw cuda_error(result, "cuCtxPushCurrent"); + pushed_ = true; + } + } + + ~context_guard() noexcept { + if (pushed_) { + CUcontext ignored = nullptr; + auto result = cuCtxPopCurrent(&ignored); + if (result != CUDA_SUCCESS) { + // Destructors cannot propagate a restoration failure. + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "cuCtxPopCurrent failed: %s\n", + description ? description : "unknown CUDA error"); + } + } + } + + context_guard(const context_guard&) = delete; + context_guard& operator=(const context_guard&) = delete; + +private: + bool pushed_ = false; +}; + +} // namespace caf::cuda::detail diff --git a/libcaf_cuda/caf/cuda/detail/context_lifetime.hpp b/libcaf_cuda/caf/cuda/detail/context_lifetime.hpp new file mode 100644 index 0000000000..ff84572abe --- /dev/null +++ b/libcaf_cuda/caf/cuda/detail/context_lifetime.hpp @@ -0,0 +1,40 @@ +#pragma once + +#include +#include + +#include + +namespace caf::cuda::detail { + +/// Identifies one lifetime of a CUDA context handle. +/// +/// CUDA may reuse the numeric value of a destroyed context handle. Objects +/// that retain this token can therefore distinguish a current context from a +/// stale context even when their raw CUcontext values happen to compare equal. +class context_lifetime { +public: + explicit context_lifetime(CUcontext context) noexcept : context_(context) { + // nop + } + + CUcontext context() const noexcept { + return context_; + } + + bool active() const noexcept { + return active_.load(std::memory_order_acquire); + } + + void deactivate() noexcept { + active_.store(false, std::memory_order_release); + } + +private: + CUcontext context_; + std::atomic active_{true}; +}; + +using context_lifetime_ptr = std::shared_ptr; + +} // namespace caf::cuda::detail diff --git a/libcaf_cuda/caf/cuda/detail/stream_order.hpp b/libcaf_cuda/caf/cuda/detail/stream_order.hpp new file mode 100644 index 0000000000..36c8d47ebf --- /dev/null +++ b/libcaf_cuda/caf/cuda/detail/stream_order.hpp @@ -0,0 +1,53 @@ +#pragma once + +#include "caf/cuda/detail/context_guard.hpp" + +#include + +#include + +namespace caf::cuda::detail { + +/// A timing-free CUDA event used only to establish stream dependencies. +class stream_transition_event { +public: + stream_transition_event() { + auto result = cuEventCreate(&event_, CU_EVENT_DISABLE_TIMING); + if (result != CUDA_SUCCESS) + throw cuda_error(result, "cuEventCreate(stream transition)"); + } + + ~stream_transition_event() noexcept { + auto result = cuEventDestroy(event_); + if (result != CUDA_SUCCESS) { + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "cuEventDestroy(stream transition) failed: %s\n", + description ? description : "unknown CUDA error"); + } + } + + stream_transition_event(const stream_transition_event&) = delete; + stream_transition_event& operator=(const stream_transition_event&) = delete; + + CUevent get() const noexcept { + return event_; + } + +private: + CUevent event_ = nullptr; +}; + +/// Makes work subsequently submitted to `target` wait for work already +/// submitted to `source`. The caller must have the streams' context current. +inline void order_stream_after(CUstream target, CUstream source) { + stream_transition_event event; + auto result = cuEventRecord(event.get(), source); + if (result != CUDA_SUCCESS) + throw cuda_error(result, "cuEventRecord(stream transition)"); + result = cuStreamWaitEvent(target, event.get(), 0); + if (result != CUDA_SUCCESS) + throw cuda_error(result, "cuStreamWaitEvent(stream transition)"); +} + +} // namespace caf::cuda::detail diff --git a/libcaf_cuda/caf/cuda/device.hpp b/libcaf_cuda/caf/cuda/device.hpp index c7192d1ce6..95145bd54b 100644 --- a/libcaf_cuda/caf/cuda/device.hpp +++ b/libcaf_cuda/caf/cuda/device.hpp @@ -1,16 +1,22 @@ #pragma once +#include +#include #include #include #include #include #include +#include #include #include #include #include +#include +#include #include #include +#include #include #include @@ -18,6 +24,9 @@ #include #include "caf/cuda/global.hpp" +#include "caf/cuda/detail/context_guard.hpp" +#include "caf/cuda/detail/context_lifetime.hpp" +#include "caf/cuda/detail/stream_order.hpp" #include "caf/cuda/types.hpp" #include "caf/cuda/streampool.hpp" #include "caf/cuda/mem_ref.hpp" @@ -25,6 +34,19 @@ namespace caf::cuda { +namespace detail { + +template +struct is_cuda_mem_ptr : std::false_type {}; + +template +struct is_cuda_mem_ptr> : std::true_type {}; + +template +inline constexpr bool is_cuda_mem_ptr_v = is_cuda_mem_ptr::value; + +} // namespace detail + class CAF_CUDA_EXPORT device : public caf::ref_counted { public: using device_ptr = caf::intrusive_ptr; @@ -32,22 +54,23 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { device(CUdevice device, CUcontext context, const char* name, int id, size_t stream_pool_size = 32) : device_(device), context_(context), + context_lifetime_( + std::make_shared(context)), id_(id), name_(name), + stream_pool_size_(stream_pool_size), stream_table_(std::make_unique(context, stream_pool_size)) { - init_device_properties(); - } - - ~device() { - check(cuCtxDestroy(context_), "cuCtxDestroy"); + init_device_properties(); } + ~device() noexcept override; + device(const device&) = delete; device& operator=(const device&) = delete; device(device&&) noexcept = default; device& operator=(device&&) noexcept = default; - const char* name() const { return name_; } + const char* name() const { return name_.c_str(); } CUdevice getDevice() const { return device_; } CUcontext getContext() const { return context_; } int getId() const { return id_; } @@ -57,17 +80,14 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { CUcontext getContext(int) { return context_; } - // Resets the CUDA context for this device. - // Destroys the old context and reinitializes the stream pool with the new one. - void reset_context(CUcontext new_ctx) { - // Destroy the old context - if (context_) { - check(cuCtxDestroy(context_), "cuCtxDestroy in device::reset_context"); - } - // Set the new context and reinitialize the stream table - context_ = new_ctx; - stream_table_ = std::make_unique(context_, stream_table_->pool_size()); - } + /// Replaces this device's CUDA context and rebuilds context-bound resources. + /// + /// Reset is a cold recovery operation. The caller must first quiesce all + /// submissions, callbacks, program construction/destruction, managed + /// memory/event operations, and raw CUDA/library-handle use for this device. + /// Existing non-scalar memory, events, and raw handles become stale. Program + /// objects reload their modules lazily when they are next used. + void reset_context(CUcontext new_ctx); // Number of streaming multiprocessors (SMs) int num_sms() const noexcept { return sm_count_; } @@ -112,962 +132,457 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { - //returns the CUStream associated with the actor id + // Returns the CUstream associated with an actor ID. An actor ID defines a + // serialized submission lane; calls using the same ID must not overlap. CUstream get_stream_for_actor(int actor_id) { return stream_table_->get_stream(actor_id); } - //releases the CUStream associated with the actor id + // Releases resources for an idle actor ID. This must not overlap with a + // submission or raw library-handle use for the same ID. void release_stream_for_actor(int actor_id) { + if (cublas_table_) + cublas_table_->release_handle(actor_id); + if (cusparse_table_) + cusparse_table_->release_handle(actor_id); stream_table_->release_stream(actor_id); } /// Creates a CUDA event on this device. event_ptr create_event(unsigned int flags = CU_EVENT_DEFAULT) { - CHECK_CUDA(cuCtxPushCurrent(context_)); + detail::context_guard context_scope{context_}; auto res = caf::make_counted(flags); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + res->bind_context_lifetime(context_lifetime_); return res; } /// Records an event on the stream associated with the actor_id. void record_event(event_ptr e, int actor_id) { + validate_event(e); CUstream stream = get_stream_for_actor(actor_id); + detail::context_guard context_scope{context_}; CHECK_CUDA(cuEventRecord(e->get(), stream)); } /// Makes a stream wait on an event. void wait_event(event_ptr e, int actor_id) { + validate_event(e); CUstream stream = get_stream_for_actor(actor_id); + detail::context_guard context_scope{context_}; CHECK_CUDA(cuStreamWaitEvent(stream, e->get(), 0)); } /// Returns true if the event has completed. bool query_event(event_ptr e) { - CHECK_CUDA(cuCtxPushCurrent(context_)); - bool res = e->query(); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return res; + validate_event(e); + detail::context_guard context_scope{context_}; + return e->query(); } /// Blocks until the event has completed. void synchronize_event(event_ptr e) { - CHECK_CUDA(cuCtxPushCurrent(context_)); + validate_event(e); + detail::context_guard context_scope{context_}; e->synchronize(); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); } /// Enable cuBLAS support. void enable_cublas() { + std::lock_guard guard{lifecycle_mutex_}; if (!cublas_table_) - cublas_table_ = std::make_unique(context_); + cublas_table_ = std::make_unique( + context_, stream_table_->pool_size()); } /// Enable cuSparse support. void enable_cusparse() { + std::lock_guard guard{lifecycle_mutex_}; if (!cusparse_table_) - cusparse_table_ = std::make_unique(context_); + cusparse_table_ = std::make_unique( + context_, stream_table_->pool_size()); } - /// Returns the cuSparse handle associated with the actor id. + /// Returns the raw cuSPARSE handle associated with an actor ID. + /// + /// The handle may be reconfigured by the caller, but must not be retained or + /// used concurrently with submissions or release for the same actor ID. cusparseHandle_t get_cusparse_handle(int stream_id) { - if (!cusparse_table_) return nullptr; - return cusparse_table_->get_handle(stream_id, get_stream_for_actor(stream_id)); + if (!cusparse_table_) + return nullptr; + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + return cusparse_table_->get_handle(stream_id, stream); } - /// Returns the cuBLAS handle associated with the actor id. + /// Returns the raw cuBLAS handle associated with an actor ID. + /// + /// The handle may be reconfigured by the caller, but must not be retained or + /// used concurrently with submissions or release for the same actor ID. cublasHandle_t get_cublas_handle(int stream_id) { - if (!cublas_table_) return nullptr; - return cublas_table_->get_handle(stream_id, get_stream_for_actor(stream_id)); + if (!cublas_table_) + return nullptr; + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + return cublas_table_->get_handle(stream_id, stream); } /// Performs single precision matrix-vector multiplication (y = alpha*A*x + beta*y). /// Assumes A is in row-major order of dimensions m x n. void sgemv(int stream_id, int m, int n, float alpha, mem_ptr A, mem_ptr x, float beta, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - // Row-major matrix A (m x n) is stored as m rows of n elements. // Viewed as column-major by cuBLAS, this is a n x m matrix. // To compute y = A * x: // Op(Memory) * x = (n x m)^T * (n x 1) = (m x n) * (n x 1) = (m x 1). // We use CUBLAS_OP_T. LDA is the 'rows' in the column-major view, which is n. - cublasStatus_t status = cublasSgemv(handle, CUBLAS_OP_T, - n, m, - &alpha, - reinterpret_cast(A->mem()), n, - reinterpret_cast(x->mem()), 1, - &beta, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSgemv failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, x, y), "cublasSgemv", + [&](cublasHandle_t handle) { + return cublasSgemv( + handle, CUBLAS_OP_T, n, m, &alpha, + reinterpret_cast(A->mem()), n, + reinterpret_cast(x->mem()), 1, &beta, + reinterpret_cast(y->mem()), 1); + }); } /// Performs element-wise multiplication of two vectors: result = x .* y. /// This is implemented using cublasSdgmm (Diagonal Matrix-Vector Multiplication). - void s_elementwise_multiply(int stream_id, int n, mem_ptr x, mem_ptr y, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); + void s_elementwise_multiply(int stream_id, int n, mem_ptr x, + mem_ptr y, mem_ptr result) { // cublasSdgmm: C = diag(X) * Y. When Y is a vector (n x 1), this is element-wise mult. - cublasStatus_t status = cublasSdgmm(handle, CUBLAS_SIDE_LEFT, - n, 1, - reinterpret_cast(y->mem()), n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(result->mem()), n); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSdgmm failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x, y, result), "cublasSdgmm", + [&](cublasHandle_t handle) { + return cublasSdgmm( + handle, CUBLAS_SIDE_LEFT, n, 1, + reinterpret_cast(y->mem()), n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(result->mem()), n); + }); } /// Performs element-wise multiplication of two vectors: result = x .* y. /// This is implemented using cublasDdgmm (Diagonal Matrix-Vector Multiplication). - void d_elementwise_multiply(int stream_id, int n, mem_ptr x, mem_ptr y, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); + void d_elementwise_multiply(int stream_id, int n, mem_ptr x, + mem_ptr y, mem_ptr result) { // cublasDdgmm: C = diag(X) * Y. When Y is a vector (n x 1), this is element-wise mult. - cublasStatus_t status = cublasDdgmm(handle, CUBLAS_SIDE_LEFT, - n, 1, - reinterpret_cast(y->mem()), n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(result->mem()), n); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDdgmm failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x, y, result), "cublasDdgmm", + [&](cublasHandle_t handle) { + return cublasDdgmm( + handle, CUBLAS_SIDE_LEFT, n, 1, + reinterpret_cast(y->mem()), n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(result->mem()), n); + }); } /// Returns the required buffer size for SpMV CSR. template - size_t spmv_csr_buffer_size(int stream_id, int m, int n, int nnz, - mem_ptr row_ptr, mem_ptr col_ind, mem_ptr values, - mem_ptr x, mem_ptr y) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCsr(&matA, m, n, nnz, reinterpret_cast(row_ptr->mem()), - reinterpret_cast(col_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matA, vecX, &beta, vecY, - type, CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmv_csr_buffer_size(int stream_id, int m, int n, int nnz, + mem_ptr row_ptr, mem_ptr col_ind, + mem_ptr values, mem_ptr x, + mem_ptr y) { + return spmv_buffer_size_impl( + sparse_format::csr, stream_id, m, n, nnz, row_ptr, col_ind, values, + x, y); } /// Returns the required buffer size for SpMV COO. template - size_t spmv_coo_buffer_size(int stream_id, int m, int n, int nnz, - mem_ptr row_ind, mem_ptr col_ind, mem_ptr values, - mem_ptr x, mem_ptr y) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCoo(&matA, m, n, nnz, reinterpret_cast(row_ind->mem()), - reinterpret_cast(col_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matA, vecX, &beta, vecY, - type, CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmv_coo_buffer_size(int stream_id, int m, int n, int nnz, + mem_ptr row_ind, mem_ptr col_ind, + mem_ptr values, mem_ptr x, + mem_ptr y) { + return spmv_buffer_size_impl( + sparse_format::coo, stream_id, m, n, nnz, row_ind, col_ind, values, + x, y); } /// Returns the required buffer size for SpMV CSC. template - size_t spmv_csc_buffer_size(int stream_id, int m, int n, int nnz, - mem_ptr col_ptr, mem_ptr row_ind, mem_ptr values, - mem_ptr x, mem_ptr y) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCsc(&matA, m, n, nnz, reinterpret_cast(col_ptr->mem()), - reinterpret_cast(row_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matA, vecX, &beta, vecY, - type, CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmv_csc_buffer_size(int stream_id, int m, int n, int nnz, + mem_ptr col_ptr, mem_ptr row_ind, + mem_ptr values, mem_ptr x, + mem_ptr y) { + return spmv_buffer_size_impl( + sparse_format::csc, stream_id, m, n, nnz, col_ptr, row_ind, values, + x, y); } /// Performs sparse matrix-vector multiplication (y = alpha*A*x + beta*y) using CSR format. template - void spmv_csr(int stream_id, int m, int n, int nnz, T alpha, - mem_ptr row_ptr, mem_ptr col_ind, mem_ptr values, - mem_ptr x, T beta, mem_ptr y, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) - throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCsr(&matA, m, n, nnz, - reinterpret_cast(row_ptr->mem()), - reinterpret_cast(col_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, - CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMV(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, d_workspace); - - if (dBuffer) - CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMV (CSR) failed on device " + std::to_string(id_)); + void spmv_csr(int stream_id, int m, int n, int nnz, T alpha, + mem_ptr row_ptr, mem_ptr col_ind, + mem_ptr values, mem_ptr x, T beta, mem_ptr y, + mem_ptr workspace = nullptr) { + spmv_impl(sparse_format::csr, stream_id, m, n, nnz, alpha, row_ptr, + col_ind, values, x, beta, y, workspace); } /// Performs sparse matrix-vector multiplication (y = alpha*A*x + beta*y) using COO format. template - void spmv_coo(int stream_id, int m, int n, int nnz, T alpha, - mem_ptr row_ind, mem_ptr col_ind, mem_ptr values, - mem_ptr x, T beta, mem_ptr y, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) - throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCoo(&matA, m, n, nnz, - reinterpret_cast(row_ind->mem()), - reinterpret_cast(col_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMV(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, d_workspace); - - if (dBuffer) - CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMV (COO) failed on device " + std::to_string(id_)); + void spmv_coo(int stream_id, int m, int n, int nnz, T alpha, + mem_ptr row_ind, mem_ptr col_ind, + mem_ptr values, mem_ptr x, T beta, mem_ptr y, + mem_ptr workspace = nullptr) { + spmv_impl(sparse_format::coo, stream_id, m, n, nnz, alpha, row_ind, + col_ind, values, x, beta, y, workspace); } /// Performs sparse matrix-vector multiplication (y = alpha*A*x + beta*y) using CSC format. template - void spmv_csc(int stream_id, int m, int n, int nnz, T alpha, - mem_ptr col_ptr, mem_ptr row_ind, mem_ptr values, - mem_ptr x, T beta, mem_ptr y, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) - throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCsc(&matA, m, n, nnz, - reinterpret_cast(col_ptr->mem()), - reinterpret_cast(row_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, - CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnVecDescr_t vecX, vecY; - cusparseCreateDnVec(&vecX, n, reinterpret_cast(x->mem()), type); - cusparseCreateDnVec(&vecY, m, reinterpret_cast(y->mem()), type); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMV(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, vecX, &beta, vecY, type, - CUSPARSE_SPMV_ALG_DEFAULT, d_workspace); - - if (dBuffer) - CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnVec(vecX); - cusparseDestroyDnVec(vecY); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMV (CSC) failed on device " + std::to_string(id_)); + void spmv_csc(int stream_id, int m, int n, int nnz, T alpha, + mem_ptr col_ptr, mem_ptr row_ind, + mem_ptr values, mem_ptr x, T beta, mem_ptr y, + mem_ptr workspace = nullptr) { + spmv_impl(sparse_format::csc, stream_id, m, n, nnz, alpha, col_ptr, + row_ind, values, x, beta, y, workspace); } /// Returns the required buffer size for SpMM CSR. template - size_t spmm_csr_buffer_size(int stream_id, int m, int n, int k, int nnz, - mem_ptr row_ptr, mem_ptr col_ind, mem_ptr values, - mem_ptr B, mem_ptr C) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCsr(&matA, m, k, nnz, reinterpret_cast(row_ptr->mem()), - reinterpret_cast(col_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmm_csr_buffer_size(int stream_id, int m, int n, int k, int nnz, + mem_ptr row_ptr, mem_ptr col_ind, + mem_ptr values, mem_ptr B, + mem_ptr C) { + return spmm_buffer_size_impl( + sparse_format::csr, stream_id, m, n, k, nnz, row_ptr, col_ind, + values, B, C); } /// Returns the required buffer size for SpMM COO. template - size_t spmm_coo_buffer_size(int stream_id, int m, int n, int k, int nnz, - mem_ptr row_ind, mem_ptr col_ind, mem_ptr values, - mem_ptr B, mem_ptr C) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCoo(&matA, m, k, nnz, reinterpret_cast(row_ind->mem()), - reinterpret_cast(col_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmm_coo_buffer_size(int stream_id, int m, int n, int k, int nnz, + mem_ptr row_ind, mem_ptr col_ind, + mem_ptr values, mem_ptr B, + mem_ptr C) { + return spmm_buffer_size_impl( + sparse_format::coo, stream_id, m, n, k, nnz, row_ind, col_ind, + values, B, C); } /// Returns the required buffer size for SpMM CSC. template - size_t spmm_csc_buffer_size(int stream_id, int m, int n, int k, int nnz, - mem_ptr col_ptr, mem_ptr row_ind, mem_ptr values, - mem_ptr B, mem_ptr C) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled"); - CHECK_CUDA(cuCtxPushCurrent(context_)); - T alpha = T{1}; T beta = T{0}; - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - cusparseSpMatDescr_t matA; - cusparseCreateCsc(&matA, m, k, nnz, reinterpret_cast(col_ptr->mem()), - reinterpret_cast(row_ind->mem()), reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return bufferSize; + size_t spmm_csc_buffer_size(int stream_id, int m, int n, int k, int nnz, + mem_ptr col_ptr, mem_ptr row_ind, + mem_ptr values, mem_ptr B, + mem_ptr C) { + return spmm_buffer_size_impl( + sparse_format::csc, stream_id, m, n, k, nnz, col_ptr, row_ind, + values, B, C); } /// Performs sparse matrix-matrix multiplication (C = alpha*A*B + beta*C) using CSR format. /// A is sparse (m x k), B is dense (k x n), C is dense (m x n). template - void spmm_csr(int stream_id, int m, int n, int k, int nnz, T alpha, - mem_ptr row_ptr, mem_ptr col_ind, mem_ptr values, - mem_ptr B, T beta, mem_ptr C, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCsr(&matA, m, k, nnz, - reinterpret_cast(row_ptr->mem()), - reinterpret_cast(col_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, - CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMM(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, d_workspace); - - if (dBuffer) CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMM (CSR) failed on device " + std::to_string(id_)); + void spmm_csr(int stream_id, int m, int n, int k, int nnz, T alpha, + mem_ptr row_ptr, mem_ptr col_ind, + mem_ptr values, mem_ptr B, T beta, mem_ptr C, + mem_ptr workspace = nullptr) { + spmm_impl(sparse_format::csr, stream_id, m, n, k, nnz, alpha, row_ptr, + col_ind, values, B, beta, C, workspace); } /// Performs sparse matrix-matrix multiplication (C = alpha*A*B + beta*C) using COO format. template - void spmm_coo(int stream_id, int m, int n, int k, int nnz, T alpha, - mem_ptr row_ind, mem_ptr col_ind, mem_ptr values, - mem_ptr B, T beta, mem_ptr C, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCoo(&matA, m, k, nnz, - reinterpret_cast(row_ind->mem()), - reinterpret_cast(col_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMM(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, d_workspace); - - if (dBuffer) CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMM (COO) failed on device " + std::to_string(id_)); + void spmm_coo(int stream_id, int m, int n, int k, int nnz, T alpha, + mem_ptr row_ind, mem_ptr col_ind, + mem_ptr values, mem_ptr B, T beta, mem_ptr C, + mem_ptr workspace = nullptr) { + spmm_impl(sparse_format::coo, stream_id, m, n, k, nnz, alpha, row_ind, + col_ind, values, B, beta, C, workspace); } /// Performs sparse matrix-matrix multiplication (C = alpha*A*B + beta*C) using CSC format. template - void spmm_csc(int stream_id, int m, int n, int k, int nnz, T alpha, - mem_ptr col_ptr, mem_ptr row_ind, mem_ptr values, - mem_ptr B, T beta, mem_ptr C, mem_ptr workspace = nullptr) { - cusparseHandle_t handle = get_cusparse_handle(stream_id); - if (!handle) throw std::runtime_error("cuSparse not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - cudaDataType type; - if constexpr (std::is_same_v) - type = CUDA_R_64F; - else - type = CUDA_R_32F; - - CUstream stream = get_stream_for_actor(stream_id); - - cusparseSpMatDescr_t matA; - cusparseCreateCsc(&matA, m, k, nnz, - reinterpret_cast(col_ptr->mem()), - reinterpret_cast(row_ind->mem()), - reinterpret_cast(values->mem()), - CUSPARSE_INDEX_32I, CUSPARSE_INDEX_32I, - CUSPARSE_INDEX_BASE_ZERO, type); - - cusparseDnMatDescr_t matB, matC; - cusparseCreateDnMat(&matB, k, n, n, reinterpret_cast(B->mem()), type, CUSPARSE_ORDER_ROW); - cusparseCreateDnMat(&matC, m, n, n, reinterpret_cast(C->mem()), type, CUSPARSE_ORDER_ROW); - - void* d_workspace = nullptr; - CUdeviceptr dBuffer = 0; - if (workspace) { - d_workspace = reinterpret_cast(workspace->mem()); - } else { - size_t bufferSize = 0; - cusparseSpMM_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, &bufferSize); - if (bufferSize > 0) { - CHECK_CUDA(cuMemAllocAsync(&dBuffer, bufferSize, stream)); - d_workspace = reinterpret_cast(dBuffer); - } - } - - cusparseStatus_t status = cusparseSpMM(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, CUSPARSE_OPERATION_NON_TRANSPOSE, - &alpha, matA, matB, &beta, matC, type, - CUSPARSE_SPMM_ALG_DEFAULT, d_workspace); - - if (dBuffer) - CHECK_CUDA(cuMemFreeAsync(dBuffer, stream)); - - cusparseDestroySpMat(matA); - cusparseDestroyDnMat(matB); - cusparseDestroyDnMat(matC); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUSPARSE_STATUS_SUCCESS) - throw std::runtime_error("cusparseSpMM (CSC) failed on device " + std::to_string(id_)); + void spmm_csc(int stream_id, int m, int n, int k, int nnz, T alpha, + mem_ptr col_ptr, mem_ptr row_ind, + mem_ptr values, mem_ptr B, T beta, mem_ptr C, + mem_ptr workspace = nullptr) { + spmm_impl(sparse_format::csc, stream_id, m, n, k, nnz, alpha, col_ptr, + row_ind, values, B, beta, C, workspace); } /// Performs double precision matrix-vector multiplication (y = alpha*A*x + beta*y). /// Assumes A is in row-major order of dimensions m x n. void dgemv(int stream_id, int m, int n, double alpha, mem_ptr A, mem_ptr x, double beta, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - // Row-major matrix A (m x n) is stored as m rows of n elements. // Viewed as column-major by cuBLAS, this is a n x m matrix. // To compute y = A * x: // Op(Memory) * x = (n x m)^T * (n x 1) = (m x n) * (n x 1) = (m x 1). // We use CUBLAS_OP_T. LDA is the 'rows' in the column-major view, which is n. - cublasStatus_t status = cublasDgemv(handle, CUBLAS_OP_T, - n, m, - &alpha, - reinterpret_cast(A->mem()), n, - reinterpret_cast(x->mem()), 1, - &beta, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDgemv failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, x, y), "cublasDgemv", + [&](cublasHandle_t handle) { + return cublasDgemv( + handle, CUBLAS_OP_T, n, m, &alpha, + reinterpret_cast(A->mem()), n, + reinterpret_cast(x->mem()), 1, &beta, + reinterpret_cast(y->mem()), 1); + }); } /// Performs symmetric rank-k update (C = alpha*A*A^T + beta*C). /// Assumes A is in row-major order of dimensions n x k, and C is n x n. void ssyrk(int stream_id, int n, int k, float alpha, mem_ptr A, float beta, mem_ptr C) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - // Row-major matrix A (n x k) viewed as column-major is k x n. // To compute C = alpha * A * A^T + beta * C: // We use CUBLAS_OP_T so that (k x n)^T * (k x n) = (n x k) * (k x n) = n x n. // Note: We use CUBLAS_FILL_MODE_UPPER because the upper triangle in // column-major maps to the lower triangle in row-major layout. - cublasStatus_t status = cublasSsyrk(handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_T, - n, k, - &alpha, - reinterpret_cast(A->mem()), k, - &beta, - reinterpret_cast(C->mem()), n); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSsyrk failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, C), "cublasSsyrk", + [&](cublasHandle_t handle) { + return cublasSsyrk( + handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_T, n, k, &alpha, + reinterpret_cast(A->mem()), k, &beta, + reinterpret_cast(C->mem()), n); + }); } /// Performs double precision symmetric rank-k update (C = alpha*A*A^T + beta*C). /// Assumes A is in row-major order of dimensions n x k, and C is n x n. void dsyrk(int stream_id, int n, int k, double alpha, mem_ptr A, double beta, mem_ptr C) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - // Row-major matrix A (n x k) viewed as column-major is k x n. // To compute C = alpha * A * A^T + beta * C: // We use CUBLAS_OP_T so that (k x n)^T * (k x n) = (n x k) * (k x n) = n x n. // Note: We use CUBLAS_FILL_MODE_UPPER because the upper triangle in // column-major maps to the lower triangle in row-major layout. - cublasStatus_t status = cublasDsyrk(handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_T, - n, k, - &alpha, - reinterpret_cast(A->mem()), k, - &beta, - reinterpret_cast(C->mem()), n); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDsyrk failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, C), "cublasDsyrk", + [&](cublasHandle_t handle) { + return cublasDsyrk( + handle, CUBLAS_FILL_MODE_UPPER, CUBLAS_OP_T, n, k, &alpha, + reinterpret_cast(A->mem()), k, &beta, + reinterpret_cast(C->mem()), n); + }); } /// Performs single precision vector-vector addition (y = alpha*x + y). - void saxpy(int stream_id, int n, float alpha, mem_ptr x, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasSaxpy(handle, n, &alpha, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSaxpy failed on device " + std::to_string(id_)); + void saxpy(int stream_id, int n, float alpha, mem_ptr x, + mem_ptr y) { + run_cublas_operation( + stream_id, std::tie(x, y), "cublasSaxpy", + [&](cublasHandle_t handle) { + return cublasSaxpy( + handle, n, &alpha, reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1); + }); } /// Performs double precision vector-vector addition (y = alpha*x + y). - void daxpy(int stream_id, int n, double alpha, mem_ptr x, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasDaxpy(handle, n, &alpha, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDaxpy failed on device " + std::to_string(id_)); + void daxpy(int stream_id, int n, double alpha, mem_ptr x, + mem_ptr y) { + run_cublas_operation( + stream_id, std::tie(x, y), "cublasDaxpy", + [&](cublasHandle_t handle) { + return cublasDaxpy( + handle, n, &alpha, reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1); + }); } /// Performs single precision vector scaling (x = alpha*x). void sscal(int stream_id, int n, float alpha, mem_ptr x) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasSscal(handle, n, &alpha, - reinterpret_cast(x->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSscal failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x), "cublasSscal", + [&](cublasHandle_t handle) { + return cublasSscal(handle, n, &alpha, + reinterpret_cast(x->mem()), 1); + }); } /// Performs double precision vector scaling (x = alpha*x). void dscal(int stream_id, int n, double alpha, mem_ptr x) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasDscal(handle, n, &alpha, - reinterpret_cast(x->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDscal failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x), "cublasDscal", + [&](cublasHandle_t handle) { + return cublasDscal(handle, n, &alpha, + reinterpret_cast(x->mem()), 1); + }); } /// Performs single precision Euclidean norm (result = ||x||2). void snrm2(int stream_id, int n, mem_ptr x, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_DEVICE); - cublasStatus_t status = cublasSnrm2(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(result->mem())); - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSnrm2 failed on device " + std::to_string(id_)); + run_cublas_device_result_operation( + stream_id, std::tie(x, result), "cublasSnrm2", + [&](cublasHandle_t handle) { + return cublasSnrm2(handle, n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(result->mem())); + }); } /// Performs double precision Euclidean norm (result = ||x||2). void dnrm2(int stream_id, int n, mem_ptr x, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_DEVICE); - cublasStatus_t status = cublasDnrm2(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(result->mem())); - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDnrm2 failed on device " + std::to_string(id_)); + run_cublas_device_result_operation( + stream_id, std::tie(x, result), "cublasDnrm2", + [&](cublasHandle_t handle) { + return cublasDnrm2(handle, n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(result->mem())); + }); } /// Performs single precision dot product (result = x^T * y). - void sdot(int stream_id, int n, mem_ptr x, mem_ptr y, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_DEVICE); - cublasStatus_t status = cublasSdot(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1, - reinterpret_cast(result->mem())); - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSdot failed on device " + std::to_string(id_)); + void sdot(int stream_id, int n, mem_ptr x, mem_ptr y, + mem_ptr result) { + run_cublas_device_result_operation( + stream_id, std::tie(x, y, result), "cublasSdot", + [&](cublasHandle_t handle) { + return cublasSdot(handle, n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1, + reinterpret_cast(result->mem())); + }); } /// Performs double precision dot product (result = x^T * y). - void ddot(int stream_id, int n, mem_ptr x, mem_ptr y, mem_ptr result) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_DEVICE); - cublasStatus_t status = cublasDdot(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1, - reinterpret_cast(result->mem())); - cublasSetPointerMode(handle, CUBLAS_POINTER_MODE_HOST); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDdot failed on device " + std::to_string(id_)); + void ddot(int stream_id, int n, mem_ptr x, mem_ptr y, + mem_ptr result) { + run_cublas_device_result_operation( + stream_id, std::tie(x, y, result), "cublasDdot", + [&](cublasHandle_t handle) { + return cublasDdot(handle, n, + reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1, + reinterpret_cast(result->mem())); + }); } /// Copies vector x to vector y (y = x). void scopy(int stream_id, int n, mem_ptr x, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasScopy(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasScopy failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x, y), "cublasScopy", + [&](cublasHandle_t handle) { + return cublasScopy( + handle, n, reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1); + }); } /// Copies double vector x to vector y (y = x). void dcopy(int stream_id, int n, mem_ptr x, mem_ptr y) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasDcopy(handle, n, - reinterpret_cast(x->mem()), 1, - reinterpret_cast(y->mem()), 1); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDcopy failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(x, y), "cublasDcopy", + [&](cublasHandle_t handle) { + return cublasDcopy( + handle, n, reinterpret_cast(x->mem()), 1, + reinterpret_cast(y->mem()), 1); + }); } /// Performs single precision matrix-matrix multiplication (C = alpha*A*B + beta*C). /// Assumes A is m x k, B is k x n, and C is m x n, all in row-major order. - void sgemm(int stream_id, int m, int n, int k, float alpha, mem_ptr A, - mem_ptr B, float beta, mem_ptr C) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - + void sgemm(int stream_id, int m, int n, int k, float alpha, + mem_ptr A, mem_ptr B, float beta, + mem_ptr C) { // For row-major matrices A(m,k), B(k,n), C(m,n) to compute C = alpha*A*B + beta*C // cuBLAS expects column-major. The equivalent column-major operation is: // C_col = alpha * B_col * A_col + beta * C_col @@ -1075,29 +590,22 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { // So, we call cublasSgemm with: // A_cublas = B (transposed), B_cublas = A (transposed) // Dimensions: m_cublas = n, n_cublas = m, k_cublas = k - cublasStatus_t status = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, - n, m, k, - &alpha, - reinterpret_cast(B->mem()), n, // B is k x n row-major, lda = n - reinterpret_cast(A->mem()), k, // A is m x k row-major, ldb = k - &beta, - reinterpret_cast(C->mem()), n); // C is m x n row-major, ldc = n - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSgemm failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, B, C), "cublasSgemm", + [&](cublasHandle_t handle) { + return cublasSgemm( + handle, CUBLAS_OP_T, CUBLAS_OP_T, n, m, k, &alpha, + reinterpret_cast(B->mem()), n, + reinterpret_cast(A->mem()), k, &beta, + reinterpret_cast(C->mem()), n); + }); } /// Performs double precision matrix-matrix multiplication (C = alpha*A*B + beta*C). /// Assumes A is m x k, B is k x n, and C is m x n, all in row-major order. - void dgemm(int stream_id, int m, int n, int k, double alpha, mem_ptr A, - mem_ptr B, double beta, mem_ptr C) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - + void dgemm(int stream_id, int m, int n, int k, double alpha, + mem_ptr A, mem_ptr B, double beta, + mem_ptr C) { // For row-major matrices A(m,k), B(k,n), C(m,n) to compute C = alpha*A*B + beta*C // cuBLAS expects column-major. The equivalent column-major operation is: // C_col = alpha * B_col * A_col + beta * C_col @@ -1105,64 +613,52 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { // So, we call cublasDgemm with: // A_cublas = B (transposed), B_cublas = A (transposed) // Dimensions: m_cublas = n, n_cublas = m, k_cublas = k - cublasStatus_t status = cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, - n, m, k, - &alpha, - reinterpret_cast(B->mem()), n, // B is k x n row-major, lda = n - reinterpret_cast(A->mem()), k, // A is m x k row-major, ldb = k - &beta, - reinterpret_cast(C->mem()), n); // C is m x n row-major, ldc = n - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDgemm failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, B, C), "cublasDgemm", + [&](cublasHandle_t handle) { + return cublasDgemm( + handle, CUBLAS_OP_T, CUBLAS_OP_T, n, m, k, &alpha, + reinterpret_cast(B->mem()), n, + reinterpret_cast(A->mem()), k, &beta, + reinterpret_cast(C->mem()), n); + }); } /// Performs single precision strided batched matrix-matrix multiplication. - void sgemm_strided_batched(int stream_id, int m, int n, int k, float alpha, mem_ptr A, long long int strideA, - mem_ptr B, long long int strideB, float beta, mem_ptr C, long long int strideC, int batchCount) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - + void sgemm_strided_batched(int stream_id, int m, int n, int k, + float alpha, mem_ptr A, + long long int strideA, mem_ptr B, + long long int strideB, float beta, + mem_ptr C, long long int strideC, + int batchCount) { // Row-major A(m,k), B(k,n), C(m,n) -> cuBLAS (col-major): C = B * A - cublasStatus_t status = cublasSgemmStridedBatched(handle, CUBLAS_OP_T, CUBLAS_OP_T, - n, m, k, - &alpha, - reinterpret_cast(B->mem()), n, strideB, - reinterpret_cast(A->mem()), k, strideA, - &beta, - reinterpret_cast(C->mem()), n, strideC, - batchCount); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasSgemmStridedBatched failed on device " + std::to_string(id_)); + run_cublas_operation( + stream_id, std::tie(A, B, C), "cublasSgemmStridedBatched", + [&](cublasHandle_t handle) { + return cublasSgemmStridedBatched( + handle, CUBLAS_OP_T, CUBLAS_OP_T, n, m, k, &alpha, + reinterpret_cast(B->mem()), n, strideB, + reinterpret_cast(A->mem()), k, strideA, &beta, + reinterpret_cast(C->mem()), n, strideC, batchCount); + }); } /// Performs double precision strided batched matrix-matrix multiplication. - void dgemm_strided_batched(int stream_id, int m, int n, int k, double alpha, mem_ptr A, long long int strideA, - mem_ptr B, long long int strideB, double beta, mem_ptr C, long long int strideC, int batchCount) { - cublasHandle_t handle = get_cublas_handle(stream_id); - if (!handle) - throw std::runtime_error("cuBLAS not enabled on device " + std::to_string(id_)); - - CHECK_CUDA(cuCtxPushCurrent(context_)); - - cublasStatus_t status = cublasDgemmStridedBatched(handle, CUBLAS_OP_T, CUBLAS_OP_T, - n, m, k, - &alpha, - reinterpret_cast(B->mem()), n, strideB, - reinterpret_cast(A->mem()), k, strideA, - &beta, - reinterpret_cast(C->mem()), n, strideC, - batchCount); - - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - if (status != CUBLAS_STATUS_SUCCESS) - throw std::runtime_error("cublasDgemm failed on device " + std::to_string(id_)); + void dgemm_strided_batched(int stream_id, int m, int n, int k, + double alpha, mem_ptr A, + long long int strideA, mem_ptr B, + long long int strideB, double beta, + mem_ptr C, long long int strideC, + int batchCount) { + run_cublas_operation( + stream_id, std::tie(A, B, C), "cublasDgemmStridedBatched", + [&](cublasHandle_t handle) { + return cublasDgemmStridedBatched( + handle, CUBLAS_OP_T, CUBLAS_OP_T, n, m, k, &alpha, + reinterpret_cast(B->mem()), n, strideB, + reinterpret_cast(A->mem()), k, strideA, &beta, + reinterpret_cast(C->mem()), n, strideC, batchCount); + }); } // Overloads for make_arg using actor_id @@ -1202,76 +698,48 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { - //handling the case that a mem_ref is passed in - //should I force synchronization onto the same stream always? + // Reuse an existing allocation, ordering a cross-stream handoff if needed. template mem_ptr make_arg(mem_ptr arg, CUstream stream) { - - if (arg -> deviceID() != id_) { - - throw std::runtime_error("Error memory on device " + std::to_string(arg->deviceID()) + - " attempted to be used on a different device, device id was " + std::to_string(id_) + "\n"); - - } - //just return the arg back - return arg; + validate_memory_argument(arg); + if (arg->requires_stream_transition(stream)) { + detail::context_guard context_scope{getContext()}; + arg->transition_to_stream_in_current_context(stream); + } + return arg; } - //given a tuple of mem_ptrs - //will copy their data back to host and place them in an output buffer + // Copies OUT and IN_OUT arguments back to host in argument order. Buffers + // produced by one launch normally share a stream, so enqueue all copies + // before waiting for that stream once. template - std::vector collect_output_buffers_helper(const std::tuple& args) { - std::vector result; - std::apply([&](auto&&... mem) { - (([&] { - if (mem && (mem->access() == OUT || mem->access() == IN_OUT)) { - //using T = typename std::decay_t::value_type; - result.emplace_back(output_buffer{buffer_variant{mem->copy_to_host()}}); - } - })(), ...); - }, args); - return result; + std::vector + collect_output_buffers_helper(const std::tuple& args) { + auto target = find_output_collection_target(args); + if (!target.has_device_buffer) + return prepare_output_buffers(args); + if (!target.has_common_stream) + return collect_output_buffers_individually(args); + return collect_output_buffers_on_stream(args, target); } - //launches a kernel using wrapper types, in, in_out and out as arguments - //and returns a tuple of mem ref's that hold device memory - template - std::tuple>...> - launch_kernel_mem_ref(CUfunction kernel, - const nd_range& range, - std::tuple args, - int actor_id, - int shared_mem = 0 //in bytes - ) { - - // Step 1: Allocate mem_ref for each wrapper type - CUstream stream = get_stream_for_actor(actor_id); - auto mem_refs = std::apply([&](auto&&... arg) { - return std::make_tuple(make_arg(std::forward(arg), stream)...); - }, args); - - // Step 2: Prepare kernel argument pointers - auto kernel_args = prepare_kernel_args(mem_refs); - - // Step 3: Launch kernel - CHECK_CUDA(cuCtxPushCurrent(getContext())); - launch_kernel_internal(kernel, range, stream, kernel_args.ptrs.data(),shared_mem); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - - // Step 4: Clean up kernel argument pointers - cleanup_kernel_args(kernel_args); - - // Step 5: Return tuple of mem_ref... - return mem_refs; -} - - + // Launches a kernel using wrapper types and returns the resulting mem_refs. + template + std::tuple>...> + launch_kernel_mem_ref(CUfunction kernel, + const nd_range& range, + std::tuple args, + int actor_id, + int shared_mem = 0) { + return launch_kernel_mem_ref_impl(kernel, range, args, actor_id, + shared_mem); + } // Launch kernel with args that have already been allocated // on the device via mem_ref @@ -1281,21 +749,21 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { std::tuple args, int actor_id) { CUstream stream = get_stream_for_actor(actor_id); - CHECK_CUDA(cuCtxPushCurrent(getContext())); - - auto kernel_args = prepare_kernel_args(args); - launch_kernel_internal(kernel, range, stream, kernel_args.ptrs.data()); - - //CHECK_CUDA(cuStreamSynchronize(stream)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + { + detail::context_guard context_scope{getContext()}; + prepare_memory_arguments_for_stream_in_current_context(args, stream); + stack_kernel_arg_pack kernel_args; + prepare_stack_kernel_args(args, kernel_args); + auto* parameters = sizeof...(Ts) == 0 ? nullptr + : kernel_args.ptrs.data(); + launch_kernel_internal(kernel, range, stream, parameters); + } auto outputs = collect_output_buffers(args); - cleanup_kernel_args(kernel_args); return outputs; } - // For testing: scalar/buffer detection and cleanup struct kernel_arg_pack { std::vector ptrs; std::vector allocated_device_ptrs; // Buffers only @@ -1341,21 +809,931 @@ class CAF_CUDA_EXPORT device : public caf::ref_counted { - // === Old method for legacy tests === template std::vector extract_kernel_args(const std::tuple& t) { return extract_kernel_args_impl(t, std::index_sequence_for{}); } private: + friend class program; + friend class scheduler; + template + friend class command_runner; + template + friend class bulk_memory_command; + + template + void run_cublas_operation(int stream_id, + const MemoryTuple& memory, + const char* operation, + Function&& function, + cublasPointerMode_t pointer_mode + = CUBLAS_POINTER_MODE_HOST) { + if (!cublas_table_) { + throw std::runtime_error("cuBLAS not enabled on device " + + std::to_string(id_)); + } + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + validate_device_buffer_arguments(memory, "cuBLAS"); + transition_memory_arguments_for_stream_in_current_context(memory, + stream); + auto handle = cublas_table_->get_handle_with_pointer_mode( + stream_id, stream, pointer_mode); + auto status = std::forward(function)(handle); + if (status != CUBLAS_STATUS_SUCCESS) { + throw std::runtime_error(std::string{operation} + " failed on device " + + std::to_string(id_) + " with cuBLAS status " + + std::to_string(static_cast(status))); + } + } + + template + void validate_device_buffer_argument(const MemPtr& memory, + const char* library) const { + validate_memory_argument(memory); + if (memory->is_scalar()) { + throw std::invalid_argument(std::string{library} + + " operations require device buffers"); + } + } + + template + void validate_device_buffer_arguments(const std::tuple& memory, + const char* library) const { + std::apply( + [&](const auto&... value) { + (validate_device_buffer_argument(value, library), ...); + }, + memory); + } + + template + void run_cublas_device_result_operation(int stream_id, + const MemoryTuple& memory, + const char* operation, + Function&& function) { + run_cublas_operation( + stream_id, memory, operation, + [&](cublasHandle_t handle) { + return std::forward(function)(handle); + }, + CUBLAS_POINTER_MODE_DEVICE); + } + + enum class sparse_format { + csr, + coo, + csc, + }; + + static const char* spmv_operation_name(sparse_format format) noexcept { + switch (format) { + case sparse_format::csr: + return "cusparseSpMV (CSR)"; + case sparse_format::coo: + return "cusparseSpMV (COO)"; + case sparse_format::csc: + return "cusparseSpMV (CSC)"; + } + return "cusparseSpMV"; + } + + static const char* spmm_operation_name(sparse_format format) noexcept { + switch (format) { + case sparse_format::csr: + return "cusparseSpMM (CSR)"; + case sparse_format::coo: + return "cusparseSpMM (COO)"; + case sparse_format::csc: + return "cusparseSpMM (CSC)"; + } + return "cusparseSpMM"; + } + + static void report_cusparse_cleanup_error(cusparseStatus_t status, + const char* operation) noexcept { + if (status != CUSPARSE_STATUS_SUCCESS) { + std::fprintf(stderr, "%s failed with cuSPARSE status %d\n", operation, + static_cast(status)); + } + } + + struct sparse_matrix_deleter { + void operator()(cusparseSpMatDescr_t descriptor) const noexcept { + report_cusparse_cleanup_error(cusparseDestroySpMat(descriptor), + "cusparseDestroySpMat"); + } + }; + + struct dense_vector_deleter { + void operator()(cusparseDnVecDescr_t descriptor) const noexcept { + report_cusparse_cleanup_error(cusparseDestroyDnVec(descriptor), + "cusparseDestroyDnVec"); + } + }; + + struct dense_matrix_deleter { + void operator()(cusparseDnMatDescr_t descriptor) const noexcept { + report_cusparse_cleanup_error(cusparseDestroyDnMat(descriptor), + "cusparseDestroyDnMat"); + } + }; + + template + class sparse_descriptor_owner { + public: + sparse_descriptor_owner() = default; + + ~sparse_descriptor_owner() noexcept { + if (descriptor_) + Deleter{}(descriptor_); + } + + sparse_descriptor_owner(const sparse_descriptor_owner&) = delete; + sparse_descriptor_owner& operator=(const sparse_descriptor_owner&) = delete; + + sparse_descriptor_owner(sparse_descriptor_owner&& other) noexcept + : descriptor_(std::exchange(other.descriptor_, nullptr)) { + // nop + } + + sparse_descriptor_owner& operator=(sparse_descriptor_owner&&) = delete; + + Descriptor* put() noexcept { + return &descriptor_; + } + + Descriptor get() const noexcept { + return descriptor_; + } + + private: + Descriptor descriptor_ = nullptr; + }; + + using sparse_matrix_descriptor = + sparse_descriptor_owner; + using dense_vector_descriptor = + sparse_descriptor_owner; + using dense_matrix_descriptor = + sparse_descriptor_owner; + + class sparse_temporary_workspace { + public: + sparse_temporary_workspace(size_t size, CUstream stream) + : stream_(stream) { + if (size > 0) + CHECK_CUDA(cuMemAllocAsync(&memory_, size, stream_)); + } + + ~sparse_temporary_workspace() noexcept { + if (!memory_) + return; + auto result = cuMemFreeAsync(memory_, stream_); + if (result != CUDA_SUCCESS) { + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "cuMemFreeAsync sparse cleanup failed: %s\n", + description ? description : "unknown CUDA error"); + } + } + + sparse_temporary_workspace(const sparse_temporary_workspace&) = delete; + sparse_temporary_workspace& + operator=(const sparse_temporary_workspace&) = delete; + sparse_temporary_workspace(sparse_temporary_workspace&&) = delete; + sparse_temporary_workspace& + operator=(sparse_temporary_workspace&&) = delete; + + void* data() const noexcept { + return reinterpret_cast(memory_); + } + + void release_checked() { + if (!memory_) + return; + CHECK_CUDA(cuMemFreeAsync(memory_, stream_)); + memory_ = 0; + } + + private: + CUdeviceptr memory_ = 0; + CUstream stream_ = nullptr; + }; + + void require_cusparse_enabled() const { + if (!cusparse_table_) { + throw std::runtime_error("cuSPARSE not enabled on device " + + std::to_string(id_)); + } + } + + void check_cusparse(cusparseStatus_t status, + const char* operation) const { + if (status != CUSPARSE_STATUS_SUCCESS) { + throw std::runtime_error( + std::string{operation} + " failed on device " + std::to_string(id_) + + " with cuSPARSE status " + + std::to_string(static_cast(status))); + } + } + + template + static constexpr cudaDataType sparse_data_type() noexcept { + static_assert(std::is_same_v || std::is_same_v, + "cuSPARSE operations support float and double values"); + if constexpr (std::is_same_v) + return CUDA_R_64F; + else + return CUDA_R_32F; + } + + template + sparse_matrix_descriptor make_sparse_matrix_descriptor( + sparse_format format, int rows, int columns, int nonzero_count, + const mem_ptr& first_indices, const mem_ptr& second_indices, + const mem_ptr& values) const { + sparse_matrix_descriptor result; + auto value_type = sparse_data_type(); + cusparseStatus_t status = CUSPARSE_STATUS_INVALID_VALUE; + const char* operation = nullptr; + switch (format) { + case sparse_format::csr: + operation = "cusparseCreateCsr"; + status = cusparseCreateCsr( + result.put(), rows, columns, nonzero_count, + reinterpret_cast(first_indices->mem()), + reinterpret_cast(second_indices->mem()), + reinterpret_cast(values->mem()), CUSPARSE_INDEX_32I, + CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, value_type); + break; + case sparse_format::coo: + operation = "cusparseCreateCoo"; + status = cusparseCreateCoo( + result.put(), rows, columns, nonzero_count, + reinterpret_cast(first_indices->mem()), + reinterpret_cast(second_indices->mem()), + reinterpret_cast(values->mem()), CUSPARSE_INDEX_32I, + CUSPARSE_INDEX_BASE_ZERO, value_type); + break; + case sparse_format::csc: + operation = "cusparseCreateCsc"; + status = cusparseCreateCsc( + result.put(), rows, columns, nonzero_count, + reinterpret_cast(first_indices->mem()), + reinterpret_cast(second_indices->mem()), + reinterpret_cast(values->mem()), CUSPARSE_INDEX_32I, + CUSPARSE_INDEX_32I, CUSPARSE_INDEX_BASE_ZERO, value_type); + break; + } + if (!operation) + throw std::logic_error("invalid cuSPARSE matrix format"); + check_cusparse(status, operation); + return result; + } + + template + void prepare_sparse_memory_arguments(const MemoryTuple& memory, + const mem_ptr& workspace, + CUstream stream) { + if (workspace) { + auto all_memory = std::tuple_cat(memory, std::tie(workspace)); + validate_device_buffer_arguments(all_memory, "cuSPARSE"); + transition_memory_arguments_for_stream_in_current_context(all_memory, + stream); + } else { + validate_device_buffer_arguments(memory, "cuSPARSE"); + transition_memory_arguments_for_stream_in_current_context(memory, + stream); + } + } + + template + size_t spmv_buffer_size_impl( + sparse_format format, int stream_id, int m, int n, int nnz, + const mem_ptr& first_indices, const mem_ptr& second_indices, + const mem_ptr& values, const mem_ptr& x, + const mem_ptr& y) { + require_cusparse_enabled(); + validate_device_buffer_arguments( + std::tie(first_indices, second_indices, values, x, y), "cuSPARSE"); + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + auto handle = cusparse_table_->get_handle_impl(stream_id, stream, false); + auto value_type = sparse_data_type(); + auto matrix = make_sparse_matrix_descriptor( + format, m, n, nnz, first_indices, second_indices, values); + dense_vector_descriptor vector_x; + check_cusparse( + cusparseCreateDnVec(vector_x.put(), n, + reinterpret_cast(x->mem()), value_type), + "cusparseCreateDnVec(x)"); + dense_vector_descriptor vector_y; + check_cusparse( + cusparseCreateDnVec(vector_y.put(), m, + reinterpret_cast(y->mem()), value_type), + "cusparseCreateDnVec(y)"); + T alpha = T{1}; + T beta = T{0}; + size_t result = 0; + check_cusparse( + cusparseSpMV_bufferSize( + handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matrix.get(), + vector_x.get(), &beta, vector_y.get(), value_type, + CUSPARSE_SPMV_ALG_DEFAULT, &result), + "cusparseSpMV_bufferSize"); + return result; + } + + template + void spmv_impl( + sparse_format format, int stream_id, int m, int n, int nnz, T alpha, + const mem_ptr& first_indices, const mem_ptr& second_indices, + const mem_ptr& values, const mem_ptr& x, T beta, + const mem_ptr& y, const mem_ptr& workspace) { + require_cusparse_enabled(); + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + auto memory = std::tie(first_indices, second_indices, values, x, y); + prepare_sparse_memory_arguments(memory, workspace, stream); + auto handle = cusparse_table_->get_handle_impl(stream_id, stream, false); + auto value_type = sparse_data_type(); + auto matrix = make_sparse_matrix_descriptor( + format, m, n, nnz, first_indices, second_indices, values); + dense_vector_descriptor vector_x; + check_cusparse( + cusparseCreateDnVec(vector_x.put(), n, + reinterpret_cast(x->mem()), value_type), + "cusparseCreateDnVec(x)"); + dense_vector_descriptor vector_y; + check_cusparse( + cusparseCreateDnVec(vector_y.put(), m, + reinterpret_cast(y->mem()), value_type), + "cusparseCreateDnVec(y)"); + + size_t workspace_size = 0; + if (!workspace) { + check_cusparse( + cusparseSpMV_bufferSize( + handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matrix.get(), + vector_x.get(), &beta, vector_y.get(), value_type, + CUSPARSE_SPMV_ALG_DEFAULT, &workspace_size), + "cusparseSpMV_bufferSize"); + } + sparse_temporary_workspace owned_workspace{ + workspace ? 0 : workspace_size, stream}; + auto workspace_data = workspace + ? reinterpret_cast(workspace->mem()) + : owned_workspace.data(); + check_cusparse( + cusparseSpMV(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, + matrix.get(), vector_x.get(), &beta, vector_y.get(), + value_type, CUSPARSE_SPMV_ALG_DEFAULT, workspace_data), + spmv_operation_name(format)); + owned_workspace.release_checked(); + } + + template + size_t spmm_buffer_size_impl( + sparse_format format, int stream_id, int m, int n, int k, int nnz, + const mem_ptr& first_indices, const mem_ptr& second_indices, + const mem_ptr& values, const mem_ptr& matrix_b, + const mem_ptr& matrix_c) { + require_cusparse_enabled(); + validate_device_buffer_arguments( + std::tie(first_indices, second_indices, values, matrix_b, matrix_c), + "cuSPARSE"); + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + auto handle = cusparse_table_->get_handle_impl(stream_id, stream, false); + auto value_type = sparse_data_type(); + auto matrix_a = make_sparse_matrix_descriptor( + format, m, k, nnz, first_indices, second_indices, values); + dense_matrix_descriptor dense_b; + check_cusparse( + cusparseCreateDnMat( + dense_b.put(), k, n, n, reinterpret_cast(matrix_b->mem()), + value_type, CUSPARSE_ORDER_ROW), + "cusparseCreateDnMat(B)"); + dense_matrix_descriptor dense_c; + check_cusparse( + cusparseCreateDnMat( + dense_c.put(), m, n, n, reinterpret_cast(matrix_c->mem()), + value_type, CUSPARSE_ORDER_ROW), + "cusparseCreateDnMat(C)"); + T alpha = T{1}; + T beta = T{0}; + size_t result = 0; + check_cusparse( + cusparseSpMM_bufferSize( + handle, CUSPARSE_OPERATION_NON_TRANSPOSE, + CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matrix_a.get(), + dense_b.get(), &beta, dense_c.get(), value_type, + CUSPARSE_SPMM_ALG_DEFAULT, &result), + "cusparseSpMM_bufferSize"); + return result; + } + + template + void spmm_impl( + sparse_format format, int stream_id, int m, int n, int k, int nnz, + T alpha, const mem_ptr& first_indices, + const mem_ptr& second_indices, const mem_ptr& values, + const mem_ptr& matrix_b, T beta, const mem_ptr& matrix_c, + const mem_ptr& workspace) { + require_cusparse_enabled(); + auto stream = get_stream_for_actor(stream_id); + detail::context_guard context_scope{context_}; + auto memory = std::tie( + first_indices, second_indices, values, matrix_b, matrix_c); + prepare_sparse_memory_arguments(memory, workspace, stream); + auto handle = cusparse_table_->get_handle_impl(stream_id, stream, false); + auto value_type = sparse_data_type(); + auto matrix_a = make_sparse_matrix_descriptor( + format, m, k, nnz, first_indices, second_indices, values); + dense_matrix_descriptor dense_b; + check_cusparse( + cusparseCreateDnMat( + dense_b.put(), k, n, n, reinterpret_cast(matrix_b->mem()), + value_type, CUSPARSE_ORDER_ROW), + "cusparseCreateDnMat(B)"); + dense_matrix_descriptor dense_c; + check_cusparse( + cusparseCreateDnMat( + dense_c.put(), m, n, n, reinterpret_cast(matrix_c->mem()), + value_type, CUSPARSE_ORDER_ROW), + "cusparseCreateDnMat(C)"); + + size_t workspace_size = 0; + if (!workspace) { + check_cusparse( + cusparseSpMM_bufferSize( + handle, CUSPARSE_OPERATION_NON_TRANSPOSE, + CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matrix_a.get(), + dense_b.get(), &beta, dense_c.get(), value_type, + CUSPARSE_SPMM_ALG_DEFAULT, &workspace_size), + "cusparseSpMM_bufferSize"); + } + sparse_temporary_workspace owned_workspace{ + workspace ? 0 : workspace_size, stream}; + auto workspace_data = workspace + ? reinterpret_cast(workspace->mem()) + : owned_workspace.data(); + check_cusparse( + cusparseSpMM( + handle, CUSPARSE_OPERATION_NON_TRANSPOSE, + CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matrix_a.get(), + dense_b.get(), &beta, dense_c.get(), value_type, + CUSPARSE_SPMM_ALG_DEFAULT, workspace_data), + spmm_operation_name(format)); + owned_workspace.release_checked(); + } + + struct output_collection_target { + CUcontext context = nullptr; + CUstream stream = nullptr; + bool has_device_buffer = false; + bool has_common_stream = true; + }; + + template + static bool is_output_memory(const MemPtr& mem) noexcept { + return mem && (mem->access() == OUT || mem->access() == IN_OUT); + } + + template + static void update_output_collection_target(output_collection_target& target, + const MemPtr& mem) noexcept { + if (!is_output_memory(mem) || mem->is_scalar()) + return; + if (!target.has_device_buffer) { + target.context = mem->get_ctx(); + target.stream = mem->stream(); + target.has_device_buffer = true; + return; + } + if (target.context != mem->get_ctx() || target.stream != mem->stream()) + target.has_common_stream = false; + } + + template + static output_collection_target + find_output_collection_target(const std::tuple& args) noexcept { + output_collection_target result; + std::apply( + [&](const auto&... mem) { + (update_output_collection_target(result, mem), ...); + }, + args); + return result; + } + + template + static void append_output_buffer(std::vector& result, + const MemPtr& mem) { + if (!is_output_memory(mem)) + return; + using value_type = typename std::remove_reference_t< + decltype(*mem)>::value_type; + if (mem->is_scalar()) { + result.emplace_back(output_buffer{ + buffer_variant{std::vector{*mem->host_scalar_ptr()}}}); + } else { + result.emplace_back(output_buffer{ + buffer_variant{std::vector(mem->size())}}); + } + } + + template + static std::vector + prepare_output_buffers(const std::tuple& args) { + std::vector result; + result.reserve(sizeof...(Ts)); + std::apply( + [&](const auto&... mem) { + (append_output_buffer(result, mem), ...); + }, + args); + return result; + } + + template + static std::vector + collect_output_buffers_individually(const std::tuple& args) { + std::vector result; + result.reserve(sizeof...(Ts)); + std::apply( + [&](const auto&... mem) { + (([&] { + if (is_output_memory(mem)) { + result.emplace_back( + output_buffer{buffer_variant{mem->copy_to_host()}}); + } + }()), + ...); + }, + args); + return result; + } + + template + static void enqueue_output_copy(std::vector& result, + size_t& output_index, + bool& copy_attempted, + const MemPtr& mem) { + if (!is_output_memory(mem)) + return; + if (!mem->is_scalar()) { + using value_type = typename std::remove_reference_t< + decltype(*mem)>::value_type; + auto& host_data = std::get>( + result[output_index].data); + if (!host_data.empty()) { + copy_attempted = true; + mem->enqueue_copy_to_host_in_current_context(host_data.data(), + host_data.size()); + } + } + ++output_index; + } + + static CUresult + synchronize_output_collection(const output_collection_target& target) { + if (target.stream) + return cuStreamSynchronize(target.stream); + return cuCtxSynchronize(); + } + + template + static std::vector + collect_output_buffers_on_stream(const std::tuple& args, + const output_collection_target& target) { + auto result = prepare_output_buffers(args); + detail::context_guard context_scope{target.context}; + bool copy_attempted = false; + try { + size_t output_index = 0; + std::apply( + [&](const auto&... mem) { + (enqueue_output_copy(result, output_index, copy_attempted, mem), + ...); + }, + args); + CHECK_CUDA(synchronize_output_collection(target)); + } catch (...) { + if (copy_attempted + && synchronize_output_collection(target) != CUDA_SUCCESS) { + // A terminal CUDA failure leaves completion uncertain. Retain the host + // destinations rather than letting a pending asynchronous copy access + // released storage. This allocation occurs only on the failure path. + auto* retained = new (std::nothrow) + std::vector{std::move(result)}; + if (!retained) + std::terminate(); + static_cast(retained); + } + throw; + } + return result; + } + + template + struct stack_kernel_arg_pack { + stack_kernel_arg_pack() = default; + stack_kernel_arg_pack(const stack_kernel_arg_pack&) = delete; + stack_kernel_arg_pack& operator=(const stack_kernel_arg_pack&) = delete; + stack_kernel_arg_pack(stack_kernel_arg_pack&&) = delete; + stack_kernel_arg_pack& operator=(stack_kernel_arg_pack&&) = delete; + + std::array device_ptr_values{}; + std::array ptrs{}; + }; + + static void release_failed_allocation(CUdeviceptr memory, + CUstream stream) noexcept { + if (!memory) + return; + auto result = cuMemFreeAsync(memory, stream); + if (result != CUDA_SUCCESS) { + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "cuMemFreeAsync cleanup failed: %s\n", + description ? description : "unknown CUDA error"); + } + } + + class pending_allocation { + public: + pending_allocation(CUdeviceptr memory, CUstream stream) noexcept + : memory_(memory), stream_(stream) { + // nop + } + + ~pending_allocation() noexcept { + release_failed_allocation(memory_, stream_); + } + + pending_allocation(const pending_allocation&) = delete; + pending_allocation& operator=(const pending_allocation&) = delete; + pending_allocation(pending_allocation&&) = delete; + pending_allocation& operator=(pending_allocation&&) = delete; + + void disarm() noexcept { + memory_ = 0; + } + + private: + CUdeviceptr memory_; + CUstream stream_; + }; + + template + mem_ptr make_arg_in_current_context(const in& arg, + CUstream stream) { + return global_argument_in_current_context(arg, stream, IN); + } + + template + mem_ptr make_arg_in_current_context(const in_out& arg, + CUstream stream) { + return global_argument_in_current_context(arg, stream, IN_OUT); + } + + template + mem_ptr make_arg_in_current_context(const out& arg, + CUstream stream) { + return scratch_argument_in_current_context(arg, stream, OUT); + } + + template + mem_ptr make_arg_in_current_context(mem_ptr arg, CUstream stream) { + prepare_memory_argument_for_stream_in_current_context(arg, stream); + return arg; + } + + template + auto make_prepared_arg_in_current_context(Arg&& arg, CUstream stream) { + using arg_type = std::remove_cvref_t; + if constexpr (detail::is_cuda_mem_ptr_v) + return std::forward(arg); + else + return make_arg_in_current_context(std::forward(arg), stream); + } + + template + static bool argument_requires_context(const in& arg, + [[maybe_unused]] CUstream stream) { + return !arg.is_scalar(); + } + + template + static bool argument_requires_context(const in_out& arg, + [[maybe_unused]] CUstream stream) { + return !arg.is_scalar(); + } + + template + static bool argument_requires_context(const out&, + [[maybe_unused]] CUstream stream) + noexcept { + return true; + } + + template + static bool argument_requires_context(const mem_ptr& arg, + CUstream stream) noexcept { + return arg && arg->requires_stream_transition(stream); + } + + template + void validate_memory_argument(const mem_ptr& arg) const { + if (!arg) + throw std::invalid_argument("CUDA memory argument must not be null"); + if (arg->deviceID() != id_) { + throw std::runtime_error( + "memory on CUDA device " + std::to_string(arg->deviceID()) + + " cannot be used on device " + std::to_string(id_)); + } + if (!arg->belongs_to_context(context_lifetime_, context_)) { + throw std::runtime_error( + "CUDA memory belongs to a different or stale device context"); + } + } + + void validate_event(const event_ptr& e) const { + if (!e) + throw std::invalid_argument("CUDA event must not be null"); + if (!e->belongs_to_context(context_lifetime_)) + throw std::runtime_error( + "CUDA event belongs to a different or stale device context"); + } + + template + void prepare_memory_argument_for_stream_in_current_context( + const mem_ptr& arg, CUstream stream) { + validate_memory_argument(arg); + arg->transition_to_stream_in_current_context(stream); + } + + template + void prepare_memory_arguments_for_stream_in_current_context( + const std::tuple& args, CUstream stream) { + std::apply( + [&](const auto&... arg) { + (validate_stream_transition_argument(arg), ...); + }, + args); + transition_memory_arguments_for_stream_in_current_context(args, stream); + } + + template + void transition_memory_arguments_for_stream_in_current_context( + const std::tuple& args, CUstream stream) { + auto source_streams = std::apply( + [&](const auto&... arg) { + return std::array{ + stream_transition_source(arg, stream)...}; + }, + args); + for (size_t index = 0; index < source_streams.size(); ++index) { + auto source = source_streams[index]; + if (source == stream) + continue; + auto already_ordered = false; + for (size_t previous = 0; previous < index; ++previous) { + if (source_streams[previous] == source) { + already_ordered = true; + break; + } + } + if (!already_ordered) + detail::order_stream_after(stream, source); + } + std::apply( + [&](const auto&... arg) { + (complete_stream_transition(arg, stream), ...); + }, + args); + } + + template + void validate_stream_transition_argument(const Arg& arg) const { + using arg_type = std::remove_cvref_t; + if constexpr (detail::is_cuda_mem_ptr_v) + validate_memory_argument(arg); + } + + template + static CUstream stream_transition_source(const Arg& arg, + CUstream target) noexcept { + using arg_type = std::remove_cvref_t; + if constexpr (detail::is_cuda_mem_ptr_v) + return arg->requires_stream_transition(target) ? arg->stream() : target; + else + return target; + } + + template + static void complete_stream_transition(const Arg& arg, + CUstream target) noexcept { + using arg_type = std::remove_cvref_t; + if constexpr (detail::is_cuda_mem_ptr_v) + arg->complete_stream_transition(target); + } + + template + std::tuple>...> + make_args_mem_ref_impl(std::tuple& args, CUstream stream) { + auto convert = [&](auto&... arg) { + return std::tuple>...>{ + make_arg_in_current_context(arg, stream)...}; + }; + auto convert_prepared = [&](auto&... arg) { + return std::tuple>...>{ + make_prepared_arg_in_current_context(arg, stream)...}; + }; + auto needs_context = std::apply( + [&](const auto&... arg) { + return (argument_requires_context(arg, stream) || ... || false); + }, + args); + if (!needs_context) + return std::apply(convert, args); + detail::context_guard context_scope{getContext()}; + prepare_memory_arguments_for_stream_in_current_context(args, stream); + return std::apply(convert_prepared, args); + } + + template + std::tuple>...> + launch_kernel_mem_ref_impl(CUfunction kernel, + const nd_range& range, + std::tuple& args, + int actor_id, + int shared_mem) { + CUstream stream = get_stream_for_actor(actor_id); + detail::context_guard context_scope{getContext()}; + prepare_memory_arguments_for_stream_in_current_context(args, stream); + auto mem_refs = std::apply( + [&](auto&&... arg) { + return std::tuple>...>{ + make_prepared_arg_in_current_context( + std::forward(arg), stream)...}; + }, + std::move(args)); + + stack_kernel_arg_pack kernel_args; + prepare_stack_kernel_args(mem_refs, kernel_args); + + auto* parameters = sizeof...(Args) == 0 ? nullptr + : kernel_args.ptrs.data(); + launch_kernel_internal(kernel, range, stream, parameters, shared_mem); + return mem_refs; + } + + template + static void store_stack_kernel_arg(stack_kernel_arg_pack& pack, + size_t index, + const MemPtr& mem) { + if (mem->is_scalar()) { + pack.ptrs[index] = const_cast( + static_cast(mem->host_scalar_ptr())); + } else { + pack.device_ptr_values[index] = mem->mem(); + pack.ptrs[index] = static_cast(&pack.device_ptr_values[index]); + } + } + + template + static void + prepare_stack_kernel_args_impl(const std::tuple& args, + stack_kernel_arg_pack& pack, + std::index_sequence) { + (store_stack_kernel_arg(pack, Indices, std::get(args)), ...); + } + + template + static void + prepare_stack_kernel_args(const std::tuple& args, + stack_kernel_arg_pack& pack) { + static_assert(Size == sizeof...(Ts)); + prepare_stack_kernel_args_impl(args, pack, + std::index_sequence_for{}); + } + CUdevice device_; CUcontext context_; + detail::context_lifetime_ptr context_lifetime_; int id_; - const char* name_; + std::string name_; + size_t stream_pool_size_; std::unique_ptr stream_table_; std::unique_ptr cublas_table_; std::unique_ptr cusparse_table_; - std::mutex stream_mutex_; + // Protects cold lifecycle mutations. Ordinary submissions deliberately do + // not acquire this mutex; reset requires the device to be quiescent. + std::mutex lifecycle_mutex_; // Cached GPU properties (queried once during construction) int sm_count_ = 0; @@ -1433,48 +1811,83 @@ mem_ptr scratch_argument(const out& arg, int actor_id, int access) { // allocate a readonly input buffer on the GPU template mem_ptr global_argument(const in& arg, CUstream stream, int access) { - if (arg.is_scalar()) { - return caf::intrusive_ptr>( - new mem_ref(arg.getscalar(), access, id_, 0, getContext(), stream)); - } - size_t bytes = arg.size() * sizeof(T); - CUdeviceptr dev_ptr; - CHECK_CUDA(cuCtxPushCurrent(getContext())); - CHECK_CUDA(cuMemAllocAsync(&dev_ptr, bytes, stream)); - CHECK_CUDA(cuMemcpyHtoDAsync(dev_ptr, arg.data(), bytes, stream)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return caf::intrusive_ptr>( - new mem_ref(arg.size(), dev_ptr, access, id_, 0, getContext(), stream)); + if (arg.is_scalar()) + return global_argument_in_current_context(arg, stream, access); + detail::context_guard context_scope{getContext()}; + return global_argument_in_current_context(arg, stream, access); } // allocate a read/write input buffer on the GPU template mem_ptr global_argument(const in_out& arg, CUstream stream, int access) { - if (arg.is_scalar()) { - return caf::intrusive_ptr>( - new mem_ref(arg.getscalar(), access, id_, 0, getContext(), stream)); - } - size_t bytes = arg.size() * sizeof(T); - CUdeviceptr dev_ptr; - CHECK_CUDA(cuCtxPushCurrent(getContext())); - CHECK_CUDA(cuMemAllocAsync(&dev_ptr, bytes, stream)); - CHECK_CUDA(cuMemcpyHtoDAsync(dev_ptr, arg.data(), bytes, stream)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return caf::intrusive_ptr>( - new mem_ref(arg.size(), dev_ptr, access, id_, 0, getContext(), stream)); + if (arg.is_scalar()) + return global_argument_in_current_context(arg, stream, access); + detail::context_guard context_scope{getContext()}; + return global_argument_in_current_context(arg, stream, access); } // allocate an output buffer on the GPU template mem_ptr scratch_argument(const out& arg, CUstream stream, int access) { - size_t size = arg.size(); - CUdeviceptr dev_ptr; - CHECK_CUDA(cuCtxPushCurrent(getContext())); - CHECK_CUDA(cuMemAllocAsync(&dev_ptr, size * sizeof(T), stream)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return caf::intrusive_ptr>( - new mem_ref(size, dev_ptr, access, id_, 0, getContext(), stream)); -} + detail::context_guard context_scope{getContext()}; + return scratch_argument_in_current_context(arg, stream, access); +} + +template +mem_ptr global_argument_in_current_context(const in& arg, + CUstream stream, + int access) { + if (arg.is_scalar()) { + return mem_ptr{ + new mem_ref(arg.getscalar(), access, id_, 0, context_, stream)}; + } + + CUdeviceptr memory = 0; + CHECK_CUDA(cuMemAllocAsync(&memory, arg.size() * sizeof(T), stream)); + pending_allocation cleanup{memory, stream}; + CHECK_CUDA(cuMemcpyHtoDAsync(memory, arg.data(), + arg.size() * sizeof(T), stream)); + auto result = mem_ptr{ + new mem_ref(arg.size(), memory, access, id_, 0, context_, stream)}; + result->bind_context_lifetime(context_lifetime_); + cleanup.disarm(); + return result; +} + +template +mem_ptr global_argument_in_current_context(const in_out& arg, + CUstream stream, + int access) { + if (arg.is_scalar()) { + return mem_ptr{ + new mem_ref(arg.getscalar(), access, id_, 0, context_, stream)}; + } + + CUdeviceptr memory = 0; + CHECK_CUDA(cuMemAllocAsync(&memory, arg.size() * sizeof(T), stream)); + pending_allocation cleanup{memory, stream}; + CHECK_CUDA(cuMemcpyHtoDAsync(memory, arg.data(), + arg.size() * sizeof(T), stream)); + auto result = mem_ptr{ + new mem_ref(arg.size(), memory, access, id_, 0, context_, stream)}; + result->bind_context_lifetime(context_lifetime_); + cleanup.disarm(); + return result; +} + +template +mem_ptr scratch_argument_in_current_context(const out& arg, + CUstream stream, + int access) { + CUdeviceptr memory = 0; + CHECK_CUDA(cuMemAllocAsync(&memory, arg.size() * sizeof(T), stream)); + pending_allocation cleanup{memory, stream}; + auto result = mem_ptr{ + new mem_ref(arg.size(), memory, access, id_, 0, context_, stream)}; + result->bind_context_lifetime(context_lifetime_); + cleanup.disarm(); + return result; +} // === Kernel launch core === void launch_kernel_internal(CUfunction kernel, @@ -1493,7 +1906,6 @@ mem_ptr scratch_argument(const out& arg, CUstream stream, int access) { (err_name ? err_name : "unknown error")); } } - // === Legacy helper === template std::vector extract_kernel_args_impl(const Tuple& t, std::index_sequence) { diff --git a/libcaf_cuda/caf/cuda/event.hpp b/libcaf_cuda/caf/cuda/event.hpp index eec93781a2..7fb5e0b868 100644 --- a/libcaf_cuda/caf/cuda/event.hpp +++ b/libcaf_cuda/caf/cuda/event.hpp @@ -1,67 +1,101 @@ #pragma once #include -#include -#include -#include +#include +#include + #include +#include + +#include "caf/cuda/detail/context_guard.hpp" +#include "caf/cuda/detail/context_lifetime.hpp" namespace caf::cuda { +class device; class event; using event_ptr = caf::intrusive_ptr; -/** - * @brief A smart-pointer managed wrapper for a CUDA event. - * Ensures that the CUevent handle is destroyed when the last reference is gone. - */ +/// A reference-counted CUDA event. class event : public caf::ref_counted { public: + /// Creates an event in the current CUDA context. explicit event(unsigned int flags = CU_EVENT_DEFAULT) { - // Note: Expects an active CUDA context for creation. - check_error(cuEventCreate(&event_, flags), "cuEventCreate"); + auto result = cuEventCreate(&event_, flags); + if (result != CUDA_SUCCESS) + throw detail::cuda_error(result, "cuEventCreate"); } - ~event() { - check_error(cuEventDestroy(event_), "cuEventDestroy"); + ~event() noexcept override { + if (!event_ || !context_is_live()) + return; + report_error(cuEventDestroy(event_), "cuEventDestroy"); } event(const event&) = delete; event& operator=(const event&) = delete; - CUevent get() const { + CUevent get() const noexcept { return event_; } - /// Returns true if the event has been recorded and the work has completed. + /// Returns true if the event has completed. bool query() const { - CUresult res = cuEventQuery(event_); - if (res == CUDA_SUCCESS) + require_live_context(); + auto result = cuEventQuery(event_); + if (result == CUDA_SUCCESS) return true; - if (res == CUDA_ERROR_NOT_READY) + if (result == CUDA_ERROR_NOT_READY) return false; - check_error(res, "cuEventQuery"); - return false; + throw detail::cuda_error(result, "cuEventQuery"); } - /// Blocks the calling thread until the event has completed. + /// Blocks until the event has completed. void synchronize() const { - check_error(cuEventSynchronize(event_), "cuEventSynchronize"); + require_live_context(); + auto result = cuEventSynchronize(event_); + if (result != CUDA_SUCCESS) + throw detail::cuda_error(result, "cuEventSynchronize"); } private: - static void check_error(CUresult result, const char* msg) { - if (result != CUDA_SUCCESS) { - const char* err_str = nullptr; - cuGetErrorString(result, &err_str); - std::cerr << "CUDA Driver API Error (" << msg << "): " - << (err_str ? err_str : "unknown error") << "\n"; - std::exit(1); - } + friend class device; + + void bind_context_lifetime( + const detail::context_lifetime_ptr& lifetime) noexcept { + context_lifetime_ = lifetime; + } + + bool belongs_to_context( + const detail::context_lifetime_ptr& lifetime) const noexcept { + return !context_lifetime_ + || (context_lifetime_.get() == lifetime.get() + && context_lifetime_->active()); + } + + bool context_is_live() const noexcept { + return !context_lifetime_ || context_lifetime_->active(); + } + + void require_live_context() const { + if (!context_is_live()) + throw std::runtime_error( + "CUDA event belongs to a stale device context"); + } + + static void report_error(CUresult result, + const char* operation) noexcept { + if (result == CUDA_SUCCESS) + return; + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "%s failed: %s\n", operation, + description ? description : "unknown CUDA error"); } - CUevent event_; + CUevent event_ = nullptr; + detail::context_lifetime_ptr context_lifetime_; }; -} // namespace caf::cuda \ No newline at end of file +} // namespace caf::cuda diff --git a/libcaf_cuda/caf/cuda/mem_ref.hpp b/libcaf_cuda/caf/cuda/mem_ref.hpp index 3ea1b04149..07528b1fe7 100644 --- a/libcaf_cuda/caf/cuda/mem_ref.hpp +++ b/libcaf_cuda/caf/cuda/mem_ref.hpp @@ -5,10 +5,17 @@ #include #include #include +#include +#include #include #include -#include +#include +#include +#include #include "caf/cuda/types.hpp" +#include "caf/cuda/detail/context_guard.hpp" +#include "caf/cuda/detail/context_lifetime.hpp" +#include "caf/cuda/detail/stream_order.hpp" #include //#include "caf/cuda/utility.hpp" #include @@ -17,7 +24,16 @@ namespace caf::cuda { -//A class that is a handle to gpu memory +class device; +class scheduler; + +template +class copy_back_command; + +/// A shared handle to GPU memory. +/// +/// Submissions that use aliases of the same handle must be serialized. In +/// particular, concurrent submissions to different streams are not supported. template class mem_ref : public caf::ref_counted { public: @@ -40,6 +56,7 @@ class mem_ref : public caf::ref_counted { ctx(context), is_scalar_(false) { + initialize_reference_count(); if (memory_ == 0) std::abort(); // unchanged } @@ -61,15 +78,48 @@ class mem_ref : public caf::ref_counted { is_scalar_(true), host_scalar_(scalar_value) { + initialize_reference_count(); // no cuMemAlloc, nothing to do } - ~mem_ref() { + ~mem_ref() noexcept override { reset(); } - mem_ref(mem_ref&&) noexcept = default; - mem_ref& operator=(mem_ref&&) noexcept = default; + mem_ref(mem_ref&& other) noexcept + : num_elements_(other.num_elements_), + memory_(other.memory_), + access_(other.access_), + device_id(other.device_id), + context_id(other.context_id), + stream_(other.stream_), + ctx(other.ctx), + context_lifetime_(std::move(other.context_lifetime_)), + token(std::move(other.token)), + is_scalar_(other.is_scalar_), + host_scalar_(std::move(other.host_scalar_)) { + initialize_reference_count(); + other.invalidate_after_move(); + } + + mem_ref& operator=(mem_ref&& other) noexcept { + if (this != &other) { + reset(); + num_elements_ = other.num_elements_; + memory_ = other.memory_; + access_ = other.access_; + device_id = other.device_id; + context_id = other.context_id; + stream_ = other.stream_; + ctx = other.ctx; + context_lifetime_ = std::move(other.context_lifetime_); + token = std::move(other.token); + is_scalar_ = other.is_scalar_; + host_scalar_ = std::move(other.host_scalar_); + other.invalidate_after_move(); + } + return *this; + } mem_ref(const mem_ref&) = delete; mem_ref& operator=(const mem_ref&) = delete; @@ -87,35 +137,32 @@ class mem_ref : public caf::ref_counted { //if it is ever needed, you can force synchronization on a mem_ptr //to ensure data on the device that the mem_ptr points to //is not in the middle of being operated on - //mostly here to avoid race conditions that can occur between streams + // Mostly here to avoid race conditions that can occur between streams. void synchronize() { - CHECK_CUDA(cuCtxPushCurrent(ctx)); + if (is_scalar_) + return; + require_live_context(); + detail::context_guard context_scope{ctx}; CUstream s = stream_ ? stream_ : nullptr; if (s) CHECK_CUDA(cuStreamSynchronize(s)); else CHECK_CUDA(cuCtxSynchronize()); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); } //Frees the memory on the gpu and //sets all its attributes to null or -1 void reset() { + // Destruction cannot report failure to the caller. Treat cleanup as + // best-effort and always invalidate the handle after reporting errors. if (!is_scalar_ && memory_) { - if (ctx) { - CUresult res = cuCtxPushCurrent(ctx); - if (res == CUDA_SUCCESS) { - CUresult free_res = cuMemFreeAsync(memory_, stream_); - if (free_res != CUDA_SUCCESS) { - const char* err_str = nullptr; - cuGetErrorString(free_res, &err_str); - std::cerr << "[ERROR] cuMemFreeAsync failed in mem_ref::reset: " - << (err_str ? err_str : "unknown error") << std::endl; - } - cuCtxPopCurrent(nullptr); - } else { - const char* err_str = nullptr; - cuGetErrorString(res, &err_str); - std::cerr << "[ERROR] cuCtxPushCurrent failed in mem_ref::reset: " - << (err_str ? err_str : "unknown error") << std::endl; + if (ctx && context_is_live()) { + try { + detail::context_guard context_scope{ctx}; + report_cuda_error(cuMemFreeAsync(memory_, stream_), + "cuMemFreeAsync in mem_ref::reset"); + } catch (const std::exception& error) { + std::fprintf(stderr, "mem_ref::reset failed: %s\n", error.what()); + } catch (...) { + std::fprintf(stderr, "mem_ref::reset failed: unknown error\n"); } } memory_ = 0; @@ -124,15 +171,17 @@ class mem_ref : public caf::ref_counted { access_ = -1; stream_ = nullptr; ctx = nullptr; + context_lifetime_.reset(); } // Enqueues a free operation on the specified stream and invalidates this reference. void free_on(CUstream s) { if (!is_scalar_ && memory_) { if (ctx) { - CHECK_CUDA(cuCtxPushCurrent(ctx)); + require_live_context(); + detail::context_guard context_scope{ctx}; + transition_to_stream_in_current_context(s); CHECK_CUDA(cuMemFreeAsync(memory_, s)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); } memory_ = 0; } @@ -140,70 +189,168 @@ class mem_ref : public caf::ref_counted { //copies gpu memory back to cpu memory in the form of an std::vector std::vector copy_to_host() const { - if (access_ == IN) - { - throw std::runtime_error("Cannt copy a read only buffer back to device\n"); - } - if (is_scalar_) { - return std::vector{host_scalar_}; - } - std::vector host_data(num_elements_); - size_t bytes = num_elements_ * sizeof(T); - CHECK_CUDA(cuCtxPushCurrent(ctx)); - CUstream s = stream_ ? stream_ : nullptr; - CHECK_CUDA(cuMemcpyDtoHAsync(host_data.data(), memory_, bytes, s)); - if (s) CHECK_CUDA(cuStreamSynchronize(s)); - else CHECK_CUDA(cuCtxSynchronize()); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); - return host_data; + validate_copy_destination(nullptr, num_elements_, false); + if (is_scalar_) + return std::vector{host_scalar_}; + + std::vector host_data(num_elements_); + require_live_context(); + detail::context_guard context_scope{ctx}; + bool copy_attempted = false; + try { + if (!host_data.empty()) { + copy_attempted = true; + enqueue_copy_to_host_in_current_context(host_data.data(), + host_data.size()); + } + CHECK_CUDA(synchronize_copy_stream()); + } catch (...) { + if (copy_attempted && synchronize_copy_stream() != CUDA_SUCCESS) { + // Retain storage if CUDA cannot confirm completion of an asynchronous + // copy. Otherwise, unwinding could expose freed memory to the driver. + auto* retained = new (std::nothrow) + std::vector{std::move(host_data)}; + if (!retained) + std::terminate(); + static_cast(retained); + } + throw; + } + return host_data; } //copies buffer back to dst buffer supplied by the user //count is number of elements the buffer has void copy_to_host(T* dst, size_t count) const { - if (access_ == IN) - throw std::runtime_error("Cannot copy a read only buffer back to host"); + validate_copy_destination(dst, count, true); + if (count == 0) + return; + if (is_scalar_) { + dst[0] = host_scalar_; + return; + } - if (is_scalar_) { - dst[0] = host_scalar_; - return; - } + require_live_context(); + detail::context_guard context_scope{ctx}; + enqueue_copy_to_host_in_current_context(dst, count); + CHECK_CUDA(synchronize_copy_stream()); + } - size_t bytes = count * sizeof(T); + // manual binding of the resource object that denotes memory + // to the pointer that contains it + void bind_token(caf::cuda::mem_token m_token) { + token = std::move(m_token); + } - CHECK_CUDA(cuCtxPushCurrent(ctx)); - CUstream s = stream_ ? stream_ : nullptr; - CHECK_CUDA(cuMemcpyDtoHAsync(dst, memory_, bytes, s)); +private: + friend class device; + friend class scheduler; - if (s) - CHECK_CUDA(cuStreamSynchronize(s)); - else - CHECK_CUDA(cuCtxSynchronize()); + template + friend class copy_back_command; - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + void bind_context_lifetime( + const detail::context_lifetime_ptr& lifetime) noexcept { + context_lifetime_ = lifetime; } - //reference counting for auto garabage collection - friend void intrusive_ptr_add_ref(const mem_ref* p) noexcept { - p->ref_count_.fetch_add(1, std::memory_order_relaxed); - } + bool belongs_to_context( + const detail::context_lifetime_ptr& lifetime, + CUcontext fallback_context) const noexcept { + if (is_scalar_) + return true; + if (context_lifetime_) + return context_lifetime_.get() == lifetime.get(); + return ctx == fallback_context; + } - friend void intrusive_ptr_release(const mem_ref* p) noexcept { - if (p->ref_count_.fetch_sub(1, std::memory_order_acq_rel) == 1) - delete p; - } + bool context_is_live() const noexcept { + return !context_lifetime_ || context_lifetime_->active(); + } + void require_live_context() const { + if (!ctx || !context_is_live()) + throw std::runtime_error( + "CUDA memory belongs to a stale device context"); + } - // manual binding of the resource object that denotes memory - // to the pointer that contains it - void bind_token(caf::cuda::mem_token m_token) { - token = std::move(m_token); + void validate_copy_destination(const T* destination, size_t count, + bool require_destination) const { + if (access_ != IN && access_ != OUT && access_ != IN_OUT) + throw std::runtime_error("Cannot copy from an invalid CUDA memory handle"); + if (access_ == IN) + throw std::runtime_error("Cannot copy a read-only buffer back to host"); + if (count > num_elements_) + throw std::out_of_range("copy count exceeds CUDA buffer size"); + if (require_destination && count > 0 && !destination) + throw std::invalid_argument("copy destination must not be null"); + } + + void enqueue_copy_to_host_in_current_context(T* destination, + size_t count) const { + if (is_scalar_) { + destination[0] = host_scalar_; + return; } + if (count == 0) + return; + CHECK_CUDA(cuMemcpyDtoHAsync(destination, memory_, count * sizeof(T), + stream_)); + } + CUresult synchronize_copy_stream() const noexcept { + if (stream_) + return cuStreamSynchronize(stream_); + return cuCtxSynchronize(); + } + + bool requires_stream_transition(CUstream target) const noexcept { + return !is_scalar_ && memory_ && stream_ != target; + } + + /// Orders the next use on `target` after all work submitted to the previous + /// stream. Callers must serialize submissions that share this memory handle. + void transition_to_stream_in_current_context(CUstream target) { + if (!requires_stream_transition(target)) + return; + detail::order_stream_after(target, stream_); + stream_ = target; + } + + void complete_stream_transition(CUstream target) noexcept { + if (requires_stream_transition(target)) + stream_ = target; + } + + void initialize_reference_count() noexcept { + // Preserve the documented intrusive_ptr(new mem_ref<...>) construction + // convention while storing the count in CAF's ref_counted base. + // As before, adopting mem_ref with add_ref=false is unsupported. + rc_.store(0, std::memory_order_relaxed); + } + + void invalidate_after_move() noexcept { + num_elements_ = 0; + memory_ = 0; + access_ = -1; + stream_ = nullptr; + ctx = nullptr; + context_lifetime_.reset(); + is_scalar_ = false; + } + + static void report_cuda_error(CUresult result, + const char* operation) noexcept { + if (result == CUDA_SUCCESS) + return; + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "%s failed: %s\n", operation, + description ? description : "unknown CUDA error"); + } -private: size_t num_elements_{0}; CUdeviceptr memory_{0}; int access_{-1}; @@ -211,7 +358,7 @@ class mem_ref : public caf::ref_counted { int context_id{0}; CUstream stream_{nullptr}; CUcontext ctx; - mutable std::atomic ref_count_{0}; + detail::context_lifetime_ptr context_lifetime_; mem_token token; diff --git a/libcaf_cuda/caf/cuda/memory_command.hpp b/libcaf_cuda/caf/cuda/memory_command.hpp index 1822ef79dc..4faefa5075 100644 --- a/libcaf_cuda/caf/cuda/memory_command.hpp +++ b/libcaf_cuda/caf/cuda/memory_command.hpp @@ -3,12 +3,17 @@ #include #include #include +#include +#include +#include +#include #include #include #include "caf/cuda/platform.hpp" #include "caf/cuda/device.hpp" +#include "caf/cuda/detail/context_guard.hpp" #include "caf/cuda/mem_ref.hpp" #include "caf/cuda/types.hpp" @@ -68,9 +73,7 @@ class bulk_memory_command : public caf::ref_counted { // Asynchronous execution result_type enqueue() { CUstream stream = dev_->get_stream_for_actor(stream_id_); - return std::apply([&](auto&&... arg) { - return std::make_tuple(dev_->make_arg(arg, stream)...); - }, args_); + return dev_->make_args_mem_ref_impl(args_, stream); } private: @@ -108,22 +111,24 @@ class copy_back_command : public caf::ref_counted { void enqueue(T* dst, size_t count) { if (ptr_->access() == IN) throw std::runtime_error("Cannot copy a read-only buffer back to host"); + validate_destination(dst, count); - CHECK_CUDA(cuCtxPushCurrent(ptr_->get_ctx())); - CUstream s = resolve_stream(); - + if (count == 0) + return; if (ptr_->is_scalar()) { // For scalars, the value is already on the host. // This assignment happens immediately on the CPU and is NOT stream-ordered. // If stream ordering is required for scalars, use run_async with a callback. dst[0] = *ptr_->host_scalar_ptr(); - } else { - size_t bytes = count * sizeof(T); - // Pure asynchronous copy with no host-side tracking. - CHECK_CUDA(cuMemcpyDtoHAsync(dst, ptr_->mem(), bytes, s)); + return; } - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + detail::context_guard context_scope{ptr_->get_ctx()}; + CUstream s = resolve_stream(); + ptr_->transition_to_stream_in_current_context(s); + size_t bytes = count * sizeof(T); + // Pure asynchronous copy with no host-side tracking. + CHECK_CUDA(cuMemcpyDtoHAsync(dst, ptr_->mem(), bytes, s)); } // Asynchronous execution with internal buffer allocation @@ -139,28 +144,42 @@ class copy_back_command : public caf::ref_counted { T host_scalar; }; - auto* state = new State{std::vector(ptr_->size()), std::move(callback), - ptr_->is_scalar(), *ptr_->host_scalar_ptr()}; + auto state = std::unique_ptr{ + new State{std::vector(ptr_->size()), std::move(callback), + ptr_->is_scalar(), *ptr_->host_scalar_ptr()}}; - CHECK_CUDA(cuCtxPushCurrent(ptr_->get_ctx())); + detail::context_guard context_scope{ptr_->get_ctx()}; CUstream s = resolve_stream(); - - if (!ptr_->is_scalar()) { - size_t bytes = ptr_->size() * sizeof(T); - CHECK_CUDA(cuMemcpyDtoHAsync(state->buffer.data(), ptr_->mem(), bytes, s)); - } - - auto host_fn = [](void* userData) { - auto* s_ptr = static_cast(userData); - if (s_ptr->is_scalar) - s_ptr->buffer[0] = s_ptr->host_scalar; - - s_ptr->user_callback(std::move(s_ptr->buffer)); - delete s_ptr; + ptr_->transition_to_stream_in_current_context(s); + + auto host_fn = [](void* user_data) noexcept { + auto callback_state = std::unique_ptr{ + static_cast(user_data)}; + try { + if (callback_state->is_scalar) + callback_state->buffer[0] = callback_state->host_scalar; + callback_state->user_callback(std::move(callback_state->buffer)); + } catch (const std::exception& error) { + std::fprintf(stderr, "CUDA host callback failed: %s\n", error.what()); + } catch (...) { + std::fprintf(stderr, "CUDA host callback failed: unknown error\n"); + } }; - CHECK_CUDA(cuLaunchHostFunc(s, host_fn, state)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + bool copy_attempted = false; + try { + if (!ptr_->is_scalar() && !state->buffer.empty()) { + copy_attempted = true; + size_t bytes = ptr_->size() * sizeof(T); + CHECK_CUDA(cuMemcpyDtoHAsync(state->buffer.data(), ptr_->mem(), + bytes, s)); + } + CHECK_CUDA(cuLaunchHostFunc(s, host_fn, state.get())); + state.release(); + } catch (...) { + synchronize_failed_submission(state, s, copy_attempted, true); + throw; + } } // Asynchronous execution with user-provided buffer @@ -168,6 +187,7 @@ class copy_back_command : public caf::ref_counted { void run_async(T* dst, size_t count, F callback) { if (ptr_->access() == IN) throw std::runtime_error("Cannot copy a read-only buffer back to host"); + validate_destination(dst, count); struct State { T* dst; @@ -177,36 +197,78 @@ class copy_back_command : public caf::ref_counted { T host_scalar; }; - auto* state = new State{dst, count, std::move(callback), - ptr_->is_scalar(), *ptr_->host_scalar_ptr()}; + auto state = std::unique_ptr{ + new State{dst, count, std::move(callback), ptr_->is_scalar(), + *ptr_->host_scalar_ptr()}}; - CHECK_CUDA(cuCtxPushCurrent(ptr_->get_ctx())); + detail::context_guard context_scope{ptr_->get_ctx()}; CUstream s = resolve_stream(); + ptr_->transition_to_stream_in_current_context(s); + + auto host_fn = [](void* user_data) noexcept { + auto callback_state = std::unique_ptr{ + static_cast(user_data)}; + try { + if (callback_state->is_scalar && callback_state->count > 0) + callback_state->dst[0] = callback_state->host_scalar; + callback_state->user_callback(callback_state->dst, + callback_state->count); + } catch (const std::exception& error) { + std::fprintf(stderr, "CUDA host callback failed: %s\n", error.what()); + } catch (...) { + std::fprintf(stderr, "CUDA host callback failed: unknown error\n"); + } + }; - if (!ptr_->is_scalar()) { - size_t bytes = count * sizeof(T); - CHECK_CUDA(cuMemcpyDtoHAsync(dst, ptr_->mem(), bytes, s)); + bool copy_attempted = false; + try { + if (!ptr_->is_scalar() && count > 0) { + copy_attempted = true; + size_t bytes = count * sizeof(T); + CHECK_CUDA(cuMemcpyDtoHAsync(dst, ptr_->mem(), bytes, s)); + } + CHECK_CUDA(cuLaunchHostFunc(s, host_fn, state.get())); + state.release(); + } catch (...) { + synchronize_failed_submission(state, s, copy_attempted, false); + throw; } + } - auto host_fn = [](void* userData) { - auto* s_ptr = static_cast(userData); - if (s_ptr->is_scalar) - s_ptr->dst[0] = s_ptr->host_scalar; - - s_ptr->user_callback(s_ptr->dst, s_ptr->count); - delete s_ptr; - }; +private: + void validate_destination(const T* destination, size_t count) const { + if (count > ptr_->size()) + throw std::out_of_range("copy count exceeds CUDA buffer size"); + if (count > 0 && !destination) + throw std::invalid_argument("copy destination must not be null"); + } - CHECK_CUDA(cuLaunchHostFunc(s, host_fn, state)); - CHECK_CUDA(cuCtxPopCurrent(nullptr)); + template + static void synchronize_failed_submission(std::unique_ptr& state, + CUstream stream, + bool copy_attempted, + bool state_owns_destination) noexcept { + if (!copy_attempted) + return; + auto result = cuStreamSynchronize(stream); + if (result == CUDA_SUCCESS) + return; + const char* description = nullptr; + cuGetErrorString(result, &description); + std::fprintf(stderr, "CUDA callback cleanup synchronization failed: %s\n", + description ? description : "unknown CUDA error"); + if (state_owns_destination) + state.release(); } -private: CUstream resolve_stream() { if (stream_id_ == -1) return ptr_->stream(); auto plat = platform::create(); auto dev = plat->schedule(stream_id_, ptr_->deviceID()); + if (!ptr_->is_scalar() && dev->getContext() != ptr_->get_ctx()) + throw std::runtime_error( + "CUDA copy stream belongs to a different or stale context"); return dev->get_stream_for_actor(stream_id_); } @@ -237,6 +299,9 @@ class free_memory_command : public caf::ref_counted { return ptr_->stream(); auto plat = platform::create(); auto dev = plat->schedule(stream_id_, ptr_->deviceID()); + if (!ptr_->is_scalar() && dev->getContext() != ptr_->get_ctx()) + throw std::runtime_error( + "CUDA free stream belongs to a different or stale context"); return dev->get_stream_for_actor(stream_id_); } diff --git a/libcaf_cuda/caf/cuda/nd_range.hpp b/libcaf_cuda/caf/cuda/nd_range.hpp index 0153d96477..67ecb6c75b 100644 --- a/libcaf_cuda/caf/cuda/nd_range.hpp +++ b/libcaf_cuda/caf/cuda/nd_range.hpp @@ -1,115 +1,145 @@ #pragma once -#include +#include #include +#include +#include +#include #include -#include #include -#include +#include namespace caf::cuda { using dim_vec = std::vector; -//class that represents the grid and block dimensions of a kernel +/// Represents the grid and block dimensions of a CUDA kernel. class nd_range { public: - // Constructor with individual dimension arguments - nd_range(int gridX, int gridY, int gridZ, int blockX, int blockY, int blockZ) - : gridDim{static_cast(gridX), + nd_range(int gridX, int gridY, int gridZ, + int blockX, int blockY, int blockZ) + : grid_dim_{static_cast(gridX), static_cast(gridY), static_cast(gridZ)}, - blockDim{static_cast(blockX), + block_dim_{static_cast(blockX), static_cast(blockY), static_cast(blockZ)} { - - computeHash(); - - } - + compute_hash(); + } - // Constructor from vectors nd_range(const dim_vec& grid, const dim_vec& block) { - if (grid.size() != 3 || block.size() != 3) { - throw std::invalid_argument("Grid and block dimensions must each be of size 3."); - } - gridDim = grid; - blockDim = block; + if (grid.size() != 3 || block.size() != 3) + throw std::invalid_argument( + "Grid and block dimensions must each be of size 3."); + grid_dim_ = {grid[0], grid[1], grid[2]}; + block_dim_ = {block[0], block[1], block[2]}; + compute_hash(); } - - //default constructor - //TODO fix issue where actors store a non in use ndrange + /// A default range has zero-valued dimensions. nd_range() = default; + nd_range(const nd_range& other) noexcept + : grid_dim_(other.grid_dim_), + block_dim_(other.block_dim_), + hash_value_(other.hash_value_) { + } - // Getters for grid dimensions - size_t getGridDimX() const { return gridDim[0]; } - size_t getGridDimY() const { return gridDim[1]; } - size_t getGridDimZ() const { return gridDim[2]; } + // Moving six scalar values is cheaper than transferring the compatibility + // caches, and leaves the source with the same value after the copy. + nd_range(nd_range&& other) noexcept + : nd_range(static_cast(other)) { + } - // Getters for block dimensions - size_t getBlockDimX() const { return blockDim[0]; } - size_t getBlockDimY() const { return blockDim[1]; } - size_t getBlockDimZ() const { return blockDim[2]; } + nd_range& operator=(const nd_range& other) { + if (this != &other) { + grid_dim_ = other.grid_dim_; + block_dim_ = other.block_dim_; + hash_value_ = other.hash_value_; + refresh_materialized_vectors(); + } + return *this; + } - // Optional: Getters for full vectors - const dim_vec& getGridDims() const { return gridDim; } - const dim_vec& getBlockDims() const { return blockDim; } + nd_range& operator=(nd_range&& other) { + return *this = static_cast(other); + } - // Returns total number of threads per block - [[nodiscard]] size_t get_num_threads() const noexcept { - return blockDim[0] * blockDim[1] * blockDim[2]; + size_t getGridDimX() const { return grid_dim_[0]; } + size_t getGridDimY() const { return grid_dim_[1]; } + size_t getGridDimZ() const { return grid_dim_[2]; } + + size_t getBlockDimX() const { return block_dim_[0]; } + size_t getBlockDimY() const { return block_dim_[1]; } + size_t getBlockDimZ() const { return block_dim_[2]; } + + // Preserve the vector-reference API without allocating in the launch path. + // These vectors are materialized only if a caller explicitly requests them. + const dim_vec& getGridDims() const { + std::call_once(grid_vector_once_, [this] { + grid_vector_.assign(grid_dim_.begin(), grid_dim_.end()); + grid_vector_initialized_ = true; + }); + return grid_vector_; } - //get number of blocks in total - size_t get_num_blocks() const noexcept { - return gridDim[0] * gridDim[1] * gridDim[2]; - } - // Returns the precomputed hash - [[nodiscard]] size_t getHash() const noexcept { - return hashValue_; + const dim_vec& getBlockDims() const { + std::call_once(block_vector_once_, [this] { + block_vector_.assign(block_dim_.begin(), block_dim_.end()); + block_vector_initialized_ = true; + }); + return block_vector_; } - ~nd_range() { - //no-op + [[nodiscard]] size_t get_num_threads() const noexcept { + return block_dim_[0] * block_dim_[1] * block_dim_[2]; } + size_t get_num_blocks() const noexcept { + return grid_dim_[0] * grid_dim_[1] * grid_dim_[2]; + } - // Returns a stable string representation of grid + block dims + [[nodiscard]] size_t getHash() const noexcept { + return hash_value_; + } [[nodiscard]] std::string to_string() const { - std::ostringstream oss; - oss << "grid(" - << gridDim[0] << "," - << gridDim[1] << "," - << gridDim[2] << ")" - << "_block(" - << blockDim[0] << "," - << blockDim[1] << "," - << blockDim[2] << ")"; - return oss.str(); + std::ostringstream oss; + oss << "grid(" + << grid_dim_[0] << "," + << grid_dim_[1] << "," + << grid_dim_[2] << ")" + << "_block(" + << block_dim_[0] << "," + << block_dim_[1] << "," + << block_dim_[2] << ")"; + return oss.str(); } - - private: - // Dimensions are stored in order of x, y, z - dim_vec gridDim{3,0}; - dim_vec blockDim{3,0}; - size_t hashValue_{0}; // store precomputed hash - + using dimensions = std::array; - // Precompute hash from to_string - void computeHash() { - std::hash hasher; - hashValue_ = hasher(to_string()); + void compute_hash() { + hash_value_ = std::hash{}(to_string()); } + void refresh_materialized_vectors() { + if (grid_vector_initialized_) + grid_vector_.assign(grid_dim_.begin(), grid_dim_.end()); + if (block_vector_initialized_) + block_vector_.assign(block_dim_.begin(), block_dim_.end()); + } + dimensions grid_dim_{}; + dimensions block_dim_{}; + size_t hash_value_ = 0; - + mutable dim_vec grid_vector_; + mutable dim_vec block_vector_; + mutable std::once_flag grid_vector_once_; + mutable std::once_flag block_vector_once_; + mutable bool grid_vector_initialized_ = false; + mutable bool block_vector_initialized_ = false; }; } // namespace caf::cuda - diff --git a/libcaf_cuda/caf/cuda/platform.hpp b/libcaf_cuda/caf/cuda/platform.hpp index a29c0ecf06..ffe77fe493 100644 --- a/libcaf_cuda/caf/cuda/platform.hpp +++ b/libcaf_cuda/caf/cuda/platform.hpp @@ -22,6 +22,7 @@ namespace caf::cuda { class CAF_CUDA_EXPORT platform : public ref_counted { public: friend class program; + friend class manager; template friend intrusive_ptr caf::make_counted(Ts&&...); @@ -49,6 +50,8 @@ class CAF_CUDA_EXPORT platform : public ref_counted { // Resets the CUDA context for a specific device. // This will destroy the existing context for the device, create a new one, // and update the device object and platform's internal records. + // The caller must first quiesce every context-bound operation on the device, + // including callbacks, program/memory/event lifecycle, and raw handle use. void reset_device_context(int device_id) { if (device_id < 0 || device_id >= static_cast(devices_.size())) { throw std::out_of_range("Invalid device_id for reset_device_context"); @@ -73,8 +76,14 @@ class CAF_CUDA_EXPORT platform : public ref_counted { cu_dev)); #endif - // Update the device object and platform's internal contexts_ vector - dev_obj->reset_context(new_ctx); + // device::reset_context adopts new_ctx only after all replacement + // resources have been constructed. Destroy it here if preparation fails. + try { + dev_obj->reset_context(new_ctx); + } catch (...) { + static_cast(cuCtxDestroy(new_ctx)); + throw; + } contexts_[device_id] = new_ctx; } @@ -83,6 +92,9 @@ class CAF_CUDA_EXPORT platform : public ref_counted { private: + /// Releases the singleton reference during manager shutdown. + static void shutdown(); + platform(); ~platform(); diff --git a/libcaf_cuda/caf/cuda/program.hpp b/libcaf_cuda/caf/cuda/program.hpp index 3ab35e1e3b..277110f2a7 100644 --- a/libcaf_cuda/caf/cuda/program.hpp +++ b/libcaf_cuda/caf/cuda/program.hpp @@ -1,13 +1,15 @@ #pragma once +#include #include #include +#include +#include #include -#include -#include #include #include "caf/cuda/global.hpp" +#include "caf/cuda/detail/context_lifetime.hpp" #include "caf/cuda/platform.hpp" #include "caf/cuda/device.hpp" @@ -23,36 +25,54 @@ class CAF_CUDA_EXPORT program : public caf::ref_counted { /// @param is_fatbin Whether the binary is a fatbinary (default: false). program(std::string name, std::vector binary, bool is_fatbin = false); + ~program() noexcept override; + + program(const program& other); + program& operator=(const program& other); + /// Returns the CUfunction for a given device. /// @throws std::runtime_error if the kernel was not loaded for the device. CUfunction get_kernel(int device_id); - friend void intrusive_ptr_add_ref([[maybe_unused]] const program* p) noexcept { - //p->ref_count_.fetch_add(1, std::memory_order_relaxed); - } - friend void intrusive_ptr_release(const program* p) noexcept { - if (p->ref_count_.fetch_sub(1, std::memory_order_acq_rel) == 1) { - //WARNING TURNING THIS ON FOR SOME REASON, CAUSES SEGFAUTLS - //I HAVE NO IDEA WHY - //TODO FIX THIS - // std::cout<< "Deleting\n"; - // delete p; - } - } - std::string getName() {return name_;} int getHash() const {return hashValue;} private: + template + friend class command_runner; + + struct module_entry { + device* owner = nullptr; + detail::context_lifetime_ptr lifetime; + CUmodule module = nullptr; + CUfunction function = nullptr; + std::atomic published_lifetime{nullptr}; + std::mutex reload_mutex; + }; + /// Internal helper to load the kernel modules on all devices. - void load_kernels(bool is_fatbin); + void load_kernels(); + + /// Loads one module and publishes it for lock-free subsequent lookups. + void load_kernel(module_entry& entry, + const detail::context_lifetime_ptr& lifetime); + + /// Unloads all successfully loaded modules without throwing. + void unload_modules() noexcept; + + /// Resolves a launch device through the platform retained by this program. + device_ptr device_for_launch(int actor_id, int device_number); + + /// Exchanges program state without touching either reference count. + void swap_state(program& other) noexcept; + platform_ptr platform_; ///< Keeps device contexts alive std::string name_; ///< Name of the kernel std::vector binary_; ///< The binary or PTX of the program - std::unordered_map kernels_; ///< Device ID -> CUfunction mapping - mutable std::atomic ref_count_{0}; + bool is_fatbin_; ///< Whether binary_ is a fatbinary + std::vector> modules_; std::hash hasher; int hashValue = 0; diff --git a/libcaf_cuda/caf/cuda/streampool.hpp b/libcaf_cuda/caf/cuda/streampool.hpp index 2841fa813a..ca66fad343 100644 --- a/libcaf_cuda/caf/cuda/streampool.hpp +++ b/libcaf_cuda/caf/cuda/streampool.hpp @@ -4,6 +4,8 @@ #include #include #include +#include +#include #include #include #include @@ -13,6 +15,8 @@ namespace caf::cuda { +class device; + /// Thread-safe pool of CUDA streams. /// /// Design goals: @@ -73,7 +77,7 @@ class CAF_CUDA_EXPORT StreamPool { mutable std::mutex pool_mutex_; ///< Protects the pool state }; -/// Pool of cuBLAS handles. Capped at 32. +/// Pool of cuBLAS handles. The default capacity is 32. class CAF_CUDA_EXPORT CublasHandlePool { public: explicit CublasHandlePool(CUcontext ctx, size_t max_size = 32); @@ -88,7 +92,7 @@ class CAF_CUDA_EXPORT CublasHandlePool { cublasHandle_t create_handle(); CUcontext ctx_; - std::deque available_handles_; + std::vector available_handles_; std::vector all_handles_; size_t max_size_; mutable std::mutex pool_mutex_; @@ -99,7 +103,9 @@ class CAF_CUDA_EXPORT DeviceCublasHandleTable { public: explicit DeviceCublasHandleTable(CUcontext ctx, size_t pool_size = 32); - /// Get the cuBLAS handle for an actor and bind it to a stream. + /// Get the raw cuBLAS handle for an actor and bind it to a stream. + /// Subsequent calls conservatively re-establish state because the caller may + /// reconfigure the returned handle. cublasHandle_t get_handle(int actor_id, CUstream stream); /// Release the handle assigned to an actor. @@ -108,12 +114,25 @@ class CAF_CUDA_EXPORT DeviceCublasHandleTable { size_t pool_size() const { return pool_.max_size(); } private: + friend class device; + + struct handle_state { + CUstream stream = nullptr; + cublasPointerMode_t pointer_mode = CUBLAS_POINTER_MODE_HOST; + bool externally_exposed = false; + }; + + cublasHandle_t get_handle_with_pointer_mode( + int actor_id, CUstream stream, cublasPointerMode_t pointer_mode, + bool expose_raw_handle = false); + CublasHandlePool pool_; std::unordered_map table_; ///< actor_id -> handle + std::unordered_map handle_states_; mutable std::shared_mutex table_mutex_; }; -/// Pool of cuSparse handles. Capped at 32. +/// Pool of cuSparse handles. The default capacity is 32. class CAF_CUDA_EXPORT CusparseHandlePool { public: explicit CusparseHandlePool(CUcontext ctx, size_t max_size = 32); @@ -128,7 +147,7 @@ class CAF_CUDA_EXPORT CusparseHandlePool { cusparseHandle_t create_handle(); CUcontext ctx_; - std::deque available_handles_; + std::vector available_handles_; std::vector all_handles_; size_t max_size_; mutable std::mutex pool_mutex_; @@ -139,7 +158,9 @@ class CAF_CUDA_EXPORT DeviceCusparseHandleTable { public: explicit DeviceCusparseHandleTable(CUcontext ctx, size_t pool_size = 32); - /// Get the cuSparse handle for an actor and bind it to a stream. + /// Get the raw cuSPARSE handle for an actor and bind it to a stream. + /// Subsequent calls conservatively re-establish state because the caller may + /// reconfigure the returned handle. cusparseHandle_t get_handle(int actor_id, CUstream stream); /// Release the handle assigned to an actor. @@ -148,16 +169,28 @@ class CAF_CUDA_EXPORT DeviceCusparseHandleTable { size_t pool_size() const { return pool_.max_size(); } private: + friend class device; + + struct handle_state { + CUstream stream = nullptr; + cusparsePointerMode_t pointer_mode = CUSPARSE_POINTER_MODE_HOST; + bool externally_exposed = false; + }; + + cusparseHandle_t get_handle_impl(int actor_id, CUstream stream, + bool expose_raw_handle); + CusparseHandlePool pool_; std::unordered_map table_; ///< actor_id -> handle + std::unordered_map handle_states_; mutable std::shared_mutex table_mutex_; }; /// Per-device stream manager. Assigns streams to actor IDs. /// /// `DeviceStreamTable` caches an assigned stream per `actor_id`. This makes the -/// common case (getting a stream for an actor who already has one) cheap: -/// a shared lock and hashmap lookup. +/// common case (repeating a lookup on the same thread) lock-free. Cache misses +/// use a shared lock and hash-map lookup. class CAF_CUDA_EXPORT DeviceStreamTable { public: explicit DeviceStreamTable(CUcontext ctx, size_t pool_size = 32); @@ -173,6 +206,8 @@ class CAF_CUDA_EXPORT DeviceStreamTable { size_t pool_size() const { return pool_.max_size(); } private: + const std::uint64_t table_id_; + std::atomic assignment_epoch_{0}; StreamPool pool_; std::unordered_map table_; ///< actor_id -> stream mutable std::shared_mutex table_mutex_; ///< Allows concurrent reads of table_ diff --git a/libcaf_cuda/sc26/Batched-Matrix-Multiply/CMakeLists.txt b/libcaf_cuda/sc26/Batched-Matrix-Multiply/CMakeLists.txt index f5e0fd114c..f163796f62 100644 --- a/libcaf_cuda/sc26/Batched-Matrix-Multiply/CMakeLists.txt +++ b/libcaf_cuda/sc26/Batched-Matrix-Multiply/CMakeLists.txt @@ -1,47 +1,17 @@ -cmake_minimum_required(VERSION 3.16.3) - -# 1) Enforce C++20 (needs modern C++ features) -set(CMAKE_CXX_STANDARD 20) -set(CMAKE_CXX_STANDARD_REQUIRED ON) -set(CMAKE_CXX_EXTENSIONS OFF) - -# 2) Set CAF source and build directories -set(CAF_SRC "${CMAKE_CURRENT_SOURCE_DIR}/../../../../actor-framework") -set(CAF_BUILD "${CAF_SRC}/build") - - -#required since cmake cant seem to find the compiler -set(CMAKE_CXX_COMPILER "/usr/bin/g++") -set(CMAKE_C_COMPILER "/usr/bin/gcc") - -project(CUDA_ACTORS) - -find_package(CUDA REQUIRED) -find_package(CUDAToolkit REQUIRED) - -include_directories( - "${CAF_SRC}/include" - "${CAF_SRC}/libcaf_io" - "${CAF_SRC}/libcaf_core" - "${CAF_SRC}/libcaf_opencl" - "${CAF_BUILD}/libcaf_core" # for generated build_config.hpp - "${CAF_SRC}/libcaf_cuda" -) - - -# 5) Declare your executable -add_executable(test main.test.cpp) -target_compile_definitions(test PRIVATE CAF_ENABLE_LOGGING) - -target_link_libraries(test - PRIVATE - "${CAF_BUILD}/libcaf_core/libcaf_core.so" - "${CAF_BUILD}/libcaf_io/libcaf_io.so" - "${CAF_BUILD}/libcaf_cuda/libcaf_cuda.so" - CUDA::nvrtc - CUDA::cuda_driver - CUDA::cusparse - CUDA::cublas -) - - +set(batched_cubin "${CMAKE_CURRENT_BINARY_DIR}/mmul.cubin") +sc26_add_cubin( + sc26_batched_mmul_cubin + "${batched_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/mmul.cu") + +add_executable(sc26_batched_matrix_multiply main.test.cpp) +target_compile_features(sc26_batched_matrix_multiply PRIVATE cxx_std_20) +target_compile_definitions(sc26_batched_matrix_multiply PRIVATE CAF_ENABLE_LOGGING) +target_link_libraries(sc26_batched_matrix_multiply PRIVATE CAF::cuda) +sc26_set_runtime_output(sc26_batched_matrix_multiply batched-matrix-multiply) +add_dependencies(sc26_batched_matrix_multiply sc26_batched_mmul_cubin) +sc26_copy_file_to_target(sc26_batched_matrix_multiply "${batched_cubin}") + +add_custom_target( + sc26_batched_mmul_benchmarks + DEPENDS sc26_batched_matrix_multiply) diff --git a/libcaf_cuda/sc26/CMakeLists.txt b/libcaf_cuda/sc26/CMakeLists.txt new file mode 100644 index 0000000000..c346463eba --- /dev/null +++ b/libcaf_cuda/sc26/CMakeLists.txt @@ -0,0 +1,67 @@ +find_package(CUDAToolkit REQUIRED) +enable_language(CUDA) + +function(sc26_set_runtime_output target suite) + set(output_dir "${CMAKE_BINARY_DIR}/bin/$/sc26/${suite}") + set_target_properties( + ${target} + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${output_dir}" + VS_DEBUGGER_WORKING_DIRECTORY "$") + + if(WIN32 AND BUILD_SHARED_LIBS) + get_target_property(target_links ${target} LINK_LIBRARIES) + if("CAF::cuda" IN_LIST target_links) + add_custom_command( + TARGET ${target} POST_BUILD + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "$" + "$" + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "$" + "$" + VERBATIM) + endif() + endif() +endfunction() + +function(sc26_add_cubin target output source) + add_custom_command( + OUTPUT "${output}" + COMMAND "${CUDAToolkit_NVCC_EXECUTABLE}" + --cubin + "-arch=${CAF_CUDA_ARCH}" + --generate-line-info + -o "${output}" + "${source}" + DEPENDS "${source}" + COMMENT "Compiling ${target}" + VERBATIM) + add_custom_target(${target} DEPENDS "${output}") +endfunction() + +function(sc26_copy_file_to_target target source) + add_custom_command( + TARGET ${target} POST_BUILD + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "${source}" + "$" + VERBATIM) +endfunction() + +add_subdirectory(Sequence-Independent-Tasks) +add_subdirectory(Runtime-Overhead) +add_subdirectory(Batched-Matrix-Multiply) +add_subdirectory(Fault-Tolerance) +add_subdirectory(Irregular-Workload/workloadA) +add_subdirectory(Irregular-Workload/workloadB) + +add_custom_target( + sc26_benchmarks + DEPENDS + sc26_sequence_benchmarks + sc26_runtime_overhead_benchmarks + sc26_batched_mmul_benchmarks + sc26_fault_tolerance_benchmarks + sc26_irregular_workload_a_benchmarks + sc26_irregular_workload_b_benchmarks) diff --git a/libcaf_cuda/sc26/Fault-Tolerance/CMakeLists.txt b/libcaf_cuda/sc26/Fault-Tolerance/CMakeLists.txt new file mode 100644 index 0000000000..e126af80c6 --- /dev/null +++ b/libcaf_cuda/sc26/Fault-Tolerance/CMakeLists.txt @@ -0,0 +1,16 @@ +set(fault_cubin "${CMAKE_CURRENT_BINARY_DIR}/monte_carlo.cubin") +sc26_add_cubin( + sc26_fault_monte_carlo_cubin + "${fault_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/monte_carlo.cu") + +add_executable(sc26_fault_tolerance main.cpp) +target_compile_features(sc26_fault_tolerance PRIVATE cxx_std_20) +target_link_libraries(sc26_fault_tolerance PRIVATE CAF::cuda) +sc26_set_runtime_output(sc26_fault_tolerance fault-tolerance) +add_dependencies(sc26_fault_tolerance sc26_fault_monte_carlo_cubin) +sc26_copy_file_to_target(sc26_fault_tolerance "${fault_cubin}") + +add_custom_target( + sc26_fault_tolerance_benchmarks + DEPENDS sc26_fault_tolerance) diff --git a/libcaf_cuda/sc26/Fault-Tolerance/main.cpp b/libcaf_cuda/sc26/Fault-Tolerance/main.cpp index 5594a65224..cc0ebececa 100644 --- a/libcaf_cuda/sc26/Fault-Tolerance/main.cpp +++ b/libcaf_cuda/sc26/Fault-Tolerance/main.cpp @@ -394,7 +394,8 @@ class Supervisor { double elapsed_s = ms_between(t_start_, t_end) / 1000.0; long long total_s = (long long)batches_done_ * samples_per_batch_; double pi_final = 4.0 * total_hits_ / (double)total_s; - double error_pct = std::abs(pi_final - M_PI) / M_PI * 100.0; + constexpr double pi = 3.14159265358979323846; + double error_pct = std::abs(pi_final - pi) / pi * 100.0; double throughput = batches_done_ / elapsed_s; self_->println("\n=== RESULT ==="); diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadA/CMakeLists.txt b/libcaf_cuda/sc26/Irregular-Workload/workloadA/CMakeLists.txt index 199d000211..db4d4de8a5 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadA/CMakeLists.txt +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadA/CMakeLists.txt @@ -1,132 +1,79 @@ -cmake_minimum_required(VERSION 3.16.3) - -# 1) Enforce C++20 (needs modern C++ features) -set(CMAKE_CXX_STANDARD 20) -set(CMAKE_CXX_STANDARD_REQUIRED ON) -set(CMAKE_CXX_EXTENSIONS OFF) - -# 2) Set CAF source and build directories -set(CAF_SRC "${CMAKE_CURRENT_SOURCE_DIR}/../../../../../actor-framework") -set(CAF_BUILD "${CAF_SRC}/build") - - -#required since cmake cant seem to find the compiler -set(CMAKE_CXX_COMPILER "/usr/bin/g++") -set(CMAKE_C_COMPILER "/usr/bin/gcc") - -# Enable CUDA as a first-class language for the project -project(CUDA_ACTORS LANGUAGES CXX CUDA) - -find_package(CUDA REQUIRED) -find_package(CUDAToolkit REQUIRED) - -include_directories( - "${CAF_SRC}/include" - "${CAF_SRC}/libcaf_io" - "${CAF_SRC}/libcaf_core" - "${CAF_SRC}/libcaf_opencl" - "${CAF_BUILD}/libcaf_core" # for generated build_config.hpp - "${CAF_SRC}/libcaf_cuda" -) - -# --- CUDA Kernel Compilation --- -set(STABILITY_KERNEL_SRC "${CMAKE_CURRENT_SOURCE_DIR}/stability_kernels.cu") -set(STABILITY_KERNEL_CUBIN "${CMAKE_CURRENT_BINARY_DIR}/stability_kernels.cubin") - -# Target CUDA architecture. 'native' targets the current machine's GPU (requires CUDA 11.6+). -# You can override this with -DCUDA_ARCH=sm_XX if needed. -set(CUDA_ARCH "native" CACHE STRING "Target CUDA architecture (e.g., native, sm_70, sm_75, sm_80, sm_86)") - -add_custom_command( - OUTPUT ${STABILITY_KERNEL_CUBIN} - COMMAND ${CUDAToolkit_NVCC_EXECUTABLE} - -cubin - -arch=${CUDA_ARCH} - -o ${STABILITY_KERNEL_CUBIN} - ${STABILITY_KERNEL_SRC} - DEPENDS ${STABILITY_KERNEL_SRC} - COMMENT "Compiling CUDA kernel ${STABILITY_KERNEL_SRC} for architecture: ${CUDA_ARCH}" - VERBATIM -) - -set(JACOBI_KERNEL_SRC "${CMAKE_CURRENT_SOURCE_DIR}/jacobi_kernels.cu") -set(JACOBI_KERNEL_CUBIN "${CMAKE_CURRENT_BINARY_DIR}/jacobi_kernels.cubin") - -add_custom_command( - OUTPUT ${JACOBI_KERNEL_CUBIN} - COMMAND ${CUDAToolkit_NVCC_EXECUTABLE} - -cubin - -arch=${CUDA_ARCH} - -o ${JACOBI_KERNEL_CUBIN} - ${JACOBI_KERNEL_SRC} - DEPENDS ${JACOBI_KERNEL_SRC} - COMMENT "Compiling CUDA kernel ${JACOBI_KERNEL_SRC} for architecture: ${CUDA_ARCH}" - VERBATIM -) - -add_custom_target(stability_kernels_cubin ALL DEPENDS ${STABILITY_KERNEL_CUBIN}) -add_custom_target(jacobi_kernels_cubin ALL DEPENDS ${JACOBI_KERNEL_CUBIN}) - - -# 5) Declare your executable and its source files -add_executable(test main.test.cpp sparse_utils.cpp) -target_sources(test PRIVATE +set(workload_cubin_dir "${CMAKE_CURRENT_BINARY_DIR}") +set(stability_cubin "${workload_cubin_dir}/stability_kernels.cubin") +set(jacobi_cubin "${workload_cubin_dir}/jacobi_kernels.cubin") + +sc26_add_cubin( + sc26_irregular_a_stability_cubin + "${stability_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/stability_kernels.cu") +sc26_add_cubin( + sc26_irregular_a_jacobi_cubin + "${jacobi_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/jacobi_kernels.cu") + +set(matrix_dir + "${CMAKE_CURRENT_SOURCE_DIR}/../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices") +file(TO_CMAKE_PATH "${matrix_dir}" matrix_dir) + +add_executable( + sc26_irregular_a_actor + main.test.cpp + sparse_utils.cpp atoms.hpp ft_cg_actor.hpp ft_cg_jacobi_actor.hpp - supervisor_actor.hpp -) - -target_compile_definitions(test PRIVATE CAF_ENABLE_LOGGING) - -target_link_libraries(test + supervisor_actor.hpp) +target_compile_features(sc26_irregular_a_actor PRIVATE cxx_std_20) +target_compile_definitions( + sc26_irregular_a_actor PRIVATE - "${CAF_BUILD}/libcaf_core/libcaf_core.so" - "${CAF_BUILD}/libcaf_io/libcaf_io.so" - "${CAF_BUILD}/libcaf_cuda/libcaf_cuda.so" - CUDA::nvrtc - CUDA::cublas - CUDA::cusparse -) - - -# # 5) Declare your executable -# add_executable(hot-potatoe hot-potatoe.cpp sparse_utils.cpp) - -# target_compile_definitions(hot-potatoe PRIVATE CAF_ENABLE_LOGGING) - -# target_link_libraries(hot-potatoe -# PRIVATE -# "${CAF_BUILD}/libcaf_core/libcaf_core.so" -# "${CAF_BUILD}/libcaf_io/libcaf_io.so" -# "${CAF_BUILD}/libcaf_cuda/libcaf_cuda.so" -# CUDA::nvrtc -# CUDA::cublas -# CUDA::cusparse -# ) + CAF_ENABLE_LOGGING + SC26_MATRIX_DIR="${matrix_dir}") +target_link_libraries( + sc26_irregular_a_actor + PRIVATE CAF::cuda CUDA::cublas CUDA::cusparse) +sc26_set_runtime_output(sc26_irregular_a_actor irregular-workload-a) +add_dependencies( + sc26_irregular_a_actor + sc26_irregular_a_stability_cubin + sc26_irregular_a_jacobi_cubin) +sc26_copy_file_to_target(sc26_irregular_a_actor "${stability_cubin}") +sc26_copy_file_to_target(sc26_irregular_a_actor "${jacobi_cubin}") - -# FindThreads is required for std::thread in the native version find_package(Threads REQUIRED) - -# 6) Declare the native benchmark executable (raw CUDA/cuBLAS/cuSPARSE) -add_executable(workload-native main.native.cpp sparse_utils.cpp native_utils.cu) - -target_link_libraries(workload-native - PRIVATE - CUDA::cudart - CUDA::cublas - CUDA::cusparse - Threads::Threads -) - -# 7) Declare the native sorted benchmark executable -add_executable(workload-native-sorted main.native_sorted.cpp sparse_utils.cpp native_utils.cu) - -target_link_libraries(workload-native-sorted - PRIVATE - CUDA::cudart - CUDA::cublas - CUDA::cusparse - Threads::Threads -) +set(sc26_cuda_architecture "${CAF_CUDA_ARCH}") +string(REGEX REPLACE "^sm_" "" sc26_cuda_architecture + "${sc26_cuda_architecture}") + +foreach(variant native native_sorted) + set(target "sc26_irregular_a_${variant}") + if(variant STREQUAL "native") + set(main_source main.native.cpp) + else() + set(main_source main.native_sorted.cpp) + endif() + add_executable( + ${target} + "${main_source}" + sparse_utils.cpp + native_utils.cu) + target_compile_features(${target} PRIVATE cxx_std_20) + set_target_properties( + ${target} + PROPERTIES + CUDA_STANDARD 20 + CUDA_STANDARD_REQUIRED ON + CUDA_ARCHITECTURES "${sc26_cuda_architecture}") + target_compile_definitions(${target} PRIVATE SC26_MATRIX_DIR="${matrix_dir}") + target_link_libraries( + ${target} + PRIVATE CUDA::cudart CUDA::cublas CUDA::cusparse Threads::Threads) + sc26_set_runtime_output(${target} irregular-workload-a) +endforeach() + +add_custom_target( + sc26_irregular_workload_a_benchmarks + DEPENDS + sc26_irregular_a_actor + sc26_irregular_a_native + sc26_irregular_a_native_sorted) diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadA/ft_cg_actor.hpp b/libcaf_cuda/sc26/Irregular-Workload/workloadA/ft_cg_actor.hpp index 60f3ebd23a..7f3615ee4b 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadA/ft_cg_actor.hpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadA/ft_cg_actor.hpp @@ -91,7 +91,7 @@ behavior fault_tolerant_cg_actor(stateful_actor>* self, st.d_err = runner.transfer_memory(st.device_id, st.stream_id, out(1)); size_t ws_sz = st.d_ptr->spmv_csr_buffer_size(st.stream_id, st.n, st.n, st.nnz, st.A_rp, st.A_ci, st.A_val, st.x, st.w); - if (ws_sz > 0) st.spmv_ws = command_runner>{}.transfer_memory(st.device_id, st.stream_id, out{static_cast(ws_sz)}); + if (ws_sz > 0) st.spmv_ws = command_runner>{}.transfer_memory(st.device_id, st.stream_id, out{ws_sz}); st.d_ptr->spmv_csr(st.stream_id, st.n, st.n, st.nnz, T{1}, st.A_rp, st.A_ci, st.A_val, st.x, T{0}, st.w, st.spmv_ws); if constexpr (std::is_same_v) st.d_ptr->dcopy(st.stream_id, st.n, st.b, st.r); @@ -263,4 +263,4 @@ behavior fault_tolerant_cg_actor(stateful_actor>* self, self->quit(); } }; -} \ No newline at end of file +} diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native.cpp index 2186aa4e93..c19c13f161 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native.cpp @@ -8,6 +8,10 @@ #include #include "native_utils.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices" +#endif + void producer(ThreadSafeQueue& queue, std::vector matrix_pool) { for (auto& task : matrix_pool) { task.enqueue_time = std::chrono::steady_clock::now(); @@ -24,7 +28,7 @@ int main(int argc, char** argv) std::cout << "[INFO] Loading matrices...\n"; //std::vector matrix_pool = scan_for_matrices("/scratch/nqr159/matrix-collection/matrices/mixed", CGS_SOLVER); - std::vector matrix_pool = scan_for_matrices("../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices", CGS_SOLVER); + std::vector matrix_pool = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); if (matrix_pool.empty()) { diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native_sorted.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native_sorted.cpp index e28fbcd7a5..269ef2287e 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native_sorted.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.native_sorted.cpp @@ -8,6 +8,10 @@ #include #include "native_utils.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices" +#endif + void producer(ThreadSafeQueue& queue, std::vector matrix_pool) { // Order tasks from lowest NNZ to highest NNZ std::sort(matrix_pool.begin(), matrix_pool.end(), [](const MatrixTask& a, const MatrixTask& b) { @@ -27,7 +31,7 @@ int main(int argc, char** argv) { // if (argc > 1) num_streams = std::max(num_streams, std::atoi(argv[1])); std::cout << "[INFO] Loading matrices...\n"; - std::vector matrix_pool = scan_for_matrices("../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices", CGS_SOLVER); + std::vector matrix_pool = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); if (matrix_pool.empty()) { diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.test.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.test.cpp index da20c1e9c5..9e3502aceb 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.test.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadA/main.test.cpp @@ -16,6 +16,10 @@ #include "ft_cg_actor.hpp" #include "supervisor_actor.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices" +#endif + using namespace caf; using namespace caf::cuda; @@ -31,7 +35,7 @@ void caf_main(actor_system& sys) { //auto tasks_vec = scan_for_matrices("/scratch/nqr159/matrix-collection/matrices/mixed", CGS_SOLVER); // auto tasks_vec = scan_for_matrices("/scratch/nqr159/matrices/workloadA", CGS_SOLVER); - auto tasks_vec = scan_for_matrices("../../scripts/Irregular-Workload/workloadA/downloaded_matrices/matrices", CGS_SOLVER); + auto tasks_vec = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadB/CMakeLists.txt b/libcaf_cuda/sc26/Irregular-Workload/workloadB/CMakeLists.txt index 199d000211..afae47fa22 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadB/CMakeLists.txt +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadB/CMakeLists.txt @@ -1,132 +1,79 @@ -cmake_minimum_required(VERSION 3.16.3) - -# 1) Enforce C++20 (needs modern C++ features) -set(CMAKE_CXX_STANDARD 20) -set(CMAKE_CXX_STANDARD_REQUIRED ON) -set(CMAKE_CXX_EXTENSIONS OFF) - -# 2) Set CAF source and build directories -set(CAF_SRC "${CMAKE_CURRENT_SOURCE_DIR}/../../../../../actor-framework") -set(CAF_BUILD "${CAF_SRC}/build") - - -#required since cmake cant seem to find the compiler -set(CMAKE_CXX_COMPILER "/usr/bin/g++") -set(CMAKE_C_COMPILER "/usr/bin/gcc") - -# Enable CUDA as a first-class language for the project -project(CUDA_ACTORS LANGUAGES CXX CUDA) - -find_package(CUDA REQUIRED) -find_package(CUDAToolkit REQUIRED) - -include_directories( - "${CAF_SRC}/include" - "${CAF_SRC}/libcaf_io" - "${CAF_SRC}/libcaf_core" - "${CAF_SRC}/libcaf_opencl" - "${CAF_BUILD}/libcaf_core" # for generated build_config.hpp - "${CAF_SRC}/libcaf_cuda" -) - -# --- CUDA Kernel Compilation --- -set(STABILITY_KERNEL_SRC "${CMAKE_CURRENT_SOURCE_DIR}/stability_kernels.cu") -set(STABILITY_KERNEL_CUBIN "${CMAKE_CURRENT_BINARY_DIR}/stability_kernels.cubin") - -# Target CUDA architecture. 'native' targets the current machine's GPU (requires CUDA 11.6+). -# You can override this with -DCUDA_ARCH=sm_XX if needed. -set(CUDA_ARCH "native" CACHE STRING "Target CUDA architecture (e.g., native, sm_70, sm_75, sm_80, sm_86)") - -add_custom_command( - OUTPUT ${STABILITY_KERNEL_CUBIN} - COMMAND ${CUDAToolkit_NVCC_EXECUTABLE} - -cubin - -arch=${CUDA_ARCH} - -o ${STABILITY_KERNEL_CUBIN} - ${STABILITY_KERNEL_SRC} - DEPENDS ${STABILITY_KERNEL_SRC} - COMMENT "Compiling CUDA kernel ${STABILITY_KERNEL_SRC} for architecture: ${CUDA_ARCH}" - VERBATIM -) - -set(JACOBI_KERNEL_SRC "${CMAKE_CURRENT_SOURCE_DIR}/jacobi_kernels.cu") -set(JACOBI_KERNEL_CUBIN "${CMAKE_CURRENT_BINARY_DIR}/jacobi_kernels.cubin") - -add_custom_command( - OUTPUT ${JACOBI_KERNEL_CUBIN} - COMMAND ${CUDAToolkit_NVCC_EXECUTABLE} - -cubin - -arch=${CUDA_ARCH} - -o ${JACOBI_KERNEL_CUBIN} - ${JACOBI_KERNEL_SRC} - DEPENDS ${JACOBI_KERNEL_SRC} - COMMENT "Compiling CUDA kernel ${JACOBI_KERNEL_SRC} for architecture: ${CUDA_ARCH}" - VERBATIM -) - -add_custom_target(stability_kernels_cubin ALL DEPENDS ${STABILITY_KERNEL_CUBIN}) -add_custom_target(jacobi_kernels_cubin ALL DEPENDS ${JACOBI_KERNEL_CUBIN}) - - -# 5) Declare your executable and its source files -add_executable(test main.test.cpp sparse_utils.cpp) -target_sources(test PRIVATE +set(workload_cubin_dir "${CMAKE_CURRENT_BINARY_DIR}") +set(stability_cubin "${workload_cubin_dir}/stability_kernels.cubin") +set(jacobi_cubin "${workload_cubin_dir}/jacobi_kernels.cubin") + +sc26_add_cubin( + sc26_irregular_b_stability_cubin + "${stability_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/stability_kernels.cu") +sc26_add_cubin( + sc26_irregular_b_jacobi_cubin + "${jacobi_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/jacobi_kernels.cu") + +set(matrix_dir + "${CMAKE_CURRENT_SOURCE_DIR}/../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices") +file(TO_CMAKE_PATH "${matrix_dir}" matrix_dir) + +add_executable( + sc26_irregular_b_actor + main.test.cpp + sparse_utils.cpp atoms.hpp ft_cg_actor.hpp ft_cg_jacobi_actor.hpp - supervisor_actor.hpp -) - -target_compile_definitions(test PRIVATE CAF_ENABLE_LOGGING) - -target_link_libraries(test + supervisor_actor.hpp) +target_compile_features(sc26_irregular_b_actor PRIVATE cxx_std_20) +target_compile_definitions( + sc26_irregular_b_actor PRIVATE - "${CAF_BUILD}/libcaf_core/libcaf_core.so" - "${CAF_BUILD}/libcaf_io/libcaf_io.so" - "${CAF_BUILD}/libcaf_cuda/libcaf_cuda.so" - CUDA::nvrtc - CUDA::cublas - CUDA::cusparse -) - - -# # 5) Declare your executable -# add_executable(hot-potatoe hot-potatoe.cpp sparse_utils.cpp) - -# target_compile_definitions(hot-potatoe PRIVATE CAF_ENABLE_LOGGING) - -# target_link_libraries(hot-potatoe -# PRIVATE -# "${CAF_BUILD}/libcaf_core/libcaf_core.so" -# "${CAF_BUILD}/libcaf_io/libcaf_io.so" -# "${CAF_BUILD}/libcaf_cuda/libcaf_cuda.so" -# CUDA::nvrtc -# CUDA::cublas -# CUDA::cusparse -# ) + CAF_ENABLE_LOGGING + SC26_MATRIX_DIR="${matrix_dir}") +target_link_libraries( + sc26_irregular_b_actor + PRIVATE CAF::cuda CUDA::cublas CUDA::cusparse) +sc26_set_runtime_output(sc26_irregular_b_actor irregular-workload-b) +add_dependencies( + sc26_irregular_b_actor + sc26_irregular_b_stability_cubin + sc26_irregular_b_jacobi_cubin) +sc26_copy_file_to_target(sc26_irregular_b_actor "${stability_cubin}") +sc26_copy_file_to_target(sc26_irregular_b_actor "${jacobi_cubin}") - -# FindThreads is required for std::thread in the native version find_package(Threads REQUIRED) - -# 6) Declare the native benchmark executable (raw CUDA/cuBLAS/cuSPARSE) -add_executable(workload-native main.native.cpp sparse_utils.cpp native_utils.cu) - -target_link_libraries(workload-native - PRIVATE - CUDA::cudart - CUDA::cublas - CUDA::cusparse - Threads::Threads -) - -# 7) Declare the native sorted benchmark executable -add_executable(workload-native-sorted main.native_sorted.cpp sparse_utils.cpp native_utils.cu) - -target_link_libraries(workload-native-sorted - PRIVATE - CUDA::cudart - CUDA::cublas - CUDA::cusparse - Threads::Threads -) +set(sc26_cuda_architecture "${CAF_CUDA_ARCH}") +string(REGEX REPLACE "^sm_" "" sc26_cuda_architecture + "${sc26_cuda_architecture}") + +foreach(variant native native_sorted) + set(target "sc26_irregular_b_${variant}") + if(variant STREQUAL "native") + set(main_source main.native.cpp) + else() + set(main_source main.native_sorted.cpp) + endif() + add_executable( + ${target} + "${main_source}" + sparse_utils.cpp + native_utils.cu) + target_compile_features(${target} PRIVATE cxx_std_20) + set_target_properties( + ${target} + PROPERTIES + CUDA_STANDARD 20 + CUDA_STANDARD_REQUIRED ON + CUDA_ARCHITECTURES "${sc26_cuda_architecture}") + target_compile_definitions(${target} PRIVATE SC26_MATRIX_DIR="${matrix_dir}") + target_link_libraries( + ${target} + PRIVATE CUDA::cudart CUDA::cublas CUDA::cusparse Threads::Threads) + sc26_set_runtime_output(${target} irregular-workload-b) +endforeach() + +add_custom_target( + sc26_irregular_workload_b_benchmarks + DEPENDS + sc26_irregular_b_actor + sc26_irregular_b_native + sc26_irregular_b_native_sorted) diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadB/ft_cg_actor.hpp b/libcaf_cuda/sc26/Irregular-Workload/workloadB/ft_cg_actor.hpp index 4fd5335cd8..8e206af3bb 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadB/ft_cg_actor.hpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadB/ft_cg_actor.hpp @@ -87,7 +87,7 @@ behavior fault_tolerant_cg_actor(stateful_actor>* self, st.d_err = runner.transfer_memory(st.device_id, st.stream_id, out(1)); size_t ws_sz = st.d_ptr->spmv_csr_buffer_size(st.stream_id, st.n, st.n, st.nnz, st.A_rp, st.A_ci, st.A_val, st.x, st.w); - if (ws_sz > 0) st.spmv_ws = command_runner>{}.transfer_memory(st.device_id, st.stream_id, out{static_cast(ws_sz)}); + if (ws_sz > 0) st.spmv_ws = command_runner>{}.transfer_memory(st.device_id, st.stream_id, out{ws_sz}); st.d_ptr->spmv_csr(st.stream_id, st.n, st.n, st.nnz, T{1}, st.A_rp, st.A_ci, st.A_val, st.x, T{0}, st.w, st.spmv_ws); if constexpr (std::is_same_v) st.d_ptr->dcopy(st.stream_id, st.n, st.b, st.r); @@ -213,4 +213,4 @@ behavior fault_tolerant_cg_actor(stateful_actor>* self, self->quit(); } }; -} \ No newline at end of file +} diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native.cpp index d3e1dafba9..5fe78d4657 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native.cpp @@ -8,6 +8,10 @@ #include #include "native_utils.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices" +#endif + void producer(ThreadSafeQueue& queue, std::vector matrix_pool) { for (auto& task : matrix_pool) { task.enqueue_time = std::chrono::steady_clock::now(); @@ -23,7 +27,7 @@ int main(int argc, char** argv) // if (argc > 1) num_streams = std::max(num_streams, std::atoi(argv[1])); std::cout << "[INFO] Loading matrices...\n"; - std::vector matrix_pool = scan_for_matrices("../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices", CGS_SOLVER); + std::vector matrix_pool = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); if (matrix_pool.empty()) { std::cerr << "No matrices found.\n"; diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native_sorted.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native_sorted.cpp index 7b53620ee1..62830548b4 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native_sorted.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.native_sorted.cpp @@ -8,6 +8,10 @@ #include #include "native_utils.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices" +#endif + void producer(ThreadSafeQueue& queue, std::vector matrix_pool) { // Order tasks from lowest NNZ to highest NNZ std::sort(matrix_pool.begin(), matrix_pool.end(), [](const MatrixTask& a, const MatrixTask& b) { @@ -27,7 +31,7 @@ int main(int argc, char** argv) { // if (argc > 1) num_streams = std::max(num_streams, std::atoi(argv[1])); std::cout << "[INFO] Loading matrices...\n"; - std::vector matrix_pool = scan_for_matrices("../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices", CGS_SOLVER); + std::vector matrix_pool = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); if (matrix_pool.empty()) { std::cerr << "No matrices found.\n"; diff --git a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.test.cpp b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.test.cpp index 54687dc1f7..4d85fb5d2d 100644 --- a/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.test.cpp +++ b/libcaf_cuda/sc26/Irregular-Workload/workloadB/main.test.cpp @@ -16,6 +16,10 @@ #include "ft_cg_actor.hpp" #include "supervisor_actor.hpp" +#ifndef SC26_MATRIX_DIR +# define SC26_MATRIX_DIR "../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices" +#endif + using namespace caf; using namespace caf::cuda; @@ -28,7 +32,7 @@ void caf_main(actor_system& sys) { //auto tasks_vec = scan_for_matrices("/scratch/nqr159/matrix-collection/matrices/spd", CGS_SOLVER); //auto tasks_vec = scan_for_matrices("/scratch/nqr159/matrix-collection/matrices/unsymmetric", CGS_SOLVER); //auto tasks_vec = scan_for_matrices("/scratch/nqr159/matrix-collection/matrix_corpus_v2/matrices/unsymmetric", CGS_SOLVER); - auto tasks_vec = scan_for_matrices("../../scripts/Irregular-Workload/workloadB/downloaded_matrices/matrices", CGS_SOLVER); + auto tasks_vec = scan_for_matrices(SC26_MATRIX_DIR, CGS_SOLVER); int num_gpus = manager::get().get_num_devices(); std::cout << "[INFO] Found " << num_gpus << " GPUs\n"; diff --git a/libcaf_cuda/sc26/README.md b/libcaf_cuda/sc26/README.md deleted file mode 100644 index e69de29bb2..0000000000 diff --git a/libcaf_cuda/sc26/Runtime-Overhead/CMakeLists.txt b/libcaf_cuda/sc26/Runtime-Overhead/CMakeLists.txt new file mode 100644 index 0000000000..5ef40dd24c --- /dev/null +++ b/libcaf_cuda/sc26/Runtime-Overhead/CMakeLists.txt @@ -0,0 +1,29 @@ +set(runtime_cubin "${CMAKE_CURRENT_BINARY_DIR}/mmul.cubin") +sc26_add_cubin( + sc26_runtime_mmul_cubin + "${runtime_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/mmul.cu") + +add_executable(sc26_runtime_native cuda_native.cpp) +target_compile_features(sc26_runtime_native PRIVATE cxx_std_20) +target_link_libraries(sc26_runtime_native PRIVATE CUDA::cuda_driver) +sc26_set_runtime_output(sc26_runtime_native runtime-overhead) +add_dependencies(sc26_runtime_native sc26_runtime_mmul_cubin) +sc26_copy_file_to_target(sc26_runtime_native "${runtime_cubin}") + +foreach(variant actor_facade command_runner) + set(target "sc26_runtime_${variant}") + add_executable(${target} "${variant}.cpp") + target_compile_features(${target} PRIVATE cxx_std_20) + target_link_libraries(${target} PRIVATE CAF::cuda) + sc26_set_runtime_output(${target} runtime-overhead) + add_dependencies(${target} sc26_runtime_mmul_cubin) + sc26_copy_file_to_target(${target} "${runtime_cubin}") +endforeach() + +add_custom_target( + sc26_runtime_overhead_benchmarks + DEPENDS + sc26_runtime_native + sc26_runtime_actor_facade + sc26_runtime_command_runner) diff --git a/libcaf_cuda/sc26/Runtime-Overhead/command_runner.cpp b/libcaf_cuda/sc26/Runtime-Overhead/command_runner.cpp index 82ddad6e09..a5501522ce 100644 --- a/libcaf_cuda/sc26/Runtime-Overhead/command_runner.cpp +++ b/libcaf_cuda/sc26/Runtime-Overhead/command_runner.cpp @@ -133,7 +133,7 @@ caf::behavior mmul_actor_fun_2(caf::stateful_actor* self) { auto output = command.run_async( program,dims, 1, 0, device, - arg1,arg2,out{N*N},in{N}); + arg1, arg2, out{static_cast(N) * N}, in{N}); auto t_response_received = clock::now(); diff --git a/libcaf_cuda/sc26/Runtime-Overhead/cuda_native.cpp b/libcaf_cuda/sc26/Runtime-Overhead/cuda_native.cpp index 19ba316789..33af35eff7 100644 --- a/libcaf_cuda/sc26/Runtime-Overhead/cuda_native.cpp +++ b/libcaf_cuda/sc26/Runtime-Overhead/cuda_native.cpp @@ -202,7 +202,14 @@ int main(int argc, char** argv) { checkCU(cuDeviceGet(&dev, 0), "cuDeviceGet(0)"); CUcontext ctx; +#if CUDA_VERSION >= 13000 + CUctxCreateParams ctx_params = {}; + checkCU(cuCtxCreate(&ctx, &ctx_params, + CU_CTX_SCHED_AUTO | CU_CTX_MAP_HOST, dev), + "cuCtxCreate"); +#else checkCU(cuCtxCreate(&ctx, CU_CTX_SCHED_AUTO | CU_CTX_MAP_HOST, dev), "cuCtxCreate"); +#endif const std::string cubinPath = "mmul.cubin"; std::string cubin; diff --git a/libcaf_cuda/sc26/Sequence-Independent-Tasks/CMakeLists.txt b/libcaf_cuda/sc26/Sequence-Independent-Tasks/CMakeLists.txt new file mode 100644 index 0000000000..5e16c08437 --- /dev/null +++ b/libcaf_cuda/sc26/Sequence-Independent-Tasks/CMakeLists.txt @@ -0,0 +1,68 @@ +set(sc26_sequence_cubin "${CMAKE_CURRENT_BINARY_DIR}/mmul.cubin") +set(sc26_sequence_output_dir "${CMAKE_BINARY_DIR}/bin/$") + +add_custom_command( + OUTPUT "${sc26_sequence_cubin}" + COMMAND "${CUDAToolkit_NVCC_EXECUTABLE}" + --cubin + "-arch=${CAF_CUDA_ARCH}" + --generate-line-info + -o "${sc26_sequence_cubin}" + "${CMAKE_CURRENT_SOURCE_DIR}/mmul.cu" + DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/mmul.cu" + COMMENT "Compiling the SC26 sequence-independent matrix kernel" + VERBATIM) + +add_custom_target(sc26_sequence_mmul_cubin + DEPENDS "${sc26_sequence_cubin}") + +function(add_sc26_sequence_benchmark target source) + add_executable(${target} "${source}") + target_compile_features(${target} PRIVATE cxx_std_20) + target_link_libraries(${target} PRIVATE CUDA::cuda_driver ${ARGN}) + set_target_properties( + ${target} + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${sc26_sequence_output_dir}" + VS_DEBUGGER_WORKING_DIRECTORY "$") + add_dependencies(${target} sc26_sequence_mmul_cubin) + + add_custom_command( + TARGET ${target} POST_BUILD + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "${sc26_sequence_cubin}" + "$/mmul.cubin" + VERBATIM) +endfunction() + +add_sc26_sequence_benchmark(sc26_sequence_native cuda_native.cpp) +add_sc26_sequence_benchmark( + sc26_sequence_actor_facade actor_facade.cpp CAF::cuda) +add_sc26_sequence_benchmark( + sc26_sequence_command_runner command_runner.cpp CAF::cuda) + +# Shared CAF builds place DLLs in their component build directories. Copy the +# two CAF runtime DLLs beside each CAF benchmark executable so they can run +# directly without modifying PATH. +if(WIN32 AND BUILD_SHARED_LIBS) + foreach(target + sc26_sequence_actor_facade + sc26_sequence_command_runner) + add_custom_command( + TARGET ${target} POST_BUILD + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "$" + "$" + COMMAND "${CMAKE_COMMAND}" -E copy_if_different + "$" + "$" + VERBATIM) + endforeach() +endif() + +add_custom_target( + sc26_sequence_benchmarks + DEPENDS + sc26_sequence_native + sc26_sequence_actor_facade + sc26_sequence_command_runner) diff --git a/libcaf_cuda/sc26/Sequence-Independent-Tasks/actor_facade.cpp b/libcaf_cuda/sc26/Sequence-Independent-Tasks/actor_facade.cpp index f8ac8d32ec..6ea704bf0d 100644 --- a/libcaf_cuda/sc26/Sequence-Independent-Tasks/actor_facade.cpp +++ b/libcaf_cuda/sc26/Sequence-Independent-Tasks/actor_facade.cpp @@ -28,8 +28,8 @@ caf::behavior throughput_mapping_manager(caf::stateful_actor #include #include + +#include +#include #include -#include +#include +#include #include -#include -#include -#include -#include -#include -#include "caf/actor_registry.hpp" -//#include - - using namespace caf; -using namespace std::chrono_literals; - - -void serial_matrix_multiply(const std::vector& a, - const std::vector& b, - std::vector& c, - int N) { - - - for (int i = 0; i < N; ++i) { - for (int j = 0; j < N; ++j) { - int sum = 0; - for (int k = 0; k < N; ++k) { - sum += a[i * N + k] * b[k * N + j]; - } - c[i * N + j] = sum; - } - } -} -using command = - caf::cuda::command_runner<>; +using command = caf::cuda::command_runner<>; command mmul_command; struct mmul_state { - caf::cuda::program_ptr program; - int total_expected = 0; - int results_received = 0; + caf::cuda::program_ptr program; + int iterations = 0; + int N = 0; + caf::actor parent; + std::shared_ptr> matrixA; + std::shared_ptr> matrixB; + std::shared_ptr> matrixC; }; -//global output buffer meant to disclude it from timing -//the other benchmark test do not include its memory allocations in it -//so its only fair that we do not either -std::vector matrixC; - - - - -caf::behavior mmul_actor_fun(caf::stateful_actor* self, - caf::cuda::program_ptr mmul_kernel, int iterations) { - - self ->state().program = mmul_kernel; - self ->state().total_expected = iterations; - -return { - [=](const std::vector& matrixA, - const std::vector& matrixB, - int N) { - - using clock = std::chrono::steady_clock; - using ms = std::chrono::duration; - - - caf::cuda::manager& mgr = caf::cuda::manager::get(); - int device = 0; - int stream = 1; - - auto t_total_start = clock::now(); - // ------------------------- - // create_in_arg A - // ------------------------- - auto t_a_inarg_start = clock::now(); - - auto inA = caf::cuda::create_in_arg(std::move(matrixA)); - - auto t_a_inarg_end = clock::now(); - - // ------------------------- - // transfer A - // ------------------------- - auto t_a_transfer_start = clock::now(); - - auto arg1 = mmul_command.transfer_memory( - device, - stream, - std::move(inA)); - - auto t_a_transfer_end = clock::now(); - - // ------------------------- - // create_in_arg B - // ------------------------- - auto t_b_inarg_start = clock::now(); - - auto inB = caf::cuda::create_in_arg(std::move(matrixB)); - - auto t_b_inarg_end = clock::now(); - - // ------------------------- - // transfer B - // ------------------------- - auto t_b_transfer_start = clock::now(); - - auto arg2 = mmul_command.transfer_memory( - device, - stream, - std::move(inB)); - - auto t_b_transfer_end = clock::now(); - - const int THREADS = 32; - const int BLOCKS = (N + THREADS - 1) / THREADS; - - caf::cuda::nd_range dims( - BLOCKS, BLOCKS, 1, - THREADS, THREADS, 1); - - caf::cuda::mmul_async_command command; - auto output = command.run_async( - self->state().program, - dims, - 1, 0, device, - arg1,arg2,out{N*N},in{N}); - - caf::cuda::mem_ptr dC = std::get<2>(output); +caf::behavior mmul_actor_fun( + caf::stateful_actor* self, + caf::cuda::program_ptr mmul_kernel, + int iterations, + int N, + caf::actor parent, + std::shared_ptr> matrixA, + std::shared_ptr> matrixB, + std::shared_ptr> matrixC) { + auto& st = self->state(); + st.program = std::move(mmul_kernel); + st.iterations = iterations; + st.N = N; + st.parent = std::move(parent); + st.matrixA = std::move(matrixA); + st.matrixB = std::move(matrixB); + st.matrixC = std::move(matrixC); + + return { + [=](int) { + auto& state = self->state(); + constexpr int device = 0; + constexpr int stream = 1; + constexpr int threads = 32; + const int blocks = (state.N + threads - 1) / threads; + const auto elements = static_cast(state.N) + * static_cast(state.N); + caf::cuda::nd_range dims(blocks, blocks, 1, + threads, threads, 1); + caf::cuda::mmul_async_command runner; + caf::cuda::mem_ptr completion; + + for (int i = 0; i < state.iterations; ++i) { + // The matrices live in actor state. create_in_arg only creates a + // non-owning view, so no matrix data is copied through a CAF message. + auto dA = mmul_command.transfer_memory( + device, stream, caf::cuda::create_in_arg(*state.matrixA)); + auto dB = mmul_command.transfer_memory( + device, stream, caf::cuda::create_in_arg(*state.matrixB)); + + auto output = runner.run_async( + state.program, dims, stream, 0, device, + dA, dB, out{elements}, in{state.N}); + auto dC = std::get<2>(output); + + mmul_command.copy_to_host_async( + dC, state.matrixC->data(), state.matrixC->size()); + + // Match native CUDA's per-iteration ordering: + // alloc A/B/C -> H2D A/B -> kernel -> D2H C -> free A/B/C. + mmul_command.free_memory(dA); + mmul_command.free_memory(dB); + mmul_command.free_memory(dC); + completion = dC; + } - auto self_hdl = caf::actor_cast(self); - mmul_command.copy_to_host_async(dC, matrixC.data(), N*N, [self_hdl](int*, size_t) { - caf::anon_mail(kernel_done_atom_v).send(self_hdl); - }); + // Match native CUDA's one stream synchronization after the full series. + // free_memory invalidates the allocation while retaining the stream + // association used here to wait for all queued work, including frees. + if (completion) + completion->synchronize(); + anon_mail(1).send(state.parent); + self->quit(); }, - [=](kernel_done_atom) { - if (++self->state().results_received == self->state().total_expected) { - self->quit(); - } - } }; } - -void run_mmul_test(caf::actor_system& sys, int matrix_size,int iterations) { - - - caf::cuda::manager::init(sys); - // ------------------------------------ - // Start timing - // ------------------------------------ - - // Spawn num_actors actors running the mmul behavior - std::vector matrixA(matrix_size * matrix_size,2); - std::vector matrixB(matrix_size * matrix_size,3); - - matrixC.resize(matrix_size*matrix_size); - - auto& mgr = caf::cuda::manager::get(); - - auto program = - mgr.create_program_from_cubin("mmul.cubin", "matrixMul"); - - - using clock = std::chrono::steady_clock; +void run_mmul_test( + caf::actor_system& sys, + caf::cuda::program_ptr program, + const std::shared_ptr>& matrixA, + const std::shared_ptr>& matrixB, + const std::shared_ptr>& matrixC, + int matrix_size, + int iterations) { + scoped_actor coordinator{sys}; + auto worker = sys.spawn(mmul_actor_fun, std::move(program), iterations, + matrix_size, actor_cast(coordinator), + matrixA, matrixB, matrixC); auto start = std::chrono::steady_clock::now(); - - caf::actor a = sys.spawn(mmul_actor_fun, program, iterations); - - for (int i = 0; i < iterations; i++) - anon_mail(matrixA,matrixB,matrix_size).send(a); - - // Wait for all actors to finish - sys.await_all_actors_done(); - - // ------------------------------------ - // Stop timing - // ------------------------------------ + anon_mail(0).send(worker); + coordinator->receive([](int) {}); auto end = std::chrono::steady_clock::now(); - auto duration_ms = - std::chrono::duration_cast(end - start).count(); - - std::cout << "[MMUL TEST] matrix_size=" << matrix_size - << " iterations = " << iterations << - ", time=" << duration_ms << " ms\n"; + coordinator->wait_for(worker); + auto duration_ms = std::chrono::duration_cast( + end - start).count(); - caf::cuda::manager::shutdown(); + if (matrixC->empty() || matrixC->front() != matrix_size) + throw std::runtime_error("matrix multiplication produced an invalid result"); + std::cout << "[MMUL TEST] matrix_size=" << matrix_size + << " iterations = " << iterations + << ", time=" << duration_ms << " ms" << std::endl; } - void caf_main(caf::actor_system& sys) { + constexpr int matrix_size = 1000; + const auto elements = static_cast(matrix_size) + * static_cast(matrix_size); - for (int i = 1000; i < 11000; i+=1000) - run_mmul_test(sys,1000,i); + caf::cuda::manager::init(sys); + auto program = caf::cuda::manager::get().create_program_from_cubin( + "mmul.cubin", "matrixMul"); -} + // Match native CUDA: host buffers persist across every timed series. + auto matrixA = std::make_shared>(elements, 1); + auto matrixB = std::make_shared>(elements, 1); + auto matrixC = std::make_shared>(elements, 0); + for (int iterations = 1000; iterations <= 10000; iterations += 1000) + run_mmul_test(sys, program, matrixA, matrixB, matrixC, + matrix_size, iterations); + + caf::cuda::manager::shutdown(); +} int main(int argc, char** argv) { - // Initialize user defined types and messages if needed. - //init_global_meta_objects(); - - // Initialize the global type information. core::init_global_meta_objects(); - // Create the config. actor_system_config cfg; - - // --- SINGLE THREAD CONFIGURATION --- cfg.set("caf.scheduler.max-threads", 1); - cfg.set("caf.scheduler.policy", "sharing"); - // ------------------------------------ + cfg.set("caf.scheduler.policy", "sharing"); - // Read CLI options. (Note: CLI flags like --caf.scheduler.max-threads=4 - // will override the hardcoded '1' above if provided by the user). auto err = cfg.parse(argc, argv); if (err) return EXIT_FAILURE; - if (cfg.helptext_printed()) return 0; - // Create the actor system (the scheduler starts here). actor_system sys{cfg}; - - // Run user-defined code. caf_main(sys); - return 0; } - -// CAF_MAIN() diff --git a/libcaf_cuda/sc26/Sequence-Independent-Tasks/cuda_native.cpp b/libcaf_cuda/sc26/Sequence-Independent-Tasks/cuda_native.cpp index db076f8860..24fe7b623b 100644 --- a/libcaf_cuda/sc26/Sequence-Independent-Tasks/cuda_native.cpp +++ b/libcaf_cuda/sc26/Sequence-Independent-Tasks/cuda_native.cpp @@ -55,7 +55,12 @@ int main() { checkCU(cuDeviceGet(&dev, 0), "cuDeviceGet(0)"); CUcontext ctx; +#if CUDA_VERSION >= 13000 + CUctxCreateParams ctx_params = {}; + checkCU(cuCtxCreate(&ctx, &ctx_params, 0, dev), "cuCtxCreate"); +#else checkCU(cuCtxCreate(&ctx, 0, dev), "cuCtxCreate"); +#endif std::string ptx = readFile("mmul.cubin"); @@ -120,9 +125,14 @@ int main() { auto end = clock::now(); double total_ms = std::chrono::duration(end - start).count(); + if (h_c.empty() || h_c.front() != N) { + std::cerr << "Matrix multiplication produced an invalid result\n"; + return EXIT_FAILURE; + } + std::cout << "[SERIES RESULT] Matrix " << N << "x" << N << ", iterations = " << iterations - << ", total GPU time = " << total_ms << " ms\n"; + << ", total GPU time = " << total_ms << " ms" << std::endl; } // ---------------------------------- diff --git a/libcaf_cuda/sc26/benchmarks/README.md b/libcaf_cuda/sc26/benchmarks/README.md new file mode 100644 index 0000000000..89d1ee68b4 --- /dev/null +++ b/libcaf_cuda/sc26/benchmarks/README.md @@ -0,0 +1,88 @@ +# SC26 benchmarks + +Run these commands from the repository root in Bash. + +## Build + +```bash +cmake -S . -B build-sc26 \ + -DCMAKE_BUILD_TYPE=Release \ + -DCAF_ENABLE_SC26_BENCHMARKS=ON \ + -DCAF_ENABLE_EXAMPLES=OFF \ + -DCAF_ENABLE_TESTING=OFF + +cmake --build build-sc26 --config Release --target sc26_benchmarks --parallel +``` + +Build one suite by replacing `sc26_benchmarks` with one of: + +```text +sc26_sequence_benchmarks +sc26_runtime_overhead_benchmarks +sc26_batched_mmul_benchmarks +sc26_fault_tolerance_benchmarks +sc26_irregular_workload_a_benchmarks +sc26_irregular_workload_b_benchmarks +``` + +## Sequence-independent tasks + +```bash +python3 libcaf_cuda/sc26/benchmarks/run_sequence_independent.py \ + --build-dir build-sc26 --config Release --runs 10 +``` + +## Runtime overhead + +```bash +python3 libcaf_cuda/sc26/benchmarks/run_runtime_overhead.py \ + --build-dir build-sc26 --config Release --runs 10 +``` + +## Other benchmarks + +Run the executables from their output directories so that their CUDA modules are found: + +```bash +( + cd build-sc26/bin/Release/sc26/batched-matrix-multiply + ./sc26_batched_matrix_multiply +) + +( + cd build-sc26/bin/Release/sc26/fault-tolerance + ./sc26_fault_tolerance +) +``` + +Download the sparse-matrix inputs before running the irregular workloads: + +```bash +python3 -m pip install numpy scipy ssgetpy + +( + cd libcaf_cuda/sc26/scripts/Irregular-Workload/workloadA + python3 download-matrices.py +) + +( + cd libcaf_cuda/sc26/scripts/Irregular-Workload/workloadB + python3 download-matrices.py +) +``` + +```bash +( + cd build-sc26/bin/Release/sc26/irregular-workload-a + ./sc26_irregular_a_actor + ./sc26_irregular_a_native + ./sc26_irregular_a_native_sorted +) + +( + cd build-sc26/bin/Release/sc26/irregular-workload-b + ./sc26_irregular_b_actor + ./sc26_irregular_b_native + ./sc26_irregular_b_native_sorted +) +``` diff --git a/libcaf_cuda/sc26/benchmarks/run_runtime_overhead.py b/libcaf_cuda/sc26/benchmarks/run_runtime_overhead.py index 65a03bafd2..9c9488e484 100755 --- a/libcaf_cuda/sc26/benchmarks/run_runtime_overhead.py +++ b/libcaf_cuda/sc26/benchmarks/run_runtime_overhead.py @@ -29,13 +29,10 @@ # Configuration # --------------------------------------------------------------------------- -BENCH_DIR = Path(__file__).resolve().parent.parent / "Runtime-Overhead" - -BINARIES = { - "cuda_native": BENCH_DIR / "cuda_native", - "actor_facade": BENCH_DIR / "actor_facade", - "command_runner": BENCH_DIR / "command_runner", -} +REPOSITORY_ROOT = Path(__file__).resolve().parents[3] +BENCH_DIR = REPOSITORY_ROOT / "libcaf_cuda" / "sc26" / "Runtime-Overhead" +DEFAULT_BUILD_DIRECTORY = REPOSITORY_ROOT / "build-sc26" +EXECUTABLE_SUFFIX = ".exe" if os.name == "nt" else "" # Patterns shared across both implementations. # cuda_native prints: N=1000 ... TOTAL: 42.37 ms @@ -43,25 +40,56 @@ RE_SIZE = re.compile(r"N=(\d+)") RE_TOTAL_NATIVE = re.compile(r"^TOTAL:\s+([\d.]+)", re.MULTILINE) RE_TOTAL_ACTOR = re.compile(r"^TOTAL end-to-end:\s+([\d.]+)", re.MULTILINE) +RE_ACTOR_LATENCY = re.compile( + r"\[LATENCY TEST\]\s+matrix_size=(\d+),\s+time=([\d.]+)\s+ms" +) # Warmup size — excluded from parsed results WARMUP_SIZE = 64 -ENV = {**os.environ, "CUDA_VISIBLE_DEVICES": "0"} - # --------------------------------------------------------------------------- # Running # --------------------------------------------------------------------------- -def run_binary(binary: Path, run_index: int, timeout: int = 600) -> str: +def locate_binaries(build_directory: Path, configuration: str) -> dict[str, Path]: + legacy = { + "cuda_native": BENCH_DIR / f"cuda_native{EXECUTABLE_SUFFIX}", + "actor_facade": BENCH_DIR / f"actor_facade{EXECUTABLE_SUFFIX}", + "command_runner": BENCH_DIR / f"command_runner{EXECUTABLE_SUFFIX}", + } + if all(binary.is_file() for binary in legacy.values()): + return legacy + + candidates = [ + build_directory / "bin" / configuration / "sc26" / "runtime-overhead", + build_directory / "bin" / "sc26" / "runtime-overhead", + ] + for candidate in candidates: + binaries = { + "cuda_native": candidate / f"sc26_runtime_native{EXECUTABLE_SUFFIX}", + "actor_facade": candidate / f"sc26_runtime_actor_facade{EXECUTABLE_SUFFIX}", + "command_runner": candidate / f"sc26_runtime_command_runner{EXECUTABLE_SUFFIX}", + } + if all(binary.is_file() for binary in binaries.values()): + return binaries + + raise FileNotFoundError( + "could not find all three runtime-overhead executables in " + f"{BENCH_DIR} or {candidates[0]}" + ) + + +def run_binary( + binary: Path, run_index: int, timeout: int, device: str +) -> str: """Run a single binary and return its stdout as a string.""" print(f" run {run_index} ...", flush=True) result = subprocess.run( [str(binary)], capture_output=True, text=True, - cwd=str(BENCH_DIR), - env=ENV, + cwd=str(binary.parent), + env={**os.environ, "CUDA_VISIBLE_DEVICES": device}, timeout=timeout, ) if result.returncode != 0: @@ -81,6 +109,13 @@ def parse_output(name: str, stdout: str) -> dict[int, float]: Tracks the most-recently seen N= header so each TOTAL is paired correctly. """ results: dict[int, float] = {} + if name == "actor_facade": + return { + int(match.group(1)): float(match.group(2)) + for match in RE_ACTOR_LATENCY.finditer(stdout) + if int(match.group(1)) != WARMUP_SIZE + } + current_size: int | None = None total_re = RE_TOTAL_NATIVE if name == "cuda_native" else RE_TOTAL_ACTOR @@ -196,7 +231,12 @@ def comparison_table(all_data: dict[str, dict[int, list[float]]]) -> str: def main() -> None: parser = argparse.ArgumentParser(description="Runtime-Overhead benchmark harness") - parser.add_argument("--runs", type=int, default=10, help="Number of times to run each binary (default: 10)") + parser.add_argument("--build-dir", type=Path, default=DEFAULT_BUILD_DIRECTORY, + help=f"CAF build directory (default: {DEFAULT_BUILD_DIRECTORY})") + parser.add_argument("--config", default="Release", help="build configuration") + parser.add_argument("--runs", type=int, default=10, help="Number of times to run each binary (default: 10)") + parser.add_argument("--timeout", type=int, default=600, help="Timeout for each run in seconds") + parser.add_argument("--device", default="0", help="CUDA_VISIBLE_DEVICES value") parser.add_argument("--output", type=str, default=str(BENCH_DIR / "results" / "benchmark_results.txt"), help="Path to the output file") args = parser.parse_args() @@ -204,13 +244,14 @@ def main() -> None: output_path = Path(args.output) output_path.parent.mkdir(parents=True, exist_ok=True) num_runs = args.runs + binaries = locate_binaries(args.build_dir.resolve(), args.config) # {impl_name: {matrix_size: [time_ms, ...]}} all_data: dict[str, dict[int, list[float]]] = {} # {impl_name: list of raw stdout strings (one per run)} raw_outputs: dict[str, list[str]] = {} - for name, binary in BINARIES.items(): + for name, binary in binaries.items(): print(f"\n[{name}]", flush=True) if not binary.exists(): print(f" SKIP — binary not found: {binary}", file=sys.stderr) @@ -220,7 +261,7 @@ def main() -> None: raw_outputs[name] = [] for i in range(1, num_runs + 1): - stdout = run_binary(binary, i) + stdout = run_binary(binary, i, args.timeout, args.device) raw_outputs[name].append(stdout) parsed = parse_output(name, stdout) for size, t in parsed.items(): @@ -236,11 +277,11 @@ def main() -> None: lines.append(header_line()) lines.append(f"Date : {datetime.now().strftime('%Y-%m-%d %H:%M:%S')}") lines.append(f"Runs : {num_runs}") - lines.append(f"GPU : CUDA_VISIBLE_DEVICES=0") - lines.append(f"Directory : {BENCH_DIR}") + lines.append(f"GPU : CUDA_VISIBLE_DEVICES={args.device}") + lines.append(f"Directory : {next(iter(binaries.values())).parent}") lines.append("") - for name in BINARIES: + for name in binaries: if name not in all_data: continue @@ -296,7 +337,7 @@ def main() -> None: lines.append(header_line()) lines.append(section("FULL RAW OUTPUT (all runs)")) lines.append(header_line()) - for name in BINARIES: + for name in binaries: if name not in raw_outputs: continue lines.append(f"\n{'='*10} {name} {'='*10}") diff --git a/libcaf_cuda/sc26/benchmarks/run_sequence_independent.py b/libcaf_cuda/sc26/benchmarks/run_sequence_independent.py index 9f630ba7f1..04af63b701 100755 --- a/libcaf_cuda/sc26/benchmarks/run_sequence_independent.py +++ b/libcaf_cuda/sc26/benchmarks/run_sequence_independent.py @@ -1,432 +1,341 @@ #!/usr/bin/env python3 -""" -Sequence-Independent-Tasks Benchmark Harness -============================================= -Runs main_cuda_native, main_actor_facade, and main_command_runner each -NUM_RUNS times. Each run is expected to execute one timed 10 000-iteration -series, emit cumulative [MILESTONE] lines every 1 000 iterations, and print -one final [SERIES RESULT] line. The output file contains: - - 1. All raw [MILESTONE] and [SERIES RESULT] lines from every run. - 2. Statistics across runs for the 10 000-iteration series at each - 1 000-iteration milestone (mean / min / max / stddev). - 3. Incremental time per 1 000-iteration step (derived from adjacent - milestones) — this shows how long each individual 1 000-iteration - "slice" took on average. - 4. A cross-implementation comparison table. - -Usage: - cd sc26/ - python3 benchmarks/run_sequence_independent.py [--runs N] [--output PATH] - -Output: - Sequence-Independent-Tasks/results/benchmark_results.txt (default) -""" +"""Run and compare the SC26 sequence-independent CUDA benchmarks.""" + +from __future__ import annotations import argparse import os import re -import subprocess import statistics +import subprocess import sys -from collections import defaultdict +from dataclasses import dataclass from datetime import datetime from pathlib import Path -# --------------------------------------------------------------------------- -# Configuration -# --------------------------------------------------------------------------- -BENCH_DIR = Path(__file__).resolve().parent.parent / "Sequence-Independent-Tasks" - -BINARIES = { - "main_cuda_native": BENCH_DIR / "main_cuda_native", - "main_actor_facade": BENCH_DIR / "main_actor_facade", - "main_command_runner": BENCH_DIR / "main_command_runner", -} +REPOSITORY_ROOT = Path(__file__).resolve().parents[3] +DEFAULT_BUILD_DIRECTORY = REPOSITORY_ROOT / "build-sc26" +BENCHMARK_DIRECTORY = ( + REPOSITORY_ROOT / "libcaf_cuda" / "sc26" / "Sequence-Independent-Tasks" +) +DEFAULT_OUTPUT = BENCHMARK_DIRECTORY / "results" / "benchmark_results.txt" +EXECUTABLE_SUFFIX = ".exe" if os.name == "nt" else "" -# [MILESTONE] 3000 / 10000 iterations, elapsed = 1564.72 ms -RE_MILESTONE = re.compile( - r"\[MILESTONE\]\s+(\d+)\s*/\s*(\d+)\s+iterations,\s+elapsed\s*=\s*([\d.]+)" +NATIVE_PATTERN = re.compile( + r"\[SERIES RESULT\].*iterations\s*=\s*(?P\d+)" + r".*total GPU time\s*=\s*(?P