really fix the perf regression in the maxwell_gpu_s example

ericniebler · ericniebler · commit f5d7bccfa25d · 2026-03-07T11:50:20.000-08:00
diff --git a/examples/nvexec/maxwell/snr.cuh b/examples/nvexec/maxwell/snr.cuh
@@ -60,19 +60,19 @@ STDEXEC_PRAGMA_IGNORE_GNU("-Wmissing-braces")
 
 namespace ex = stdexec;
 
-namespace repeat_n_detail
+namespace _repeat_n
 {
   template <class Sender, class Closure>
-  struct repeat_n_sender_t;
-}  // namespace repeat_n_detail
+  struct sender;
+}  // namespace _repeat_n
 
 struct repeat_n_t
 {
   template <ex::sender Sender, ex::__sender_adaptor_closure Closure>
   auto operator()(Sender __sndr, std::size_t n, Closure closure) const noexcept
-    -> repeat_n_detail::repeat_n_sender_t<Sender, Closure>
+    -> _repeat_n::sender<Sender, Closure>
   {
-    return repeat_n_detail::repeat_n_sender_t<Sender, Closure>{
+    return _repeat_n::sender<Sender, Closure>{
       {},
       {closure, n},
       std::move(__sndr)
@@ -88,9 +88,8 @@ struct repeat_n_t
 
 inline constexpr repeat_n_t repeat_n{};
 
-namespace repeat_n_detail
+namespace _repeat_n
 {
-
   template <class OpT>
   class receiver_2_t
   {
@@ -245,11 +244,11 @@ namespace repeat_n_detail
   };
 
   template <class Sender, class Closure>
-  struct repeat_n_sender_t
+  struct sender
   {
     using sender_concept = ex::sender_t;
 
-    template <class>
+    template <class, class Env>
     static consteval auto get_completion_signatures() noexcept
     {
       return ex::completion_signatures<
@@ -260,11 +259,11 @@ namespace repeat_n_detail
         >();
     }
 
-    template <ex::__decays_to<repeat_n_sender_t> Self, ex::receiver Receiver>
+    template <ex::__decays_to<sender> Self, ex::receiver Receiver>
     STDEXEC_EXPLICIT_THIS_BEGIN(auto connect)(this Self&& self, Receiver r)
-      -> repeat_n_detail::operation_state_t<Sender, Closure, Receiver>
+      -> _repeat_n::operation_state_t<Sender, Closure, Receiver>
     {
-      return repeat_n_detail::operation_state_t<Sender, Closure, Receiver>(
+      return _repeat_n::operation_state_t<Sender, Closure, Receiver>(
         static_cast<Self&&>(self).sender_,
         static_cast<Self&&>(self).data_.first,
         static_cast<Receiver&&>(r),
@@ -283,26 +282,25 @@ namespace repeat_n_detail
     std::pair<Closure, std::size_t> data_;
     Sender                          sender_;
   };
-}  // namespace repeat_n_detail
+}  // namespace _repeat_n
 
 namespace STDEXEC
 {
   template <class Sender, class Closure>
-  inline constexpr std::size_t
-    __structured_binding_size_v<repeat_n_detail::repeat_n_sender_t<Sender, Closure>> = 3;
+  inline constexpr std::size_t __structured_binding_size_v<_repeat_n::sender<Sender, Closure>> = 3;
 }  // namespace STDEXEC
 
 #if STDEXEC_CUDA_COMPILATION()
 // A CUDA stream implementation of repeat_n
 namespace nv::execution::_strm
 {
-  namespace repeat_n
+  namespace _repeat_n
   {
     template <class OpT>
     class receiver_2_t : public stream_receiver_base
     {
-      using Sender   = OpT::PredSender;
-      using Receiver = OpT::Receiver;
+      using Sender   = OpT::child_t;
+      using Receiver = OpT::receiver_t;
 
       OpT& op_state_;
 
@@ -353,11 +351,15 @@ namespace nv::execution::_strm
     template <class OpT>
     class receiver_1_t : public stream_receiver_base
     {
-      using Receiver = OpT::Receiver;
+      using Receiver = OpT::receiver_t;
 
       OpT& op_state_;
 
      public:
+      explicit receiver_1_t(OpT& op_state)
+        : op_state_(op_state)
+      {}
+
       void set_value() noexcept
       {
         using inner_op_state_t = OpT::inner_op_state_t;
@@ -394,15 +396,13 @@ namespace nv::execution::_strm
       {
         return op_state_.make_env();
       }
-
-      explicit receiver_1_t(OpT& op_state)
-        : op_state_(op_state)
-      {}
     };
 
     template <class PredSender, class Closure, class Receiver>
     struct operation_state_t : _strm::opstate_base<Receiver>
     {
+      using receiver_t  = Receiver;
+      using child_t     = PredSender;
       using Scheduler   = std::invoke_result_t<ex::get_completion_scheduler_t<ex::set_value_t>,
                                                ex::env_of_t<PredSender>,
                                                ex::env_of_t<Receiver>>;
@@ -445,7 +445,7 @@ namespace nv::execution::_strm
             static_cast<Receiver&&>(rcvr),
             ex::get_completion_scheduler<ex::set_value_t>(ex::get_env(pred_sender),
                                                           ex::get_env(rcvr))
-              .context_state_)
+              .ctx_)
         , scheduler_(ex::get_completion_scheduler<ex::set_value_t>(ex::get_env(pred_sender),
                                                                    ex::get_env(rcvr)))
         , closure_(closure)
@@ -458,7 +458,7 @@ namespace nv::execution::_strm
     };
 
     template <class Sender, class Closure>
-    struct sender_t
+    struct sender
     {
       using sender_concept = ex::sender_t;
 
@@ -467,12 +467,12 @@ namespace nv::execution::_strm
                                                               ex::set_error_t(std::exception_ptr),
                                                               ex::set_error_t(cudaError_t)>;
 
-      template <ex::__decays_to<sender_t> Self, ex::receiver Receiver>
+      template <ex::__decays_to<sender> Self, ex::receiver Receiver>
         requires(ex::sender_to<Sender, Receiver>)
       STDEXEC_EXPLICIT_THIS_BEGIN(auto connect)(this Self&& self, Receiver r)
-        -> nvexec::_strm::repeat_n::operation_state_t<Sender, Closure, Receiver>
+        -> nvexec::_strm::_repeat_n::operation_state_t<Sender, Closure, Receiver>
       {
-        return nvexec::_strm::repeat_n::operation_state_t<Sender, Closure, Receiver>(
+        return nvexec::_strm::_repeat_n::operation_state_t<Sender, Closure, Receiver>(
           static_cast<Self&&>(self).sender_,
           static_cast<Self&&>(self).closure_,
           static_cast<Receiver&&>(r),
@@ -490,21 +490,20 @@ namespace nv::execution::_strm
       Closure     closure_;
       std::size_t n_{};
     };
-  }  // namespace repeat_n
+  }  // namespace _repeat_n
 
   template <>
-  struct transform_sender_for<repeat_n_t>
+  struct transform_sender_for<::repeat_n_t>
   {
     template <class Env, class Data, class Sender>
-    auto operator()(Env const &, ex::__ignore, Data&& data, Sender sndr) const
+    auto operator()(Env const &, ::repeat_n_t, Data&& data, Sender sndr) const
     {
-      static_assert(sizeof(Env) == 0);
-      auto& [closure, n] = data;
-      using closure_t    = decltype(closure);
+      auto& [closure, count] = data;
+      using closure_t        = decltype(closure);
 
-      return repeat_n::sender_t<Sender, closure_t>(static_cast<Sender&&>(sndr),
-                                                   ex::__forward_like<Data>(closure),
-                                                   n);
+      return _strm::_repeat_n::sender<Sender, closure_t>(static_cast<Sender&&>(sndr),
+                                                         ex::__forward_like<Data>(closure),
+                                                         count);
     }
   };
 }  // namespace nv::execution::_strm
@@ -528,8 +527,8 @@ auto maxwell_eqs_snr(float                dt,
                      ex::scheduler auto&& computer)
 {
   return ex::just()
-       | repeat_n(n_iterations,
-                  ex::on(computer,
+       | ex::on(computer,
+                repeat_n(n_iterations,
                          ex::bulk(ex::par, accessor.cells, update_h(accessor))
                            | ex::bulk(ex::par, accessor.cells, update_e(time, dt, accessor))))
        | ex::then(dump_vtk(write_results, accessor));
diff --git a/include/exec/detail/basic_sequence.hpp b/include/exec/detail/basic_sequence.hpp
@@ -113,6 +113,9 @@ namespace exec = experimental::execution;
 
 namespace STDEXEC::__detail
 {
+  template <auto _DescriptorFn>
+  extern __result_of<_DescriptorFn> __desc_of_v<exec::__seqexpr<_DescriptorFn>>;
+
   template <auto _DescriptorFn>
   extern __declfn_t<__minvoke<__result_of<_DescriptorFn>, __q<exec::__basic_sequence_sender_t>>>
     __demangle_v<exec::__seqexpr<_DescriptorFn>>;
diff --git a/include/exec/static_thread_pool.hpp b/include/exec/static_thread_pool.hpp
@@ -1281,6 +1281,14 @@ namespace experimental::execution
       using _bulk_opstate_t =
         _static_thread_pool::_bulk_opstate<Parallelize, Shape, Fun, Sender, Receiver>;
 
+      explicit _bulk_sender(_static_thread_pool& pool, Sender sndr, Shape shape, Fun fun)
+        noexcept(__nothrow_move_constructible<Sender, Fun>)
+        : pool_(pool)
+        , sndr_(static_cast<Sender&&>(sndr))
+        , shape_(shape)
+        , fun_(static_cast<Fun&&>(fun))
+      {}
+
       template <__decays_to<_bulk_sender> Self, receiver Receiver>
         requires receiver_of<Receiver, _completions_t<Self, env_of_t<Receiver>>>
       STDEXEC_EXPLICIT_THIS_BEGIN(auto connect)(this Self&& self, Receiver rcvr)
@@ -1310,6 +1318,7 @@ namespace experimental::execution
         return STDEXEC::get_env(sndr_);
       }
 
+     private:
       _static_thread_pool& pool_;
       Sender               sndr_;
       Shape                shape_;
diff --git a/include/nvexec/stream/common.cuh b/include/nvexec/stream/common.cuh
@@ -98,10 +98,20 @@ namespace nv::execution
   struct stream_domain : STDEXEC::default_domain
   {
     template <::exec::sender_for Sender, class Tag = STDEXEC::tag_of_t<Sender>, class Env>
-      requires STDEXEC::__applicable<_strm::transform_sender_for<Tag>, Sender, Env const &>
+      requires STDEXEC::__callable<STDEXEC::__structured_apply_t,
+                                   _strm::transform_sender_for<Tag>,
+                                   Sender,
+                                   Env const &>
     static auto transform_sender(STDEXEC::set_value_t, Sender&& sndr, Env const & env)
+      noexcept(STDEXEC::__nothrow_callable<STDEXEC::__structured_apply_t,
+                                           _strm::transform_sender_for<Tag>,
+                                           Sender,
+                                           Env const &>)
+
     {
-      return STDEXEC::__apply(_strm::transform_sender_for<Tag>{}, static_cast<Sender&&>(sndr), env);
+      return STDEXEC::__structured_apply(_strm::transform_sender_for<Tag>{},
+                                         static_cast<Sender&&>(sndr),
+                                         env);
     }
 
     template <class Tag, STDEXEC::sender Sender, class... Args>
@@ -115,7 +125,6 @@ namespace nv::execution
 
   namespace _strm
   {
-
 #if STDEXEC_HAS_BUILTIN(__is_reference)
     template <class... Ts>
     concept trivially_copyable = ((STDEXEC_IS_TRIVIALLY_COPYABLE(Ts) || __is_reference(Ts)) && ...);
@@ -701,8 +710,9 @@ namespace nv::execution
         else
         {
           // pass a cudaError_t by value:
-          continuation_kernel<OuterReceiver, Error>
-            <<<1, 1, 0, get_stream()>>>(static_cast<OuterReceiver&&>(rcvr_), set_error_t(), status);
+          continuation_kernel<<<1, 1, 0, get_stream()>>>(static_cast<OuterReceiver&&>(rcvr_),
+                                                         set_error_t(),
+                                                         cudaError_t(status));
         }
       }
 
diff --git a/include/nvexec/stream/schedule_from.cuh b/include/nvexec/stream/schedule_from.cuh
@@ -38,8 +38,8 @@ namespace nv::execution
       template <class Sender, class Receiver>
       struct opstate : _strm::opstate_base<Receiver>
       {
-        using env_t = _strm::opstate_base<Receiver>::env_t;
         struct receiver;
+        using env_t              = _strm::opstate_base<Receiver>::env_t;
         using variant_t          = variant_storage_t<Sender, env_t>;
         using task_t             = continuation_task<receiver, variant_t>;
         using enqueue_receiver_t = stream_enqueue_receiver<env_t, variant_t>;
@@ -75,7 +75,7 @@ namespace nv::execution
           opstate& opstate_;
         };
 
-        opstate(Sender&& sender, Receiver&& rcvr, context ctx)
+        opstate(Sender&& sndr, Receiver&& rcvr, context ctx)
           : _strm::opstate_base<Receiver>(static_cast<Receiver&&>(rcvr), ctx)
           , ctx_(ctx)
           , storage_(host_allocate<variant_t>(this->status_, ctx.pinned_resource_))
@@ -88,7 +88,7 @@ namespace nv::execution
                     .release())
           , env_(host_allocate(this->status_, ctx_.pinned_resource_, this->make_env()))
           , inner_op_{
-              connect(static_cast<Sender&&>(sender),
+              connect(static_cast<Sender&&>(sndr),
                       enqueue_receiver_t{env_.get(), storage_.get(), task_, ctx_.hub_->producer()})}
         {
           if (this->status_ == cudaSuccess)
diff --git a/include/stdexec/__detail/__config.hpp b/include/stdexec/__detail/__config.hpp
@@ -801,16 +801,16 @@ namespace STDEXEC
                                   _NAME))
 
 #  define STDEXEC_EXPLICIT_THIS_EAT_this
-#  define STDEXEC_EXPLICIT_THIS_MANGLE_auto auto STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN
-#  define STDEXEC_EXPLICIT_THIS_MANGLE_void void STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN
-#  define STDEXEC_EXPLICIT_THIS_MANGLE_bool bool STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN
+#  define STDEXEC_EXPLICIT_THIS_MANGLE_auto auto STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN()
+#  define STDEXEC_EXPLICIT_THIS_MANGLE_void void STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN()
+#  define STDEXEC_EXPLICIT_THIS_MANGLE_bool bool STDEXEC_EXPLICIT_THIS_MANGLE STDEXEC_PP_LPAREN()
 
 #  define STDEXEC_EXPLICIT_THIS_ARGS(...)                                                       \
-    STDEXEC_PP_CAT(STDEXEC_EXPLICIT_THIS_EAT_, __VA_ARGS__) STDEXEC_PP_RPAREN
+    STDEXEC_PP_CAT(STDEXEC_EXPLICIT_THIS_EAT_, __VA_ARGS__) STDEXEC_PP_RPAREN()
 
 #  define STDEXEC_EXPLICIT_THIS_BEGIN(...)                                                      \
     static STDEXEC_PP_EXPAND(STDEXEC_PP_CAT(STDEXEC_EXPLICIT_THIS_MANGLE_, __VA_ARGS__)         \
-                             STDEXEC_PP_RPAREN) STDEXEC_PP_LPAREN STDEXEC_EXPLICIT_THIS_ARGS
+                             STDEXEC_PP_RPAREN()) STDEXEC_PP_LPAREN() STDEXEC_EXPLICIT_THIS_ARGS
 
 #  define STDEXEC_EXPLICIT_THIS_END(_NAME)                                                      \
     template <class... _Ts>                                                                     \
diff --git a/include/stdexec/__detail/__preprocessor.hpp b/include/stdexec/__detail/__preprocessor.hpp
@@ -18,10 +18,10 @@
 #define STDEXEC_PP_STRINGIZE_I(...) #__VA_ARGS__
 #define STDEXEC_PP_STRINGIZE(...)   STDEXEC_PP_STRINGIZE_I(__VA_ARGS__)
 
-#define STDEXEC_PP_LPAREN (
-#define STDEXEC_PP_RPAREN )
-#define STDEXEC_PP_PARENS            ()
-#define STDEXEC_PP_COMMA             ,
+#define STDEXEC_PP_LPAREN() (
+#define STDEXEC_PP_RPAREN() )
+#define STDEXEC_PP_PARENS() ()
+#define STDEXEC_PP_COMMA()  ,
 
 #define STDEXEC_PP_CAT_I(_XP, ...)   _XP##__VA_ARGS__
 #define STDEXEC_PP_CAT(_XP, ...)     STDEXEC_PP_CAT_I(_XP, __VA_ARGS__)
@@ -83,7 +83,7 @@
 
 #define STDEXEC_PP_FOR_EACH_AGAIN() STDEXEC_PP_FOR_EACH_HELPER
 #define STDEXEC_PP_FOR_EACH_HELPER(_MACRO, _A1, ...)                                               \
-  _MACRO(_A1) __VA_OPT__(STDEXEC_PP_FOR_EACH_AGAIN STDEXEC_PP_PARENS(_MACRO, __VA_ARGS__)) /**/
+  _MACRO(_A1) __VA_OPT__(STDEXEC_PP_FOR_EACH_AGAIN STDEXEC_PP_PARENS()(_MACRO, __VA_ARGS__)) /**/
 #define STDEXEC_PP_FOR_EACH(_MACRO, ...)                                                           \
   __VA_OPT__(STDEXEC_PP_EXPAND_R(STDEXEC_PP_FOR_EACH_HELPER(_MACRO, __VA_ARGS__)))
 
@@ -94,7 +94,7 @@
 #define STDEXEC_PP_BACK_AGAIN()      STDEXEC_PP_BACK_I
 #define STDEXEC_PP_BACK_I(_A1, ...)                                                                \
   STDEXEC_PP_FRONT(__VA_OPT__(, ) _A1, )                                                           \
-  __VA_OPT__(STDEXEC_PP_BACK_AGAIN STDEXEC_PP_PARENS(__VA_ARGS__))
+  __VA_OPT__(STDEXEC_PP_BACK_AGAIN STDEXEC_PP_PARENS()(__VA_ARGS__))
 #define STDEXEC_PP_BACK(...)                 __VA_OPT__(STDEXEC_PP_EXPAND_R(STDEXEC_PP_BACK_I(__VA_ARGS__)))
 
 #define STDEXEC_PP_TAIL(_IGN, ...)           __VA_ARGS__
diff --git a/include/stdexec/__detail/__sender_introspection.hpp b/include/stdexec/__detail/__sender_introspection.hpp

Original file line number	Diff line number	Diff line change
`@@ -113,6 +113,9 @@ namespace exec = experimental::execution;`
`113`	`113`
`114`	`114`	`namespace STDEXEC::__detail`
`115`	`115`	`{`
	`116`	`+ template <auto _DescriptorFn>`
	`117`	`+ extern __result_of<_DescriptorFn> __desc_of_v<exec::__seqexpr<_DescriptorFn>>;`
	`118`	`+`
`116`	`119`	`template <auto _DescriptorFn>`
`117`	`120`	`extern __declfn_t<__minvoke<__result_of<_DescriptorFn>, __q<exec::__basic_sequence_sender_t>>>`
`118`	`121`	`__demangle_v<exec::__seqexpr<_DescriptorFn>>;`
Original file line number	Diff line number	Diff line change
`@@ -98,10 +98,20 @@ namespace nv::execution`
`98`	`98`	`struct stream_domain : STDEXEC::default_domain`
`99`	`99`	`{`
`100`	`100`	`template <::exec::sender_for Sender, class Tag = STDEXEC::tag_of_t<Sender>, class Env>`
`101`		`- requires STDEXEC::__applicable<_strm::transform_sender_for<Tag>, Sender, Env const &>`
	`101`	`+ requires STDEXEC::__callable<STDEXEC::__structured_apply_t,`
	`102`	`+ _strm::transform_sender_for<Tag>,`
	`103`	`+ Sender,`
	`104`	`+ Env const &>`
`102`	`105`	`static auto transform_sender(STDEXEC::set_value_t, Sender&& sndr, Env const & env)`
	`106`	`+ noexcept(STDEXEC::__nothrow_callable<STDEXEC::__structured_apply_t,`
	`107`	`+ _strm::transform_sender_for<Tag>,`
	`108`	`+ Sender,`
	`109`	`+ Env const &>)`
	`110`	`+`
`103`	`111`	`{`
`104`		`- return STDEXEC::__apply(_strm::transform_sender_for<Tag>{}, static_cast<Sender&&>(sndr), env);`
	`112`	`+ return STDEXEC::__structured_apply(_strm::transform_sender_for<Tag>{},`
	`113`	`+ static_cast<Sender&&>(sndr),`
	`114`	`+ env);`
`105`	`115`	`}`
`106`	`116`
`107`	`117`	`template <class Tag, STDEXEC::sender Sender, class... Args>`
`@@ -115,7 +125,6 @@ namespace nv::execution`
`115`	`125`
`116`	`126`	`namespace _strm`
`117`	`127`	`{`
`118`		`-`
`119`	`128`	`#if STDEXEC_HAS_BUILTIN(__is_reference)`
`120`	`129`	`template <class... Ts>`
`121`	`130`	`concept trivially_copyable = ((STDEXEC_IS_TRIVIALLY_COPYABLE(Ts) \|\| __is_reference(Ts)) && ...);`
`@@ -701,8 +710,9 @@ namespace nv::execution`
`701`	`710`	`else`
`702`	`711`	`{`
`703`	`712`	`// pass a cudaError_t by value:`
`704`		`- continuation_kernel<OuterReceiver, Error>`
`705`		`- <<<1, 1, 0, get_stream()>>>(static_cast<OuterReceiver&&>(rcvr_), set_error_t(), status);`
	`713`	`+ continuation_kernel<<<1, 1, 0, get_stream()>>>(static_cast<OuterReceiver&&>(rcvr_),`
	`714`	`+ set_error_t(),`
	`715`	`+ cudaError_t(status));`
`706`	`716`	`}`
`707`	`717`	`}`
`708`	`718`