Merge remote-tracking branch 'origin/main' into tmp_cuda_reduce

ericniebler · ericniebler · commit 76f863d93ccc · 2023-09-23T16:57:09.000-07:00
diff --git a/examples/nvexec/reduce.cpp b/examples/nvexec/reduce.cpp
@@ -23,52 +23,19 @@
 #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::just(std::span{first, last}) | nvexec::reduce(42.0f);
+  // 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();
 
-  auto on_snd = ex::on(sched, std::move(snd));
-  //::print<stdexec::__detail::__name_of<decltype(on_snd)>>();
+  // nvexec::stream_context stream_ctx{};
 
-  // recursively transforms the sender using the stream domain
-  // auto stream_on_snd = nvexec::_strm::stream_domain().transform_sender(
-  //   std::move(on_snd), ex::empty_env());
+  // auto snd = ex::transfer_just(stream_ctx.get_scheduler(), std::span{first, last})
+  //          | nvexec::reduce(42.0f);
 
-  // the name of the transformed sender shows that the reduce node
-  // in the tree was transformed from a basic_sender<> to a nvexec::reduce_::sender_t<>
-  //::print<stdexec::__detail::__name_of<decltype(stream_on_snd)>>();
+  // auto [result] = stdexec::sync_wait(std::move(snd)).value();
 
-  //auto [result] = stdexec::sync_wait(std::move(on_snd)).value();
+  // std::cout << "result: " << result << std::endl;
 }
diff --git a/include/exec/env.hpp b/include/exec/env.hpp
@@ -23,8 +23,8 @@
 #endif
 
 namespace exec {
-  template <class... _TagValue>
-  using with_t = stdexec::__with<_TagValue...>;
+  template <class _Tag, class _Value = stdexec::__none_such>
+  using with_t = stdexec::__with<_Tag, _Value>;
 
   namespace __detail {
     struct __with_t {
diff --git a/include/exec/linux/io_uring_context.hpp b/include/exec/linux/io_uring_context.hpp
@@ -74,8 +74,13 @@ namespace exec {
       unsigned int __to_submit,
       unsigned int __min_complete,
       unsigned int __flags) {
-      return (int) ::syscall(
+      int rc = (int) ::syscall(
         __NR_io_uring_enter, __ring_fd, __to_submit, __min_complete, __flags, nullptr, 0);
+      if (rc == -1) {
+        return -errno;
+      } else {
+        return rc;
+      }
     }
 
     inline memory_mapped_region __map_region(int __fd, ::off_t __offset, std::size_t __size) {
diff --git a/include/exec/task.hpp b/include/exec/task.hpp
@@ -31,8 +31,8 @@
 #include "scope.hpp"
 
 STDEXEC_PRAGMA_PUSH()
-STDEXEC_PRAGMA_IGNORE("-Wpragmas")
-STDEXEC_PRAGMA_IGNORE("-Wundefined-inline")
+STDEXEC_PRAGMA_IGNORE_GNU("-Wpragmas")
+STDEXEC_PRAGMA_IGNORE_GNU("-Wundefined-inline")
 
 namespace exec {
   namespace __task {
diff --git a/include/nvexec/multi_gpu_context.cuh b/include/nvexec/multi_gpu_context.cuh
@@ -66,7 +66,7 @@ namespace nvexec {
             if (op.status_ == cudaSuccess) {
               continuation_kernel<<<1, 1, 0, op.stream_>>>(std::move(op.rec_), stdexec::set_value);
             } else {
-              continuation_kernel<cudaError_t><<<1, 1, 0, op.stream_>>>(
+              continuation_kernel<<<1, 1, 0, op.stream_>>>(
                 std::move(op.rec_), stdexec::set_error, std::move(op.status_));
             }
           }
diff --git a/include/nvexec/stream/bulk.cuh b/include/nvexec/stream/bulk.cuh
@@ -202,7 +202,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
                 if (begin < end) {
                   cudaSetDevice(dev);
-                  cudaStreamWaitEvent(stream, op_state.ready_to_launch_);
+                  cudaStreamWaitEvent(stream, op_state.ready_to_launch_, 0);
                   kernel<block_threads, As&...>
                     <<<grid_blocks, block_threads, 0, stream>>>(begin, end, self.f_, as...);
                   cudaEventRecord(op_state.ready_to_complete_[dev], op_state.streams_[dev]);
@@ -225,7 +225,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
             for (int dev = 0; dev < op_state.num_devices_; dev++) {
               if (dev != op_state.current_device_) {
-                cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev]);
+                cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev], 0);
               }
             }
           }
diff --git a/include/nvexec/stream/ensure_started.cuh b/include/nvexec/stream/ensure_started.cuh
@@ -246,7 +246,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
           if (status == cudaSuccess) {
             if constexpr (stream_sender<Sender, env_t>) {
               status = STDEXEC_DBG_ERR(
-                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));
+                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));
             }
 
             visit(
diff --git a/include/nvexec/stream/split.cuh b/include/nvexec/stream/split.cuh
@@ -236,7 +236,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
           if (status == cudaSuccess) {
             if constexpr (stream_sender<Sender, env_t>) {
               status = STDEXEC_DBG_ERR(
-                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));
+                cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));
             }
 
             visit(
diff --git a/include/nvexec/stream/when_all.cuh b/include/nvexec/stream/when_all.cuh
@@ -259,7 +259,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {
 
               for (int i = 0; i < sizeof...(SenderIds); i++) {
                 if (status_ == cudaSuccess) {
-                  status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i]));
+                  status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i], 0));
                 }
               }
             } else {
diff --git a/include/nvexec/stream_context.cuh b/include/nvexec/stream_context.cuh
@@ -101,7 +101,7 @@ namespace nvexec {
 
       // Lazy algorithm customizations require a recursive tree transformation
       template <sender_expr Sender, class Env>
-        requires _non_stream_sender<Sender> // no need to transform it a second time
+        requires _non_stream_sender<Sender, Env> // no need to transform it a second time
       auto transform_sender(Sender&& sndr, const Env& env) const noexcept {
         return stdexec::apply_sender(
           (Sender&&) sndr,
@@ -113,7 +113,7 @@ namespace nvexec {
 
       // reduce senders get a special transformation
       template <sender_expr_for<reduce_t> Sender, class Env>
-        requires _non_stream_sender<Sender> // no need to transform it a second time
+        requires _non_stream_sender<Sender, Env> // no need to transform it a second time
       auto transform_sender(Sender&& sndr, const Env& env) const noexcept {
         return stdexec::apply_sender(
           (Sender&&) sndr,
diff --git a/include/stdexec/__detail/__config.hpp b/include/stdexec/__detail/__config.hpp
@@ -54,6 +54,8 @@
 #define STDEXEC_NVCC() 1
 #elif defined(__NVCOMPILER)
 #define STDEXEC_NVHPC() 1
+#elif defined(__EDG__)
+#define LEGATE_EDG() 1
 #elif defined(__clang__)
 #define STDEXEC_CLANG() 1
 #elif defined(__GNUC__)
@@ -68,6 +70,9 @@
 #ifndef STDEXEC_NVHPC
 #define STDEXEC_NVHPC() 0
 #endif
+#ifndef STDEXEC_EDG
+#define STDEXEC_EDG() 0
+#endif
 #ifndef STDEXEC_CLANG
 #define STDEXEC_CLANG() 0
 #endif
@@ -78,15 +83,31 @@
 #define STDEXEC_MSVC() 0
 #endif
 
-#if STDEXEC_CLANG() || STDEXEC_GCC()
 #define STDEXEC_STRINGIZE(_ARG) #_ARG
+
+#if STDEXEC_NVCC()
+#define STDEXEC_PRAGMA_PUSH() _Pragma("nv_diagnostic push")
+#define STDEXEC_PRAGMA_POP() _Pragma("nv_diagnostic pop")
+#define STDEXEC_PRAGMA_IGNORE_EDG(...) _Pragma(STDEXEC_STRINGIZE(nv_diag_suppress __VA_ARGS__))
+#elif STDEXEC_NVHPC() || STDEXEC_EDG()
+#define STDEXEC_PRAGMA_PUSH() \
+  _Pragma("diagnostic push") STDEXEC_PRAGMA_IGNORE_EDG(invalid_error_number)
+#define STDEXEC_PRAGMA_POP() _Pragma("diagnostic pop")
+#define STDEXEC_PRAGMA_IGNORE_EDG(...) _Pragma(STDEXEC_STRINGIZE(diag_suppress __VA_ARGS__))
+#elif STDEXEC_CLANG() || STDEXEC_GCC()
 #define STDEXEC_PRAGMA_PUSH() _Pragma("GCC diagnostic push")
 #define STDEXEC_PRAGMA_POP() _Pragma("GCC diagnostic pop")
-#define STDEXEC_PRAGMA_IGNORE(_ARG) _Pragma(STDEXEC_STRINGIZE(GCC diagnostic ignored _ARG))
+#define STDEXEC_PRAGMA_IGNORE_GNU(_ARG) _Pragma(STDEXEC_STRINGIZE(GCC diagnostic ignored _ARG))
 #else
 #define STDEXEC_PRAGMA_PUSH()
 #define STDEXEC_PRAGMA_POP()
-#define STDEXEC_PRAGMA_IGNORE(_ARG)
+#endif
+
+#ifndef STDEXEC_PRAGMA_IGNORE_GNU
+#define STDEXEC_PRAGMA_IGNORE_GNU(_ARG)
+#endif
+#ifndef STDEXEC_PRAGMA_IGNORE_EDG
+#define STDEXEC_PRAGMA_IGNORE_EDG(_ARG)
 #endif
 
 #if !STDEXEC_MSVC() && defined(__has_builtin)
diff --git a/include/stdexec/__detail/__execution_fwd.hpp b/include/stdexec/__detail/__execution_fwd.hpp
@@ -123,9 +123,6 @@ namespace stdexec {
   using __get_completion_signatures::get_completion_signatures_t;
   extern const get_completion_signatures_t get_completion_signatures;
 
-  template <class _Sender, class _Env>
-  using __completion_signatures_of_t = __call_result_t< get_completion_signatures_t, _Sender, _Env>;
-
   //////////////////////////////////////////////////////////////////////////////////////////////////
   namespace __connect {
     struct connect_t;
diff --git a/include/stdexec/__detail/__meta.hpp b/include/stdexec/__detail/__meta.hpp
@@ -351,6 +351,12 @@ namespace stdexec {
     requires(sizeof...(_False) <= 1)
   using __if_c = __minvoke<__if_::__<_Pred>, _True, _False...>;
 
+  template <class _Pred, class _True, class... _False>
+  using __minvoke_if = __minvoke<__if<_Pred, _True, _False...>>;
+
+  template <bool _Pred, class _True, class... _False>
+  using __minvoke_if_c = __minvoke<__if_c<_Pred, _True, _False...>>;
+
   template <class _Tp>
   struct __mconst {
     template <class...>
@@ -372,6 +378,9 @@ namespace stdexec {
   template <class _Fn, class _Default>
   using __with_default = __mtry_catch<_Fn, __mconst<_Default>>;
 
+  template <template <class...> class _Fn, class _Default>
+  using __with_default_q = __mtry_catch_q<_Fn, __mconst<_Default>>;
+
   inline constexpr __mstring __mbad_substitution =
     "The specified meta-function could not be evaluated with the types provided."__csz;
 
diff --git a/include/stdexec/execution.hpp b/include/stdexec/execution.hpp
diff --git a/test/nvexec/CMakeLists.txt b/test/nvexec/CMakeLists.txt
diff --git a/test/stdexec/queries/test_get_forward_progress_guarantee.cpp b/test/stdexec/queries/test_get_forward_progress_guarantee.cpp

Original file line number	Diff line number	Diff line change
`@@ -66,7 +66,7 @@ namespace nvexec {`
`66`	`66`	`if (op.status_ == cudaSuccess) {`
`67`	`67`	`continuation_kernel<<<1, 1, 0, op.stream_>>>(std::move(op.rec_), stdexec::set_value);`
`68`	`68`	`} else {`
`69`		`- continuation_kernel<cudaError_t><<<1, 1, 0, op.stream_>>>(`
	`69`	`+ continuation_kernel<<<1, 1, 0, op.stream_>>>(`
`70`	`70`	`std::move(op.rec_), stdexec::set_error, std::move(op.status_));`
`71`	`71`	`}`
`72`	`72`	`}`
Original file line number	Diff line number	Diff line change
`@@ -202,7 +202,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {`
`202`	`202`
`203`	`203`	`if (begin < end) {`
`204`	`204`	`cudaSetDevice(dev);`
`205`		`- cudaStreamWaitEvent(stream, op_state.ready_to_launch_);`
	`205`	`+ cudaStreamWaitEvent(stream, op_state.ready_to_launch_, 0);`
`206`	`206`	`kernel<block_threads, As&...>`
`207`	`207`	`<<<grid_blocks, block_threads, 0, stream>>>(begin, end, self.f_, as...);`
`208`	`208`	`cudaEventRecord(op_state.ready_to_complete_[dev], op_state.streams_[dev]);`
`@@ -225,7 +225,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {`
`225`	`225`
`226`	`226`	`for (int dev = 0; dev < op_state.num_devices_; dev++) {`
`227`	`227`	`if (dev != op_state.current_device_) {`
`228`		`- cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev]);`
	`228`	`+ cudaStreamWaitEvent(baseline_stream, op_state.ready_to_complete_[dev], 0);`
`229`	`229`	`}`
`230`	`230`	`}`
`231`	`231`	`}`
Original file line number	Diff line number	Diff line change
`@@ -246,7 +246,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {`
`246`	`246`	`if (status == cudaSuccess) {`
`247`	`247`	`if constexpr (stream_sender<Sender, env_t>) {`
`248`	`248`	`status = STDEXEC_DBG_ERR(`
`249`		`- cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));`
	`249`	`+ cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));`
`250`	`250`	`}`
`251`	`251`
`252`	`252`	`visit(`
Original file line number	Diff line number	Diff line change
`@@ -236,7 +236,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {`
`236`	`236`	`if (status == cudaSuccess) {`
`237`	`237`	`if constexpr (stream_sender<Sender, env_t>) {`
`238`	`238`	`status = STDEXEC_DBG_ERR(`
`239`		`- cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_));`
	`239`	`+ cudaStreamWaitEvent(op->get_stream(), op->shared_state_->event_, 0));`
`240`	`240`	`}`
`241`	`241`
`242`	`242`	`visit(`
Original file line number	Diff line number	Diff line change
`@@ -259,7 +259,7 @@ namespace nvexec::STDEXEC_STREAM_DETAIL_NS {`
`259`	`259`
`260`	`260`	`for (int i = 0; i < sizeof...(SenderIds); i++) {`
`261`	`261`	`if (status_ == cudaSuccess) {`
`262`		`- status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i]));`
	`262`	`+ status_ = STDEXEC_DBG_ERR(cudaStreamWaitEvent(stream, events_[i], 0));`
`263`	`263`	`}`
`264`	`264`	`}`
`265`	`265`	`} else {`