[Mirror] Introduction of Kokkos::Event interface (based on and includes 9122)

core/src/Cuda/Kokkos_Cuda_Event.hpp

@@ -0,0 +1,113 @@
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright Contributors to the Kokkos project
+
+#ifndef KOKKOS_IMPL_PUBLIC_INCLUDE
+#include <Kokkos_Macros.hpp>
+static_assert(false,
+              "Including non-public Kokkos header files is not allowed.");
+#endif
+
+#ifndef KOKKOS_CUDA_EVENT_HPP
+#define KOKKOS_CUDA_EVENT_HPP
+
+#include <Kokkos_Macros.hpp>
+#if defined(KOKKOS_ENABLE_CUDA)
+
+#include <Kokkos_Event.hpp>
+
+#include <Cuda/Kokkos_Cuda.hpp>
+#include <Cuda/Kokkos_Cuda_Error.hpp>
+
+#include <memory>
+#include <string>
+
+namespace Kokkos {
+namespace Impl {
+
+template <>
+struct EventResource<Kokkos::Cuda> {
+  std::string label   = "unknown";
+  cudaEvent_t m_event = nullptr;
+  int m_cudaDev       = -1;
+
+  explicit EventResource(const std::string& label,
+                         const Kokkos::Cuda& exec_space)
+      : m_label(label), m_cudaDev(exec_space.cuda_device()) {
+    KOKKOS_IMPL_CUDA_SAFE_CALL(cudaSetDevice(m_cudaDev));
+    KOKKOS_IMPL_CUDA_SAFE_CALL(
+        cudaEventCreateWithFlags(&m_event, cudaEventDisableTiming));
+  }
+
+  ~EventResource() {
+    if (m_event != nullptr) {
+      KOKKOS_IMPL_CUDA_SAFE_CALL(cudaEventDestroy(m_event));
+    }
+  }
+
+  EventResource(const EventResource&)            = delete;
+  EventResource& operator=(const EventResource&) = delete;
+};
+
+}  // namespace Impl
+
+namespace Experimental {
+
+//============================================================================
+// CUDA specialization — native cudaEvent_t implementation
+//============================================================================
+
+/// CUDA specialization of Event.
+///
+/// Copyable: copies share the underlying cudaEvent_t via reference
+/// counting.  The last copy standing destroys the event.  Re-recording
+/// through any copy affects all copies (same semantics as sharing a
+/// raw cudaEvent_t).
+template <>
+class Event<Kokkos::Cuda> {
+ public:
+  Event(const std::string& label)
+      : m_handle(std::make_shared<Kokkos::Impl::EventResource<Kokkos::Cuda>>(
+            label, Kokkos::Cuda())) {}
+
+  Event(const std::string& label, const Kokkos::Cuda& exec_space)
+      : m_handle(std::make_shared<Kokkos::Impl::EventResource<Kokkos::Cuda>>(
+            label, exec_space)) {
+    record(exec_space);
+  }
+
+  void record(const Kokkos::Cuda& exec_space) {
+    KOKKOS_IMPL_CUDA_SAFE_CALL(
+        cudaEventRecord(m_handle->m_event, exec_space.cuda_stream()));
+  }
+
+  void fence() const {
+    KOKKOS_IMPL_CUDA_SAFE_CALL(cudaEventSynchronize(m_handle->m_event));
+  }
+
+  bool is_complete() const {
+    cudaError_t err = cudaEventQuery(m_handle->m_event);
+    if (err == cudaSuccess) return true;
+    if (err == cudaErrorNotReady) return false;
+    KOKKOS_IMPL_CUDA_SAFE_CALL(err);
+    return false;
+  }
+
+  const std::string& label() const { return m_handle->m_label; }
+  cudaEvent_t cuda_event() const noexcept { return m_handle->m_event; }
+
+ private:
+  std::shared_ptr<Kokkos::Impl::EventResource<Kokkos::Cuda>> m_handle;
+};
+
+/// CUDA: insert a stream wait for the recorded event (non-blocking on host).
+inline void space_depends_on(const Kokkos::Cuda& exec_space,
+                             const Event<Kokkos::Cuda>& event) {
+  KOKKOS_IMPL_CUDA_SAFE_CALL(
+      cudaStreamWaitEvent(exec_space.cuda_stream(), event.cuda_event(), 0));
+}
+
+}  // namespace Experimental
+}  // namespace Kokkos
+
+#endif  // KOKKOS_ENABLE_CUDA
+#endif  // KOKKOS_CUDA_EVENT_HPP

core/src/Kokkos_Core.cppm

@@ -52,7 +52,9 @@ export {
 #endif
   }  // namespace Experimental
   namespace Experimental {
+  using ::Kokkos::Experimental::Event;
   using ::Kokkos::Experimental::partition_space;
+  using ::Kokkos::Experimental::space_depends_on;
   }  // namespace Experimental
   using ::Kokkos::AnonymousSpace;
   using ::Kokkos::DefaultExecutionSpace;

core/src/Kokkos_Core.hpp

@@ -61,6 +61,7 @@
 //----------------------------------------------------------------------------
 //----------------------------------------------------------------------------
 
+#include <Kokkos_Event.hpp>
 #include <Kokkos_Crs.hpp>
 #include <Kokkos_WorkGraphPolicy.hpp>
 // Including this in Kokkos_Parallel_Reduce.hpp led to a circular dependency

core/src/Kokkos_Event.hpp

@@ -0,0 +1,139 @@
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright Contributors to the Kokkos project
+
+/// \file Kokkos_Event.hpp
+/// \brief Experimental event API for fine-grained stream dependencies.
+///
+/// Events capture a point in an execution space's asynchronous timeline.
+/// They enable cross-stream dependencies without a full fence, and
+/// selective host synchronisation.
+///
+/// API:
+///   - space_depends_on(exec_space, event) — GPU-side dependency
+///   (non-blocking on host)
+///   - event.fence()                       — host-side blocking synchronisation
+///   - event.is_complete()                 — non-blocking query
+///
+/// Currently only the CUDA backend provides a native implementation.
+/// For other backends the fallback records a fence on record() and
+/// space_depends_on / fence / is_complete are no-ops or trivially satisfied.
+
+#ifndef KOKKOS_EVENT_HPP
+#define KOKKOS_EVENT_HPP
+#ifndef KOKKOS_IMPL_PUBLIC_INCLUDE
+#define KOKKOS_IMPL_PUBLIC_INCLUDE
+#define KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_EVENT
+#endif
+
+#include <Kokkos_Macros.hpp>
+#include <Kokkos_Core.hpp>
+#include <thread>
+
+namespace Kokkos {
+namespace Impl {
+template <class ExecutionSpace>
+struct EventResource {
+  EventResource(const std::string& label_,
+                const Kokkos::View<int, Kokkos::SharedHostPinnedSpace>& flag_,
+                const ExecutionSpace& exec_)
+      : label(label_), flag(flag_), exec(exec_) {}
+  std::string label;
+  Kokkos::View<int, Kokkos::SharedHostPinnedSpace> flag;
+  ExecutionSpace exec;
+};
+}  // namespace Impl
+
+namespace Experimental {
+
+//============================================================================
+// Backend-agnostic Event — fallback for non-native-event backends
+//============================================================================
+
+/// Portable fallback event for backends without native event support.
+///
+/// On record(), a fence is issued so that subsequent space_depends_on()
+/// and fence() are trivially satisfied.  This preserves correctness at
+/// the cost of synchronisation -- the same trade-off existing Kokkos
+/// code already makes.
+///
+/// Backends that provide a native implementation (e.g. CUDA) specialize
+/// this template in backend-specific headers.
+
+// forward declare the class and the friend function space_depends_on
+// so that we can make namespace qualified call work
+template <Kokkos::ExecutionSpace Exec = DefaultExecutionSpace>
+struct Event;
+
+// Device-side dependency: the given execution space waits until the event
+// has occured.
+template <Kokkos::ExecutionSpace Exec>
+void space_depends_on(const Exec& exec_space, const Event<Exec>& event);
+
+template <Kokkos::ExecutionSpace Exec>
+struct Event {
+  using execution_space = Exec;
+
+ private:
+  using resource_t = Kokkos::Impl::EventResource<execution_space>;
+  using handle_t   = std::shared_ptr<resource_t>;
+  using flag_t     = Kokkos::View<int, Kokkos::SharedHostPinnedSpace>;
+
+ public:
+  Event(const std::string& label_)
+      : m_handle(std::make_shared<resource_t>(
+            label_, flag_t(std::string("Kokkos::Event::flag:" + label_)),
+            execution_space())) {
+    m_handle->flag() = 1;
+  };
+
+  Event(const std::string& label_, const execution_space& exec_space)
+      : m_handle(std::make_shared<resource_t>(
+            label_,
+            Kokkos::View<int, Kokkos::SharedHostPinnedSpace>(
+                std::string("Kokkos::Event::flag:") + label_),
+            execution_space())) {
+    record(exec_space);
+  };
+
+  // Create an event at the current spot in the execution space queue
+  void record(const execution_space& exec_space) {
+    m_handle->flag() = 0;
+    m_handle->exec   = execution_space();
+    auto flag        = m_handle->flag;
+    Kokkos::parallel_for(
+        std::string("Kokkos::Event::record:" + m_handle->label), 1,
+        KOKKOS_LAMBDA(int) { flag() = 1; });
+  }
+
+  // Wait untile the even occurs
+  void fence() const {
+    while (m_handle->flag() != 1) std::this_thread::yield();
+  }
+
+  // Check whether the even has occured
+  bool is_complete() const { return m_handle->flag() == 1; }
+
+  const std::string& label() const { return m_handle->label; }
+
+  // Enqueue a dependency on the event in an execution space instance
+  friend void space_depends_on<execution_space>(
+      const execution_space& exec_space, const Event<execution_space>& event);
+
+ private:
+  handle_t m_handle;
+};
+
+template <Kokkos::ExecutionSpace Exec>
+void space_depends_on(const Exec& exec_space, const Event<Exec>& event) {
+  // Only need to wait if its not the same execution space instance
+  // Otherwise any work issues to
+  if (exec_space != event.m_handle->exec) event.fence();
+}
+}  // namespace Experimental
+}  // namespace Kokkos
+
+#ifdef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_EVENT
+#undef KOKKOS_IMPL_PUBLIC_INCLUDE
+#undef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_EVENT
+#endif
+#endif  // KOKKOS_EVENT_HPP

core/src/decl/Kokkos_Declare_CUDA.hpp

@@ -15,6 +15,7 @@
 #include <Cuda/Kokkos_Cuda_KernelLaunch.hpp>
 #include <Cuda/Kokkos_Cuda_Instance.hpp>
 #include <Cuda/Kokkos_Cuda_View.hpp>
+#include <Cuda/Kokkos_Cuda_Event.hpp>
 #include <Cuda/Kokkos_Cuda_Team.hpp>
 #include <Cuda/Kokkos_Cuda_MDRangePolicy.hpp>
 #include <Cuda/Kokkos_Cuda_UniqueToken.hpp>

core/unit_test/CMakeLists.txt

@@ -125,6 +125,7 @@ foreach(Tag Threads;Serial;OpenMP;Cuda;HPX;OpenACC;HIP;SYCL)
       DeepCopy_Assignment
       DeepCopy_Narrowing
       DeepCopyAlignment
+      Event
       ExecSpacePartitioning
       ExecSpaceThreadSafety
       ExecutionSpace
@@ -603,6 +604,7 @@ if(Kokkos_ENABLE_CUDA)
   kokkos_add_executable_and_test(
     CoreUnitTest_CudaInterOpGraphMultiGPU SOURCES UnitTestMainInit.cpp cuda/TestCuda_InterOp_GraphMultiGPU.cpp
   )
+
 endif()
 
 if(Kokkos_ENABLE_HIP)