From 8c72e52f0e8c61acd423a875201d00c47b27a21f Mon Sep 17 00:00:00 2001
From: Mikael Simberg <mikael.simberg@iki.fi>
Date: Tue, 6 Feb 2024 18:54:24 +0000
Subject: [PATCH] Merge pull request #1017 from
 msimberg/fix-cublas-handles-cuda-pool

Unrevert #872
---
 libs/pika/CMakeLists.txt                      |   1 +
 libs/pika/async_cuda/CMakeLists.txt           |  31 +----
 .../include/pika/async_cuda/cuda_event.hpp    |  72 -----------
 .../include/pika/async_cuda/cuda_pool.hpp     |  77 +++++++++++-
 .../pika/async_cuda/cuda_scheduler.hpp        |   5 +-
 .../pika/async_cuda/cuda_scheduler_bulk.hpp   |   2 +-
 .../async_cuda/detail/cuda_event_callback.hpp |   2 +-
 .../pika/async_cuda/then_with_stream.hpp      |  72 ++++++-----
 .../async_cuda/src/cuda_event_callback.cpp    |   8 +-
 libs/pika/async_cuda/src/cuda_pool.cpp        | 114 +++++++++++++++++-
 libs/pika/async_cuda/src/cuda_scheduler.cpp   |  11 ++
 libs/pika/async_cuda/src/then_with_stream.cpp |  37 ------
 .../pika/async_cuda/tests/unit/CMakeLists.txt |   7 +-
 .../tests/unit/cuda_device_reset.cpp          |  81 +++++++++++++
 libs/pika/async_cuda_base/CMakeLists.txt      |  54 +++++++++
 libs/pika/async_cuda_base/README.rst          |  13 ++
 libs/pika/async_cuda_base/docs/index.rst      |  14 +++
 .../async_cuda_base/examples/CMakeLists.txt   |  18 +++
 .../async_cuda_base}/cublas_exception.hpp     |   2 +-
 .../pika/async_cuda_base}/cublas_handle.hpp   |   4 +-
 .../async_cuda_base}/cuda_device_scope.hpp    |   0
 .../pika/async_cuda_base/cuda_event.hpp       |  44 +++++++
 .../pika/async_cuda_base}/cuda_stream.hpp     |   0
 .../async_cuda_base}/cusolver_exception.hpp   |   2 +-
 .../pika/async_cuda_base}/cusolver_handle.hpp |   4 +-
 .../pika/async_cuda_base}/custom_blas_api.hpp |   0
 .../async_cuda_base}/custom_lapack_api.hpp    |   0
 .../async_cuda_base}/detail/cuda_debug.hpp    |   0
 .../src/cublas_exception.cpp                  |   4 +-
 .../src/cublas_handle.cpp                     |  10 +-
 .../src/cuda_device_scope.cpp                 |   2 +-
 libs/pika/async_cuda_base/src/cuda_event.cpp  |  73 +++++++++++
 .../src/cuda_stream.cpp                       |   4 +-
 .../src/cusolver_exception.cpp                |   4 +-
 .../src/cusolver_handle.cpp                   |  10 +-
 .../pika/async_cuda_base/tests/CMakeLists.txt |  44 +++++++
 .../tests/performance/CMakeLists.txt          |   5 +
 .../tests/regressions/CMakeLists.txt          |   5 +
 .../async_cuda_base/tests/unit/CMakeLists.txt |  34 ++++++
 .../tests/unit/cublas_handle.cu               |   0
 .../tests/unit/cuda_stream.cu                 |   0
 .../tests/unit/cusolver_handle.cu             |   0
 libs/pika/include/include/pika/cuda.hpp       |   1 +
 libs/pika/runtime/CMakeLists.txt              |   4 +
 libs/pika/runtime/src/runtime.cpp             |   8 ++
 45 files changed, 674 insertions(+), 209 deletions(-)
 delete mode 100644 libs/pika/async_cuda/include/pika/async_cuda/cuda_event.hpp
 delete mode 100644 libs/pika/async_cuda/src/then_with_stream.cpp
 create mode 100644 libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp
 create mode 100644 libs/pika/async_cuda_base/CMakeLists.txt
 create mode 100644 libs/pika/async_cuda_base/README.rst
 create mode 100644 libs/pika/async_cuda_base/docs/index.rst
 create mode 100644 libs/pika/async_cuda_base/examples/CMakeLists.txt
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cublas_exception.hpp (94%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cublas_handle.hpp (95%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cuda_device_scope.hpp (100%)
 create mode 100644 libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cuda_stream.hpp (100%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cusolver_exception.hpp (94%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/cusolver_handle.hpp (95%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/custom_blas_api.hpp (100%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/custom_lapack_api.hpp (100%)
 rename libs/pika/{async_cuda/include/pika/async_cuda => async_cuda_base/include/pika/async_cuda_base}/detail/cuda_debug.hpp (100%)
 rename libs/pika/{async_cuda => async_cuda_base}/src/cublas_exception.cpp (97%)
 rename libs/pika/{async_cuda => async_cuda_base}/src/cublas_handle.cpp (91%)
 rename libs/pika/{async_cuda => async_cuda_base}/src/cuda_device_scope.cpp (92%)
 create mode 100644 libs/pika/async_cuda_base/src/cuda_event.cpp
 rename libs/pika/{async_cuda => async_cuda_base}/src/cuda_stream.cpp (96%)
 rename libs/pika/{async_cuda => async_cuda_base}/src/cusolver_exception.cpp (97%)
 rename libs/pika/{async_cuda => async_cuda_base}/src/cusolver_handle.cpp (90%)
 create mode 100644 libs/pika/async_cuda_base/tests/CMakeLists.txt
 create mode 100644 libs/pika/async_cuda_base/tests/performance/CMakeLists.txt
 create mode 100644 libs/pika/async_cuda_base/tests/regressions/CMakeLists.txt
 create mode 100644 libs/pika/async_cuda_base/tests/unit/CMakeLists.txt
 rename libs/pika/{async_cuda => async_cuda_base}/tests/unit/cublas_handle.cu (100%)
 rename libs/pika/{async_cuda => async_cuda_base}/tests/unit/cuda_stream.cu (100%)
 rename libs/pika/{async_cuda => async_cuda_base}/tests/unit/cusolver_handle.cu (100%)

diff --git a/libs/pika/CMakeLists.txt b/libs/pika/CMakeLists.txt
index 9db14e4d9..1d16d4a76 100644
--- a/libs/pika/CMakeLists.txt
+++ b/libs/pika/CMakeLists.txt
@@ -16,6 +16,7 @@ set(_pika_modules
     assertion
     async_base
     async_cuda
+    async_cuda_base
     async_mpi
     command_line_handling
     concepts
diff --git a/libs/pika/async_cuda/CMakeLists.txt b/libs/pika/async_cuda/CMakeLists.txt
index e9c349a2d..530080a36 100644
--- a/libs/pika/async_cuda/CMakeLists.txt
+++ b/libs/pika/async_cuda/CMakeLists.txt
@@ -10,39 +10,17 @@ endif()
 
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 
-# Default location is $PIKA_ROOT/libs/async_cuda/include
 set(async_cuda_headers
-    pika/async_cuda/cublas_exception.hpp
-    pika/async_cuda/cublas_handle.hpp
-    pika/async_cuda/cuda_device_scope.hpp
-    pika/async_cuda/cuda_event.hpp
     pika/async_cuda/cuda_polling_helper.hpp
     pika/async_cuda/cuda_pool.hpp
     pika/async_cuda/cuda_scheduler.hpp
     pika/async_cuda/cuda_scheduler_bulk.hpp
-    pika/async_cuda/cuda_stream.hpp
-    pika/async_cuda/cusolver_exception.hpp
-    pika/async_cuda/cusolver_handle.hpp
-    pika/async_cuda/custom_blas_api.hpp
-    pika/async_cuda/custom_lapack_api.hpp
-    pika/async_cuda/detail/cuda_debug.hpp
     pika/async_cuda/detail/cuda_event_callback.hpp
     pika/async_cuda/then_on_host.hpp
     pika/async_cuda/then_with_stream.hpp
 )
 
-set(async_cuda_sources
-    cublas_exception.cpp
-    cublas_handle.cpp
-    cuda_device_scope.cpp
-    cuda_event_callback.cpp
-    cuda_pool.cpp
-    cuda_scheduler.cpp
-    cuda_stream.cpp
-    cusolver_exception.cpp
-    cusolver_handle.cpp
-    then_with_stream.cpp
-)
+set(async_cuda_sources cuda_event_callback.cpp cuda_pool.cpp cuda_scheduler.cpp)
 
 if(PIKA_WITH_HIP)
   set(async_cuda_extra_deps roc::rocblas roc::rocsolver)
@@ -59,17 +37,18 @@ pika_add_module(
   SOURCES ${async_cuda_sources}
   HEADERS ${async_cuda_headers}
   MODULE_DEPENDENCIES
-    pika_allocator_support
     pika_assertion
     pika_async_base
+    pika_async_cuda_base
     pika_concurrency
     pika_config
-    pika_debugging
+    pika_coroutines
     pika_errors
+    pika_execution
     pika_execution_base
-    pika_memory
     pika_runtime
     pika_threading_base
+    pika_topology
   DEPENDENCIES ${async_cuda_extra_deps}
   CMAKE_SUBDIRS examples tests
 )
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_event.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_event.hpp
deleted file mode 100644
index 2477f8d51..000000000
--- a/libs/pika/async_cuda/include/pika/async_cuda/cuda_event.hpp
+++ /dev/null
@@ -1,72 +0,0 @@
-//  Copyright (c) 2020 John Biddiscombe
-//  Copyright (c) 2020 Teodor Nikolov
-//
-//  SPDX-License-Identifier: BSL-1.0
-//  Distributed under the Boost Software License, Version 1.0. (See accompanying
-//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
-
-#pragma once
-
-#include <boost/lockfree/stack.hpp>
-#include <whip.hpp>
-
-namespace pika::cuda::experimental::detail {
-
-    // a pool of cudaEvent_t objects.
-    // Since allocation of a cuda event passes into the cuda runtime
-    // it might be an expensive operation, so we pre-allocate a pool
-    // of them at startup.
-    struct cuda_event_pool
-    {
-        static constexpr int initial_events_in_pool = 128;
-
-        static cuda_event_pool& get_event_pool()
-        {
-            static cuda_event_pool event_pool_;
-            return event_pool_;
-        }
-
-        // create a bunch of events on initialization
-        cuda_event_pool()
-          : free_list_(initial_events_in_pool)
-        {
-            for (int i = 0; i < initial_events_in_pool; ++i) { add_event_to_pool(); }
-        }
-
-        // on destruction, all objects in stack will be freed
-        ~cuda_event_pool()
-        {
-            whip::event_t event;
-            bool ok = true;
-            while (ok)
-            {
-                ok = free_list_.pop(event);
-                if (ok) whip::event_destroy(event);
-            }
-        }
-
-        inline bool pop(whip::event_t& event)
-        {
-            // pop an event off the pool, if that fails, create a new one
-            while (!free_list_.pop(event)) { add_event_to_pool(); }
-            return true;
-        }
-
-        inline bool push(whip::event_t event) { return free_list_.push(event); }
-
-    private:
-        void add_event_to_pool()
-        {
-            whip::event_t event;
-            // Create an cuda_event to query a CUDA/CUBLAS kernel for completion.
-            // Timing is disabled for performance. [1]
-            //
-            // [1]: CUDA Runtime API, section 5.5 cuda_event Management
-            whip::event_create_with_flags(&event, whip::event_disable_timing);
-            free_list_.push(event);
-        }
-
-        // pool is dynamically sized and can grow if needed
-        boost::lockfree::stack<whip::event_t, boost::lockfree::fixed_sized<false>> free_list_;
-    };
-}    // namespace pika::cuda::experimental::detail
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_pool.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_pool.hpp
index fe8459d2d..86e1cabd8 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cuda_pool.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cuda_pool.hpp
@@ -7,7 +7,9 @@
 #pragma once
 
 #include <pika/assert.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda_base/cublas_handle.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
+#include <pika/async_cuda_base/cusolver_handle.hpp>
 #include <pika/concurrency/cache_line_data.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 
@@ -16,10 +18,43 @@
 #include <atomic>
 #include <cstddef>
 #include <memory>
+#include <mutex>
 #include <string>
 #include <vector>
 
 namespace pika::cuda::experimental {
+    class locked_cublas_handle
+    {
+        cublas_handle& handle;
+        std::unique_lock<std::mutex> mutex;
+
+    public:
+        PIKA_EXPORT locked_cublas_handle(
+            cublas_handle& handle, std::unique_lock<std::mutex>&& mutex);
+        locked_cublas_handle(locked_cublas_handle&&) = delete;
+        locked_cublas_handle(locked_cublas_handle const&) = delete;
+        locked_cublas_handle& operator=(locked_cublas_handle&&) = delete;
+        locked_cublas_handle& operator=(locked_cublas_handle const&) = delete;
+
+        PIKA_EXPORT cublas_handle const& get() noexcept;
+    };
+
+    class locked_cusolver_handle
+    {
+        cusolver_handle& handle;
+        std::unique_lock<std::mutex> mutex;
+
+    public:
+        PIKA_EXPORT locked_cusolver_handle(
+            cusolver_handle& handle, std::unique_lock<std::mutex>&& mutex);
+        locked_cusolver_handle(locked_cusolver_handle&&) = delete;
+        locked_cusolver_handle(locked_cusolver_handle const&) = delete;
+        locked_cusolver_handle& operator=(locked_cusolver_handle&&) = delete;
+        locked_cusolver_handle& operator=(locked_cusolver_handle const&) = delete;
+
+        PIKA_EXPORT cusolver_handle const& get() noexcept;
+    };
+
     /// A pool of CUDA streams, used for scheduling work on a CUDA device.
     ///
     /// The pool initializes a set of CUDA (thread-local) streams on
@@ -47,11 +82,48 @@ namespace pika::cuda::experimental {
             PIKA_EXPORT cuda_stream const& get_next_stream();
         };
 
+        struct cublas_handles_holder
+        {
+            std::size_t const concurrency;
+            std::vector<cublas_handle> unsynchronized_handles;
+            std::atomic<std::size_t> synchronized_handle_index;
+            std::vector<cublas_handle> synchronized_handles;
+            std::vector<std::mutex> handle_mutexes;
+
+            PIKA_EXPORT cublas_handles_holder();
+            cublas_handles_holder(cublas_handles_holder&&) = delete;
+            cublas_handles_holder(cublas_handles_holder const&) = delete;
+            cublas_handles_holder& operator=(cublas_handles_holder&&) = delete;
+            cublas_handles_holder& operator=(cublas_handles_holder const&) = delete;
+
+            PIKA_EXPORT locked_cublas_handle get_locked_handle(
+                cuda_stream const& stream, cublasPointerMode_t pointer_mode);
+        };
+
+        struct cusolver_handles_holder
+        {
+            std::size_t const concurrency;
+            std::vector<cusolver_handle> unsynchronized_handles;
+            std::atomic<std::size_t> synchronized_handle_index;
+            std::vector<cusolver_handle> synchronized_handles;
+            std::vector<std::mutex> handle_mutexes;
+
+            PIKA_EXPORT cusolver_handles_holder();
+            cusolver_handles_holder(cusolver_handles_holder&&) = delete;
+            cusolver_handles_holder(cusolver_handles_holder const&) = delete;
+            cusolver_handles_holder& operator=(cusolver_handles_holder&&) = delete;
+            cusolver_handles_holder& operator=(cusolver_handles_holder const&) = delete;
+
+            PIKA_EXPORT locked_cusolver_handle get_locked_handle(cuda_stream const& stream);
+        };
+
         struct pool_data
         {
             int device;
             streams_holder normal_priority_streams;
             streams_holder high_priority_streams;
+            cublas_handles_holder cublas_handles;
+            cusolver_handles_holder cusolver_handles;
 
             PIKA_EXPORT pool_data(int device, std::size_t num_normal_priority_streams_per_thread,
                 std::size_t num_high_priority_streams_per_thread, unsigned int flags);
@@ -80,6 +152,9 @@ namespace pika::cuda::experimental {
         PIKA_EXPORT explicit operator bool() noexcept;
         PIKA_EXPORT cuda_stream const& get_next_stream(
             pika::execution::thread_priority priority = pika::execution::thread_priority::normal);
+        PIKA_EXPORT locked_cublas_handle get_cublas_handle(
+            cuda_stream const& stream, cublasPointerMode_t pointer_mode);
+        PIKA_EXPORT locked_cusolver_handle get_cusolver_handle(cuda_stream const& stream);
 
         /// \cond NOINTERNAL
         friend bool operator==(cuda_pool const& lhs, cuda_pool const& rhs) noexcept
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler.hpp
index 33fc052a4..4d0c787cc 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler.hpp
@@ -9,7 +9,7 @@
 #include <pika/assert.hpp>
 #include <pika/async_base/scheduling_properties.hpp>
 #include <pika/async_cuda/cuda_pool.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
 #include <pika/concepts/concepts.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 #include <pika/execution/algorithms/execute.hpp>
@@ -40,6 +40,9 @@ namespace pika::cuda::experimental {
 
         PIKA_EXPORT cuda_pool const& get_pool() const noexcept;
         PIKA_EXPORT cuda_stream const& get_next_stream();
+        PIKA_EXPORT locked_cublas_handle get_cublas_handle(
+            cuda_stream const& stream, cublasPointerMode_t pointer_mode);
+        PIKA_EXPORT locked_cusolver_handle get_cusolver_handle(cuda_stream const& stream);
 
         /// \cond NOINTERNAL
         friend bool operator==(cuda_scheduler const& lhs, cuda_scheduler const& rhs)
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler_bulk.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler_bulk.hpp
index 80c281282..20cc78950 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler_bulk.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cuda_scheduler_bulk.hpp
@@ -8,8 +8,8 @@
 
 #include <pika/assert.hpp>
 #include <pika/async_cuda/cuda_pool.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
 #include <pika/async_cuda/then_with_stream.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
 #include <pika/concepts/concepts.hpp>
 #include <pika/execution/algorithms/bulk.hpp>
 #include <pika/execution/algorithms/execute.hpp>
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_event_callback.hpp b/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_event_callback.hpp
index 7719feace..d7fb0f6e4 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_event_callback.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_event_callback.hpp
@@ -15,7 +15,7 @@
 #pragma once
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
 #include <pika/functional/unique_function.hpp>
 #include <pika/threading_base/thread_pool_base.hpp>
 
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/then_with_stream.hpp b/libs/pika/async_cuda/include/pika/async_cuda/then_with_stream.hpp
index 4633a5bb9..58329a67f 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/then_with_stream.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/then_with_stream.hpp
@@ -8,12 +8,12 @@
 
 #include <pika/config.hpp>
 #include <pika/assert.hpp>
-#include <pika/async_cuda/cublas_handle.hpp>
 #include <pika/async_cuda/cuda_scheduler.hpp>
-#include <pika/async_cuda/cusolver_handle.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
-#include <pika/async_cuda/custom_lapack_api.hpp>
 #include <pika/async_cuda/detail/cuda_event_callback.hpp>
+#include <pika/async_cuda_base/cublas_handle.hpp>
+#include <pika/async_cuda_base/cusolver_handle.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
+#include <pika/async_cuda_base/custom_lapack_api.hpp>
 #include <pika/datastructures/variant.hpp>
 #include <pika/execution/algorithms/detail/helpers.hpp>
 #include <pika/execution/algorithms/detail/partial_algorithm.hpp>
@@ -35,29 +35,23 @@
 #include <utility>
 
 namespace pika::cuda::experimental::then_with_stream_detail {
-    PIKA_EXPORT pika::cuda::experimental::cublas_handle const& get_thread_local_cublas_handle(
-        cuda_stream const& stream, cublasPointerMode_t pointer_mode);
-
     template <typename F, typename... Ts>
-    auto invoke_with_thread_local_cublas_handle(
-        cuda_stream const& stream, cublasPointerMode_t pointer_mode, F&& f, Ts&&... ts)
+    auto invoke_with_thread_local_cublas_handle(cuda_scheduler& sched, cuda_stream const& stream,
+        cublasPointerMode_t pointer_mode, F&& f, Ts&&... ts)
         -> decltype(PIKA_INVOKE(
             PIKA_FORWARD(F, f), std::declval<cublasHandle_t>(), PIKA_FORWARD(Ts, ts)...))
     {
-        return PIKA_INVOKE(PIKA_FORWARD(F, f),
-            get_thread_local_cublas_handle(stream, pointer_mode).get(), PIKA_FORWARD(Ts, ts)...);
+        auto locked_handle = sched.get_cublas_handle(stream, pointer_mode);
+        return PIKA_INVOKE(PIKA_FORWARD(F, f), locked_handle.get().get(), PIKA_FORWARD(Ts, ts)...);
     }
 
-    PIKA_EXPORT pika::cuda::experimental::cusolver_handle const& get_thread_local_cusolver_handle(
-        cuda_stream const& stream);
-
     template <typename F, typename... Ts>
-    auto invoke_with_thread_local_cusolver_handle(cuda_stream const& stream, F&& f, Ts&&... ts)
-        -> decltype(PIKA_INVOKE(
-            PIKA_FORWARD(F, f), std::declval<cusolverDnHandle_t>(), PIKA_FORWARD(Ts, ts)...))
+    auto invoke_with_thread_local_cusolver_handle(cuda_scheduler& sched, cuda_stream const& stream,
+        F&& f, Ts&&... ts) -> decltype(PIKA_INVOKE(PIKA_FORWARD(F, f),
+        std::declval<cusolverDnHandle_t>(), PIKA_FORWARD(Ts, ts)...))
     {
-        return PIKA_INVOKE(PIKA_FORWARD(F, f), get_thread_local_cusolver_handle(stream).get(),
-            PIKA_FORWARD(Ts, ts)...);
+        auto locked_handle = sched.get_cusolver_handle(stream);
+        return PIKA_INVOKE(PIKA_FORWARD(F, f), locked_handle.get().get(), PIKA_FORWARD(Ts, ts)...);
     }
 
     template <typename R, typename... Ts>
@@ -157,11 +151,11 @@ namespace pika::cuda::experimental::then_with_stream_detail {
 
 #if defined(PIKA_HAVE_STDEXEC)
         template <typename... Ts>
-        requires std::is_invocable_v<F, cuda_stream const&,
+        requires std::is_invocable_v<F, cuda_scheduler&, cuda_stream const&,
             std::add_lvalue_reference_t<std::decay_t<Ts>>...>
         using invoke_result_helper =
             pika::execution::experimental::completion_signatures<pika::execution::experimental::
-                    detail::result_type_signature_helper_t<std::invoke_result_t<F,
+                    detail::result_type_signature_helper_t<std::invoke_result_t<F, cuda_scheduler&,
                         cuda_stream const&, std::add_lvalue_reference_t<std::decay_t<Ts>>...>>>;
 
         using completion_signatures =
@@ -177,7 +171,7 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         template <template <typename...> class Tuple, typename... Ts>
         struct invoke_result_helper<Tuple<Ts...>>
         {
-            using result_type = std::invoke_result_t<F, cuda_stream const&,
+            using result_type = std::invoke_result_t<F, cuda_scheduler&, cuda_stream const&,
                 std::add_lvalue_reference_t<std::decay_t<Ts>>...>;
             using type =
                 std::conditional_t<std::is_void_v<result_type>, Tuple<>, Tuple<result_type>>;
@@ -256,7 +250,8 @@ namespace pika::cuda::experimental::then_with_stream_detail {
 
                 template <typename... Ts>
                 auto set_value(Ts&&... ts) noexcept
-                    -> decltype(PIKA_INVOKE(PIKA_MOVE(f), stream.value(), ts...), void())
+                    -> decltype(PIKA_INVOKE(PIKA_MOVE(f), op_state.sched, stream.value(), ts...),
+                        void())
                 {
                     pika::detail::try_catch_exception_ptr(
                         [&]() mutable {
@@ -287,16 +282,16 @@ namespace pika::cuda::experimental::then_with_stream_detail {
                                 }
                             }
 
-                            using invoke_result_type =
-                                std::decay_t<std::invoke_result_t<F, cuda_stream const&,
+                            using invoke_result_type = std::decay_t<
+                                std::invoke_result_t<F, cuda_scheduler&, cuda_stream const&,
                                     std::add_lvalue_reference_t<std::decay_t<Ts>>...>>;
                             constexpr bool is_void_result = std::is_void_v<invoke_result_type>;
                             if constexpr (is_void_result)
                             {
                                 std::apply(
                                     [&](auto&... ts) mutable {
-                                        PIKA_INVOKE(
-                                            PIKA_MOVE(op_state.f), op_state.stream.value(), ts...);
+                                        PIKA_INVOKE(PIKA_MOVE(op_state.f), op_state.sched,
+                                            op_state.stream.value(), ts...);
                                     },
                                     t);
 
@@ -336,7 +331,7 @@ namespace pika::cuda::experimental::then_with_stream_detail {
                                 std::apply(
                                     [&](auto&... ts) mutable {
                                         op_state.result.template emplace<invoke_result_type>(
-                                            PIKA_INVOKE(PIKA_MOVE(op_state.f),
+                                            PIKA_INVOKE(PIKA_MOVE(op_state.f), op_state.sched,
                                                 op_state.stream.value(), ts...));
                                     },
                                     t);
@@ -548,7 +543,7 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         std::decay_t<F> f;
 
         template <typename... Ts>
-        auto operator()(cuda_stream const& stream, Ts&&... ts)
+        auto operator()(cuda_scheduler&, cuda_stream const& stream, Ts&&... ts)
         // nvcc does not compile this correctly with noexcept(...)
 #if defined(PIKA_CLANG_VERSION)
             noexcept(noexcept(PIKA_INVOKE(f, PIKA_FORWARD(Ts, ts)..., stream.get())))
@@ -568,17 +563,17 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         cublasPointerMode_t pointer_mode;
 
         template <typename... Ts>
-        auto operator()(cuda_stream const& stream, Ts&&... ts)
+        auto operator()(cuda_scheduler& sched, cuda_stream const& stream, Ts&&... ts)
         // nvcc does not compile this correctly with noexcept(...)
 #if defined(PIKA_CLANG_VERSION)
             noexcept(noexcept(invoke_with_thread_local_cublas_handle(
-                stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...)))
+                sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...)))
 #endif
                 -> decltype(invoke_with_thread_local_cublas_handle(
-                    stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...))
+                    sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...))
         {
             return invoke_with_thread_local_cublas_handle(
-                stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...);
+                sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...);
         }
     };
 
@@ -590,16 +585,17 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         std::decay_t<F> f;
 
         template <typename... Ts>
-        auto operator()(cuda_stream const& stream, Ts&&... ts)
+        auto operator()(cuda_scheduler& sched, cuda_stream const& stream, Ts&&... ts)
         // nvcc does not compile this correctly with noexcept(...)
 #if defined(PIKA_CLANG_VERSION)
-            noexcept(noexcept(
-                invoke_with_thread_local_cusolver_handle(stream, f, PIKA_FORWARD(Ts, ts)...)))
+            noexcept(noexcept(invoke_with_thread_local_cusolver_handle(
+                sched, stream, f, PIKA_FORWARD(Ts, ts)...)))
 #endif
                 -> decltype(invoke_with_thread_local_cusolver_handle(
-                    stream, f, PIKA_FORWARD(Ts, ts)...))
+                    sched, stream, f, PIKA_FORWARD(Ts, ts)...))
         {
-            return invoke_with_thread_local_cusolver_handle(stream, f, PIKA_FORWARD(Ts, ts)...);
+            return invoke_with_thread_local_cusolver_handle(
+                sched, stream, f, PIKA_FORWARD(Ts, ts)...);
         }
     };
 }    // namespace pika::cuda::experimental::then_with_stream_detail
diff --git a/libs/pika/async_cuda/src/cuda_event_callback.cpp b/libs/pika/async_cuda/src/cuda_event_callback.cpp
index 2d2e29974..7d1c4ae78 100644
--- a/libs/pika/async_cuda/src/cuda_event_callback.cpp
+++ b/libs/pika/async_cuda/src/cuda_event_callback.cpp
@@ -9,9 +9,9 @@
 
 #include <pika/config.hpp>
 #include <pika/assert.hpp>
-#include <pika/async_cuda/cuda_event.hpp>
-#include <pika/async_cuda/detail/cuda_debug.hpp>
 #include <pika/async_cuda/detail/cuda_event_callback.hpp>
+#include <pika/async_cuda_base/cuda_event.hpp>
+#include <pika/async_cuda_base/detail/cuda_debug.hpp>
 #include <pika/concurrency/concurrentqueue.hpp>
 #include <pika/concurrency/spinlock.hpp>
 #include <pika/datastructures/detail/small_vector.hpp>
@@ -309,6 +309,10 @@ namespace pika::cuda::experimental::detail {
 #if defined(PIKA_DEBUG)
         ++get_register_polling_count();
 #endif
+
+        // pre-create CUDA/HIP events
+        cuda_event_pool::get_event_pool().grow(cuda_event_pool::initial_events_in_pool);
+
         PIKA_DETAIL_DP(cud_debug<2>, debug(str<>("enable polling"), pool.get_pool_name()));
         auto* sched = pool.get_scheduler();
         sched->set_cuda_polling_functions(&pika::cuda::experimental::detail::poll, &get_work_count);
diff --git a/libs/pika/async_cuda/src/cuda_pool.cpp b/libs/pika/async_cuda/src/cuda_pool.cpp
index 1592130da..953baf027 100644
--- a/libs/pika/async_cuda/src/cuda_pool.cpp
+++ b/libs/pika/async_cuda/src/cuda_pool.cpp
@@ -6,7 +6,7 @@
 
 #include <pika/assert.hpp>
 #include <pika/async_cuda/cuda_pool.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
 #include <pika/concurrency/cache_line_data.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 #include <pika/runtime/runtime_fwd.hpp>
@@ -15,6 +15,7 @@
 
 #include <cstddef>
 #include <memory>
+#include <utility>
 #include <vector>
 
 namespace pika::cuda::experimental {
@@ -40,6 +41,102 @@ namespace pika::cuda::experimental {
         return streams[global_stream_index];
     }
 
+    cuda_pool::cublas_handles_holder::cublas_handles_holder()
+      : concurrency(pika::get_runtime_ptr() ? pika::get_num_worker_threads() :
+                                              pika::threads::detail::hardware_concurrency())
+      , unsynchronized_handles(concurrency, cublas_handle{})
+      , synchronized_handle_index{0}
+      , synchronized_handles(concurrency, cublas_handle{})
+      , handle_mutexes(concurrency)
+    {
+    }
+
+    locked_cublas_handle::locked_cublas_handle(
+        cublas_handle& handle, std::unique_lock<std::mutex>&& mutex)
+      : handle(handle)
+      , mutex(std::move(mutex))
+    {
+    }
+
+    cublas_handle const& locked_cublas_handle::get() noexcept { return handle; }
+
+    locked_cublas_handle cuda_pool::cublas_handles_holder::get_locked_handle(
+        cuda_stream const& stream, cublasPointerMode_t pointer_mode)
+    {
+        auto const t = pika::threads::detail::get_global_thread_num_tss();
+
+        // If we are on a pika runtime worker thread we use one of the unsynchronized handles since
+        // this is the only thread that will access this handle
+        if (t < unsynchronized_handles.size())
+        {
+            auto& handle = unsynchronized_handles[t];
+            handle.set_stream(stream);
+            handle.set_pointer_mode(pointer_mode);
+
+            return locked_cublas_handle(handle, std::unique_lock<std::mutex>{});
+        }
+        // We use synchronized (locked) handles in a round-robin fashion for all other threads
+        else
+        {
+            auto const t = synchronized_handle_index++ % synchronized_handles.size();
+
+            std::unique_lock lock{handle_mutexes[t]};
+
+            auto& handle = synchronized_handles[t];
+            handle.set_stream(stream);
+            handle.set_pointer_mode(pointer_mode);
+
+            return locked_cublas_handle(handle, std::move(lock));
+        }
+    }
+
+    cuda_pool::cusolver_handles_holder::cusolver_handles_holder()
+      : concurrency(pika::get_runtime_ptr() ? pika::get_num_worker_threads() :
+                                              pika::threads::detail::hardware_concurrency())
+      , unsynchronized_handles(concurrency, cusolver_handle{})
+      , synchronized_handle_index{0}
+      , synchronized_handles(concurrency, cusolver_handle{})
+      , handle_mutexes(concurrency)
+    {
+    }
+
+    locked_cusolver_handle::locked_cusolver_handle(
+        cusolver_handle& handle, std::unique_lock<std::mutex>&& mutex)
+      : handle(handle)
+      , mutex(std::move(mutex))
+    {
+    }
+
+    cusolver_handle const& locked_cusolver_handle::get() noexcept { return handle; }
+
+    locked_cusolver_handle cuda_pool::cusolver_handles_holder::get_locked_handle(
+        cuda_stream const& stream)
+    {
+        auto const t = pika::threads::detail::get_global_thread_num_tss();
+
+        // If we are on a pika runtime worker thread we use one of the unsynchronized handles since
+        // this is the only thread that will access this handle
+        if (t < unsynchronized_handles.size())
+        {
+            auto& handle = unsynchronized_handles[t];
+            handle.set_stream(stream);
+
+            return locked_cusolver_handle(handle, std::unique_lock<std::mutex>{});
+        }
+        // We use synchronized (locked) handles in a round-robin fashion for all other threads
+        else
+        {
+            auto const t = synchronized_handle_index++ % synchronized_handles.size();
+
+            std::unique_lock lock{handle_mutexes[t]};
+
+            auto& handle = synchronized_handles[t];
+            handle.set_stream(stream);
+
+            return {handle, std::move(lock)};
+        }
+    }
+
     cuda_pool::pool_data::pool_data(int device, std::size_t num_normal_priority_streams_per_thread,
         std::size_t num_high_priority_streams_per_thread, unsigned int flags)
       : device(device)
@@ -47,6 +144,8 @@ namespace pika::cuda::experimental {
             pika::execution::thread_priority::normal, flags)
       , high_priority_streams(device, num_high_priority_streams_per_thread,
             pika::execution::thread_priority::high, flags)
+      , cublas_handles()
+      , cusolver_handles()
     {
     }
 
@@ -71,4 +170,17 @@ namespace pika::cuda::experimental {
         }
         else { return data->high_priority_streams.get_next_stream(); }
     }
+
+    locked_cublas_handle cuda_pool::get_cublas_handle(
+        cuda_stream const& stream, cublasPointerMode_t pointer_mode)
+    {
+        PIKA_ASSERT(data);
+        return data->cublas_handles.get_locked_handle(stream, pointer_mode);
+    }
+
+    locked_cusolver_handle cuda_pool::get_cusolver_handle(cuda_stream const& stream)
+    {
+        PIKA_ASSERT(data);
+        return data->cusolver_handles.get_locked_handle(stream);
+    }
 }    // namespace pika::cuda::experimental
diff --git a/libs/pika/async_cuda/src/cuda_scheduler.cpp b/libs/pika/async_cuda/src/cuda_scheduler.cpp
index c96b1bc00..1f371a21c 100644
--- a/libs/pika/async_cuda/src/cuda_scheduler.cpp
+++ b/libs/pika/async_cuda/src/cuda_scheduler.cpp
@@ -18,6 +18,17 @@ namespace pika::cuda::experimental {
 
     cuda_stream const& cuda_scheduler::get_next_stream() { return pool.get_next_stream(priority); }
 
+    locked_cublas_handle cuda_scheduler::get_cublas_handle(
+        cuda_stream const& stream, cublasPointerMode_t pointer_mode)
+    {
+        return pool.get_cublas_handle(stream, pointer_mode);
+    }
+
+    locked_cusolver_handle cuda_scheduler::get_cusolver_handle(cuda_stream const& stream)
+    {
+        return pool.get_cusolver_handle(stream);
+    }
+
     namespace detail {
         cuda_scheduler_sender::cuda_scheduler_sender(cuda_scheduler scheduler)
           : scheduler(PIKA_MOVE(scheduler))
diff --git a/libs/pika/async_cuda/src/then_with_stream.cpp b/libs/pika/async_cuda/src/then_with_stream.cpp
deleted file mode 100644
index 37913126b..000000000
--- a/libs/pika/async_cuda/src/then_with_stream.cpp
+++ /dev/null
@@ -1,37 +0,0 @@
-//  Copyright (c) 2022 ETH Zurich
-//
-//  SPDX-License-Identifier: BSL-1.0
-//  Distributed under the Boost Software License, Version 1.0. (See accompanying
-//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
-
-#pragma once
-
-#include <pika/config.hpp>
-#include <pika/async_cuda/cublas_handle.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
-#include <pika/async_cuda/cusolver_handle.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
-#include <pika/async_cuda/then_with_stream.hpp>
-
-namespace pika::cuda::experimental::then_with_stream_detail {
-    pika::cuda::experimental::cublas_handle const& get_thread_local_cublas_handle(
-        cuda_stream const& stream, cublasPointerMode_t pointer_mode)
-    {
-        static thread_local pika::cuda::experimental::cublas_handle handle{stream};
-
-        handle.set_stream(stream);
-        handle.set_pointer_mode(pointer_mode);
-
-        return handle;
-    }
-
-    pika::cuda::experimental::cusolver_handle const& get_thread_local_cusolver_handle(
-        cuda_stream const& stream)
-    {
-        static thread_local pika::cuda::experimental::cusolver_handle handle{stream};
-
-        handle.set_stream(stream);
-
-        return handle;
-    }
-}    // namespace pika::cuda::experimental::then_with_stream_detail
diff --git a/libs/pika/async_cuda/tests/unit/CMakeLists.txt b/libs/pika/async_cuda/tests/unit/CMakeLists.txt
index 60b43f322..9ad34cf4a 100644
--- a/libs/pika/async_cuda/tests/unit/CMakeLists.txt
+++ b/libs/pika/async_cuda/tests/unit/CMakeLists.txt
@@ -5,14 +5,12 @@
 # file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 set(tests
-    cublas_handle
     cublas_matmul
     cuda_bulk
+    cuda_device_reset
     cuda_pool
     cuda_scheduler
     cuda_sender
-    cuda_stream
-    cusolver_handle
     then_with_stream
 )
 
@@ -21,13 +19,10 @@ set(cuda_sender_PARAMETERS THREADS 4)
 set(then_with_stream_PARAMETERS THREADS 4)
 set(then_with_stream_DEPENDENCIES pika_execution_test_utilities)
 
-set(cublas_handle_GPU ON)
 set(cuda_bulk_GPU ON)
 set(cuda_pool_GPU ON)
 set(cuda_scheduler_GPU ON)
 set(cuda_sender_GPU ON)
-set(cuda_stream_GPU ON)
-set(cusolver_handle_GPU ON)
 set(then_with_stream_GPU ON)
 
 foreach(test ${tests})
diff --git a/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp b/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp
new file mode 100644
index 000000000..10025143c
--- /dev/null
+++ b/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp
@@ -0,0 +1,81 @@
+//  Copyright (c) 2023 ETH Zurich
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+#include <pika/cuda.hpp>
+#include <pika/execution.hpp>
+#include <pika/init.hpp>
+#include <pika/testing.hpp>
+
+#include <whip.hpp>
+
+#include <atomic>
+#include <cstddef>
+#include <cstdlib>
+#include <exception>
+#include <functional>
+#include <string>
+#include <utility>
+
+namespace cu = pika::cuda::experimental;
+namespace ex = pika::execution::experimental;
+namespace tt = pika::this_thread::experimental;
+
+int main(int argc, char* argv[])
+{
+    pika::start(argc, argv);
+
+    {
+        std::size_t const n = 1000;
+
+        cu::cuda_pool pool{};
+
+        // Explicitly do not use the RAII version here to test that pika::stop also clears events
+        cu::detail::register_polling(pika::resource::get_thread_pool("default"));
+
+        // cuBLAS handle on a non-pika thread
+        for (std::size_t i = 0; i < n; ++i)
+        {
+            tt::sync_wait(ex::schedule(cu::cuda_scheduler{pool}) |
+                cu::then_with_cublas([](cublasHandle_t) {}, CUBLAS_POINTER_MODE_HOST));
+        }
+
+        // cuBLAS handle on a pika thread
+        for (std::size_t i = 0; i < n; ++i)
+        {
+            tt::sync_wait(ex::schedule(ex::thread_pool_scheduler{}) |
+                ex::transfer(cu::cuda_scheduler{pool}) |
+                cu::then_with_cublas([](cublasHandle_t) {}, CUBLAS_POINTER_MODE_HOST));
+        }
+
+        // cuSOLVER handle on a non-pika thread
+        for (std::size_t i = 0; i < n; ++i)
+        {
+            tt::sync_wait(ex::schedule(cu::cuda_scheduler{pool}) |
+                cu::then_with_cusolver([](cusolverDnHandle_t) {}));
+        }
+
+        // cuSOLVER handle on a pika thread
+        for (std::size_t i = 0; i < n; ++i)
+        {
+            tt::sync_wait(ex::schedule(ex::thread_pool_scheduler{}) |
+                ex::transfer(cu::cuda_scheduler{pool}) |
+                cu::then_with_cusolver([](cusolverDnHandle_t) {}));
+        }
+    }
+
+    pika::finalize();
+    pika::stop();
+
+    whip::check_error(
+#if defined(PIKA_HAVE_CUDA)
+        cudaDeviceReset()
+#elif defined(PIKA_HAVE_HIP)
+        hipDeviceReset()
+#endif
+    );
+
+    PIKA_TEST(true);
+}
diff --git a/libs/pika/async_cuda_base/CMakeLists.txt b/libs/pika/async_cuda_base/CMakeLists.txt
new file mode 100644
index 000000000..fcf95e41c
--- /dev/null
+++ b/libs/pika/async_cuda_base/CMakeLists.txt
@@ -0,0 +1,54 @@
+# Copyright (c) 2019 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+if(NOT PIKA_WITH_GPU_SUPPORT)
+  return()
+endif()
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+
+set(async_cuda_base_headers
+    pika/async_cuda_base/cublas_exception.hpp
+    pika/async_cuda_base/cublas_handle.hpp
+    pika/async_cuda_base/cuda_device_scope.hpp
+    pika/async_cuda_base/cuda_event.hpp
+    pika/async_cuda_base/cuda_stream.hpp
+    pika/async_cuda_base/cusolver_exception.hpp
+    pika/async_cuda_base/cusolver_handle.hpp
+    pika/async_cuda_base/custom_blas_api.hpp
+    pika/async_cuda_base/custom_lapack_api.hpp
+    pika/async_cuda_base/detail/cuda_debug.hpp
+)
+
+set(async_cuda_base_sources
+    cublas_exception.cpp
+    cublas_handle.cpp
+    cuda_device_scope.cpp
+    cuda_event.cpp
+    cuda_stream.cpp
+    cusolver_exception.cpp
+    cusolver_handle.cpp
+)
+
+if(PIKA_WITH_HIP)
+  set(async_cuda_base_extra_deps roc::rocblas roc::rocsolver)
+elseif(PIKA_WITH_CUDA)
+  set(async_cuda_base_extra_deps pika_internal::cuda ${CUDA_CUBLAS_LIBRARIES}
+                                 ${CUDA_CUSOLVER_LIBRARIES}
+  )
+endif()
+
+include(pika_add_module)
+pika_add_module(
+  pika async_cuda_base
+  GLOBAL_HEADER_GEN ON
+  SOURCES ${async_cuda_base_sources}
+  HEADERS ${async_cuda_base_headers}
+  MODULE_DEPENDENCIES pika_assertion pika_config pika_coroutines pika_debugging
+                      pika_errors
+  DEPENDENCIES ${async_cuda_base_extra_deps}
+  CMAKE_SUBDIRS examples tests
+)
diff --git a/libs/pika/async_cuda_base/README.rst b/libs/pika/async_cuda_base/README.rst
new file mode 100644
index 000000000..552a7aa84
--- /dev/null
+++ b/libs/pika/async_cuda_base/README.rst
@@ -0,0 +1,13 @@
+
+..
+    Copyright (c) 2019 The STE||AR-Group
+
+    SPDX-License-Identifier: BSL-1.0
+    Distributed under the Boost Software License, Version 1.0. (See accompanying
+    file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+===============
+async_cuda_base
+===============
+
+This library is part of pika.
diff --git a/libs/pika/async_cuda_base/docs/index.rst b/libs/pika/async_cuda_base/docs/index.rst
new file mode 100644
index 000000000..48fe14084
--- /dev/null
+++ b/libs/pika/async_cuda_base/docs/index.rst
@@ -0,0 +1,14 @@
+..
+    Copyright (c) 2019 The STE||AR-Group
+
+    SPDX-License-Identifier: BSL-1.0
+    Distributed under the Boost Software License, Version 1.0. (See accompanying
+    file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+.. _modules_async_cuda_base:
+
+===============
+async_cuda_base
+===============
+
+This module defines basic CUDA and HIP utilities that are independent of the pika runtime.
diff --git a/libs/pika/async_cuda_base/examples/CMakeLists.txt b/libs/pika/async_cuda_base/examples/CMakeLists.txt
new file mode 100644
index 000000000..e7c52167b
--- /dev/null
+++ b/libs/pika/async_cuda_base/examples/CMakeLists.txt
@@ -0,0 +1,18 @@
+# Copyright (c) 2019 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+if(PIKA_WITH_EXAMPLES)
+  pika_add_pseudo_target(examples.modules.async_cuda_base)
+  pika_add_pseudo_dependencies(
+    examples.modules examples.modules.async_cuda_base
+  )
+  if(PIKA_WITH_TESTS AND PIKA_WITH_TESTS_EXAMPLES)
+    pika_add_pseudo_target(tests.examples.modules.async_cuda_base)
+    pika_add_pseudo_dependencies(
+      tests.examples.modules tests.examples.modules.async_cuda_base
+    )
+  endif()
+endif()
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_exception.hpp
similarity index 94%
rename from libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_exception.hpp
index bf0c10a6d..67beb782c 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp
+++ b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_exception.hpp
@@ -7,7 +7,7 @@
 #pragma once
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
 #include <pika/errors/exception.hpp>
 
 #include <string>
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp
similarity index 95%
rename from libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp
index 31aa1ba34..48b0eda5c 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp
+++ b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp
@@ -7,8 +7,8 @@
 #pragma once
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
 
 #include <fmt/format.h>
 #include <whip.hpp>
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_device_scope.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_device_scope.hpp
similarity index 100%
rename from libs/pika/async_cuda/include/pika/async_cuda/cuda_device_scope.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_device_scope.hpp
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp
new file mode 100644
index 000000000..654201ece
--- /dev/null
+++ b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp
@@ -0,0 +1,44 @@
+//  Copyright (c) 2020 John Biddiscombe
+//  Copyright (c) 2020 Teodor Nikolov
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+#pragma once
+
+#include <pika/config.hpp>
+
+#include <boost/lockfree/stack.hpp>
+#include <whip.hpp>
+
+#include <cstddef>
+
+namespace pika::cuda::experimental::detail {
+
+    // a pool of cudaEvent_t objects.
+    // Since allocation of a cuda event passes into the cuda runtime
+    // it might be an expensive operation, so we pre-allocate a pool
+    // of them at startup.
+    struct PIKA_EXPORT cuda_event_pool
+    {
+        static constexpr std::size_t initial_events_in_pool = 128;
+
+        static cuda_event_pool& get_event_pool();
+
+        cuda_event_pool();
+        ~cuda_event_pool();
+
+        bool pop(whip::event_t& event);
+        bool push(whip::event_t event);
+
+        void clear();
+        void grow(std::size_t n);
+
+    private:
+        void add_event_to_pool();
+
+        // pool is dynamically sized and can grow if needed
+        boost::lockfree::stack<whip::event_t, boost::lockfree::fixed_sized<false>> free_list_;
+    };
+}    // namespace pika::cuda::experimental::detail
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cuda_stream.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_stream.hpp
similarity index 100%
rename from libs/pika/async_cuda/include/pika/async_cuda/cuda_stream.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_stream.hpp
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp
similarity index 94%
rename from libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp
index 5e7b37d3e..85add9d52 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp
+++ b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp
@@ -8,7 +8,7 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda/custom_lapack_api.hpp>
+# include <pika/async_cuda_base/custom_lapack_api.hpp>
 # include <pika/errors/exception.hpp>
 
 # include <string>
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp
similarity index 95%
rename from libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp
index fe62be58e..302492501 100644
--- a/libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp
+++ b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp
@@ -8,8 +8,8 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda/cuda_stream.hpp>
-# include <pika/async_cuda/custom_lapack_api.hpp>
+# include <pika/async_cuda_base/cuda_stream.hpp>
+# include <pika/async_cuda_base/custom_lapack_api.hpp>
 
 # include <fmt/format.h>
 # include <whip.hpp>
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/custom_blas_api.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_blas_api.hpp
similarity index 100%
rename from libs/pika/async_cuda/include/pika/async_cuda/custom_blas_api.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_blas_api.hpp
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/custom_lapack_api.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_lapack_api.hpp
similarity index 100%
rename from libs/pika/async_cuda/include/pika/async_cuda/custom_lapack_api.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_lapack_api.hpp
diff --git a/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_debug.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/detail/cuda_debug.hpp
similarity index 100%
rename from libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_debug.hpp
rename to libs/pika/async_cuda_base/include/pika/async_cuda_base/detail/cuda_debug.hpp
diff --git a/libs/pika/async_cuda/src/cublas_exception.cpp b/libs/pika/async_cuda_base/src/cublas_exception.cpp
similarity index 97%
rename from libs/pika/async_cuda/src/cublas_exception.cpp
rename to libs/pika/async_cuda_base/src/cublas_exception.cpp
index 9ca16f0eb..2c4bcf98d 100644
--- a/libs/pika/async_cuda/src/cublas_exception.cpp
+++ b/libs/pika/async_cuda_base/src/cublas_exception.cpp
@@ -5,8 +5,8 @@
 //  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cublas_exception.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
+#include <pika/async_cuda_base/cublas_exception.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
 #include <pika/errors/exception.hpp>
 
 #include <string>
diff --git a/libs/pika/async_cuda/src/cublas_handle.cpp b/libs/pika/async_cuda_base/src/cublas_handle.cpp
similarity index 91%
rename from libs/pika/async_cuda/src/cublas_handle.cpp
rename to libs/pika/async_cuda_base/src/cublas_handle.cpp
index e5bc6a253..a2b51d408 100644
--- a/libs/pika/async_cuda/src/cublas_handle.cpp
+++ b/libs/pika/async_cuda_base/src/cublas_handle.cpp
@@ -5,11 +5,11 @@
 //  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cublas_exception.hpp>
-#include <pika/async_cuda/cublas_handle.hpp>
-#include <pika/async_cuda/cuda_device_scope.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
+#include <pika/async_cuda_base/cublas_exception.hpp>
+#include <pika/async_cuda_base/cublas_handle.hpp>
+#include <pika/async_cuda_base/cuda_device_scope.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
 
 #include <whip.hpp>
 
diff --git a/libs/pika/async_cuda/src/cuda_device_scope.cpp b/libs/pika/async_cuda_base/src/cuda_device_scope.cpp
similarity index 92%
rename from libs/pika/async_cuda/src/cuda_device_scope.cpp
rename to libs/pika/async_cuda_base/src/cuda_device_scope.cpp
index 6dbbf1f45..6f829ff2e 100644
--- a/libs/pika/async_cuda/src/cuda_device_scope.cpp
+++ b/libs/pika/async_cuda_base/src/cuda_device_scope.cpp
@@ -5,7 +5,7 @@
 //  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cuda_device_scope.hpp>
+#include <pika/async_cuda_base/cuda_device_scope.hpp>
 
 #include <whip.hpp>
 
diff --git a/libs/pika/async_cuda_base/src/cuda_event.cpp b/libs/pika/async_cuda_base/src/cuda_event.cpp
new file mode 100644
index 000000000..42e8b4f38
--- /dev/null
+++ b/libs/pika/async_cuda_base/src/cuda_event.cpp
@@ -0,0 +1,73 @@
+//  Copyright (c) 2023 ETH Zurich
+//  Copyright (c) 2020 John Biddiscombe
+//  Copyright (c) 2020 Teodor Nikolov
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+#include <pika/async_cuda_base/cuda_event.hpp>
+
+#include <whip.hpp>
+
+#include <cstddef>
+
+namespace pika::cuda::experimental::detail {
+    cuda_event_pool& cuda_event_pool::get_event_pool()
+    {
+        static cuda_event_pool event_pool_;
+        return event_pool_;
+    }
+
+    // reserve space for a bunch of events on initialization
+    cuda_event_pool::cuda_event_pool()
+      : free_list_(initial_events_in_pool)
+    {
+    }
+
+    // on destruction, all objects in stack will be freed
+    cuda_event_pool::~cuda_event_pool() { clear(); }
+
+    bool cuda_event_pool::pop(whip::event_t& event)
+    {
+        // pop an event off the pool, if that fails, create a new one
+        while (!free_list_.pop(event)) { add_event_to_pool(); }
+        return true;
+    }
+
+    bool cuda_event_pool::push(whip::event_t event) { return free_list_.push(event); }
+
+    void cuda_event_pool::clear()
+    {
+        whip::event_t event;
+        while (free_list_.pop(event))
+        {
+            if (!whip::event_ready(event))
+            {
+                PIKA_THROW_EXCEPTION(pika::error::invalid_status,
+                    "pika::cuda::experimental::detail::cuda_event_pool::clear",
+                    "Clearing pool of CUDA/HIP events, but found an event that is not yet "
+                    "ready. Are you disabling event polling before all kernels have "
+                    "completed?");
+            }
+
+            whip::event_destroy(event);
+        }
+    }
+
+    void cuda_event_pool::grow(std::size_t n)
+    {
+        for (std::size_t i = 0; i < n; ++i) { add_event_to_pool(); }
+    }
+
+    void cuda_event_pool::add_event_to_pool()
+    {
+        whip::event_t event;
+        // Create an cuda_event to query a CUDA/CUBLAS kernel for completion.
+        // Timing is disabled for performance. [1]
+        //
+        // [1]: CUDA Runtime API, section 5.5 cuda_event Management
+        whip::event_create_with_flags(&event, whip::event_disable_timing);
+        free_list_.push(event);
+    }
+}    // namespace pika::cuda::experimental::detail
diff --git a/libs/pika/async_cuda/src/cuda_stream.cpp b/libs/pika/async_cuda_base/src/cuda_stream.cpp
similarity index 96%
rename from libs/pika/async_cuda/src/cuda_stream.cpp
rename to libs/pika/async_cuda_base/src/cuda_stream.cpp
index 944e70bd3..9521b41e8 100644
--- a/libs/pika/async_cuda/src/cuda_stream.cpp
+++ b/libs/pika/async_cuda_base/src/cuda_stream.cpp
@@ -5,8 +5,8 @@
 //  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cuda_device_scope.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda_base/cuda_device_scope.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 
 #include <whip.hpp>
diff --git a/libs/pika/async_cuda/src/cusolver_exception.cpp b/libs/pika/async_cuda_base/src/cusolver_exception.cpp
similarity index 97%
rename from libs/pika/async_cuda/src/cusolver_exception.cpp
rename to libs/pika/async_cuda_base/src/cusolver_exception.cpp
index be2b7bb52..4bebaeeb6 100644
--- a/libs/pika/async_cuda/src/cusolver_exception.cpp
+++ b/libs/pika/async_cuda_base/src/cusolver_exception.cpp
@@ -8,8 +8,8 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda/cusolver_exception.hpp>
-# include <pika/async_cuda/custom_lapack_api.hpp>
+# include <pika/async_cuda_base/cusolver_exception.hpp>
+# include <pika/async_cuda_base/custom_lapack_api.hpp>
 # include <pika/errors/exception.hpp>
 
 # include <string>
diff --git a/libs/pika/async_cuda/src/cusolver_handle.cpp b/libs/pika/async_cuda_base/src/cusolver_handle.cpp
similarity index 90%
rename from libs/pika/async_cuda/src/cusolver_handle.cpp
rename to libs/pika/async_cuda_base/src/cusolver_handle.cpp
index 8a491e01c..7b54c0df1 100644
--- a/libs/pika/async_cuda/src/cusolver_handle.cpp
+++ b/libs/pika/async_cuda_base/src/cusolver_handle.cpp
@@ -5,11 +5,11 @@
 //  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
 
 #include <pika/config.hpp>
-#include <pika/async_cuda/cuda_device_scope.hpp>
-#include <pika/async_cuda/cuda_stream.hpp>
-#include <pika/async_cuda/cusolver_exception.hpp>
-#include <pika/async_cuda/cusolver_handle.hpp>
-#include <pika/async_cuda/custom_blas_api.hpp>
+#include <pika/async_cuda_base/cuda_device_scope.hpp>
+#include <pika/async_cuda_base/cuda_stream.hpp>
+#include <pika/async_cuda_base/cusolver_exception.hpp>
+#include <pika/async_cuda_base/cusolver_handle.hpp>
+#include <pika/async_cuda_base/custom_blas_api.hpp>
 
 #include <whip.hpp>
 
diff --git a/libs/pika/async_cuda_base/tests/CMakeLists.txt b/libs/pika/async_cuda_base/tests/CMakeLists.txt
new file mode 100644
index 000000000..f81df4d8a
--- /dev/null
+++ b/libs/pika/async_cuda_base/tests/CMakeLists.txt
@@ -0,0 +1,44 @@
+# Copyright (c) 2019 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+include(pika_message)
+include(pika_option)
+
+if(PIKA_WITH_TESTS)
+  if(PIKA_WITH_TESTS_UNIT)
+    pika_add_pseudo_target(tests.unit.modules.async_cuda_base)
+    pika_add_pseudo_dependencies(
+      tests.unit.modules tests.unit.modules.async_cuda_base
+    )
+    add_subdirectory(unit)
+  endif()
+
+  if(PIKA_WITH_TESTS_REGRESSIONS)
+    pika_add_pseudo_target(tests.regressions.modules.async_cuda_base)
+    pika_add_pseudo_dependencies(
+      tests.regressions.modules tests.regressions.modules.async_cuda_base
+    )
+    add_subdirectory(regressions)
+  endif()
+
+  if(PIKA_WITH_TESTS_BENCHMARKS)
+    pika_add_pseudo_target(tests.performance.modules.async_cuda_base)
+    pika_add_pseudo_dependencies(
+      tests.performance.modules tests.performance.modules.async_cuda_base
+    )
+    add_subdirectory(performance)
+  endif()
+
+  if(PIKA_WITH_TESTS_HEADERS)
+    pika_add_header_tests(
+      modules.async_cuda_base
+      HEADERS ${async_cuda_base_headers}
+      HEADER_ROOT ${PROJECT_SOURCE_DIR}/include
+      NOLIBS
+      DEPENDENCIES pika_async_cuda_base
+    )
+  endif()
+endif()
diff --git a/libs/pika/async_cuda_base/tests/performance/CMakeLists.txt b/libs/pika/async_cuda_base/tests/performance/CMakeLists.txt
new file mode 100644
index 000000000..e5e7cb287
--- /dev/null
+++ b/libs/pika/async_cuda_base/tests/performance/CMakeLists.txt
@@ -0,0 +1,5 @@
+# Copyright (c) 2020 John Biddiscombe
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
diff --git a/libs/pika/async_cuda_base/tests/regressions/CMakeLists.txt b/libs/pika/async_cuda_base/tests/regressions/CMakeLists.txt
new file mode 100644
index 000000000..e05062746
--- /dev/null
+++ b/libs/pika/async_cuda_base/tests/regressions/CMakeLists.txt
@@ -0,0 +1,5 @@
+# Copyright (c) 2019 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
diff --git a/libs/pika/async_cuda_base/tests/unit/CMakeLists.txt b/libs/pika/async_cuda_base/tests/unit/CMakeLists.txt
new file mode 100644
index 000000000..866183ad5
--- /dev/null
+++ b/libs/pika/async_cuda_base/tests/unit/CMakeLists.txt
@@ -0,0 +1,34 @@
+# Copyright (c) 2019 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+set(tests cublas_handle cuda_stream cusolver_handle)
+
+set(cublas_handle_GPU ON)
+set(cuda_stream_GPU ON)
+set(cusolver_handle_GPU ON)
+
+foreach(test ${tests})
+  if(${${test}_GPU})
+    set(sources ${test}.cu)
+  else()
+    set(sources ${test}.cpp)
+  endif()
+
+  source_group("Source Files" FILES ${sources})
+
+  # add test executable
+  pika_add_executable(
+    ${test}_test INTERNAL_FLAGS
+    SOURCES ${sources} ${${test}_FLAGS}
+    DEPENDENCIES ${${test}_DEPENDENCIES}
+    EXCLUDE_FROM_ALL
+    FOLDER "Tests/Unit/Modules/AsyncCudaBase"
+  )
+
+  pika_add_unit_test(
+    "modules.async_cuda_base" ${test} ${${test}_PARAMETERS} RUN_SERIAL
+  )
+endforeach()
diff --git a/libs/pika/async_cuda/tests/unit/cublas_handle.cu b/libs/pika/async_cuda_base/tests/unit/cublas_handle.cu
similarity index 100%
rename from libs/pika/async_cuda/tests/unit/cublas_handle.cu
rename to libs/pika/async_cuda_base/tests/unit/cublas_handle.cu
diff --git a/libs/pika/async_cuda/tests/unit/cuda_stream.cu b/libs/pika/async_cuda_base/tests/unit/cuda_stream.cu
similarity index 100%
rename from libs/pika/async_cuda/tests/unit/cuda_stream.cu
rename to libs/pika/async_cuda_base/tests/unit/cuda_stream.cu
diff --git a/libs/pika/async_cuda/tests/unit/cusolver_handle.cu b/libs/pika/async_cuda_base/tests/unit/cusolver_handle.cu
similarity index 100%
rename from libs/pika/async_cuda/tests/unit/cusolver_handle.cu
rename to libs/pika/async_cuda_base/tests/unit/cusolver_handle.cu
diff --git a/libs/pika/include/include/pika/cuda.hpp b/libs/pika/include/include/pika/cuda.hpp
index 60a64d0ca..b7f970524 100644
--- a/libs/pika/include/include/pika/cuda.hpp
+++ b/libs/pika/include/include/pika/cuda.hpp
@@ -9,4 +9,5 @@
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
 # include <pika/modules/async_cuda.hpp>
+# include <pika/modules/async_cuda_base.hpp>
 #endif
diff --git a/libs/pika/runtime/CMakeLists.txt b/libs/pika/runtime/CMakeLists.txt
index 8b28198b8..8a9bffea0 100644
--- a/libs/pika/runtime/CMakeLists.txt
+++ b/libs/pika/runtime/CMakeLists.txt
@@ -40,6 +40,10 @@ set(runtime_sources
     thread_stacktrace.cpp
 )
 
+if(PIKA_WITH_GPU_SUPPORT)
+  list(APPEND runtime_additional_module_dependencies pika_async_cuda_base)
+endif()
+
 if(PIKA_WITH_MPI)
   list(APPEND runtime_additional_module_dependencies pika_mpi_base)
 endif()
diff --git a/libs/pika/runtime/src/runtime.cpp b/libs/pika/runtime/src/runtime.cpp
index 0051a73ca..b5c1fd8c2 100644
--- a/libs/pika/runtime/src/runtime.cpp
+++ b/libs/pika/runtime/src/runtime.cpp
@@ -38,6 +38,10 @@
 #include <pika/topology/topology.hpp>
 #include <pika/version.hpp>
 
+#if defined(PIKA_HAVE_GPU_SUPPORT)
+# include <pika/async_cuda_base/cuda_event.hpp>
+#endif
+
 #if defined(PIKA_HAVE_MPI)
 # include <pika/mpi_base/mpi.hpp>
 #endif
@@ -364,6 +368,10 @@ namespace pika {
         if (0 == instance_number_counter_) --instance_number_counter_;
 
         resource::detail::delete_partitioner();
+
+#if defined(PIKA_HAVE_GPU_SUPPORT)
+        pika::cuda::experimental::detail::cuda_event_pool::get_event_pool().clear();
+#endif
     }
 
     void runtime::on_exit(util::detail::function<void()> const& f)