[cudax -> libcu++] Move memory resources to libcu++ (#6384)

* Move memory resources to libcu++

* Add moved headers to the main header

* Unused parameter

Author: pciolkosz

cudax/examples/async_buffer_add.cu

@@ -55,7 +55,7 @@ int main()
   // The execution policy we want to use to run all work on the same stream
   auto policy = thrust::cuda::par_nosync.on(stream.get());
 
-  cudax::device_memory_pool_ref device_resource = cudax::device_default_memory_pool(cuda::device_ref{0});
+  cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});
 
   // Allocate the two inputs and output, but do not zero initialize via `cudax::no_init`
   cudax::async_device_buffer<float> A{stream, device_resource, numElements, cudax::no_init};
@@ -69,7 +69,7 @@ int main()
   // Add the vectors together
   thrust::transform(policy, A.begin(), A.end(), B.begin(), C.begin(), cuda::std::plus<>{});
 
-  cudax::pinned_memory_pool_ref pinned_resource = cudax::pinned_default_memory_pool();
+  cuda::pinned_memory_pool_ref pinned_resource = cuda::pinned_default_memory_pool();
 
   // Verify that the result vector is correct, by copying it to host
   cudax::async_host_buffer<float> h_A{stream, pinned_resource, A};

cudax/examples/cub_reduce.cu

@@ -25,8 +25,8 @@ int main()
   constexpr int num_items = 50000;
 
   // A CUDA stream on which to execute the reduction
-  cudax::stream stream{cuda::devices[0]};
-  cudax::device_memory_pool_ref mr = cudax::device_default_memory_pool(cuda::devices[0]);
+  cuda::stream stream{cuda::devices[0]};
+  cuda::device_memory_pool_ref mr = cuda::device_default_memory_pool(cuda::devices[0]);
 
   // Allocate input and output, but do not zero initialize output (`cudax::no_init`)
   auto d_in  = cudax::make_async_buffer<int>(stream, mr, num_items, 1);
@@ -41,7 +41,7 @@ int main()
     exit(EXIT_FAILURE);
   }
 
-  auto h_out = cudax::make_async_buffer<float>(stream, cudax::pinned_default_memory_pool(), d_out);
+  auto h_out = cudax::make_async_buffer<float>(stream, cuda::pinned_default_memory_pool(), d_out);
 
   stream.sync();
 

cudax/examples/simple_p2p.cu

@@ -120,7 +120,7 @@ void test_cross_device_access_from_kernel(
 
   // This will be a pinned memory vector once available
   cudax::uninitialized_buffer<float, cuda::mr::host_accessible> host_buffer(
-    cudax::legacy_pinned_memory_resource(), dev0_buffer.size());
+    cuda::legacy_pinned_memory_resource(), dev0_buffer.size());
   std::generate(host_buffer.begin(), host_buffer.end(), []() {
     static int i = 0;
     return static_cast<float>((i++) % 4096);
@@ -215,13 +215,13 @@ try
     return 0;
   }
 
-  cudax::stream dev0_stream(peers[0]);
-  cudax::stream dev1_stream(peers[1]);
+  cuda::stream dev0_stream(peers[0]);
+  cuda::stream dev1_stream(peers[1]);
 
   printf("Enabling peer access between GPU%d and GPU%d...\n", peers[0].get(), peers[1].get());
-  cudax::device_memory_pool_ref dev0_resource = cudax::device_default_memory_pool(peers[0]);
+  cuda::device_memory_pool_ref dev0_resource = cuda::device_default_memory_pool(peers[0]);
   dev0_resource.enable_access_from(peers[1]);
-  cudax::device_memory_pool_ref dev1_resource = cudax::device_default_memory_pool(peers[1]);
+  cuda::device_memory_pool_ref dev1_resource = cuda::device_default_memory_pool(peers[1]);
   dev1_resource.enable_access_from(peers[0]);
 
   // Allocate buffers

cudax/include/cuda/experimental/__memory_resource/resource.cuh

@@ -23,14 +23,14 @@
 #endif // !LIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE
 
 #include <cuda/__memory_resource/any_resource.h>
+#include <cuda/__memory_resource/device_memory_pool.h>
+#include <cuda/__memory_resource/legacy_managed_memory_resource.h>
+#include <cuda/__memory_resource/legacy_pinned_memory_resource.h>
+#include <cuda/__memory_resource/managed_memory_pool.h>
+#include <cuda/__memory_resource/pinned_memory_pool.h>
 #include <cuda/__memory_resource/properties.h>
+#include <cuda/__memory_resource/resource.h>
 
-#include <cuda/experimental/__memory_resource/device_memory_pool.cuh>
-#include <cuda/experimental/__memory_resource/legacy_managed_memory_resource.cuh>
-#include <cuda/experimental/__memory_resource/legacy_pinned_memory_resource.cuh>
-#include <cuda/experimental/__memory_resource/managed_memory_pool.cuh>
-#include <cuda/experimental/__memory_resource/pinned_memory_pool.cuh>
-#include <cuda/experimental/__memory_resource/resource.cuh>
 #include <cuda/experimental/__memory_resource/shared_resource.cuh>
 
 #endif // __CUDAX_MEMORY_RESOURCE___

cudax/include/cuda/experimental/memory_resource.cuh

@@ -135,10 +135,6 @@ foreach(cn_target IN LISTS cudax_TARGETS)
   )
 
   cudax_add_catch2_test(test_target memory_resource ${cn_target}
-    memory_resource/memory_pools.cu
-    memory_resource/device_memory_resource.cu
-    memory_resource/managed_memory_resource.cu
-    memory_resource/pinned_memory_resource.cu
     memory_resource/shared_resource.cu
   )
 

cudax/test/CMakeLists.txt

@@ -47,7 +47,7 @@ void check_result_and_erase(cudax::stream_ref stream, Result&& result, uint8_t p
 template <typename Layout = cuda::std::layout_right, typename Extents>
 auto make_buffer_for_mdspan(Extents extents, char value = 0)
 {
-  cudax::legacy_pinned_memory_resource host_resource;
+  cuda::legacy_pinned_memory_resource host_resource;
   auto mapping = typename Layout::template mapping<decltype(extents)>{extents};
 
   cudax::uninitialized_buffer<int, cuda::mr::host_accessible> buffer(host_resource, mapping.required_span_size());

cudax/test/algorithm/common.cuh

@@ -12,11 +12,11 @@
 
 C2H_TEST("1d Copy", "[data_manipulation]")
 {
-  cudax::stream _stream{cuda::device_ref{0}};
+  cuda::stream _stream{cuda::device_ref{0}};
 
   SECTION("Device resource")
   {
-    cudax::device_memory_pool_ref device_resource = cudax::device_default_memory_pool(cuda::device_ref{0});
+    cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});
     std::vector<int> host_vector(buffer_size);
 
     {
@@ -46,8 +46,8 @@ C2H_TEST("1d Copy", "[data_manipulation]")
 
   SECTION("Host and managed resource")
   {
-    cudax::legacy_managed_memory_resource managed_resource;
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_managed_memory_resource managed_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
 
     {
       cudax::uninitialized_buffer<int, cuda::mr::host_accessible> host_buffer(host_resource, buffer_size);
@@ -78,7 +78,7 @@ C2H_TEST("1d Copy", "[data_manipulation]")
   }
   SECTION("Launch transform")
   {
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
     cudax::weird_buffer input(host_resource, buffer_size);
     cudax::weird_buffer output(host_resource, buffer_size);
 
@@ -90,7 +90,7 @@ C2H_TEST("1d Copy", "[data_manipulation]")
 
   SECTION("Asymmetric size")
   {
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
     cudax::uninitialized_buffer<int, cuda::mr::host_accessible> host_buffer(host_resource, 1);
     cudax::fill_bytes(_stream, host_buffer, fill_byte);
 
@@ -133,7 +133,7 @@ void test_mdspan_copy_bytes(
 
 C2H_TEST("Mdspan copy", "[data_manipulation]")
 {
-  cudax::stream stream{cuda::device_ref{0}};
+  cuda::stream stream{cuda::device_ref{0}};
 
   SECTION("Different extents")
   {
@@ -154,7 +154,7 @@ C2H_TEST("Mdspan copy", "[data_manipulation]")
 
   SECTION("Launch transform")
   {
-    auto host_resource = cudax::legacy_pinned_memory_resource{};
+    auto host_resource = cuda::legacy_pinned_memory_resource{};
     auto mixed_extents =
       cuda::std::extents<size_t, 1024, cuda::std::dynamic_extent, 2, cuda::std::dynamic_extent>(1024, 2);
     [[maybe_unused]] auto static_extents = cuda::std::extents<size_t, 1024, 1024, 2, 2>();
@@ -171,7 +171,7 @@ C2H_TEST("Mdspan copy", "[data_manipulation]")
 
 C2H_TEST("Non exhaustive mdspan copy_bytes", "[data_manipulation]")
 {
-  cudax::stream stream{cuda::device_ref{0}};
+  cuda::stream stream{cuda::device_ref{0}};
   {
     auto fake_strided_mdspan = create_fake_strided_mdspan();
 

cudax/test/algorithm/copy.cu

@@ -12,10 +12,10 @@
 
 C2H_TEST("Fill", "[data_manipulation]")
 {
-  cudax::stream _stream{cuda::device_ref{0}};
+  cuda::stream _stream{cuda::device_ref{0}};
   SECTION("Host resource")
   {
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
     cudax::uninitialized_buffer<int, cuda::mr::device_accessible> buffer(host_resource, buffer_size);
 
     cudax::fill_bytes(_stream, buffer, fill_byte);
@@ -25,7 +25,7 @@ C2H_TEST("Fill", "[data_manipulation]")
 
   SECTION("Device resource")
   {
-    cudax::device_memory_pool_ref device_resource = cudax::device_default_memory_pool(cuda::device_ref{0});
+    cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});
     cudax::uninitialized_buffer<int, cuda::mr::device_accessible> buffer(device_resource, buffer_size);
     cudax::fill_bytes(_stream, buffer, fill_byte);
 
@@ -37,7 +37,7 @@ C2H_TEST("Fill", "[data_manipulation]")
   }
   SECTION("Launch transform")
   {
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
     cudax::weird_buffer buffer(host_resource, buffer_size);
 
     cudax::fill_bytes(_stream, buffer, fill_byte);
@@ -47,7 +47,7 @@ C2H_TEST("Fill", "[data_manipulation]")
 
 C2H_TEST("Mdspan Fill", "[data_manipulation]")
 {
-  cudax::stream stream{cuda::device_ref{0}};
+  cuda::stream stream{cuda::device_ref{0}};
   {
     cuda::std::dextents<size_t, 3> dynamic_extents{1, 2, 3};
     auto buffer = make_buffer_for_mdspan(dynamic_extents, 0);
@@ -65,7 +65,7 @@ C2H_TEST("Mdspan Fill", "[data_manipulation]")
     check_result_and_erase(stream, cuda::std::span(buffer.data(), buffer.size()));
   }
   {
-    cudax::legacy_pinned_memory_resource host_resource;
+    cuda::legacy_pinned_memory_resource host_resource;
     using static_extents = cuda::std::extents<size_t, 2, 3, 4>;
     auto size            = cuda::std::layout_left::mapping<static_extents>().required_span_size();
     cudax::weird_buffer<cuda::std::mdspan<int, static_extents>> buffer(host_resource, size);
@@ -77,7 +77,7 @@ C2H_TEST("Mdspan Fill", "[data_manipulation]")
 
 C2H_TEST("Non exhaustive mdspan fill_bytes", "[data_manipulation]")
 {
-  cudax::stream stream{cuda::device_ref{0}};
+  cuda::stream stream{cuda::device_ref{0}};
   {
     auto fake_strided_mdspan = create_fake_strided_mdspan();
 

cudax/test/algorithm/fill.cu

@@ -160,12 +160,12 @@ C2H_CCCLRT_TEST("make_async_buffer variants", "[container][async_buffer]")
   cudax::stream stream{cuda::device_ref{0}};
   const cudax::async_buffer<int, cuda::mr::device_accessible, other_property> input{
     stream,
-    cudax::device_default_memory_pool(cuda::device_ref{0}),
+    cuda::device_default_memory_pool(cuda::device_ref{0}),
     {int(1), int(42), int(1337), int(0), int(12), int(-1)}};
 
   // straight from a resource
-  auto buf = cuda::experimental::make_async_buffer(
-    input.stream(), cudax::device_default_memory_pool(cuda::device_ref{0}), input);
+  auto buf =
+    cuda::experimental::make_async_buffer(input.stream(), cuda::device_default_memory_pool(cuda::device_ref{0}), input);
   CUDAX_CHECK(equal_range(buf));
   static_assert(
     ::cuda::mr::synchronous_resource_with<typename decltype(buf)::__resource_t, cuda::mr::device_accessible>);
@@ -174,7 +174,7 @@ C2H_CCCLRT_TEST("make_async_buffer variants", "[container][async_buffer]")
   static_assert(!::cuda::mr::synchronous_resource_with<typename decltype(buf)::__resource_t, other_property>);
 
   auto buf2 = cuda::experimental::make_async_buffer<int, cuda::mr::device_accessible>(
-    input.stream(), {cudax::device_default_memory_pool(cuda::device_ref{0})}, input);
+    input.stream(), {cuda::device_default_memory_pool(cuda::device_ref{0})}, input);
   CUDAX_CHECK(equal_range(buf2));
   static_assert(
     ::cuda::mr::synchronous_resource_with<typename decltype(buf2)::__resource_t, cuda::mr::device_accessible>);
@@ -184,7 +184,7 @@ C2H_CCCLRT_TEST("make_async_buffer variants", "[container][async_buffer]")
 
   // from any resource
   auto any_res = cuda::mr::any_resource<cuda::mr::device_accessible, other_property>(
-    cudax::device_default_memory_pool(cuda::device_ref{0}));
+    cuda::device_default_memory_pool(cuda::device_ref{0}));
   auto buf3 = cudax::make_async_buffer(input.stream(), any_res, input);
   CUDAX_CHECK(equal_range(buf3));
   static_assert(
@@ -220,7 +220,7 @@ C2H_CCCLRT_TEST("make_async_buffer variants", "[container][async_buffer]")
     !::cuda::mr::synchronous_resource_with<typename decltype(buf6)::__resource_t, cuda::mr::host_accessible>);
 
   auto shared_res =
-    cudax::make_shared_resource<cudax::device_memory_pool_ref>(cudax::device_default_memory_pool(cuda::device_ref{0}));
+    cudax::make_shared_resource<cuda::device_memory_pool_ref>(cuda::device_default_memory_pool(cuda::device_ref{0}));
   auto buf7 = cudax::make_async_buffer(input.stream(), shared_res, input);
   CUDAX_CHECK(equal_range(buf7));
   static_assert(