Merge branch 'main' into deprecate_cuda_stream_ref_header

Author: davebayer

cudax/examples/stdexec_stream.cu

@@ -17,48 +17,54 @@
 #include <cuda_runtime_api.h>
 
 namespace cudax = cuda::experimental;
-namespace task  = cudax::execution;
+namespace ex    = cudax::execution;
 
 // This example demonstrates how to use the experimental CUDA implementation of
 // C++26's std::execution async tasking framework.
 
-struct say_hello
+int main()
 {
-  __device__ int operator()() const
+  try
   {
-    printf("Hello from lambda on device!\n");
-    return value;
-  }
+    auto tctx = ex::thread_context{};
+    auto sctx = ex::stream_context{cuda::device_ref{0}};
+    auto gpu  = sctx.get_scheduler();
 
-  int value;
-};
+    const auto bulk_shape = 10;
+    const auto bulk_fn    = [] __device__(const int index, int i) noexcept {
+      const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+      if (tid < bulk_shape)
+      {
+        printf("Hello from bulk task on device! index = %d, i = %d\n", index, i);
+      }
+    };
 
-__host__ void run()
-{
-  /*
-  try
-  {
-    task::thread_context tctx;
-    task::stream_context sctx{cuda::device_ref{0}};
-    auto sch = sctx.get_scheduler();
+    auto start =
+      // begin work on the GPU:
+      ex::schedule(gpu)
 
-    auto start = //
-      task::schedule(sch) // begin work on the GPU
-      | task::then(say_hello{42}) // enqueue a function object on the GPU
-      | task::then([] __device__(int i) noexcept -> int { // enqueue a lambda on the GPU
-          printf("Hello again from lambda on device! i = %d\n", i);
-          return i + 1;
+      // execute a device lambda on the GPU:
+      | ex::then([] __device__() noexcept -> int {
+          printf("Hello from lambda on device!\n");
+          return 42;
         })
-      | task::continues_on(tctx.get_scheduler()) // continue work on the CPU
-      | task::then([] __host__ __device__(int i) noexcept -> int { // run a lambda on the CPU
+
+      // do some parallel work on the GPU:
+      | ex::bulk(ex::par, bulk_shape, bulk_fn) //
+
+      // transfer execution back to the CPU:
+      | ex::continues_on(tctx.get_scheduler())
+
+      // execute a host/device lambda on the CPU:
+      | ex::then([] __host__ __device__(int i) noexcept -> int {
           NV_IF_TARGET(NV_IS_HOST,
                        (printf("Hello from lambda on host! i = %d\n", i);),
                        (printf("OOPS! still on the device! i = %d\n", i);))
-          return i;
+          return i + 1;
         });
 
     // run the task, wait for it to finish, and get the result
-    auto [i] = task::sync_wait(std::move(start)).value();
+    auto [i] = ex::sync_wait(std::move(start)).value();
     printf("All done on the host! result = %d\n", i);
   }
   catch (cuda::cuda_error const& e)
@@ -73,10 +79,4 @@ __host__ void run()
   {
     std::printf("Unknown exception\n");
   }
-  */
-}
-
-int main()
-{
-  run();
 }

cudax/include/cuda/experimental/__container/async_buffer.cuh

@@ -27,7 +27,6 @@
 
 #include <cuda/__memory_resource/get_memory_resource.h>
 #include <cuda/__memory_resource/properties.h>
-#include <cuda/__memory_resource/resource_ref.h>
 #include <cuda/__stream/get_stream.h>
 #include <cuda/std/__execution/env.h>
 #include <cuda/std/__iterator/concepts.h>
@@ -50,7 +49,6 @@
 #include <cuda/experimental/__execution/policy.cuh>
 #include <cuda/experimental/__launch/host_launch.cuh>
 #include <cuda/experimental/__memory_resource/any_resource.cuh>
-#include <cuda/experimental/__memory_resource/properties.cuh>
 #include <cuda/experimental/__utility/ensure_current_device.cuh>
 
 #include <cuda/std/__cccl/prologue.h>
@@ -96,7 +94,7 @@ public:
   using const_reverse_iterator = ::cuda::std::reverse_iterator<const_iterator>;
   using size_type              = ::cuda::std::size_t;
   using difference_type        = ::cuda::std::ptrdiff_t;
-  using properties_list        = ::cuda::experimental::properties_list<_Properties...>;
+  using properties_list        = ::cuda::mr::properties_list<_Properties...>;
 
   using __buffer_t       = ::cuda::experimental::uninitialized_async_buffer<_Tp, _Properties...>;
   using __resource_t     = ::cuda::experimental::any_resource<_Properties...>;
@@ -533,7 +531,7 @@ public:
 
   //! @brief Causes the buffer to be treated as a span when passed to cudax::launch.
   //! @pre The buffer must have the cuda::mr::device_accessible property.
-  template <class _DeviceAccessible = device_accessible>
+  template <class _DeviceAccessible = ::cuda::mr::device_accessible>
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, async_buffer& __self) noexcept
     _CCCL_TRAILING_REQUIRES(::cuda::std::span<_Tp>)(::cuda::std::__is_included_in_v<_DeviceAccessible, _Properties...>)
@@ -544,7 +542,7 @@ public:
 
   //! @brief Causes the buffer to be treated as a span when passed to cudax::launch
   //! @pre The buffer must have the cuda::mr::device_accessible property.
-  template <class _DeviceAccessible = device_accessible>
+  template <class _DeviceAccessible = ::cuda::mr::device_accessible>
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, const async_buffer& __self) noexcept _CCCL_TRAILING_REQUIRES(
     ::cuda::std::span<const _Tp>)(::cuda::std::__is_included_in_v<_DeviceAccessible, _Properties...>)
@@ -634,7 +632,7 @@ async_buffer<_Tp, _TargetProperties...> make_async_buffer(
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class... _SourceProperties, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND __has_default_queries<_Resource>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>)
 auto make_async_buffer(
   stream_ref __stream, _Resource&& __mr, const async_buffer<_Tp, _SourceProperties...>& __source, const _Env& __env = {})
 {
@@ -656,8 +654,8 @@ make_async_buffer(stream_ref __stream, any_resource<_Properties...> __mr, const
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(
-  ::cuda::mr::resource<_Resource> _CCCL_AND __has_default_queries<_Resource> _CCCL_AND __buffer_compatible_env<_Env>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource> _CCCL_AND
+                 __buffer_compatible_env<_Env>)
 auto make_async_buffer(stream_ref __stream, _Resource&& __mr, const _Env& __env = {})
 {
   using __buffer_type = __buffer_type_for_props<_Tp, typename ::cuda::std::decay_t<_Resource>::default_queries>;
@@ -669,7 +667,11 @@ _CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT
 // Size and value make function
 template <class _Tp, class... _Properties, class _Env = ::cuda::std::execution::env<>>
 async_buffer<_Tp, _Properties...> make_async_buffer(
-  stream_ref __stream, any_resource<_Properties...> __mr, size_t __size, const _Tp& __value, const _Env& __env = {})
+  stream_ref __stream,
+  any_resource<_Properties...> __mr,
+  size_t __size,
+  const _Tp& __value,
+  [[maybe_unused]] const _Env& __env = {})
 {
   auto __res = async_buffer<_Tp, _Properties...>{__stream, __mr, __size, no_init};
   __fill_n<_Tp, !::cuda::mr::__is_device_accessible<_Properties...>>(
@@ -678,7 +680,7 @@ async_buffer<_Tp, _Properties...> make_async_buffer(
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND __has_default_queries<_Resource>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>)
 auto make_async_buffer(
   stream_ref __stream, _Resource&& __mr, size_t __size, const _Tp& __value, [[maybe_unused]] const _Env& __env = {})
 {
@@ -705,7 +707,7 @@ async_buffer<_Tp, _Properties...> make_async_buffer(
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND __has_default_queries<_Resource>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>)
 auto make_async_buffer(
   stream_ref __stream, _Resource&& __mr, size_t __size, ::cuda::experimental::no_init_t, const _Env& __env = {})
 {
@@ -723,8 +725,8 @@ async_buffer<_Tp, _Properties...> make_async_buffer(
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Iter, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND
-                 __has_default_queries<_Resource> _CCCL_AND ::cuda::std::__has_forward_traversal<_Iter>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>
+                 _CCCL_AND ::cuda::std::__has_forward_traversal<_Iter>)
 auto make_async_buffer(stream_ref __stream, _Resource&& __mr, _Iter __first, _Iter __last, const _Env& __env = {})
 {
   using __buffer_type = __buffer_type_for_props<_Tp, typename ::cuda::std::decay_t<_Resource>::default_queries>;
@@ -743,7 +745,7 @@ async_buffer<_Tp, _Properties...> make_async_buffer(
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND __has_default_queries<_Resource>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>)
 auto make_async_buffer(
   stream_ref __stream, _Resource&& __mr, ::cuda::std::initializer_list<_Tp> __ilist, const _Env& __env = {})
 {
@@ -761,8 +763,8 @@ make_async_buffer(stream_ref __stream, any_resource<_Properties...> __mr, _Range
 }
 
 _CCCL_TEMPLATE(class _Tp, class _Resource, class _Range, class _Env = ::cuda::std::execution::env<>)
-_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND
-                 __has_default_queries<_Resource> _CCCL_AND ::cuda::std::ranges::forward_range<_Range>)
+_CCCL_REQUIRES(::cuda::mr::resource<_Resource> _CCCL_AND ::cuda::mr::__has_default_queries<_Resource>
+                 _CCCL_AND ::cuda::std::ranges::forward_range<_Range>)
 auto make_async_buffer(stream_ref __stream, _Resource&& __mr, _Range&& __range, const _Env& __env = {})
 {
   using __buffer_type = __buffer_type_for_props<_Tp, typename ::cuda::std::decay_t<_Resource>::default_queries>;

cudax/include/cuda/experimental/__container/uninitialized_async_buffer.cuh

@@ -33,7 +33,6 @@
 #include <cuda/std/span>
 
 #include <cuda/experimental/__memory_resource/any_resource.cuh>
-#include <cuda/experimental/__memory_resource/properties.cuh>
 
 #include <cuda/std/__cccl/prologue.h>
 
@@ -119,7 +118,7 @@ private:
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, uninitialized_async_buffer& __self) noexcept
     _CCCL_TRAILING_REQUIRES(::cuda::std::span<_Tp>)(
-      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<device_accessible, _Properties...>)
+      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<::cuda::mr::device_accessible, _Properties...>)
   {
     // TODO add auto synchronization
     return {__self.__get_data(), __self.size()};
@@ -131,7 +130,7 @@ private:
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, const uninitialized_async_buffer& __self) noexcept
     _CCCL_TRAILING_REQUIRES(::cuda::std::span<const _Tp>)(
-      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<device_accessible, _Properties...>)
+      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<::cuda::mr::device_accessible, _Properties...>)
   {
     // TODO add auto synchronization
     return {__self.__get_data(), __self.size()};
@@ -390,7 +389,7 @@ public:
 };
 
 template <class _Tp>
-using uninitialized_async_device_buffer = uninitialized_async_buffer<_Tp, device_accessible>;
+using uninitialized_async_device_buffer = uninitialized_async_buffer<_Tp, ::cuda::mr::device_accessible>;
 
 } // namespace cuda::experimental
 

cudax/include/cuda/experimental/__container/uninitialized_buffer.cuh

@@ -31,7 +31,6 @@
 #include <cuda/std/span>
 
 #include <cuda/experimental/__memory_resource/any_resource.cuh>
-#include <cuda/experimental/__memory_resource/properties.cuh>
 
 #include <cuda/std/__cccl/prologue.h>
 
@@ -108,7 +107,7 @@ private:
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, uninitialized_buffer& __self) noexcept
     _CCCL_TRAILING_REQUIRES(::cuda::std::span<_Tp>)(
-      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<device_accessible, _Properties...>)
+      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<::cuda::mr::device_accessible, _Properties...>)
   {
     return {__self.__get_data(), __self.size()};
   }
@@ -119,7 +118,7 @@ private:
   [[nodiscard]] _CCCL_HIDE_FROM_ABI friend auto
   transform_device_argument(::cuda::stream_ref, const uninitialized_buffer& __self) noexcept
     _CCCL_TRAILING_REQUIRES(::cuda::std::span<const _Tp>)(
-      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<device_accessible, _Properties...>)
+      ::cuda::std::same_as<_Tp, _Tp2>&& ::cuda::std::__is_included_in_v<::cuda::mr::device_accessible, _Properties...>)
   {
     return {__self.__get_data(), __self.size()};
   }
@@ -289,7 +288,7 @@ public:
 };
 
 template <class _Tp>
-using uninitialized_device_buffer = uninitialized_buffer<_Tp, device_accessible>;
+using uninitialized_device_buffer = uninitialized_buffer<_Tp, ::cuda::mr::device_accessible>;
 
 } // namespace cuda::experimental
 

cudax/include/cuda/experimental/__execution/apply_sender.cuh

@@ -63,7 +63,7 @@ public:
   //! @throws Any exception thrown by the underlying domain's `apply_sender`.
   _CCCL_EXEC_CHECK_DISABLE
   template <class _Domain, class _Tag, class _Sndr, class... _Args>
-  _CCCL_NODEBUG_API constexpr auto operator()(_Domain, _Tag, _Sndr&& __sndr, _Args&&... __args) const
+  _CCCL_API constexpr auto operator()(_Domain, _Tag, _Sndr&& __sndr, _Args&&... __args) const
     noexcept(noexcept(__apply_domain_t<_Domain, _Tag, _Sndr, _Args...>{}.apply_sender(
       _Tag{}, static_cast<_Sndr&&>(__sndr), static_cast<_Args&&>(__args)...)))
       -> __apply_sender_result_t<__apply_domain_t<_Domain, _Tag, _Sndr, _Args...>, _Tag, _Sndr, _Args...>

cudax/include/cuda/experimental/__execution/bulk.cuh

@@ -68,13 +68,22 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __state_t
 template <class _Sndr, class _Shape>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT __attrs_t
 {
-  [[nodiscard]] _CCCL_HOST_API constexpr auto query(get_launch_config_t) const noexcept
+  [[nodiscard]] _CCCL_HOST_API static constexpr auto __get_launch_config(_Shape __shape) noexcept
   {
     constexpr int __block_threads = 256;
-    const int __grid_blocks       = ::cuda::ceil_div(static_cast<int>(__shape_), __block_threads);
+    const int __grid_blocks       = ::cuda::ceil_div(static_cast<int>(__shape), __block_threads);
     return experimental::make_config(block_dims<__block_threads>(), grid_dims(__grid_blocks));
   }
 
+  using __launch_config_t = decltype(__get_launch_config(_Shape()));
+
+  [[nodiscard]] _CCCL_API constexpr auto query(get_launch_config_t) const noexcept -> __launch_config_t
+  {
+    NV_IF_TARGET(NV_IS_HOST,
+                 (return __get_launch_config(__shape_);),
+                 (_CCCL_ASSERT(false, "cannot get a launch configuration from device"); ::cuda::std::terminate();))
+  }
+
   _CCCL_EXEC_CHECK_DISABLE
   _CCCL_TEMPLATE(class _Query, class... _Args)
   _CCCL_REQUIRES(__forwarding_query<_Query> _CCCL_AND __queryable_with<env_of_t<_Sndr>, _Query, _Args...>)
@@ -152,7 +161,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __bulk_t
       execution::set_stopped(static_cast<_Rcvr&&>(__state_->__rcvr_));
     }
 
-    [[nodiscard]] _CCCL_NODEBUG_API constexpr auto get_env() const noexcept -> __fwd_env_t<env_of_t<_Rcvr>>
+    [[nodiscard]] _CCCL_API constexpr auto get_env() const noexcept -> __fwd_env_t<env_of_t<_Rcvr>>
     {
       return __fwd_env(execution::get_env(__state_->__rcvr_));
     }
@@ -188,7 +197,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __bulk_t
   struct _CCCL_TYPE_VISIBILITY_DEFAULT __closure_base_t
   {
     template <class _Sndr>
-    [[nodiscard]] _CCCL_NODEBUG_API friend constexpr auto operator|(_Sndr&& __sndr, __closure_base_t __self)
+    [[nodiscard]] _CCCL_API friend constexpr auto operator|(_Sndr&& __sndr, __closure_base_t __self)
     {
       static_assert(__is_sender<_Sndr>);
 
@@ -202,7 +211,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __bulk_t
         {{}, static_cast<__closure_base_t&&>(__self), static_cast<_Sndr&&>(__sndr)}};
     }
 
-    _CCCL_NO_UNIQUE_ADDRESS _Policy __policy_;
+    /*_CCCL_NO_UNIQUE_ADDRESS*/ _Policy __policy_;
     _Shape __shape_;
     _Fn __fn_;
   };
@@ -250,7 +259,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __bulk_t
       return {__state_.__shape_, __sndr_};
     }
 
-    _CCCL_NO_UNIQUE_ADDRESS _BulkTag __tag_;
+    /*_CCCL_NO_UNIQUE_ADDRESS*/ _BulkTag __tag_;
     __closure_base_t<_Policy, _Shape, _Fn> __state_;
     _Sndr __sndr_;
   };
@@ -267,7 +276,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT __bulk_t
   // This function call operator creates a sender adaptor closure object that can appear
   // on the right-hand side of a pipe operator, like: sndr | bulk(par, shape, fn).
   template <class _Policy, class _Shape, class _Fn>
-  [[nodiscard]] _CCCL_NODEBUG_API auto operator()(_Policy __policy, _Shape __shape, _Fn __fn) const
+  [[nodiscard]] _CCCL_API auto operator()(_Policy __policy, _Shape __shape, _Fn __fn) const
   {
     static_assert(::cuda::std::integral<_Shape>);
     static_assert(::cuda::std::is_execution_policy_v<_Policy>);
@@ -389,7 +398,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT bulk_t : __bulk_t<bulk_t>
   {
     _CCCL_EXEC_CHECK_DISABLE
     template <class... _Ts>
-    _CCCL_NODEBUG_API void operator()(_Shape __begin, _Shape __end, _Ts&&... __values) noexcept(
+    _CCCL_API void operator()(_Shape __begin, _Shape __end, _Ts&&... __values) noexcept(
       __nothrow_callable<_Fn&, _Shape, decltype(__values)&...>)
     {
       for (; __begin != __end; ++__begin)

cudax/include/cuda/experimental/__execution/completion_behavior.cuh

@@ -139,7 +139,7 @@ struct get_completion_behavior_t
     return __attrs.query(*this, __env);
   }
 
-  [[nodiscard]] _CCCL_NODEBUG_API static constexpr auto query(forwarding_query_t) noexcept -> bool
+  [[nodiscard]] _CCCL_API static constexpr auto query(forwarding_query_t) noexcept -> bool
   {
     return true;
   }