diff --git a/libs/pika/CMakeLists.txt b/libs/pika/CMakeLists.txt
index 1d16d4a76..9db14e4d9 100644
--- a/libs/pika/CMakeLists.txt
+++ b/libs/pika/CMakeLists.txt
@@ -16,7 +16,6 @@ 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 530080a36..e9c349a2d 100644
--- a/libs/pika/async_cuda/CMakeLists.txt
+++ b/libs/pika/async_cuda/CMakeLists.txt
@@ -10,17 +10,39 @@ 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 cuda_event_callback.cpp cuda_pool.cpp cuda_scheduler.cpp)
+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
+)
 
 if(PIKA_WITH_HIP)
   set(async_cuda_extra_deps roc::rocblas roc::rocsolver)
@@ -37,18 +59,17 @@ 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_coroutines
+    pika_debugging
     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_base/include/pika/async_cuda_base/cublas_exception.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp
similarity index 94%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_exception.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp
index 67beb782c..bf0c10a6d 100644
--- a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_exception.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cublas_exception.hpp
@@ -7,7 +7,7 @@
 #pragma once
 
 #include <pika/config.hpp>
-#include <pika/async_cuda_base/custom_blas_api.hpp>
+#include <pika/async_cuda/custom_blas_api.hpp>
 #include <pika/errors/exception.hpp>
 
 #include <string>
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp
similarity index 95%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp
index 48b0eda5c..31aa1ba34 100644
--- a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cublas_handle.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cublas_handle.hpp
@@ -7,8 +7,8 @@
 #pragma once
 
 #include <pika/config.hpp>
-#include <pika/async_cuda_base/cuda_stream.hpp>
-#include <pika/async_cuda_base/custom_blas_api.hpp>
+#include <pika/async_cuda/cuda_stream.hpp>
+#include <pika/async_cuda/custom_blas_api.hpp>
 
 #include <fmt/format.h>
 #include <whip.hpp>
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_device_scope.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_device_scope.hpp
similarity index 100%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_device_scope.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cuda_device_scope.hpp
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
new file mode 100644
index 000000000..2477f8d51
--- /dev/null
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cuda_event.hpp
@@ -0,0 +1,72 @@
+//  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 86e1cabd8..fe8459d2d 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,9 +7,7 @@
 #pragma once
 
 #include <pika/assert.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/async_cuda/cuda_stream.hpp>
 #include <pika/concurrency/cache_line_data.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 
@@ -18,43 +16,10 @@
 #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
@@ -82,48 +47,11 @@ 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);
@@ -152,9 +80,6 @@ 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 4d0c787cc..33fc052a4 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_base/cuda_stream.hpp>
+#include <pika/async_cuda/cuda_stream.hpp>
 #include <pika/concepts/concepts.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 #include <pika/execution/algorithms/execute.hpp>
@@ -40,9 +40,6 @@ 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 20cc78950..80c281282 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_base/include/pika/async_cuda_base/cuda_stream.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cuda_stream.hpp
similarity index 100%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_stream.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cuda_stream.hpp
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp
similarity index 94%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp
index 85add9d52..5e7b37d3e 100644
--- a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_exception.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cusolver_exception.hpp
@@ -8,7 +8,7 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda_base/custom_lapack_api.hpp>
+# include <pika/async_cuda/custom_lapack_api.hpp>
 # include <pika/errors/exception.hpp>
 
 # include <string>
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp b/libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp
similarity index 95%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp
index 302492501..fe62be58e 100644
--- a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cusolver_handle.hpp
+++ b/libs/pika/async_cuda/include/pika/async_cuda/cusolver_handle.hpp
@@ -8,8 +8,8 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda_base/cuda_stream.hpp>
-# include <pika/async_cuda_base/custom_lapack_api.hpp>
+# include <pika/async_cuda/cuda_stream.hpp>
+# include <pika/async_cuda/custom_lapack_api.hpp>
 
 # include <fmt/format.h>
 # include <whip.hpp>
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_blas_api.hpp b/libs/pika/async_cuda/include/pika/async_cuda/custom_blas_api.hpp
similarity index 100%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_blas_api.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/custom_blas_api.hpp
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_lapack_api.hpp b/libs/pika/async_cuda/include/pika/async_cuda/custom_lapack_api.hpp
similarity index 100%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/custom_lapack_api.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/custom_lapack_api.hpp
diff --git a/libs/pika/async_cuda_base/include/pika/async_cuda_base/detail/cuda_debug.hpp b/libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_debug.hpp
similarity index 100%
rename from libs/pika/async_cuda_base/include/pika/async_cuda_base/detail/cuda_debug.hpp
rename to libs/pika/async_cuda/include/pika/async_cuda/detail/cuda_debug.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 d7fb0f6e4..7719feace 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_base/cuda_stream.hpp>
+#include <pika/async_cuda/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 58329a67f..4633a5bb9 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,23 +35,29 @@
 #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_scheduler& sched, cuda_stream const& stream,
-        cublasPointerMode_t pointer_mode, F&& f, Ts&&... ts)
+    auto invoke_with_thread_local_cublas_handle(
+        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)...))
     {
-        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)...);
+        return PIKA_INVOKE(PIKA_FORWARD(F, f),
+            get_thread_local_cublas_handle(stream, pointer_mode).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_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)...))
+    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 locked_handle = sched.get_cusolver_handle(stream);
-        return PIKA_INVOKE(PIKA_FORWARD(F, f), locked_handle.get().get(), PIKA_FORWARD(Ts, ts)...);
+        return PIKA_INVOKE(PIKA_FORWARD(F, f), get_thread_local_cusolver_handle(stream).get(),
+            PIKA_FORWARD(Ts, ts)...);
     }
 
     template <typename R, typename... Ts>
@@ -151,11 +157,11 @@ namespace pika::cuda::experimental::then_with_stream_detail {
 
 #if defined(PIKA_HAVE_STDEXEC)
         template <typename... Ts>
-        requires std::is_invocable_v<F, cuda_scheduler&, cuda_stream const&,
+        requires std::is_invocable_v<F, 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, cuda_scheduler&,
+                    detail::result_type_signature_helper_t<std::invoke_result_t<F,
                         cuda_stream const&, std::add_lvalue_reference_t<std::decay_t<Ts>>...>>>;
 
         using completion_signatures =
@@ -171,7 +177,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_scheduler&, cuda_stream const&,
+            using result_type = std::invoke_result_t<F, 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>>;
@@ -250,8 +256,7 @@ namespace pika::cuda::experimental::then_with_stream_detail {
 
                 template <typename... Ts>
                 auto set_value(Ts&&... ts) noexcept
-                    -> decltype(PIKA_INVOKE(PIKA_MOVE(f), op_state.sched, stream.value(), ts...),
-                        void())
+                    -> decltype(PIKA_INVOKE(PIKA_MOVE(f), stream.value(), ts...), void())
                 {
                     pika::detail::try_catch_exception_ptr(
                         [&]() mutable {
@@ -282,16 +287,16 @@ namespace pika::cuda::experimental::then_with_stream_detail {
                                 }
                             }
 
-                            using invoke_result_type = std::decay_t<
-                                std::invoke_result_t<F, cuda_scheduler&, cuda_stream const&,
+                            using invoke_result_type =
+                                std::decay_t<std::invoke_result_t<F, 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.sched,
-                                            op_state.stream.value(), ts...);
+                                        PIKA_INVOKE(
+                                            PIKA_MOVE(op_state.f), op_state.stream.value(), ts...);
                                     },
                                     t);
 
@@ -331,7 +336,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), op_state.sched,
+                                            PIKA_INVOKE(PIKA_MOVE(op_state.f),
                                                 op_state.stream.value(), ts...));
                                     },
                                     t);
@@ -543,7 +548,7 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         std::decay_t<F> f;
 
         template <typename... Ts>
-        auto operator()(cuda_scheduler&, cuda_stream const& stream, Ts&&... ts)
+        auto operator()(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())))
@@ -563,17 +568,17 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         cublasPointerMode_t pointer_mode;
 
         template <typename... Ts>
-        auto operator()(cuda_scheduler& sched, cuda_stream const& stream, Ts&&... ts)
+        auto operator()(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(
-                sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...)))
+                stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...)))
 #endif
                 -> decltype(invoke_with_thread_local_cublas_handle(
-                    sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...))
+                    stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...))
         {
             return invoke_with_thread_local_cublas_handle(
-                sched, stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...);
+                stream, pointer_mode, f, PIKA_FORWARD(Ts, ts)...);
         }
     };
 
@@ -585,17 +590,16 @@ namespace pika::cuda::experimental::then_with_stream_detail {
         std::decay_t<F> f;
 
         template <typename... Ts>
-        auto operator()(cuda_scheduler& sched, cuda_stream const& stream, Ts&&... ts)
+        auto operator()(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(
-                sched, stream, f, PIKA_FORWARD(Ts, ts)...)))
+            noexcept(noexcept(
+                invoke_with_thread_local_cusolver_handle(stream, f, PIKA_FORWARD(Ts, ts)...)))
 #endif
                 -> decltype(invoke_with_thread_local_cusolver_handle(
-                    sched, stream, f, PIKA_FORWARD(Ts, ts)...))
+                    stream, f, PIKA_FORWARD(Ts, ts)...))
         {
-            return invoke_with_thread_local_cusolver_handle(
-                sched, stream, f, PIKA_FORWARD(Ts, ts)...);
+            return invoke_with_thread_local_cusolver_handle(stream, f, PIKA_FORWARD(Ts, ts)...);
         }
     };
 }    // namespace pika::cuda::experimental::then_with_stream_detail
diff --git a/libs/pika/async_cuda_base/src/cublas_exception.cpp b/libs/pika/async_cuda/src/cublas_exception.cpp
similarity index 97%
rename from libs/pika/async_cuda_base/src/cublas_exception.cpp
rename to libs/pika/async_cuda/src/cublas_exception.cpp
index 2c4bcf98d..9ca16f0eb 100644
--- a/libs/pika/async_cuda_base/src/cublas_exception.cpp
+++ b/libs/pika/async_cuda/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_base/cublas_exception.hpp>
-#include <pika/async_cuda_base/custom_blas_api.hpp>
+#include <pika/async_cuda/cublas_exception.hpp>
+#include <pika/async_cuda/custom_blas_api.hpp>
 #include <pika/errors/exception.hpp>
 
 #include <string>
diff --git a/libs/pika/async_cuda_base/src/cublas_handle.cpp b/libs/pika/async_cuda/src/cublas_handle.cpp
similarity index 91%
rename from libs/pika/async_cuda_base/src/cublas_handle.cpp
rename to libs/pika/async_cuda/src/cublas_handle.cpp
index a2b51d408..e5bc6a253 100644
--- a/libs/pika/async_cuda_base/src/cublas_handle.cpp
+++ b/libs/pika/async_cuda/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_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 <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 <whip.hpp>
 
diff --git a/libs/pika/async_cuda_base/src/cuda_device_scope.cpp b/libs/pika/async_cuda/src/cuda_device_scope.cpp
similarity index 92%
rename from libs/pika/async_cuda_base/src/cuda_device_scope.cpp
rename to libs/pika/async_cuda/src/cuda_device_scope.cpp
index 6f829ff2e..6dbbf1f45 100644
--- a/libs/pika/async_cuda_base/src/cuda_device_scope.cpp
+++ b/libs/pika/async_cuda/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_base/cuda_device_scope.hpp>
+#include <pika/async_cuda/cuda_device_scope.hpp>
 
 #include <whip.hpp>
 
diff --git a/libs/pika/async_cuda/src/cuda_event_callback.cpp b/libs/pika/async_cuda/src/cuda_event_callback.cpp
index 7d1c4ae78..2d2e29974 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,10 +309,6 @@ 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 0b548ff21..1592130da 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_base/cuda_stream.hpp>
+#include <pika/async_cuda/cuda_stream.hpp>
 #include <pika/concurrency/cache_line_data.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 #include <pika/runtime/runtime_fwd.hpp>
@@ -15,7 +15,6 @@
 
 #include <cstddef>
 #include <memory>
-#include <utility>
 #include <vector>
 
 namespace pika::cuda::experimental {
@@ -41,102 +40,6 @@ 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++;
-
-            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++;
-
-            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)
@@ -144,8 +47,6 @@ 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()
     {
     }
 
@@ -170,17 +71,4 @@ 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 1f371a21c..c96b1bc00 100644
--- a/libs/pika/async_cuda/src/cuda_scheduler.cpp
+++ b/libs/pika/async_cuda/src/cuda_scheduler.cpp
@@ -18,17 +18,6 @@ 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_base/src/cuda_stream.cpp b/libs/pika/async_cuda/src/cuda_stream.cpp
similarity index 96%
rename from libs/pika/async_cuda_base/src/cuda_stream.cpp
rename to libs/pika/async_cuda/src/cuda_stream.cpp
index 9521b41e8..944e70bd3 100644
--- a/libs/pika/async_cuda_base/src/cuda_stream.cpp
+++ b/libs/pika/async_cuda/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_base/cuda_device_scope.hpp>
-#include <pika/async_cuda_base/cuda_stream.hpp>
+#include <pika/async_cuda/cuda_device_scope.hpp>
+#include <pika/async_cuda/cuda_stream.hpp>
 #include <pika/coroutines/thread_enums.hpp>
 
 #include <whip.hpp>
diff --git a/libs/pika/async_cuda_base/src/cusolver_exception.cpp b/libs/pika/async_cuda/src/cusolver_exception.cpp
similarity index 97%
rename from libs/pika/async_cuda_base/src/cusolver_exception.cpp
rename to libs/pika/async_cuda/src/cusolver_exception.cpp
index 4bebaeeb6..be2b7bb52 100644
--- a/libs/pika/async_cuda_base/src/cusolver_exception.cpp
+++ b/libs/pika/async_cuda/src/cusolver_exception.cpp
@@ -8,8 +8,8 @@
 
 #include <pika/config.hpp>
 #if defined(PIKA_HAVE_GPU_SUPPORT)
-# include <pika/async_cuda_base/cusolver_exception.hpp>
-# include <pika/async_cuda_base/custom_lapack_api.hpp>
+# include <pika/async_cuda/cusolver_exception.hpp>
+# include <pika/async_cuda/custom_lapack_api.hpp>
 # include <pika/errors/exception.hpp>
 
 # include <string>
diff --git a/libs/pika/async_cuda_base/src/cusolver_handle.cpp b/libs/pika/async_cuda/src/cusolver_handle.cpp
similarity index 90%
rename from libs/pika/async_cuda_base/src/cusolver_handle.cpp
rename to libs/pika/async_cuda/src/cusolver_handle.cpp
index 7b54c0df1..8a491e01c 100644
--- a/libs/pika/async_cuda_base/src/cusolver_handle.cpp
+++ b/libs/pika/async_cuda/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_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 <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 <whip.hpp>
 
diff --git a/libs/pika/async_cuda/src/then_with_stream.cpp b/libs/pika/async_cuda/src/then_with_stream.cpp
new file mode 100644
index 000000000..37913126b
--- /dev/null
+++ b/libs/pika/async_cuda/src/then_with_stream.cpp
@@ -0,0 +1,37 @@
+//  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 9ad34cf4a..60b43f322 100644
--- a/libs/pika/async_cuda/tests/unit/CMakeLists.txt
+++ b/libs/pika/async_cuda/tests/unit/CMakeLists.txt
@@ -5,12 +5,14 @@
 # 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
 )
 
@@ -19,10 +21,13 @@ 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_base/tests/unit/cublas_handle.cu b/libs/pika/async_cuda/tests/unit/cublas_handle.cu
similarity index 100%
rename from libs/pika/async_cuda_base/tests/unit/cublas_handle.cu
rename to libs/pika/async_cuda/tests/unit/cublas_handle.cu
diff --git a/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp b/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp
deleted file mode 100644
index 32d68f358..000000000
--- a/libs/pika/async_cuda/tests/unit/cuda_device_reset.cpp
+++ /dev/null
@@ -1,66 +0,0 @@
-//  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);
-
-    {
-        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
-        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
-        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
-        tt::sync_wait(ex::schedule(cu::cuda_scheduler{pool}) |
-            cu::then_with_cusolver([](cusolverDnHandle_t) {}));
-
-        // cuSOLVER handle on a pika thread
-        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/tests/unit/cuda_stream.cu b/libs/pika/async_cuda/tests/unit/cuda_stream.cu
similarity index 100%
rename from libs/pika/async_cuda_base/tests/unit/cuda_stream.cu
rename to libs/pika/async_cuda/tests/unit/cuda_stream.cu
diff --git a/libs/pika/async_cuda_base/tests/unit/cusolver_handle.cu b/libs/pika/async_cuda/tests/unit/cusolver_handle.cu
similarity index 100%
rename from libs/pika/async_cuda_base/tests/unit/cusolver_handle.cu
rename to libs/pika/async_cuda/tests/unit/cusolver_handle.cu
diff --git a/libs/pika/async_cuda_base/CMakeLists.txt b/libs/pika/async_cuda_base/CMakeLists.txt
deleted file mode 100644
index fcf95e41c..000000000
--- a/libs/pika/async_cuda_base/CMakeLists.txt
+++ /dev/null
@@ -1,54 +0,0 @@
-# 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
deleted file mode 100644
index 552a7aa84..000000000
--- a/libs/pika/async_cuda_base/README.rst
+++ /dev/null
@@ -1,13 +0,0 @@
-
-..
-    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
deleted file mode 100644
index 48fe14084..000000000
--- a/libs/pika/async_cuda_base/docs/index.rst
+++ /dev/null
@@ -1,14 +0,0 @@
-..
-    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
deleted file mode 100644
index e7c52167b..000000000
--- a/libs/pika/async_cuda_base/examples/CMakeLists.txt
+++ /dev/null
@@ -1,18 +0,0 @@
-# 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_base/include/pika/async_cuda_base/cuda_event.hpp b/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp
deleted file mode 100644
index 654201ece..000000000
--- a/libs/pika/async_cuda_base/include/pika/async_cuda_base/cuda_event.hpp
+++ /dev/null
@@ -1,44 +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 <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_base/src/cuda_event.cpp b/libs/pika/async_cuda_base/src/cuda_event.cpp
deleted file mode 100644
index 42e8b4f38..000000000
--- a/libs/pika/async_cuda_base/src/cuda_event.cpp
+++ /dev/null
@@ -1,73 +0,0 @@
-//  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_base/tests/CMakeLists.txt b/libs/pika/async_cuda_base/tests/CMakeLists.txt
deleted file mode 100644
index f81df4d8a..000000000
--- a/libs/pika/async_cuda_base/tests/CMakeLists.txt
+++ /dev/null
@@ -1,44 +0,0 @@
-# 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
deleted file mode 100644
index e5e7cb287..000000000
--- a/libs/pika/async_cuda_base/tests/performance/CMakeLists.txt
+++ /dev/null
@@ -1,5 +0,0 @@
-# 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
deleted file mode 100644
index e05062746..000000000
--- a/libs/pika/async_cuda_base/tests/regressions/CMakeLists.txt
+++ /dev/null
@@ -1,5 +0,0 @@
-# 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
deleted file mode 100644
index 866183ad5..000000000
--- a/libs/pika/async_cuda_base/tests/unit/CMakeLists.txt
+++ /dev/null
@@ -1,34 +0,0 @@
-# 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/include/include/pika/cuda.hpp b/libs/pika/include/include/pika/cuda.hpp
index b7f970524..60a64d0ca 100644
--- a/libs/pika/include/include/pika/cuda.hpp
+++ b/libs/pika/include/include/pika/cuda.hpp
@@ -9,5 +9,4 @@
 #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 8a9bffea0..8b28198b8 100644
--- a/libs/pika/runtime/CMakeLists.txt
+++ b/libs/pika/runtime/CMakeLists.txt
@@ -40,10 +40,6 @@ 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 b5c1fd8c2..0051a73ca 100644
--- a/libs/pika/runtime/src/runtime.cpp
+++ b/libs/pika/runtime/src/runtime.cpp
@@ -38,10 +38,6 @@
 #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
@@ -368,10 +364,6 @@ 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)