Feature/cuda reduce

CMakeLists.txt

@@ -176,7 +176,7 @@ target_compile_options(stdexec_executable_flags INTERFACE
 # Template backtrace limit
 target_compile_options(stdexec_executable_flags INTERFACE
                        $<$<OR:$<CXX_COMPILER_ID:Clang>,$<CXX_COMPILER_ID:AppleClang>>:
-                       -ferror-limit=0
+                       -ferror-limit=1
                        -fmacro-backtrace-limit=0
                        -ftemplate-backtrace-limit=0>
                        )

examples/nvexec/reduce.cpp

@@ -23,19 +23,42 @@
 #include <span>
 
 namespace ex = stdexec;
+using stdexec::__tag_invoke::tag_invoke;
+struct sink_receiver {
+  using is_receiver = void;
+  friend void tag_invoke(stdexec::set_value_t, sink_receiver, auto&&...) noexcept {}
+  friend void tag_invoke(stdexec::set_error_t, sink_receiver, auto&&) noexcept {}
+  friend void tag_invoke(stdexec::set_stopped_t, sink_receiver) noexcept {}
+  friend stdexec::empty_env tag_invoke(stdexec::get_env_t, sink_receiver) noexcept { return {}; }
+};
 
+struct empty_environment {
+};
+
+template <class...>
+[[deprecated]] void print() {}
+
+// unqualified call to tag_invoke:
 int main() {
   const int n = 2 * 1024;
   thrust::device_vector<float> input(n, 1.0f);
   float* first = thrust::raw_pointer_cast(input.data());
   float* last = thrust::raw_pointer_cast(input.data()) + input.size();
-
   nvexec::stream_context stream_ctx{};
+  auto sched = stream_ctx.get_scheduler();
 
-  auto snd = ex::transfer_just(stream_ctx.get_scheduler(), std::span{first, last})
+ // auto domain = ex::get_domain(sched);
+ // print<decltype(domain)>();
+
+  auto snd = ex::just(std::span{first, last})
            | nvexec::reduce(42.0f);
 
-  auto [result] = stdexec::sync_wait(std::move(snd)).value();
+ // ::print<stdexec::__detail::__name_of<decltype(snd)>>();
+  // nvexec::stream_scheduler gpu = stream_ctx.get_scheduler();         
+  // using stdexec::__tag_invoke::tag_invoke;  
+  // tag_invoke(stdexec::get_completion_signatures, snd, empty_environment{});
+
+  auto [result] =
+    stdexec::sync_wait(ex::on(sched, std::move(snd))).value();
 
-  std::cout << "result: " << result << std::endl;
-}
+ }

include/nvexec/stream/algorithm_base.cuh

@@ -91,22 +91,58 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS::__algo_range_init_fun {
     };
   };
 
-  template <class SenderId, class InitT, class Fun, class DerivedSender>
+  template <class Tag, class SenderId, class InitT, class Fun, class DerivedSender>
   struct sender_t {
+
     struct __t : stream_sender_base {
+
+
       using Sender = stdexec::__t<SenderId>;
       using __id = sender_t;
-
+      using is_sender = void;
       template <class Receiver>
       using receiver_t = typename DerivedSender::template receiver_t<Receiver>;
 
       template <class Range>
       using _set_value_t = typename DerivedSender::template _set_value_t<Range>;
 
       Sender sndr_;
+      // why is this called initT, anyway? If other algorithms will use this in the future im not sure initT is a good name
       STDEXEC_NO_UNIQUE_ADDRESS InitT init_;
       STDEXEC_NO_UNIQUE_ADDRESS Fun fun_;
 
+      template <std::size_t Index, typename... Types>
+      using nth_type_of = std::tuple_element_t<Index, std::tuple<Types...>>;
+
+      template <typename T>
+      void print_type_name() const {
+        stdexec::print(std::declval<stdexec::__detail::__name_of<T>>());
+      }
+      template <typename T>
+      struct print_the_type;
+
+      // This shouldn't be here. Imo I think algorithm_base should 
+      // have a __data struct that each inheritor is responsible for providing. I put this here to get things to compile. 
+      template <class _InitT, class _Fun>
+      struct __data {
+        _InitT __initT_;
+        STDEXEC_NO_UNIQUE_ADDRESS _Fun __fun_;
+        static constexpr auto __mbrs_ = __mliterals<&__data::__initT_, &__data::__fun_>();
+      };
+      template <class _InitT, class _Fun>
+      __data(_InitT, _Fun) -> __data<_InitT, _Fun>;
+
+      // this is basically the apply function that sender_apply is looking for. 
+      template <class S, class Fn>
+      auto plscompile(S s, Fn f) {
+        auto inside = s.sndr_; 
+        auto data = __data(s.init_, s.fun_);
+        // stdexec::__detail::__name_of<decltype(inside)> hi;
+        auto invoked = f(std::declval<Tag>(), data, inside);
+        return invoked;
+      }
+
+
       template <class Self, class Env>
       using completion_signatures = //
         __try_make_completion_signatures<

include/nvexec/stream/common.cuh

@@ -33,7 +33,8 @@
 
 namespace nvexec {
   using stdexec::operator""__csz;
-
+  [[gnu::deprecated]]
+  void print(auto&&...) {}
   enum class stream_priority {
     high,
     normal,
@@ -312,6 +313,8 @@ namespace nvexec {
         set_stopped_t>>;
 
     inline constexpr get_stream_provider_t get_stream_provider{};
+[[gnu::deprecated]]
+  void prints(auto&&...) {}
 
     struct get_stream_t {
       template <class Env>