@@ -208,10 +208,10 @@ namespace sycl::ext::oneapi::experimental::property{
208
208
template<sycl::memory_scope Scope>
209
209
struct access_scope {};
210
210
211
- inline constexpr auto access_scope_local =
211
+ inline constexpr auto access_scope_work_group =
212
212
access_scope<memory_scope_work_group>;
213
213
214
- inline constexpr auto access_scope_private =
214
+ inline constexpr auto access_scope_work_item =
215
215
access_scope<memory_scope_work_item>;
216
216
217
217
} // namespace sycl::ext::oneapi::experimental::property
@@ -227,11 +227,11 @@ the property is specified on an accessor) or in any kernel in the graph (in case
227
227
the property is specified on a buffer or an USM pointer).
228
228
229
229
More concretely, the two shortcuts express the following semantics:
230
- * `access_scope_local `: Applying this specialization asserts that each element
231
- in the buffer or allocated device memory is accessed by no more than one
230
+ * `access_scope_work_group `: Applying this specialization asserts that each
231
+ element in the buffer or allocated device memory is accessed by no more than one
232
232
work-group.
233
- * `access_scope_private `: Applying this specialization asserts that each element
234
- in the buffer or allocated device memory is accessed by no more than one
233
+ * `access_scope_work_item `: Applying this specialization asserts that each
234
+ element in the buffer or allocated device memory is accessed by no more than one
235
235
work-item.
236
236
237
237
Implementations may treat specializations of the access scope property as a
@@ -478,31 +478,31 @@ properties must be combined as follows:
478
478
|None
479
479
|None
480
480
481
- |Local
482
- |Local
481
+ |Work Group
482
+ |Work Group
483
483
484
- |Private
485
- |Private
484
+ |Work Item
485
+ |Work Item
486
486
487
- .3+.^|Local
487
+ .3+.^|Work Group
488
488
|None
489
- |Local
489
+ |Work Group
490
490
491
- |Local
492
- |Local
491
+ |Work Group
492
+ |Work Group
493
493
494
- |Private
494
+ |Work Item
495
495
|*Error*
496
496
497
- .3+.^|Private
497
+ .3+.^|Work Item
498
498
|None
499
- |Private
499
+ |Work Item
500
500
501
- |Local
501
+ |Work Group
502
502
|*Error*
503
503
504
- |Private
505
- |Private
504
+ |Work Item
505
+ |Work Item
506
506
|===
507
507
508
508
In case different internalization targets are used for accessors to the same
@@ -516,16 +516,16 @@ buffer, the following (commutative and associative) rules are followed:
516
516
|_Any_
517
517
|None
518
518
519
- .2+.^|Local
520
- |Local
521
- |Local
519
+ .2+.^|Work Group
520
+ |Work Group
521
+ |Work Group
522
522
523
- |Private
523
+ |Work Item
524
524
|None
525
525
526
- |Private
527
- |Private
528
- |Private
526
+ |Work Item
527
+ |Work Item
528
+ |Work Item
529
529
|===
530
530
531
531
If no work-group size is specified or two accessors specify different
@@ -569,13 +569,13 @@ int main() {
569
569
// Internalization specified on the buffer
570
570
buffer<int> bTmp2{
571
571
range{dataSize},
572
- {sycl::ext::oneapi::experimental::property::access_scope_private {},
572
+ {sycl::ext::oneapi::experimental::property::access_scope_work_item {},
573
573
sycl::ext::oneapi::experimental::property::fusion_internal_memory{},
574
574
no_init}};
575
575
// Internalization specified on the buffer
576
576
buffer<int> bTmp3{
577
577
range{dataSize},
578
- {sycl::ext::oneapi::experimental::property::access_scope_private {},
578
+ {sycl::ext::oneapi::experimental::property::access_scope_work_item {},
579
579
sycl::ext::oneapi::experimental::property::fusion_internal_memory{},
580
580
no_init}};
581
581
buffer<int> bOut{out, range{dataSize}};
@@ -592,17 +592,17 @@ int main() {
592
592
auto accIn1 = bIn1.get_access(cgh);
593
593
auto accIn2 = bIn2.get_access(cgh);
594
594
// Internalization specified on each accessor.
595
- auto accTmp1 = bTmp1.get_access(
596
- cgh, sycl::ext::oneapi::experimental::property::access_scope_private {}
595
+ auto accTmp1 = bTmp1.get_access(cgh,
596
+ sycl::ext::oneapi::experimental::property::access_scope_work_item {}
597
597
sycl::ext::oneapi::experimental::property::fusion_internal_memory{},
598
598
no_init);
599
599
cgh.parallel_for<AddKernel>(dataSize, AddKernel{accIn1, accIn2, accTmp1});
600
600
});
601
601
602
602
q.submit([&](handler &cgh) {
603
603
// Internalization specified on each accessor.
604
- auto accTmp1 = bTmp1.get_access(
605
- cgh, sycl::ext::oneapi::experimental::property::access_scope_private {}
604
+ auto accTmp1 = bTmp1.get_access(cgh,
605
+ sycl::ext::oneapi::experimental::property::access_scope_work_item {}
606
606
sycl::ext::oneapi::experimental::property::fusion_internal_memory{},
607
607
no_init);
608
608
auto accIn3 = bIn3.get_access(cgh);
@@ -613,8 +613,8 @@ int main() {
613
613
614
614
q.submit([&](handler &cgh) {
615
615
// Internalization specified on each accessor.
616
- auto accTmp1 = bTmp1.get_access(
617
- cgh, sycl::ext::oneapi::experimental::property::access_scope_private {}
616
+ auto accTmp1 = bTmp1.get_access(cgh,
617
+ sycl::ext::oneapi::experimental::property::access_scope_work_item {}
618
618
sycl::ext::oneapi::experimental::property::fusion_internal_memory{},
619
619
no_init);
620
620
auto accTmp3 = bTmp3.get_access(cgh);
@@ -674,7 +674,7 @@ int main() {
674
674
// Specify internalization for an USM pointer
675
675
dTmp = malloc_device<int>(
676
676
q, dataSize,
677
- {sycl_ext::property::access_scope_private {},
677
+ {sycl_ext::property::access_scope_work_item {},
678
678
sycl_ext::property::fusion_internal_memory{}, no_init});
679
679
680
680
// This explicit memory operation is compatible with fusion, as it can be
0 commit comments