save work before merge main squash

.history/examples/nvexec/bulk_20230821124034.cpp

@@ -0,0 +1,52 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <cstdio>
+
+namespace ex = stdexec;
+
+int main() {
+  using nvexec::is_on_gpu;
+
+  nvexec::stream_context stream_ctx{};
+  ex::scheduler auto sch = stream_ctx.get_scheduler();
+
+  auto bulk_fn = [](int lbl) {
+    return [=](int i) {
+      std::printf("B%d: i = %d\n", lbl, i);
+    };
+  };
+
+  auto then_fn = [](int lbl) {
+    return [=] {
+      std::printf("T%d\n", lbl);
+    };
+  };
+
+  auto fork = ex::schedule(sch) | ex::then(then_fn(0)) | ex::split();
+
+  auto snd = ex::transfer_when_all( //
+               sch,
+               fork | ex::bulk(4, bulk_fn(1)),
+               fork | ex::then(then_fn(1)),
+               fork | ex::bulk(4, bulk_fn(2)))
+           | ex::then(then_fn(2));
+
+  stdexec::sync_wait(std::move(snd));
+}

.history/examples/nvexec/bulk_20230824201208.cpp

@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <cstdio>
+
+namespace ex = stdexec;
+
+int main() {
+  using nvexec::is_on_gpu;
+
+  nvexec::stream_context stream_ctx{};
+  ex::scheduler auto sch = stream_ctx.get_scheduler();
+
+  auto bulk_fn = [](int lbl) {
+    return [=](int i) {
+      std::printf("B%d: i = %d\n", lbl, i);
+    };
+  };
+
+  auto then_fn = [](int lbl) {
+    return [=] {
+      std::printf("T%d\n", lbl);
+    };
+  };
+
+  auto fork = ex::schedule(sch) | ex::then(then_fn(0)) | ex::split();
+
+  auto snd = ex::transfer_when_all( //
+               sch,
+               fork 
+               | ex::bulk(4, bulk_fn(1)),
+               fork | ex::then(then_fn(1)),
+               fork | ex::bulk(4, bulk_fn(2)))
+           | ex::then(then_fn(2));
+
+  stdexec::sync_wait(std::move(snd));
+}

.history/examples/nvexec/bulk_20230824201210.cpp

@@ -0,0 +1,52 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <cstdio>
+
+namespace ex = stdexec;
+
+int main() {
+  using nvexec::is_on_gpu;
+
+  nvexec::stream_context stream_ctx{};
+  ex::scheduler auto sch = stream_ctx.get_scheduler();
+
+  auto bulk_fn = [](int lbl) {
+    return [=](int i) {
+      std::printf("B%d: i = %d\n", lbl, i);
+    };
+  };
+
+  auto then_fn = [](int lbl) {
+    return [=] {
+      std::printf("T%d\n", lbl);
+    };
+  };
+
+  auto fork = ex::schedule(sch) | ex::then(then_fn(0)) | ex::split();
+
+  auto snd = ex::transfer_when_all( //
+               sch,
+               fork | ex::bulk(4, bulk_fn(1)),
+               fork | ex::then(then_fn(1)),
+               fork | ex::bulk(4, bulk_fn(2)))
+           | ex::then(then_fn(2));
+
+  stdexec::sync_wait(std::move(snd));
+}

.history/examples/nvexec/reduce_20230823223415.cpp

@@ -0,0 +1,51 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <thrust/device_vector.h>
+
+#include <cstdio>
+#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 {}; }
+};
+// 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 snd = ex::just(std::span{first, last})
+           | nvexec::reduce(42.0f);
+  auto [result] =
+    stdexec::sync_wait(ex::on(stream_ctx.get_scheduler(), std::move(snd))).value();
+
+
+
+
+}

.history/examples/nvexec/reduce_20230824090732.cpp

@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <thrust/device_vector.h>
+
+#include <cstdio>
+#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 {}; }
+};
+// 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 snd = ex::just(std::span{first, last})
+           | nvexec::reduce(42.0f);
+
+
+  auto [result] =
+    stdexec::sync_wait(ex::on(stream_ctx.get_scheduler(), std::move(snd))).value();
+
+
+
+
+}

.history/examples/nvexec/reduce_20230824090737.cpp

@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <thrust/device_vector.h>
+
+#include <cstdio>
+#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 {}; }
+};
+// 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 snd = ex::just(std::span{first, last})
+           | nvexec::reduce(42.0f);
+  using tag_
+
+  auto [result] =
+    stdexec::sync_wait(ex::on(stream_ctx.get_scheduler(), std::move(snd))).value();
+
+
+
+
+}

.history/examples/nvexec/reduce_20230824090809.cpp

@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2022 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.
+ */
+
+#include <nvexec/stream_context.cuh>
+#include <stdexec/execution.hpp>
+
+#include <thrust/device_vector.h>
+
+#include <cstdio>
+#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 {}; }
+};
+// 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 snd = ex::just(std::span{first, last})
+           | nvexec::reduce(42.0f);
+using stdexec::__tag_invoke::tag_invoke;
+tag_invoke(stdexec::get_completion_signatures, my-sender, the-environment);
+  auto [result] =
+    stdexec::sync_wait(ex::on(stream_ctx.get_scheduler(), std::move(snd))).value();
+
+
+
+
+}