Merge branch 'main' into patch-3

Author: ericniebler

CMakeLists.txt

@@ -427,18 +427,22 @@ if (STDEXEC_ENABLE_NUMA)
   target_compile_definitions(stdexec INTERFACE STDEXEC_ENABLE_NUMA)
 endif()
 
-set(SYSTEM_CONTEXT_SOURCES src/system_context/system_context.cpp)
-add_library(system_context ${SYSTEM_CONTEXT_SOURCES})
-target_compile_features(system_context PUBLIC cxx_std_20)
-set_target_properties(system_context PROPERTIES
-    CXX_STANDARD 20
-    CXX_STANDARD_REQUIRED ON
-    CXX_EXTENSIONS OFF)
-target_compile_options(system_context PUBLIC
-    $<$<COMPILE_LANG_AND_ID:CXX,MSVC>:/Zc:__cplusplus /Zc:preprocessor /Zc:externConstexpr>
-    )
-add_library(STDEXEC::system_context ALIAS system_context)
-target_link_libraries(system_context PUBLIC stdexec)
+option(STDEXEC_BUILD_SYSTEM_CONTEXT "Build the system_context compiled library" OFF)
+
+if(STDEXEC_BUILD_SYSTEM_CONTEXT)
+  set(SYSTEM_CONTEXT_SOURCES src/system_context/system_context.cpp)
+  add_library(system_context ${SYSTEM_CONTEXT_SOURCES})
+  target_compile_features(system_context PUBLIC cxx_std_20)
+  set_target_properties(system_context PROPERTIES
+      CXX_STANDARD 20
+      CXX_STANDARD_REQUIRED ON
+      CXX_EXTENSIONS OFF)
+  target_compile_options(system_context PUBLIC
+      $<$<COMPILE_LANG_AND_ID:CXX,MSVC>:/Zc:__cplusplus /Zc:preprocessor /Zc:externConstexpr>
+      )
+  add_library(STDEXEC::system_context ALIAS system_context)
+  target_link_libraries(system_context PUBLIC stdexec)
+endif()
 
 
 option(STDEXEC_ENABLE_IO_URING "Enable the use of the io_uring scheduler on Linux" OFF)
@@ -504,8 +508,13 @@ endif()
 if(STDEXEC_INSTALL)
 	include(CPack)
 
-	install(TARGETS stdexec system_context
-		EXPORT stdexec-exports
+  set(stdexec_install_targets stdexec)
+  if(STDEXEC_BUILD_SYSTEM_CONTEXT)
+    list(APPEND stdexec_install_targets system_context)
+  endif()
+
+  install(TARGETS ${stdexec_install_targets}
+    EXPORT stdexec-exports
 		FILE_SET headers
 		FILE_SET version_config)
 

conanfile.py

@@ -14,6 +14,12 @@ class StdexecPackage(ConanFile):
   license = "Apache 2.0"
 
   settings = "os", "arch", "compiler", "build_type"
+  options = {
+    "system_context": [True, False],
+  }
+  default_options = {
+    "system_context": False,
+  }
   exports_sources = (
     "include/*",
     "src/*",
@@ -24,6 +30,12 @@ class StdexecPackage(ConanFile):
   )
   generators = "CMakeToolchain"
 
+  def configure(self):
+    if self.options.system_context:
+      self.package_type = "static-library"
+    else:
+      self.package_type = "header-library"
+
   def validate(self):
     check_min_cppstd(self, "20")
 
@@ -37,18 +49,21 @@ def layout(self):
 
   def build(self):
     tests = "OFF" if self.conf.get("tools.build:skip_test", default=False) else "ON"
+    system_context = "ON" if self.options.system_context else "OFF"
 
     cmake = CMake(self)
     cmake.configure(variables={
       "STDEXEC_BUILD_TESTS": tests,
       "STDEXEC_BUILD_EXAMPLES": tests,
+      "STDEXEC_BUILD_SYSTEM_CONTEXT": system_context,
     })
     cmake.build()
     cmake.test()
 
   def package_id(self):
-    # Clear settings because this package is header-only.
-    self.info.clear()
+    if not self.info.options.system_context:
+      # Clear settings because this package is header-only.
+      self.info.clear()
 
   def package(self):
     cmake = CMake(self)
@@ -58,4 +73,8 @@ def package_info(self):
     self.cpp_info.set_property("cmake_file_name", "P2300")
     self.cpp_info.set_property("cmake_target_name", "P2300::P2300")
     self.cpp_info.set_property("cmake_target_aliases", ["STDEXEC::stdexec"])
-    self.cpp_info.libs = ["system_context"]
+    if self.options.system_context:
+      self.cpp_info.components["system_context"].libs = ["system_context"]
+      self.cpp_info.components["system_context"].set_property(
+        "cmake_target_name", "STDEXEC::system_context"
+      )

include/nvexec/stream/common.cuh

@@ -532,28 +532,21 @@ namespace nv::execution
       {}
 
       template <class... Args>
-      STDEXEC_ATTRIBUTE(host, device)
+      STDEXEC_ATTRIBUTE(device)
       void set_value(Args&&... args) noexcept
       {
         using tuple_t = decayed_tuple_t<set_value_t, Args...>;
         variant_->template emplace<tuple_t>(set_value_t(), static_cast<Args&&>(args)...);
         producer_(task_);
       }
 
-      STDEXEC_ATTRIBUTE(host, device) void set_stopped() noexcept
-      {
-        using tuple_t = decayed_tuple_t<set_stopped_t>;
-        variant_->template emplace<tuple_t>(set_stopped_t());
-        producer_(task_);
-      }
-
       template <class Error>
-      STDEXEC_ATTRIBUTE(host, device)
+      STDEXEC_ATTRIBUTE(device)
       void set_error(Error&& err) noexcept
       {
         if constexpr (__decays_to<Error, std::exception_ptr>)
         {
-          // What is `exception_ptr` but death pending
+          // What is `exception_ptr` but death pending?
           using tuple_t = decayed_tuple_t<set_error_t, cudaError_t>;
           variant_->template emplace<tuple_t>(STDEXEC::set_error, cudaErrorUnknown);
         }
@@ -565,6 +558,15 @@ namespace nv::execution
         producer_(task_);
       }
 
+      STDEXEC_ATTRIBUTE(device)
+      void set_stopped() noexcept
+      {
+        using tuple_t = decayed_tuple_t<set_stopped_t>;
+        variant_->template emplace<tuple_t>(set_stopped_t());
+        producer_(task_);
+      }
+
+      [[nodiscard]]
       auto get_env() const noexcept -> Env const &
       {
         return *env_;

include/nvexec/stream/repeat_n.cuh

@@ -68,13 +68,8 @@ namespace nv::execution::_strm
     {
       using operation_state_concept = STDEXEC::operation_state_t;
 
-      using scheduler_t =
-        STDEXEC::__result_of<STDEXEC::get_completion_scheduler<STDEXEC::set_value_t>,
-                             STDEXEC::env_of_t<Sender>,
-                             STDEXEC::env_of_t<Receiver>>;
-
-      using inner_sender_t =
-        STDEXEC::__result_of<exec::sequence, STDEXEC::schedule_result_t<scheduler_t&>, Sender&>;
+      using scheduler_t     = __completion_scheduler_of_t<set_value_t, Sender, env_of_t<Receiver>>;
+      using inner_sender_t  = STDEXEC::__result_of<STDEXEC::starts_on, scheduler_t, Sender&>;
       using inner_opstate_t = STDEXEC::connect_result_t<inner_sender_t, receiver<opstate>>;
 
       explicit opstate(Sender&& sndr, Receiver rcvr, std::size_t count, scheduler_t sched)
@@ -91,9 +86,9 @@ namespace nv::execution::_strm
 
       auto& _connect()
       {
-        inner_opstate_.__emplace_from(STDEXEC::connect,
-                                      exec::sequence(STDEXEC::schedule(sched_), sndr_),
-                                      receiver{*this});
+        return inner_opstate_.__emplace_from(STDEXEC::connect,
+                                             STDEXEC::starts_on(sched_, sndr_),
+                                             receiver{*this});
       }
 
       template <class Tag, class... Args>

include/nvexec/stream/starts_on.cuh

@@ -0,0 +1,39 @@
+/*
+ * Copyright (c) 2026 NVIDIA Corporation
+ *
+ * Licensed under the Apache License Version 2.0 with LLVM Exceptions
+ * (the "License"); you may not use this file except in compliance with
+ * the License. You may obtain a copy of the License at
+ *
+ *   https://llvm.org/LICENSE.txt
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+// clang-format Language: Cpp
+
+#pragma once
+
+#include "../../stdexec/execution.hpp"
+#include "../../stdexec/functional.hpp"
+
+#include "let_xxx.cuh"
+
+#include "common.cuh"
+
+namespace nv::execution::_strm
+{
+  template <>
+  struct transform_sender_for<STDEXEC::starts_on_t>
+  {
+    template <class Env, STDEXEC::scheduler Scheduler, STDEXEC::sender Sender>
+    auto operator()(Env const &, STDEXEC::starts_on_t, Scheduler&& sched, Sender&& sndr) const
+    {
+      return STDEXEC::let_value(STDEXEC::schedule(sched),
+                                STDEXEC::__always(static_cast<Sender&&>(sndr)));
+    }
+  };
+}  // namespace nv::execution::_strm

include/nvexec/stream/then.cuh

@@ -71,10 +71,9 @@ namespace nv::execution::_strm
         requires std::invocable<Fun, __decay_t<Args>...>
       void set_value(Args&&... args) noexcept
       {
-        using result_t = std::invoke_result_t<Fun, __decay_t<Args>...>;
-        constexpr bool                 does_not_return_a_value = std::is_same_v<void, result_t>;
-        _strm::opstate_base<Receiver>& opstate                 = opstate_;
-        cudaStream_t                   stream                  = opstate.get_stream();
+        using result_t                         = std::invoke_result_t<Fun, __decay_t<Args>...>;
+        constexpr bool does_not_return_a_value = std::is_same_v<void, result_t>;
+        cudaStream_t   stream                  = opstate_.get_stream();
 
         if constexpr (does_not_return_a_value)
         {
@@ -83,29 +82,29 @@ namespace nv::execution::_strm
           if (cudaError_t status = STDEXEC_LOG_CUDA_API(cudaPeekAtLastError());
               status == cudaSuccess)
           {
-            opstate.propagate_completion_signal(STDEXEC::set_value);
+            opstate_.propagate_completion_signal(STDEXEC::set_value);
           }
           else
           {
-            opstate.propagate_completion_signal(STDEXEC::set_error, std::move(status));
+            opstate_.propagate_completion_signal(STDEXEC::set_error, std::move(status));
           }
         }
         else
         {
           using decayed_result_t = __decay_t<result_t>;
-          auto* d_result         = static_cast<decayed_result_t*>(opstate.temp_storage_);
+          auto* d_result         = static_cast<decayed_result_t*>(opstate_.temp_storage_);
           _then_kernel_with_result<Args&&...>
             <<<1, 1, 0, stream>>>(std::move(f_), d_result, static_cast<Args&&>(args)...);
-          opstate.defer_temp_storage_destruction(d_result);
+          opstate_.defer_temp_storage_destruction(d_result);
 
           if (cudaError_t status = STDEXEC_LOG_CUDA_API(cudaPeekAtLastError());
               status == cudaSuccess)
           {
-            opstate.propagate_completion_signal(STDEXEC::set_value, std::move(*d_result));
+            opstate_.propagate_completion_signal(STDEXEC::set_value, std::move(*d_result));
           }
           else
           {
-            opstate.propagate_completion_signal(STDEXEC::set_error, std::move(status));
+            opstate_.propagate_completion_signal(STDEXEC::set_error, std::move(status));
           }
         }
       }
@@ -185,7 +184,7 @@ namespace nv::execution::_strm
         static_cast<Self&&>(self).sndr_,
         static_cast<Receiver&&>(rcvr),
         [&](_strm::opstate_base<Receiver>& stream_provider) -> receiver_t<Receiver>
-        { return receiver_t<Receiver>(self.fun_, stream_provider); });
+        { return receiver_t<Receiver>(static_cast<Self&&>(self).fun_, stream_provider); });
     }
     STDEXEC_EXPLICIT_THIS_END(connect)
 
@@ -209,11 +208,11 @@ namespace nv::execution::_strm
   struct transform_sender_for<STDEXEC::then_t>
   {
     template <class Env, class Fn, class CvSender>
-    auto operator()(Env const &, __ignore, Fn fun, CvSender&& sndr) const
+    auto operator()(Env const &, __ignore, Fn&& fun, CvSender&& sndr) const
     {
       if constexpr (stream_completing_sender<CvSender, Env>)
       {
-        using _sender_t = then_sender<__decay_t<CvSender>, Fn>;
+        using _sender_t = then_sender<__decay_t<CvSender>, __decay_t<Fn>>;
         return _sender_t{static_cast<CvSender&&>(sndr), static_cast<Fn&&>(fun)};
       }
       else

include/nvexec/stream_context.cuh

@@ -35,6 +35,7 @@
 #include "stream/schedule_from.cuh"        // IWYU pragma: export
 #include "stream/split.cuh"                // IWYU pragma: export
 #include "stream/start_detached.cuh"       // IWYU pragma: export
+#include "stream/starts_on.cuh"            // IWYU pragma: export
 #include "stream/sync_wait.cuh"            // IWYU pragma: export
 #include "stream/then.cuh"                 // IWYU pragma: export
 #include "stream/upon_error.cuh"           // IWYU pragma: export

include/stdexec/__detail/__bulk.hpp

@@ -242,6 +242,24 @@ namespace STDEXEC
     template <class _Fun>
     STDEXEC_HOST_DEVICE_DEDUCTION_GUIDE __as_bulk_chunked_fn(_Fun) -> __as_bulk_chunked_fn<_Fun>;
 
+    template <class _Child>
+    struct __attrs : env<__fwd_env_t<env_of_t<_Child>>>
+    {
+      using __base_t = env<__fwd_env_t<env_of_t<_Child>>>;
+      using __base_t::query;
+
+      constexpr explicit __attrs(_Child const & __child) noexcept
+        : __base_t{__fwd_env(STDEXEC::get_env(__child))}
+      {}
+
+      template <class... _Env>
+      STDEXEC_ATTRIBUTE(nodiscard, always_inline, host, device)
+      constexpr auto query(__get_completion_behavior_t<set_value_t>, _Env&&...) const noexcept
+      {
+        return STDEXEC::__get_completion_behavior<set_value_t, _Child, _Env...>();
+      }
+    };
+
     template <class _AlgoTag>
     struct __impl_base : __sexpr_defaults
     {
@@ -252,9 +270,10 @@ namespace STDEXEC
       using __shape_t = decltype(__decay_t<__data_of<_Sender>>::__shape_);
 
       // Forward the child sender's environment (which contains completion scheduler)
-      static constexpr auto __get_attrs = [](__ignore, __ignore, auto const & __child) noexcept
+      static constexpr auto __get_attrs =  //
+        []<class _Child>(__ignore, __ignore, _Child const & __child) noexcept
       {
-        return __fwd_env(STDEXEC::get_env(__child));
+        return __attrs{__child};
       };
 
       template <class _Sender, class... _Env>