diff --git a/cudax/include/cuda/experimental/__detail/utility.cuh b/cudax/include/cuda/experimental/__detail/utility.cuh
new file mode 100644
index 0000000000..874075b107
--- /dev/null
+++ b/cudax/include/cuda/experimental/__detail/utility.cuh
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CUDAX_DETAIL_UTILITY_H
+#define __CUDAX_DETAIL_UTILITY_H
+
+namespace cuda::experimental
+{
+struct uninit_t
+{
+  explicit uninit_t() = default;
+};
+
+inline constexpr uninit_t uninit{};
+} // namespace cuda::experimental
+
+#endif // __CUDAX_DETAIL_UTILITY_H
diff --git a/cudax/include/cuda/experimental/__event/event.cuh b/cudax/include/cuda/experimental/__event/event.cuh
new file mode 100644
index 0000000000..0b6b7802b2
--- /dev/null
+++ b/cudax/include/cuda/experimental/__event/event.cuh
@@ -0,0 +1,156 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDAX_EVENT_DETAIL_H
+#define _CUDAX_EVENT_DETAIL_H
+
+#include <cuda_runtime_api.h>
+// cuda_runtime_api needs to come first
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__cuda/api_wrapper.h>
+#include <cuda/std/cstddef>
+#include <cuda/std/utility>
+
+#include <cuda/experimental/__detail/utility.cuh>
+#include <cuda/experimental/__event/event_ref.cuh>
+
+namespace cuda::experimental
+{
+class timed_event;
+
+//! @brief An owning wrapper for an untimed `cudaEvent_t`.
+class event : public event_ref
+{
+  friend class timed_event;
+
+public:
+  //! @brief Flags to use when creating the event.
+  enum class flags : unsigned int
+  {
+    none          = cudaEventDefault,
+    blocking_sync = cudaEventBlockingSync,
+    interprocess  = cudaEventInterprocess
+  };
+
+  //! @brief Construct a new `event` object with timing disabled, and record
+  //!        the event in the specified stream.
+  //!
+  //! @throws cuda_error if the event creation fails.
+  explicit event(stream_ref __stream, flags __flags = flags::none)
+      : event(static_cast<unsigned int>(__flags) | cudaEventDisableTiming)
+  {
+    record(__stream);
+  }
+
+  //! @brief Construct a new `event` object into the moved-from state.
+  //!
+  //! @post `get()` returns `cudaEvent_t()`.
+  explicit constexpr event(uninit_t) noexcept
+      : event_ref(::cudaEvent_t{})
+  {}
+
+  //! @brief Move-construct a new `event` object
+  //!
+  //! @param __other
+  //!
+  //! @post `__other` is in a moved-from state.
+  constexpr event(event&& __other) noexcept
+      : event_ref(_CUDA_VSTD::exchange(__other.__event_, {}))
+  {}
+
+  // Disallow copy construction.
+  event(const event&) = delete;
+
+  //! @brief Destroy the `event` object
+  //!
+  //! @note If the event fails to be destroyed, the error is silently ignored.
+  ~event()
+  {
+    if (__event_ != nullptr)
+    {
+      [[maybe_unused]] auto __status = ::cudaEventDestroy(__event_);
+    }
+  }
+
+  //! @brief Move-assign an `event` object
+  //!
+  //! @param __other
+  //!
+  //! @post `__other` is in a moved-from state.
+  event& operator=(event&& __other) noexcept
+  {
+    event __tmp(_CUDA_VSTD::move(__other));
+    _CUDA_VSTD::swap(__event_, __tmp.__event_);
+    return *this;
+  }
+
+  // Disallow copy assignment.
+  event& operator=(const event&) = delete;
+
+  //! @brief Construct an `event` object from a native `cudaEvent_t` handle.
+  //!
+  //! @param __evnt The native handle
+  //!
+  //! @return event The constructed `event` object
+  //!
+  //! @note The constructed `event` object takes ownership of the native handle.
+  _CCCL_NODISCARD static event from_native_handle(::cudaEvent_t __evnt) noexcept
+  {
+    return event(__evnt);
+  }
+
+  // Disallow construction from an `int`, e.g., `0`.
+  static event from_native_handle(int) = delete;
+
+  // Disallow construction from `nullptr`.
+  static event from_native_handle(_CUDA_VSTD::nullptr_t) = delete;
+
+  //! @brief Retrieve the native `cudaEvent_t` handle and give up ownership.
+  //!
+  //! @return cudaEvent_t The native handle being held by the `event` object.
+  //!
+  //! @post The event object is in a moved-from state.
+  _CCCL_NODISCARD constexpr ::cudaEvent_t release() noexcept
+  {
+    return _CUDA_VSTD::exchange(__event_, {});
+  }
+
+  _CCCL_NODISCARD_FRIEND constexpr flags operator|(flags __lhs, flags __rhs) noexcept
+  {
+    return static_cast<flags>(static_cast<unsigned int>(__lhs) | static_cast<unsigned int>(__rhs));
+  }
+
+private:
+  // Use `event::from_native_handle(e)` to construct an owning `event`
+  // object from a `cudaEvent_t` handle.
+  explicit constexpr event(::cudaEvent_t __evnt) noexcept
+      : event_ref(__evnt)
+  {}
+
+  explicit event(unsigned int __flags)
+      : event_ref(::cudaEvent_t{})
+  {
+    _CCCL_TRY_CUDA_API(
+      ::cudaEventCreateWithFlags, "Failed to create CUDA event", &__event_, static_cast<unsigned int>(__flags));
+  }
+};
+} // namespace cuda::experimental
+
+#endif // _CUDAX_EVENT_DETAIL_H
diff --git a/cudax/include/cuda/experimental/__event/event_ref.cuh b/cudax/include/cuda/experimental/__event/event_ref.cuh
new file mode 100644
index 0000000000..b795d46a77
--- /dev/null
+++ b/cudax/include/cuda/experimental/__event/event_ref.cuh
@@ -0,0 +1,134 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDAX_EVENT_REF_DETAIL_H
+#define _CUDAX_EVENT_REF_DETAIL_H
+
+#include <cuda_runtime_api.h>
+// cuda_runtime_api needs to come first
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__cuda/api_wrapper.h>
+#include <cuda/std/cassert>
+#include <cuda/std/cstddef>
+#include <cuda/std/utility>
+#include <cuda/stream_ref>
+
+namespace cuda::experimental
+{
+class event;
+class timed_event;
+
+//! @brief An non-owning wrapper for an untimed `cudaEvent_t`.
+class event_ref
+{
+private:
+  friend class event;
+  friend class timed_event;
+
+  ::cudaEvent_t __event_{};
+
+public:
+  using value_type = ::cudaEvent_t;
+
+  //! @brief Construct a new `event_ref` object from a `cudaEvent_t`
+  //!
+  //! This constructor provides an implicit conversion from `cudaEvent_t`
+  //!
+  //! @post `get() == __evnt`
+  //!
+  //! @note: It is the callers responsibilty to ensure the `event_ref` does not
+  //! outlive the event denoted by the `cudaEvent_t` handle.
+  constexpr event_ref(::cudaEvent_t __evnt) noexcept
+      : __event_(__evnt)
+  {}
+
+  /// Disallow construction from an `int`, e.g., `0`.
+  event_ref(int) = delete;
+
+  /// Disallow construction from `nullptr`.
+  event_ref(_CUDA_VSTD::nullptr_t) = delete;
+
+  //! @brief Records an event on the specified stream
+  //!
+  //! @param __stream
+  //!
+  //! @throws cuda_error if the event record fails
+  void record(stream_ref __stream) const
+  {
+    assert(__event_ != nullptr);
+    assert(__stream.get() != nullptr);
+    _CCCL_TRY_CUDA_API(::cudaEventRecord, "Failed to record CUDA event", __event_, __stream.get());
+  }
+
+  //! @brief Waits until all the work in the stream prior to the record of the
+  //!        event has completed.
+  //!
+  //! @throws cuda_error if waiting for the event fails
+  void wait() const
+  {
+    assert(__event_ != nullptr);
+    _CCCL_TRY_CUDA_API(::cudaEventSynchronize, "Failed to wait for CUDA event", __event_);
+  }
+
+  //! @brief Retrieve the native `cudaEvent_t` handle.
+  //!
+  //! @return cudaEvent_t The native handle being held by the event_ref object.
+  _CCCL_NODISCARD constexpr ::cudaEvent_t get() const noexcept
+  {
+    return __event_;
+  }
+
+  //! @brief Checks if the `event_ref` is valid
+  //!
+  //! @return true if the `event_ref` is valid, false otherwise.
+  _CCCL_NODISCARD explicit constexpr operator bool() const noexcept
+  {
+    return __event_ != nullptr;
+  }
+
+  //! @brief Compares two `event_ref`s for equality
+  //!
+  //! @note Allows comparison with `cudaEvent_t` due to implicit conversion to
+  //! `event_ref`.
+  //!
+  //! @param lhs The first `event_ref` to compare
+  //! @param rhs The second `event_ref` to compare
+  //! @return true if `lhs` and `rhs` refer to the same `cudaEvent_t` object.
+  _CCCL_NODISCARD_FRIEND constexpr bool operator==(event_ref __lhs, event_ref __rhs) noexcept
+  {
+    return __lhs.__event_ == __rhs.__event_;
+  }
+
+  //! @brief Compares two `event_ref`s for inequality
+  //!
+  //! @note Allows comparison with `cudaEvent_t` due to implicit conversion to
+  //! `event_ref`.
+  //!
+  //! @param lhs The first `event_ref` to compare
+  //! @param rhs The second `event_ref` to compare
+  //! @return true if `lhs` and `rhs` refer to different `cudaEvent_t` objects.
+  _CCCL_NODISCARD_FRIEND constexpr bool operator!=(event_ref __lhs, event_ref __rhs) noexcept
+  {
+    return __lhs.__event_ != __rhs.__event_;
+  }
+};
+} // namespace cuda::experimental
+
+#endif // _CUDAX_EVENT_REF_DETAIL_H
diff --git a/cudax/include/cuda/experimental/__event/timed_event.cuh b/cudax/include/cuda/experimental/__event/timed_event.cuh
new file mode 100644
index 0000000000..debcbcd26e
--- /dev/null
+++ b/cudax/include/cuda/experimental/__event/timed_event.cuh
@@ -0,0 +1,107 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDAX_TIMED_EVENT_DETAIL_H
+#define _CUDAX_TIMED_EVENT_DETAIL_H
+
+#include <cuda_runtime_api.h>
+// cuda_runtime_api needs to come first
+
+#include <cuda/std/detail/__config>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cuda/std/__cuda/api_wrapper.h>
+#include <cuda/std/chrono>
+#include <cuda/std/cstddef>
+
+#include <cuda/experimental/__detail/utility.cuh>
+#include <cuda/experimental/__event/event.cuh>
+
+namespace cuda::experimental
+{
+//! @brief An owning wrapper for a `cudaEvent_t` with timing enabled.
+class timed_event : public event
+{
+public:
+  //! @brief Construct a new `timed_event` object with the specified flags
+  //!        and record the event on the specified stream.
+  //!
+  //! @throws cuda_error if the event creation fails.
+  explicit timed_event(stream_ref __stream, flags __flags = flags::none)
+      : event(static_cast<unsigned int>(__flags))
+  {
+    record(__stream);
+  }
+
+  //! @brief Construct a new `timed_event` object into the moved-from state.
+  //!
+  //! @post `get()` returns `cudaEvent_t()`.
+  explicit constexpr timed_event(uninit_t) noexcept
+      : event(uninit)
+  {}
+
+  timed_event(timed_event&&) noexcept            = default;
+  timed_event(const timed_event&)                = delete;
+  timed_event& operator=(timed_event&&) noexcept = default;
+  timed_event& operator=(const timed_event&)     = delete;
+
+  //! @brief Construct a `timed_event` object from a native `cudaEvent_t` handle.
+  //!
+  //! @param __evnt The native handle
+  //!
+  //! @return timed_event The constructed `timed_event` object
+  //!
+  //! @note The constructed `timed_event` object takes ownership of the native handle.
+  _CCCL_NODISCARD static timed_event from_native_handle(::cudaEvent_t __evnt) noexcept
+  {
+    return timed_event(__evnt);
+  }
+
+  // Disallow construction from an `int`, e.g., `0`.
+  static timed_event from_native_handle(int) = delete;
+
+  // Disallow construction from `nullptr`.
+  static timed_event from_native_handle(_CUDA_VSTD::nullptr_t) = delete;
+
+  //! @brief Compute the time elapsed between two `timed_event` objects.
+  //!
+  //! @throws cuda_error if the query for the elapsed time fails.
+  //!
+  //! @param __end The `timed_event` object representing the end time.
+  //! @param __start The `timed_event` object representing the start time.
+  //!
+  //! @return cuda::std::chrono::nanoseconds The elapsed time in nanoseconds.
+  //!
+  //! @note The elapsed time has a resolution of approximately 0.5 microseconds.
+  _CCCL_NODISCARD_FRIEND _CUDA_VSTD::chrono::nanoseconds operator-(const timed_event& __end, const timed_event& __start)
+  {
+    float __ms = 0.0f;
+    _CCCL_TRY_CUDA_API(
+      ::cudaEventElapsedTime, "Failed to get CUDA event elapsed time", &__ms, __start.get(), __end.get());
+    return _CUDA_VSTD::chrono::nanoseconds(static_cast<_CUDA_VSTD::chrono::nanoseconds::rep>(__ms * 1'000'000.0));
+  }
+
+private:
+  // Use `timed_event::from_native_handle(e)` to construct an owning `timed_event`
+  // object from a `cudaEvent_t` handle.
+  explicit constexpr timed_event(::cudaEvent_t __evnt) noexcept
+      : event(__evnt)
+  {}
+};
+} // namespace cuda::experimental
+
+#endif // _CUDAX_TIMED_EVENT_DETAIL_H
diff --git a/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh b/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh
index 21e298ca40..48d4b38b1d 100644
--- a/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh
+++ b/cudax/include/cuda/experimental/__hierarchy/hierarchy_dimensions.cuh
@@ -11,10 +11,11 @@
 #ifndef _CUDAX__HIERARCHY_HIERARCHY_DIMENSIONS
 #define _CUDAX__HIERARCHY_HIERARCHY_DIMENSIONS
 
-#include <cuda/experimental/__hierarchy/level_dimensions.cuh>
+#include <cuda/std/__utility/declval.h>
 #include <cuda/std/tuple>
 
-#include "cuda/std/__utility/declval.h"
+#include <cuda/experimental/__hierarchy/level_dimensions.cuh>
+
 #include <nv/target>
 
 #if _CCCL_STD_VER >= 2017
diff --git a/cudax/include/cuda/experimental/__hierarchy/level_dimensions.cuh b/cudax/include/cuda/experimental/__hierarchy/level_dimensions.cuh
index 13d04ca24a..c685a2db6a 100644
--- a/cudax/include/cuda/experimental/__hierarchy/level_dimensions.cuh
+++ b/cudax/include/cuda/experimental/__hierarchy/level_dimensions.cuh
@@ -11,9 +11,10 @@
 #ifndef _CUDAX__HIERARCHY_LEVEL_DIMENSIONS
 #define _CUDAX__HIERARCHY_LEVEL_DIMENSIONS
 
-#include <cuda/experimental/__hierarchy/hierarchy_levels.cuh>
 #include <cuda/std/type_traits>
 
+#include <cuda/experimental/__hierarchy/hierarchy_levels.cuh>
+
 #if _CCCL_STD_VER >= 2017
 namespace cuda::experimental
 {
diff --git a/cudax/include/cuda/experimental/__launch/configuration.cuh b/cudax/include/cuda/experimental/__launch/configuration.cuh
index 99ff65f7c3..d85a6ff5b9 100644
--- a/cudax/include/cuda/experimental/__launch/configuration.cuh
+++ b/cudax/include/cuda/experimental/__launch/configuration.cuh
@@ -10,10 +10,12 @@
 
 #ifndef _CUDAX__LAUNCH_CONFIGURATION
 #define _CUDAX__LAUNCH_CONFIGURATION
-#include <cuda/experimental/hierarchy.cuh>
+
 #include <cuda/std/span>
 #include <cuda/std/tuple>
 
+#include <cuda/experimental/hierarchy.cuh>
+
 #if _CCCL_STD_VER >= 2017
 namespace cuda::experimental
 {
diff --git a/cudax/include/cuda/experimental/__launch/launch.cuh b/cudax/include/cuda/experimental/__launch/launch.cuh
index 21d9d3377b..790af2a9d5 100644
--- a/cudax/include/cuda/experimental/__launch/launch.cuh
+++ b/cudax/include/cuda/experimental/__launch/launch.cuh
@@ -12,10 +12,11 @@
 #define _CUDAX__LAUNCH_LAUNCH
 #include <cuda_runtime.h>
 
-#include <cuda/experimental/__launch/configuration.cuh>
 #include <cuda/std/__exception/cuda_error.h>
 #include <cuda/stream_ref>
 
+#include <cuda/experimental/__launch/configuration.cuh>
+
 #if _CCCL_STD_VER >= 2017
 namespace cuda::experimental
 {
diff --git a/cudax/include/cuda/experimental/event.cuh b/cudax/include/cuda/experimental/event.cuh
new file mode 100644
index 0000000000..ecb31597dc
--- /dev/null
+++ b/cudax/include/cuda/experimental/event.cuh
@@ -0,0 +1,18 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDAX_EVENT_H
+#define _CUDAX_EVENT_H
+
+#include <cuda/experimental/__event/event.cuh>
+#include <cuda/experimental/__event/event_ref.cuh>
+#include <cuda/experimental/__event/timed_event.cuh>
+
+#endif // _CUDAX_EVENT_H
diff --git a/cudax/test/CMakeLists.txt b/cudax/test/CMakeLists.txt
index 5cf3b67c84..2f337668dd 100644
--- a/cudax/test/CMakeLists.txt
+++ b/cudax/test/CMakeLists.txt
@@ -47,14 +47,19 @@ foreach(cn_target IN LISTS cudax_TARGETS)
   add_dependencies(${config_prefix}.all ${config_meta_target})
 
   # Add tests:
-  Cudax_add_catch2_test(test_target hierarchy_tests ${cn_target}
+  cudax_add_catch2_test(test_target hierarchy_tests ${cn_target}
     hierarchy/hierarchy_smoke.cu
     hierarchy/hierarchy_custom_types.cu
   )
 
-  Cudax_add_catch2_test(test_target launch_tests ${cn_target}
+  cudax_add_catch2_test(test_target launch_tests ${cn_target}
     launch/launch_smoke.cu
     launch/configuration.cu
   )
   target_compile_options(${test_target} PRIVATE $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:--extended-lambda>)
+
+  cudax_add_catch2_test(test_target event_tests ${cn_target}
+    event/event_smoke.cu
+  )
+  target_compile_options(${test_target} PRIVATE $<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:--extended-lambda>)
 endforeach()
diff --git a/cudax/test/common/utility.cuh b/cudax/test/common/utility.cuh
new file mode 100644
index 0000000000..d259b64d58
--- /dev/null
+++ b/cudax/test/common/utility.cuh
@@ -0,0 +1,118 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda_runtime_api.h>
+// cuda_runtime_api needs to come first
+
+#include <cuda/std/__cuda/api_wrapper.h>
+#include <cuda/std/utility>
+#include <cuda/stream_ref>
+
+#include <new> // IWYU pragma: keep (needed for placement new)
+
+namespace
+{
+namespace test
+{
+struct stream : cuda::stream_ref
+{
+  stream()
+      : cuda::stream_ref(::cudaStream_t{})
+  {
+    ::cudaStream_t stream{};
+    _CCCL_TRY_CUDA_API(::cudaStreamCreate, "failed to create a CUDA stream", &stream);
+    static_cast<cuda::stream_ref&>(*this) = cuda::stream_ref(stream);
+  }
+
+  cuda::stream_ref ref() const noexcept
+  {
+    return *this;
+  }
+
+  void wait() const
+  {
+    _CCCL_TRY_CUDA_API(::cudaStreamSynchronize, "failed to synchronize a CUDA stream", get());
+  }
+
+  ~stream()
+  {
+    [[maybe_unused]] auto status = ::cudaStreamDestroy(get());
+  }
+};
+
+struct _malloc_managed
+{
+private:
+  void* pv = nullptr;
+
+public:
+  explicit _malloc_managed(std::size_t size)
+  {
+    _CCCL_TRY_CUDA_API(::cudaMallocManaged, "failed to allocate managed memory", &pv, size);
+  }
+
+  ~_malloc_managed()
+  {
+    [[maybe_unused]] auto status = ::cudaFree(pv);
+  }
+
+  template <class T>
+  T* get_as() const noexcept
+  {
+    return static_cast<T*>(pv);
+  }
+};
+
+template <class T>
+struct managed
+{
+private:
+  _malloc_managed _mem;
+
+public:
+  explicit managed(T t)
+      : _mem(sizeof(T))
+  {
+    ::new (_mem.get_as<void>()) T(_CUDA_VSTD::move(t));
+  }
+
+  ~managed()
+  {
+    get()->~T();
+  }
+
+  T* get() noexcept
+  {
+    return _mem.get_as<T>();
+  }
+  const T* get() const noexcept
+  {
+    return _mem.get_as<T>();
+  }
+
+  T& operator*() noexcept
+  {
+    return *get();
+  }
+  const T& operator*() const noexcept
+  {
+    return *get();
+  }
+};
+
+/// A kernel that takes a callable object and invokes it with a set of arguments
+template <class Fn, class... Args>
+__global__ void invokernel(Fn fn, Args... args)
+{
+  fn(args...);
+}
+
+} // namespace test
+} // namespace
diff --git a/cudax/test/event/event_smoke.cu b/cudax/test/event/event_smoke.cu
new file mode 100644
index 0000000000..bc61775dab
--- /dev/null
+++ b/cudax/test/event/event_smoke.cu
@@ -0,0 +1,125 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of CUDA Experimental in CUDA C++ Core Libraries,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/experimental/event.cuh>
+
+#include "../common/utility.cuh"
+#include "../hierarchy/testing_common.cuh"
+#include <catch2/catch.hpp>
+
+namespace
+{
+namespace test
+{
+cudax::event_ref fn_takes_event_ref(cudax::event_ref ref)
+{
+  return ref;
+}
+} // namespace test
+} // namespace
+
+static_assert(!_CUDA_VSTD::is_default_constructible_v<cudax::event_ref>);
+static_assert(!_CUDA_VSTD::is_default_constructible_v<cudax::event>);
+static_assert(!_CUDA_VSTD::is_default_constructible_v<cudax::timed_event>);
+
+TEST_CASE("can construct an event_ref from a cudaEvent_t", "[event]")
+{
+  ::cudaEvent_t ev;
+  CUDAX_REQUIRE(::cudaEventCreate(&ev) == ::cudaSuccess);
+  cudax::event_ref ref(ev);
+  CUDAX_REQUIRE(ref.get() == ev);
+  CUDAX_REQUIRE(!!ref);
+  // test implicit converstion from cudaEvent_t:
+  cudax::event_ref ref2 = ::test::fn_takes_event_ref(ev);
+  CUDAX_REQUIRE(ref2.get() == ev);
+  CUDAX_REQUIRE(::cudaEventDestroy(ev) == ::cudaSuccess);
+  // test an empty event_ref:
+  cudax::event_ref ref3(::cudaEvent_t{});
+  CUDAX_REQUIRE(ref3.get() == ::cudaEvent_t{});
+  CUDAX_REQUIRE(!ref3);
+}
+
+TEST_CASE("can copy construct an event_ref and compare for equality", "[event]")
+{
+  ::cudaEvent_t ev;
+  CUDAX_REQUIRE(::cudaEventCreate(&ev) == ::cudaSuccess);
+  const cudax::event_ref ref(ev);
+  const cudax::event_ref ref2 = ref;
+  CUDAX_REQUIRE(ref2 == ref);
+  CUDAX_REQUIRE(!(ref != ref2));
+  CUDAX_REQUIRE((ref ? true : false)); // test contextual convertibility to bool
+  CUDAX_REQUIRE(!!ref);
+  CUDAX_REQUIRE(::cudaEvent_t{} != ref);
+  CUDAX_REQUIRE(::cudaEventDestroy(ev) == ::cudaSuccess);
+  // copy from empty event_ref:
+  const cudax::event_ref ref3(::cudaEvent_t{});
+  const cudax::event_ref ref4 = ref3;
+  CUDAX_REQUIRE(ref4 == ref3);
+  CUDAX_REQUIRE(!(ref3 != ref4));
+  CUDAX_REQUIRE(!ref4);
+}
+
+struct assign_42
+{
+  __device__ constexpr void operator()(int* pi) const noexcept
+  {
+    *pi = 42;
+  }
+};
+
+TEST_CASE("can use event_ref to record and wait on an event", "[event]")
+{
+  ::cudaEvent_t ev;
+  CUDAX_REQUIRE(::cudaEventCreate(&ev) == ::cudaSuccess);
+  const cudax::event_ref ref(ev);
+
+  test::managed<int> i(0);
+  test::stream stream;
+  ::test::invokernel<<<1, 1, 0, stream.get()>>>(assign_42{}, i.get());
+  ref.record(stream);
+  ref.wait();
+  CUDAX_REQUIRE(*i == 42);
+
+  stream.wait();
+  CUDAX_REQUIRE(::cudaEventDestroy(ev) == ::cudaSuccess);
+}
+
+TEST_CASE("can construct an event with a stream_ref", "[event]")
+{
+  test::stream stream;
+  cudax::event ev(stream.ref());
+  CUDAX_REQUIRE(ev.get() != ::cudaEvent_t{});
+}
+
+TEST_CASE("can wait on an event", "[event]")
+{
+  test::stream stream;
+  ::test::managed<int> i(0);
+  ::test::invokernel<<<1, 1, 0, stream.get()>>>(assign_42{}, i.get());
+  cudax::event ev(stream);
+  ev.wait();
+  CUDAX_REQUIRE(*i == 42);
+  stream.wait();
+}
+
+TEST_CASE("can take the difference of two timed_event objects", "[event]")
+{
+  test::stream stream;
+  ::test::managed<int> i(0);
+  cudax::timed_event start(stream);
+  ::test::invokernel<<<1, 1, 0, stream.get()>>>(assign_42{}, i.get());
+  cudax::timed_event end(stream);
+  end.wait();
+  CUDAX_REQUIRE(*i == 42);
+  auto elapsed = end - start;
+  CUDAX_REQUIRE(elapsed.count() >= 0);
+  STATIC_REQUIRE(_CUDA_VSTD::is_same_v<decltype(elapsed), _CUDA_VSTD::chrono::nanoseconds>);
+  stream.wait();
+}
diff --git a/cudax/test/launch/launch_smoke.cu b/cudax/test/launch/launch_smoke.cu
index d7624f3a4e..554cabd015 100644
--- a/cudax/test/launch/launch_smoke.cu
+++ b/cudax/test/launch/launch_smoke.cu
@@ -9,11 +9,8 @@
 //===----------------------------------------------------------------------===//
 #define LIBCUDACXX_ENABLE_EXCEPTIONS
 #include <cuda/atomic>
-#include <cuda/experimental/launch.cuh>
 
-#include <functional>
-#include <iostream>
-#include <type_traits>
+#include <cuda/experimental/launch.cuh>
 
 #include "../hierarchy/testing_common.cuh"