@@ -158,9 +158,9 @@ The following property can be passed to three different APIs, namely:
158
158
* The `accessor` constructor, giving a more granular control
159
159
* The `buffer` constructor, in which case all the `accessors` derived from
160
160
this 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.
164
164
165
165
```c++
166
166
sycl::ext::oneapi::experimental::property::promote_local
@@ -195,9 +195,9 @@ The following property can be passed to three different APIs, namely:
195
195
* The `accessor` constructor, giving a more granular control
196
196
* The `buffer` constructor, in which case all the `accessors` derived from
197
197
this 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.
201
201
202
202
```c++
203
203
sycl::ext::oneapi::experimental::property::promote_private
@@ -552,40 +552,36 @@ int main() {
552
552
553
553
int *dIn1, dIn2, dIn3, dTmp, dOut;
554
554
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 );
559
559
560
560
// Specify internalization for an USM pointer
561
- auto node_tmp = graph.add_malloc_device(
562
- dTmp, numBytes,
561
+ dTmp = malloc_device<int>(q, dataSize,
563
562
{sycl::ext::oneapi::experimental::property::promote_private});
564
563
565
564
// This explicit memory operation is compatible with fusion, as it can be
566
565
// linearized before any device kernel in the graph.
567
566
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); });
570
568
571
569
// This explicit memory operation is compatible with fusion, as it can be
572
570
// linearized before any device kernel in the graph.
573
571
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); });
576
573
577
574
auto kernel1 = graph.add(
578
575
[&](handler &cgh) {
579
576
cgh.parallel_for<class KernelOne>(
580
577
dataSize, [=](id<1> i) { tmp[i] = in1[i] + in2[i]; });
581
578
},
582
- {sycl_ext::property::node::depends_on(copy_in1, copy_in2, node_tmp )});
579
+ {sycl_ext::property::node::depends_on(copy_in1, copy_in2)});
583
580
584
581
// This explicit memory operation is compatible with fusion, as it can be
585
582
// linearized before any device kernel in the graph.
586
583
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); });
589
585
590
586
auto kernel2 = graph.add(
591
587
[&](handler &cgh) {
@@ -600,12 +596,6 @@ int main() {
600
596
graph.add([&](handler &cgh) { cgh.memcpy(out, dOut, numBytes); },
601
597
{sycl_ext::property::node::depends_on(kernel2)});
602
598
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
-
609
599
// Trigger fusion during finalization.
610
600
auto exec = graph.finalize(q.get_context(),
611
601
{sycl::ext::oneapi::experimental::property::
@@ -614,6 +604,12 @@ int main() {
614
604
// use queue shortcut for graph submission
615
605
q.ext_oneapi_graph(exec).wait();
616
606
607
+ free(dIn1, q);
608
+ free(dIn2, q);
609
+ free(dIn3, q);
610
+ free(dOut, q);
611
+ free(dTmp, q);
612
+
617
613
return 0;
618
614
}
619
615
```
@@ -633,4 +629,5 @@ Ewan Crawford, Codeplay +
633
629
|========================================
634
630
|Rev|Date|Authors|Changes
635
631
|1|2023-02-16|Lukas Sommer|*Initial draft*
632
+ |2|2023-03-16|Lukas Sommer|*Remove reference to outdated `add_malloc_device` API*
636
633
|========================================
0 commit comments