-
Notifications
You must be signed in to change notification settings - Fork 236
Expand file tree
/
Copy pathreduce.cpp
More file actions
64 lines (52 loc) · 2.08 KB
/
reduce.cpp
File metadata and controls
64 lines (52 loc) · 2.08 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
/*
* 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 {}; }
};
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 domain = ex::get_domain(sched);
// print<decltype(domain)>();
auto snd = ex::just(std::span{first, last})
| nvexec::reduce(42.0f);
// ::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();
}