[CUDAX->libcu++] Move device APIs to libcu++ (#5279)

* Move device APIs to libcudacxx

* Missing semicolon

* Exclude all device APIs for NVRTC

Author: pciolkosz

cudax/examples/async_buffer_add.cu

@@ -48,13 +48,13 @@ struct generator
 int main()
 {
   // A CUDA stream on which to execute the vector addition kernel
-  cudax::stream stream{cudax::device_ref{0}};
+  cudax::stream stream{cuda::device_ref{0}};
 
   // The execution policy we want to use to run all work on the same stream
   auto policy = thrust::cuda::par_nosync.on(stream.get());
 
   // An environment we use to pass all necessary information to the containers
-  cudax::env_t<cuda::mr::device_accessible> env{cudax::device_memory_resource{cudax::device_ref{0}}, stream};
+  cudax::env_t<cuda::mr::device_accessible> env{cudax::device_memory_resource{cuda::device_ref{0}}, stream};
 
   // Allocate the two inputs and output, but do not zero initialize via `cudax::no_init`
   cudax::async_device_buffer<float> A{env, numElements, cudax::no_init};

cudax/examples/cub_reduce.cu

@@ -25,10 +25,10 @@ int main()
   constexpr int num_items = 50000;
 
   // A CUDA stream on which to execute the reduction
-  cudax::stream stream{cudax::devices[0]};
+  cudax::stream stream{cuda::devices[0]};
 
   // An environment we use to pass all necessary information to the containers
-  cudax::env_t<cuda::mr::device_accessible> env{cudax::device_memory_resource{cudax::devices[0]}, stream};
+  cudax::env_t<cuda::mr::device_accessible> env{cudax::device_memory_resource{cuda::devices[0]}, stream};
 
   // Allocate input and output, but do not zero initialize output (`cudax::no_init`)
   cudax::async_device_buffer<int> d_in{env, num_items, 1};

cudax/examples/simple_p2p.cu

@@ -30,11 +30,11 @@
  * Unified Virtual Address Space (UVA) features.
  */
 
+#include <cuda/devices>
 #include <cuda/memory_resource>
 
 #include <cuda/experimental/algorithm.cuh>
 #include <cuda/experimental/container.cuh>
-#include <cuda/experimental/device.cuh>
 #include <cuda/experimental/launch.cuh>
 #include <cuda/experimental/memory_resource.cuh>
 
@@ -61,9 +61,9 @@ void print_peer_accessibility()
   // Check possibility for peer access
   printf("\nChecking GPU(s) for support of peer to peer memory access...\n");
 
-  for (auto& dev_i : cudax::devices)
+  for (auto& dev_i : cuda::devices)
   {
-    for (auto& dev_j : cudax::devices)
+    for (auto& dev_j : cuda::devices)
     {
       if (dev_i != dev_j)
       {
@@ -112,8 +112,8 @@ template <typename BufferType>
 void test_cross_device_access_from_kernel(
   cudax::stream_ref dev0_stream, cudax::stream_ref dev1_stream, BufferType& dev0_buffer, BufferType& dev1_buffer)
 {
-  cudax::device_ref dev0 = dev0_stream.device();
-  cudax::device_ref dev1 = dev1_stream.device();
+  cuda::device_ref dev0 = dev0_stream.device();
+  cuda::device_ref dev1 = dev1_stream.device();
 
   // Prepare host buffer and copy to GPU 0
   printf("Preparing host buffer and copy to GPU%d...\n", dev0.get());
@@ -184,9 +184,9 @@ try
 
   // Number of GPUs
   printf("Checking for multiple GPUs...\n");
-  printf("CUDA-capable device count: %zu\n", cudax::devices.size());
+  printf("CUDA-capable device count: %zu\n", cuda::devices.size());
 
-  if (cudax::devices.size() < 2)
+  if (cuda::devices.size() < 2)
   {
     printf("Two or more GPUs with Peer-to-Peer access capability are required for %s.\n", argv[0]);
     printf("Waiving test.\n");
@@ -197,8 +197,8 @@ try
   print_peer_accessibility();
 
   // But use a shorthand to find all peers of a device
-  std::vector<cudax::device_ref> peers;
-  for (auto& dev : cudax::devices)
+  std::vector<cuda::device_ref> peers;
+  for (auto& dev : cuda::devices)
   {
     peers = dev.peer_devices();
     if (peers.size() != 0)

cudax/examples/stdexec_stream.cu

@@ -38,7 +38,7 @@ __host__ void run()
   try
   {
     task::thread_context tctx;
-    task::stream_context sctx{cuda::experimental::device_ref{0}};
+    task::stream_context sctx{cuda::device_ref{0}};
     auto sch = sctx.get_scheduler();
 
     auto start = //

cudax/examples/vector_add.cu

@@ -72,7 +72,7 @@ int main(void)
 try
 {
   // A CUDA stream on which to execute the vector addition kernel
-  cudax::stream stream(cudax::devices[0]);
+  cudax::stream stream(cuda::devices[0]);
 
   // Print the vector length to be used, and compute its size
   int numElements = 50000;

cudax/include/cuda/experimental/__device/logical_device.cuh

@@ -21,7 +21,8 @@
 #  pragma system_header
 #endif // no system header
 
-#include <cuda/experimental/__device/all_devices.cuh>
+#include <cuda/__device/all_devices.h>
+
 #include <cuda/experimental/__green_context/green_ctx.cuh>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -82,7 +83,7 @@ public:
   // More of a micro-optimization, we can also remove this (depending if we keep device_ref)
   //!
   //! Constructing a logical_device for a given device has a side effect of initializing that device
-  logical_device(const ::cuda::experimental::physical_device& __dev)
+  logical_device(const ::cuda::physical_device& __dev)
       : __dev_id(__dev.get())
       , __kind(kinds::device)
       , __ctx(__dev.primary_context())

cudax/include/cuda/experimental/__execution/stream/context.cuh

@@ -21,9 +21,9 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/__device/device_ref.h>
 #include <cuda/__stream/get_stream.h>
 
-#include <cuda/experimental/__device/device_ref.cuh>
 #include <cuda/experimental/__execution/stream/scheduler.cuh>
 #include <cuda/experimental/stream.cuh>
 

cudax/include/cuda/experimental/__green_context/green_ctx.cuh

@@ -21,18 +21,16 @@
 #  pragma system_header
 #endif // no system header
 
+#include <cuda/__device/all_devices.h>
 #include <cuda/__driver/driver_api.h>
 #include <cuda/std/__cuda/api_wrapper.h>
 #include <cuda/std/utility>
 
-#include <cuda/experimental/__device/all_devices.cuh>
-
 #include <cuda/std/__cccl/prologue.h>
 
 #if _CCCL_CTK_AT_LEAST(12, 5)
 namespace cuda::experimental
 {
-struct device_ref;
 
 struct green_context
 {

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

@@ -62,8 +62,7 @@ public:
   //! @throws cuda_error if the CUDA version does not support ``cudaMallocAsync``.
   //! @param __device_id The device id of the device the stream pool is constructed on.
   //! @param __pool_properties Optional, additional properties of the pool to be created.
-  explicit device_memory_pool(const ::cuda::experimental::device_ref __device_id,
-                              memory_pool_properties __properties = {})
+  explicit device_memory_pool(const ::cuda::device_ref __device_id, memory_pool_properties __properties = {})
       : __memory_pool_base(__memory_location_type::__device, __properties, __device_id.get())
   {}
 

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

@@ -77,7 +77,7 @@ private:
 public:
   //! @brief Constructs a device_memory_resource using the default \c cudaMemPool_t of a given device.
   //! @throws cuda_error if retrieving the default \c cudaMemPool_t fails.
-  explicit device_memory_resource(::cuda::experimental::device_ref __device)
+  explicit device_memory_resource(::cuda::device_ref __device)
       : __memory_resource_base(__get_default_device_mem_pool(__device.get()))
   {}