@@ -158,9 +158,9 @@ The following property can be passed to three different APIs, namely:
158158* The `accessor` constructor, giving a more granular control
159159* The `buffer` constructor, in which case all the `accessors` derived from
160160this buffer will inherit this property (unless overridden).
161- * The `property_list` parameter of
162- `command_graph<graph_state::modifiable>::add_malloc_device ()` to apply the
163- property to an USM pointer.
161+ * The `property_list` parameter of `sycl::malloc_device()`,
162+ `sycl::aligned_alloc_device()`, `sycl::malloc_shared ()`, or
163+ `sycl::aligned_alloc_shared()` to apply the property to an USM pointer.
164164
165165```c++
166166sycl::ext::oneapi::experimental::property::promote_local
@@ -195,9 +195,9 @@ The following property can be passed to three different APIs, namely:
195195* The `accessor` constructor, giving a more granular control
196196* The `buffer` constructor, in which case all the `accessors` derived from
197197this buffer will inherit this property (unless overridden).
198- * The `property_list` parameter of
199- `command_graph<graph_state::modifiable>::add_malloc_device ()` to apply the
200- property to an USM pointer.
198+ * The `property_list` parameter of `sycl::malloc_device()`,
199+ `sycl::aligned_alloc_device()`, `sycl::malloc_shared ()`, or
200+ `sycl::aligned_alloc_shared()` to apply the property to an USM pointer.
201201
202202```c++
203203sycl::ext::oneapi::experimental::property::promote_private
@@ -552,40 +552,36 @@ int main() {
552552
553553 int *dIn1, dIn2, dIn3, dTmp, dOut;
554554
555- auto node_in1 = graph.add_malloc_device(dIn1, numBytes );
556- auto node_in2 = graph.add_malloc_device(dIn2, numBytes );
557- auto node_in3 = graph.add_malloc_device(dIn3, numBytes );
558- auto node_out = graph.add_malloc_device(dOut, numBytes );
555+ dIn1 = malloc_device<int>(q, dataSize );
556+ dIn2 = malloc_device<int>(q, dataSize );
557+ dIn3 = malloc_device<int>(q, dataSize );
558+ dOut = malloc_device<int>(q, dataSize );
559559
560560 // Specify internalization for an USM pointer
561- auto node_tmp = graph.add_malloc_device(
562- dTmp, numBytes,
561+ dTmp = malloc_device<int>(q, dataSize,
563562 {sycl::ext::oneapi::experimental::property::promote_private});
564563
565564 // This explicit memory operation is compatible with fusion, as it can be
566565 // linearized before any device kernel in the graph.
567566 auto copy_in1 =
568- graph.add([&](handler &cgh) { cgh.memcpy(dIn1, in1, numBytes); },
569- {sycl_ext::property::node::depends_on(node_in1)});
567+ graph.add([&](handler &cgh) { cgh.memcpy(dIn1, in1, numBytes); });
570568
571569 // This explicit memory operation is compatible with fusion, as it can be
572570 // linearized before any device kernel in the graph.
573571 auto copy_in2 =
574- graph.add([&](handler &cgh) { cgh.memcpy(dIn2, in2, numBytes); },
575- {sycl_ext::property::node::depends_on(node_in2)});
572+ graph.add([&](handler &cgh) { cgh.memcpy(dIn2, in2, numBytes); });
576573
577574 auto kernel1 = graph.add(
578575 [&](handler &cgh) {
579576 cgh.parallel_for<class KernelOne>(
580577 dataSize, [=](id<1> i) { tmp[i] = in1[i] + in2[i]; });
581578 },
582- {sycl_ext::property::node::depends_on(copy_in1, copy_in2, node_tmp )});
579+ {sycl_ext::property::node::depends_on(copy_in1, copy_in2)});
583580
584581 // This explicit memory operation is compatible with fusion, as it can be
585582 // linearized before any device kernel in the graph.
586583 auto copy_in3 =
587- graph.add([&](handler &cgh) { cgh.memcpy(dIn3, in3, numBytes); },
588- {sycl_ext::property::node::depends_on(node_in3)});
584+ graph.add([&](handler &cgh) { cgh.memcpy(dIn3, in3, numBytes); });
589585
590586 auto kernel2 = graph.add(
591587 [&](handler &cgh) {
@@ -600,12 +596,6 @@ int main() {
600596 graph.add([&](handler &cgh) { cgh.memcpy(out, dOut, numBytes); },
601597 {sycl_ext::property::node::depends_on(kernel2)});
602598
603- graph.add_free(dIn1, {sycl_ext::property::node::depends_on(copy_out)});
604- graph.add_free(dIn2, {sycl_ext::property::node::depends_on(copy_out)});
605- graph.add_free(dIn3, {sycl_ext::property::node::depends_on(copy_out)});
606- graph.add_free(dTmp, {sycl_ext::property::node::depends_on(copy_out)});
607- graph.add_free(dOut, {sycl_ext::property::node::depends_on(copy_out)});
608-
609599 // Trigger fusion during finalization.
610600 auto exec = graph.finalize(q.get_context(),
611601 {sycl::ext::oneapi::experimental::property::
@@ -614,6 +604,12 @@ int main() {
614604 // use queue shortcut for graph submission
615605 q.ext_oneapi_graph(exec).wait();
616606
607+ free(dIn1, q);
608+ free(dIn2, q);
609+ free(dIn3, q);
610+ free(dOut, q);
611+ free(dTmp, q);
612+
617613 return 0;
618614}
619615```
@@ -633,4 +629,5 @@ Ewan Crawford, Codeplay +
633629|========================================
634630|Rev|Date|Authors|Changes
635631|1|2023-02-16|Lukas Sommer|*Initial draft*
632+ |2|2023-03-16|Lukas Sommer|*Remove reference to outdated `add_malloc_device` API*
636633|========================================
0 commit comments