Skip to content

Commit e9a0a19

Browse files
reblereble
authored
Update sycl_ext_oneapi_graph.asciidoc
Renaming functions, update example and removing USM shortcuts
1 parent e07d7f8 commit e9a0a19

File tree

1 file changed

+32
-167
lines changed

1 file changed

+32
-167
lines changed

‎sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc‎

Lines changed: 32 additions & 167 deletions
Original file line numberDiff line numberDiff line change
@@ -126,8 +126,6 @@ namespace sycl::ext::oneapi::experimental {
126126
command_graph() = delete;
127127
};
128128
129-
command_graph<graph_state::executable> compile(const command_graph<graph_state::modifiable> Graph);
130-
131129
}
132130
133131
----
@@ -176,182 +174,47 @@ Table 4. Member functions of the `command_graph` class.
176174
|
177175
[source,c++]
178176
----
179-
node add_node(const std::vector<node>& dep = {});
177+
node add(const std::vector<node>& dep = {});
180178
----
181179
|This creates an empty node which is associated to no task. Its intended use is either a connection point inside a graph between groups of nodes, and can significantly reduce the number of edges ( O(n) vs. O(n^2) ). Another use-case is building the structure of a graph first and adding tasks later.
182180

183181
|
184182
[source,c++]
185183
----
186184
template<typename T>
187-
node add_node(T cgf, const std::vector<node>& dep = {});
185+
node add(T cgf, const std::vector<node>& dep = {});
188186
----
189-
|This node captures a command group function object containing host task which is scheduled by the SYCL runtime or a SYCL function for invoking kernels with all restrictions that apply as described in the spec.
190-
191-
|===
192-
193-
Memory that is allocated by the following functions is owned by the specific graph. When freed inside the graph, the memory is only accessible before the `free` node is executed and after the `malloc` node is executed.
194-
195-
Table 5. Member functions of the `command_graph` class (memory operations).
196-
[cols="2a,a"]
197-
|===
198-
|Member function|Description
187+
|This function adds a command group function object to a graph. The function object can contain single or multiple commands such as a host task which is scheduled by the SYCL runtime or a SYCL function for invoking kernels with all restrictions that apply as described in the spec.
199188

200189
|
201190
[source,c++]
202191
----
203-
node memcpy(void* dest, const void* src, size_t numBytes, const std::vector<node>& dep = {});
192+
command_graph<graph_state::executable> finalize(context &syclContext) const;
204193
----
205-
|Adding a node that encapsulates a `memcpy` operation.
194+
| This function creates an executable graph object with an immutable topology that can be executed on a queue that matches the given context.
206195

207-
|
208-
[source,c++]
209-
----
210-
template<typename T> node
211-
copy(const T* src, T* dest, size_t count, const std::vector<node>& dep = {});
212-
----
213-
|Adding a node that encapsulates a `copy` operation.
214-
215-
|
216-
[source,c++]
217-
----
218-
node memset(void* ptr, int value, size_t numBytes, const std::vector<node>& dep = {});
219-
----
220-
|Adding a node that encapsulates a `memset` operation.
221-
222-
|
223-
[source,c++]
224-
----
225-
template<typename T>
226-
node fill(void* ptr, const T& pattern, size_t count, const std::vector<node>& dep = {});
227-
----
228-
|Adding a node that encapsulates a `fill` operation.
229-
230-
|
231-
[source,c++]
232-
----
233-
node malloc(void *data, size_t numBytes, usm::alloc kind, const std::vector<node>& dep = {});
234-
----
235-
|Adding a node that encapsulates a `malloc` operation.
236-
237-
|
238-
[source,c++]
239-
----
240-
node malloc_shared(void *data, size_t numBytes, const std::vector<node>& dep = {});
241-
----
242-
|Adding a node that encapsulates a `malloc` operation.
243-
244-
|
245-
[source,c++]
246-
----
247-
node malloc_host(void *data, size_t numBytes, const std::vector<node>& dep = {});
248-
----
249-
|Adding a node that encapsulates a `malloc` operation.
250-
251-
|
252-
[source,c++]
253-
----
254-
node malloc_device(void *data, size_t numBytes, const std::vector<node>& dep = {});
255-
----
256-
|Adding a node that encapsulates a `malloc` operation.
257-
258-
|
259-
[source,c++]
260-
----
261-
node free(void *data, const std::vector<node>& dep = {});
262-
----
263-
|Adding a node that encapsulates a `free` operation.
264-
265-
|===
266-
267-
Table 6. Member functions of the `command_graph` class (convenience shortcuts).
268-
[cols="2a,a"]
269-
|===
270-
|Member function|Description
271-
272-
|
273-
[source,c++]
274-
----
275-
template <typename KernelName, typename KernelType>
276-
node single_task(const KernelType &kernelFunc, const std::vector<node>& dep = {});
277-
----
278-
|Adding a node that encapsulates a `single_task` operation.
279-
280-
|
281-
[source,c++]
282-
----
283-
template <typename KernelName, int Dims, typename... Rest>
284-
node parallel_for(range<Dims> numWorkItems, Rest&& rest, const std::vector<node>& dep = {});
285-
----
286-
|Adding a node that encapsulates a `parallel_for` operation.
287-
288-
|
289-
[source,c++]
290-
----
291-
template <typename KernelName, int Dims, typename... Rest>
292-
node parallel_for(nd_range<Dims> executionRange, Rest&& rest, const std::vector<node>& dep = {});
293-
----
294-
|Adding a node that encapsulates a `parallel_for` operation.
295-
296-
// |===
297-
//
298-
// Table 7. Helper functions of the `graph` class.
299-
// [cols="a,a"]
300-
// |===
301-
// |Function name|Description
302-
//
303-
// |
304-
// [source,c++]
305-
// ----
306-
// graph<graph_state::modifiable> make_graph();
307-
// ----
308-
// |Creates a `graph` object in the `graph_state::modifiable` state.
309-
310-
|===
311-
312-
=== Node member functions
313-
314-
Table 8. Constructor of the `node` class.
315-
[cols="a,a"]
316196
|===
317-
|Constructor|Description
318-
319-
|
320-
[source,c++]
321-
----
322-
node();
323-
----
324-
|Creates an empty `node` object. That encapsulates no tasks and is not assigned to a graph. Prior to execution it has to be assigned to a graph.
325197

326-
|===
198+
Memory that is allocated by the following functions is owned by the specific graph. When freed inside the graph, the memory is only accessible before the `free` node is executed and after the `malloc` node is executed.
327199

328-
Table 9. Member functions of the `node` class.
200+
Table 5. Member functions of the `command_graph` class (memory operations).
329201
[cols="2a,a"]
330202
|===
331-
|Function name|Description
332-
333-
|
334-
[source,c++]
335-
----
336-
void set_graph(command_graph<graph_state::modifiable>& Graph);
337-
----
338-
|Assigns a `node` object to a `graph`.
203+
|Member function|Description
339204

340205
|
341206
[source,c++]
342207
----
343-
template<typename T>
344-
void update(T cgf);
208+
node add_malloc_device(void *&data, size_t numBytes, const std::vector<node>& dep = {});
345209
----
346-
|Update a `node` object.
210+
|Adding a node that encapsulates a `malloc` operation.
347211

348212
|
349213
[source,c++]
350214
----
351-
template<typename T>
352-
void update(T cgf, graph<graph_state::modifiable>& Graph);
215+
node add_free(void *data, const std::vector<node>& dep = {});
353216
----
354-
|Update a `node` object and assign it to a task.
217+
|Adding a node that encapsulates a `free` operation.
355218

356219
|===
357220

@@ -378,36 +241,38 @@ int main() {
378241
sycl::ext::oneapi::experimental::command_graph g;
379242
380243
float *x , *y, *z;
381-
382-
auto n_x = g.malloc_shared<float>(x, n, q);
383-
auto n_y = g.malloc_shared<float>(y, n, q);
384-
auto n_z = g.malloc_shared<float>(z, n, q);
385-
244+
386245
float *dotp = sycl::malloc_shared<float>(1, q);
387246
388-
/* init data by using usm shortcut */
389-
auto n_i = g.parallel_for(n, [=](sycl::id<1> it){
390-
const size_t i = it[0];
391-
x[i] = 1.0f;
392-
y[i] = 2.0f;
393-
z[i] = 3.0f;
247+
auto n_x = g.add_malloc_device<float>(x, n);
248+
auto n_y = g.add_malloc_device<float>(y, n);
249+
auto n_z = g.add_malloc_device<float>(z, n);
250+
251+
/* init data on the device */
252+
auto n_i = g.add([&](sycl::handler &h) {
253+
h.parallel_for(n, [=](sycl::id<1> it){
254+
const size_t i = it[0];
255+
x[i] = 1.0f;
256+
y[i] = 2.0f;
257+
z[i] = 3.0f;
258+
});
394259
}, {n_x, n_y, n_z});
395260
396-
auto node_a = g.add_node([&](sycl::handler &h) {
261+
auto node_a = g.add([&](sycl::handler &h) {
397262
h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) {
398263
const size_t i = it[0];
399264
x[i] = alpha * x[i] + beta * y[i];
400265
});
401266
}, {n_i});
402267
403-
auto node_b = g.add_node([&](sycl::handler &h) {
268+
auto node_b = g.add([&](sycl::handler &h) {
404269
h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) {
405270
const size_t i = it[0];
406271
z[i] = gamma * z[i] + beta * y[i];
407272
});
408273
}, {n_i});
409274
410-
auto node_c = g.add_node(
275+
auto node_c = g.add(
411276
[&](sycl::handler &h) {
412277
h.parallel_for(sycl::range<1>{n},
413278
sycl::reduction(dotp, 0.0f, std::plus()),
@@ -418,15 +283,15 @@ int main() {
418283
},
419284
{node_a, node_b});
420285
421-
auto node_f1 = g.free(x, {node_c});
422-
auto node_f1 = g.free(y, {node_b});
286+
auto node_f1 = g.add_free(x, {node_c});
287+
auto node_f2 = g.add_free(y, {node_b});
423288
424-
auto exec = g.compile(q);
289+
auto exec = g.finalize(q.get_context());
425290
426291
q.submit(exec).wait();
427292
428293
// memory can be freed inside or outside the graph
429-
sycl::free(z, q);
294+
sycl::free(z, q.get_context());
430295
sycl::free(dotp, q);
431296
432297
return 0;

0 commit comments

Comments
 (0)