10
10
11
11
#include < CL/sycl/detail/defines_elementary.hpp>
12
12
13
- #include < set>
14
13
#include < list>
14
+ #include < set>
15
15
16
16
__SYCL_INLINE_NAMESPACE (cl) {
17
17
namespace sycl {
@@ -29,90 +29,98 @@ using node_ptr = std::shared_ptr<node_impl>;
29
29
using graph_ptr = std::shared_ptr<graph_impl>;
30
30
31
31
class wrapper {
32
- using T = std::function<void (sycl::handler&)>;
33
- T my_func;
34
- std::vector<sycl::event> my_deps;
32
+ using T = std::function<void (sycl::handler &)>;
33
+ T my_func;
34
+ std::vector<sycl::event> my_deps;
35
+
35
36
public:
36
- wrapper (T t, const std::vector<sycl::event>& deps) : my_func(t), my_deps(deps) {};
37
+ wrapper (T t, const std::vector<sycl::event> &deps)
38
+ : my_func(t), my_deps(deps){};
37
39
38
- void operator ()(sycl::handler& cgh) {
39
- cgh.depends_on (my_deps);
40
- std::invoke (my_func,cgh);
41
- }
40
+ void operator ()(sycl::handler & cgh) {
41
+ cgh.depends_on (my_deps);
42
+ std::invoke (my_func, cgh);
43
+ }
42
44
};
43
45
44
46
struct node_impl {
45
- bool is_scheduled;
47
+ bool is_scheduled;
46
48
47
- graph_ptr my_graph;
48
- sycl::event my_event;
49
+ graph_ptr my_graph;
50
+ sycl::event my_event;
49
51
50
- std::vector<node_ptr> my_successors;
51
- std::vector<node_ptr> my_predecessors;
52
+ std::vector<node_ptr> my_successors;
53
+ std::vector<node_ptr> my_predecessors;
52
54
53
- std::function<void (sycl::handler&)> my_body;
55
+ std::function<void (sycl::handler &)> my_body;
54
56
55
- void exec ( sycl::queue q ) {
56
- std::vector<sycl::event> __deps;
57
- for (auto i:my_predecessors) __deps.push_back (i->get_event ());
58
- my_event = q.submit (wrapper{my_body,__deps});
59
- }
57
+ void exec (sycl::queue q) {
58
+ std::vector<sycl::event> __deps;
59
+ for (auto i : my_predecessors)
60
+ __deps.push_back (i->get_event ());
61
+ my_event = q.submit (wrapper{my_body, __deps});
62
+ }
60
63
61
- void register_successor (node_ptr n) {
62
- my_successors.push_back (n);
63
- n->register_predecessor (node_ptr (this ));
64
- }
64
+ void register_successor (node_ptr n) {
65
+ my_successors.push_back (n);
66
+ n->register_predecessor (node_ptr (this ));
67
+ }
65
68
66
- void register_predecessor (node_ptr n) { my_predecessors.push_back (n); }
69
+ void register_predecessor (node_ptr n) { my_predecessors.push_back (n); }
67
70
68
- sycl::event get_event (void ) {return my_event;}
71
+ sycl::event get_event (void ) { return my_event; }
69
72
70
- template <typename T>
71
- node_impl (graph_ptr g, T cgf) : is_scheduled(false ), my_graph(g), my_body(cgf) {}
73
+ template <typename T>
74
+ node_impl (graph_ptr g, T cgf)
75
+ : is_scheduled(false ), my_graph(g), my_body(cgf) {}
72
76
73
- // Recursively adding nodes to execution stack:
74
- void topology_sort (std::list<node_ptr>& schedule) {
75
- is_scheduled = true ;
76
- for (auto i:my_successors) {
77
- if (!i->is_scheduled ) i->topology_sort (schedule);
78
- }
79
- schedule.push_front (node_ptr (this ));
77
+ // Recursively adding nodes to execution stack:
78
+ void topology_sort (std::list<node_ptr> &schedule) {
79
+ is_scheduled = true ;
80
+ for (auto i : my_successors) {
81
+ if (!i->is_scheduled )
82
+ i->topology_sort (schedule);
80
83
}
84
+ schedule.push_front (node_ptr (this ));
85
+ }
81
86
};
82
87
83
88
struct graph_impl {
84
- std::set<node_ptr> my_roots;
85
- std::list<node_ptr> my_schedule;
86
-
87
- graph_ptr parent;
88
-
89
- void exec ( sycl::queue q ) {
90
- if ( my_schedule.empty () ) {
91
- for (auto n : my_roots) {
92
- n->topology_sort (my_schedule);
93
- }
94
- }
95
- for (auto n : my_schedule) n->exec (q);
96
- }
89
+ std::set<node_ptr> my_roots;
90
+ std::list<node_ptr> my_schedule;
97
91
98
- void exec_and_wait ( sycl::queue q ) {
99
- exec (q);
100
- q.wait ();
101
- }
92
+ graph_ptr parent;
102
93
103
- void add_root (node_ptr n) {
104
- my_roots.insert (n);
105
- for (auto n : my_schedule) n->is_scheduled =false ;
106
- my_schedule.clear ();
94
+ void exec (sycl::queue q) {
95
+ if (my_schedule.empty ()) {
96
+ for (auto n : my_roots) {
97
+ n->topology_sort (my_schedule);
98
+ }
107
99
}
108
-
109
- void remove_root (node_ptr n) {
110
- my_roots.erase (n);
111
- for (auto n : my_schedule) n->is_scheduled =false ;
112
- my_schedule.clear ();
113
- }
114
-
115
- graph_impl () {}
100
+ for (auto n : my_schedule)
101
+ n->exec (q);
102
+ }
103
+
104
+ void exec_and_wait (sycl::queue q) {
105
+ exec (q);
106
+ q.wait ();
107
+ }
108
+
109
+ void add_root (node_ptr n) {
110
+ my_roots.insert (n);
111
+ for (auto n : my_schedule)
112
+ n->is_scheduled = false ;
113
+ my_schedule.clear ();
114
+ }
115
+
116
+ void remove_root (node_ptr n) {
117
+ my_roots.erase (n);
118
+ for (auto n : my_schedule)
119
+ n->is_scheduled = false ;
120
+ my_schedule.clear ();
121
+ }
122
+
123
+ graph_impl () {}
116
124
};
117
125
118
126
} // namespace detail
@@ -124,89 +132,92 @@ class graph;
124
132
class executable_graph ;
125
133
126
134
struct node {
127
- // TODO: add properties to distinguish between empty, host, device nodes.
128
- detail::node_ptr my_node;
129
- detail::graph_ptr my_graph;
135
+ // TODO: add properties to distinguish between empty, host, device nodes.
136
+ detail::node_ptr my_node;
137
+ detail::graph_ptr my_graph;
130
138
131
- template <typename T>
132
- node (detail::graph_ptr g, T cgf) : my_graph(g), my_node(new detail::node_impl(g,cgf)) {};
133
- void register_successor (node n) { my_node->register_successor (n.my_node ); }
134
- void exec ( sycl::queue q, sycl::event = sycl::event() ) { my_node->exec (q); }
139
+ template <typename T>
140
+ node (detail::graph_ptr g, T cgf)
141
+ : my_graph(g), my_node(new detail::node_impl(g, cgf)){};
142
+ void register_successor (node n) { my_node->register_successor (n.my_node ); }
143
+ void exec (sycl::queue q, sycl::event = sycl::event()) { my_node->exec (q); }
135
144
136
- void set_root () { my_graph->add_root (my_node);}
145
+ void set_root () { my_graph->add_root (my_node); }
137
146
138
- // TODO: Add query functions: is_root, ...
147
+ // TODO: Add query functions: is_root, ...
139
148
};
140
149
141
150
class executable_graph {
142
151
public:
143
- int my_tag;
144
- sycl::queue my_queue;
152
+ int my_tag;
153
+ sycl::queue my_queue;
145
154
146
- void exec_and_wait ();// { my_queue.wait(); }
155
+ void exec_and_wait (); // { my_queue.wait(); }
147
156
148
- executable_graph (detail::graph_ptr g, sycl::queue q) : my_queue(q), my_tag(rand()) {
149
- g->exec (my_queue);
150
- }
157
+ executable_graph (detail::graph_ptr g, sycl::queue q)
158
+ : my_queue(q), my_tag(rand()) {
159
+ g->exec (my_queue);
160
+ }
151
161
};
152
162
153
163
class graph {
154
164
public:
155
- // Adding empty node with [0..n] predecessors:
156
- node add_empty_node (const std::vector<node>& dep = {});
165
+ // Adding empty node with [0..n] predecessors:
166
+ node add_empty_node (const std::vector<node> & dep = {});
157
167
158
- // Adding node for host task
159
- template <typename T>
160
- node add_host_node (T hostTaskCallable, const std::vector<node>& dep = {});
168
+ // Adding node for host task
169
+ template <typename T>
170
+ node add_host_node (T hostTaskCallable, const std::vector<node> & dep = {});
161
171
162
- // Adding device node:
163
- template <typename T>
164
- node add_device_node (T cgf, const std::vector<node>& dep = {});
172
+ // Adding device node:
173
+ template <typename T>
174
+ node add_device_node (T cgf, const std::vector<node> & dep = {});
165
175
166
- // Adding dependency between two nodes.
167
- void make_edge (node sender, node receiver);
176
+ // Adding dependency between two nodes.
177
+ void make_edge (node sender, node receiver);
168
178
169
- // TODO: Extend queue to directly submit graph
170
- void exec_and_wait ( sycl::queue q );
179
+ // TODO: Extend queue to directly submit graph
180
+ void exec_and_wait (sycl::queue q);
171
181
172
- executable_graph exec ( sycl::queue q ) { return executable_graph{my_graph,q};};
182
+ executable_graph exec (sycl::queue q) {
183
+ return executable_graph{my_graph, q};
184
+ };
173
185
174
- graph () : my_graph(new detail::graph_impl()) {}
186
+ graph () : my_graph(new detail::graph_impl()) {}
175
187
176
- // Creating a subgraph (with predecessors)
177
- graph (graph& parent, const std::vector<node>& dep = {}) {}
188
+ // Creating a subgraph (with predecessors)
189
+ graph (graph & parent, const std::vector<node> & dep = {}) {}
178
190
179
- bool is_subgraph ();
191
+ bool is_subgraph ();
180
192
181
193
private:
182
- detail::graph_ptr my_graph;
194
+ detail::graph_ptr my_graph;
183
195
};
184
196
185
197
void executable_graph::exec_and_wait () { my_queue.wait (); }
186
198
187
- template <typename T>
188
- node graph::add_device_node (T cgf , const std::vector<node>& dep) {
189
- node _node (my_graph,cgf);
190
- if ( !dep.empty () ) {
191
- for (auto n : dep) this ->make_edge (n,_node);
192
- } else {
193
- _node.set_root ();
194
- }
195
- return _node;
199
+ template <typename T>
200
+ node graph::add_device_node (T cgf, const std::vector<node> &dep) {
201
+ node _node (my_graph, cgf);
202
+ if (!dep.empty ()) {
203
+ for (auto n : dep)
204
+ this ->make_edge (n, _node);
205
+ } else {
206
+ _node.set_root ();
207
+ }
208
+ return _node;
196
209
}
197
210
198
211
void graph::make_edge (node sender, node receiver) {
199
- sender.register_successor (receiver);// register successor
200
- my_graph->remove_root (receiver.my_node ); // remove receiver from root node list
212
+ sender.register_successor (receiver); // register successor
213
+ my_graph->remove_root (receiver.my_node ); // remove receiver from root node
214
+ // list
201
215
}
202
216
203
- void graph::exec_and_wait ( sycl::queue q ) {
204
- my_graph->exec_and_wait (q);
205
- };
217
+ void graph::exec_and_wait (sycl::queue q) { my_graph->exec_and_wait (q); };
206
218
207
219
} // namespace experimental
208
220
} // namespace oneapi
209
221
} // namespace ext
210
222
} // namespace sycl
211
223
} // __SYCL_INLINE_NAMESPACE(cl)
212
-
0 commit comments