Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 7 additions & 11 deletions include/ghex/device/cuda/event.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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
{
Expand Down
2 changes: 1 addition & 1 deletion include/ghex/device/cuda/event_pool.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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];
}
Expand Down
41 changes: 18 additions & 23 deletions include/ghex/device/cuda/future.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -10,9 +10,9 @@
#pragma once

#include <ghex/config.hpp>
#include <ghex/util/c_managed_struct.hpp>
#include <ghex/device/cuda/error.hpp>
#ifdef GHEX_CUDACC
#include <ghex/device/cuda/event.hpp>
#include <ghex/device/cuda/error.hpp>
#include <ghex/device/cuda/runtime.hpp>
#endif
#include <memory>
Expand All @@ -28,31 +28,29 @@ namespace device
template<typename T>
struct future
{
GHEX_C_MANAGED_STRUCT(
event_type, cudaEvent_t, [](auto&&... args)
{ GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(std::forward<decltype(args)>(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;
future& operator=(const future&) = delete;
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()
Expand All @@ -65,30 +63,27 @@ struct future
template<>
struct future<void>
{
GHEX_C_MANAGED_STRUCT(
event_type, cudaEvent_t, [](auto&&... args)
{ GHEX_CHECK_CUDA_RESULT(cudaEventCreateWithFlags(std::forward<decltype(args)>(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;
future& operator=(const future&) = delete;
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(); }
Expand Down
10 changes: 1 addition & 9 deletions include/ghex/device/cuda/stream.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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
{
Expand Down
Loading