diff --git a/include/ghex/device/cuda/event.hpp b/include/ghex/device/cuda/event.hpp index 4e0305df..b35c8ee8 100644 --- a/include/ghex/device/cuda/event.hpp +++ b/include/ghex/device/cuda/event.hpp @@ -27,8 +27,12 @@ struct cuda_event cudaEvent_t m_event; ghex::util::moved_bit m_moved; - cuda_event() { - GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(&m_event, cudaEventDisableTiming)) + cuda_event() + : cuda_event(cudaEventDisableTiming) + { + } + explicit cuda_event(unsigned int flags) { + GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(&m_event, flags)) }; cuda_event(const cuda_event&) = delete; cuda_event& operator=(const cuda_event&) = delete; @@ -40,15 +44,7 @@ struct cuda_event if (!m_moved) { GHEX_CHECK_CUDA_RESULT_NO_THROW(cudaEventDestroy(m_event)) } } - /** - * @brief Returns `true` if `*this` has been moved, i.e. can no longer be used. - * - * @todo The semantic of this function is a bit confusing as a valid object returns - * `false`. It should be changed such that a valid object returns `true` and an - * invalid one returns `false`. This is the behaviour for `GHEX_C_STRUCT` and - * `GHEX_C_MANAGED_STRUCT` but not for `stream` and `cuda_event`. - */ - operator bool() const noexcept { return m_moved; } + operator bool() const noexcept { return !m_moved; } cudaEvent_t& get() noexcept { diff --git a/include/ghex/device/cuda/event_pool.hpp b/include/ghex/device/cuda/event_pool.hpp index f65a2b67..ce1b77bf 100644 --- a/include/ghex/device/cuda/event_pool.hpp +++ b/include/ghex/device/cuda/event_pool.hpp @@ -70,7 +70,7 @@ struct event_pool while (!(m_next_event < m_events.size())) { m_events.emplace_back(cuda_event()); } const std::size_t event_to_use = m_next_event; - assert(!bool(m_events[event_to_use])); + assert(bool(m_events[event_to_use])); m_next_event += 1; return m_events[event_to_use]; } diff --git a/include/ghex/device/cuda/future.hpp b/include/ghex/device/cuda/future.hpp index bdb0965f..c06b290b 100644 --- a/include/ghex/device/cuda/future.hpp +++ b/include/ghex/device/cuda/future.hpp @@ -10,9 +10,9 @@ #pragma once #include -#include -#include #ifdef GHEX_CUDACC +#include +#include #include #endif #include @@ -28,19 +28,14 @@ namespace device template struct future { - GHEX_C_MANAGED_STRUCT( - event_type, cudaEvent_t, [](auto&&... args) - { GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(std::forward(args)...)) }, - [](auto& e) { GHEX_CHECK_CUDA_RESULT_NO_THROW(cudaEventDestroy(e)) }) - - event_type m_event; + cuda_event m_event; T m_data; future(T&& data, stream& stream) - : m_event{cudaEventDisableTiming} //: m_event{cudaEventDisableTiming | cudaEventBlockingSync} + : m_event{} , m_data{std::move(data)} { - GHEX_CHECK_CUDA_RESULT(cudaEventRecord(m_event, stream)); + GHEX_CHECK_CUDA_RESULT(cudaEventRecord(m_event.get(), stream)); } future(const future&) = delete; @@ -48,11 +43,14 @@ struct future future(future&& other) = default; future& operator=(future&&) = default; - bool test() noexcept { return (m_event ? (cudaSuccess == cudaEventQuery(m_event)) : true); } + bool test() noexcept + { + return (m_event ? (cudaSuccess == cudaEventQuery(m_event.get())) : true); + } void wait() { - if (m_event) GHEX_CHECK_CUDA_RESULT(cudaEventSynchronize(m_event)); + if (m_event) GHEX_CHECK_CUDA_RESULT(cudaEventSynchronize(m_event.get())); } [[nodiscard]] T get() @@ -65,18 +63,12 @@ struct future template<> struct future { - GHEX_C_MANAGED_STRUCT( - event_type, cudaEvent_t, [](auto&&... args) - { GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(std::forward(args)...)) }, - [](auto& e) { GHEX_CHECK_CUDA_RESULT_NO_THROW(cudaEventDestroy(e)) }) - - event_type m_event; + cuda_event m_event; future(stream& stream) - : m_event{cudaEventDisableTiming} - //: m_event{cudaEventDisableTiming | cudaEventBlockingSync} + : m_event{} { - GHEX_CHECK_CUDA_RESULT(cudaEventRecord(m_event, stream)); + GHEX_CHECK_CUDA_RESULT(cudaEventRecord(m_event.get(), stream)); } future(const future&) = delete; @@ -84,11 +76,14 @@ struct future future(future&& other) = default; future& operator=(future&&) = default; - bool test() noexcept { return (m_event ? (cudaSuccess == cudaEventQuery(m_event)) : true); } + bool test() noexcept + { + return (m_event ? (cudaSuccess == cudaEventQuery(m_event.get())) : true); + } void wait() { - if (m_event) GHEX_CHECK_CUDA_RESULT(cudaEventSynchronize(m_event)); + if (m_event) GHEX_CHECK_CUDA_RESULT(cudaEventSynchronize(m_event.get())); } void get() { wait(); } diff --git a/include/ghex/device/cuda/stream.hpp b/include/ghex/device/cuda/stream.hpp index 0c93ed4b..7ade4771 100644 --- a/include/ghex/device/cuda/stream.hpp +++ b/include/ghex/device/cuda/stream.hpp @@ -39,15 +39,7 @@ struct stream if (!m_moved) { GHEX_CHECK_CUDA_RESULT_NO_THROW(cudaStreamDestroy(m_stream)) } } - /** - * @brief Returns `true` if `*this` has been moved, i.e. can no longer be used. - * - * @todo The semantic of this function is a bit confusing as a valid object returns - * `false`. It should be changed such that a valid object returns `true` and an - * invalid one returns `false`. This is the behaviour for `GHEX_C_STRUCT` and - * `GHEX_C_MANAGED_STRUCT` but not for `stream` and `cuda_event`. - */ - operator bool() const noexcept { return m_moved; } + operator bool() const noexcept { return !m_moved; } operator cudaStream_t() const noexcept {