33//
44// SPDX-License-Identifier: MIT
55// =============================================================
6- #include < CL/sycl.hpp>
76#include < iostream>
7+ #include < sycl/sycl.hpp>
88
9- template <typename T>
10- auto get_multi_ptr (T *raw_ptr) {
9+ template <typename T> auto get_multi_ptr (T *raw_ptr) {
1110 auto multi_ptr =
12- sycl::address_space_cast<
13- sycl::access::address_space::global_space,
14- sycl::access::decorated::yes>(raw_ptr);
11+ sycl::address_space_cast<sycl::access::address_space::global_space,
12+ sycl::access::decorated::yes>(raw_ptr);
1513
1614 return multi_ptr;
1715}
1816
19-
2017int main () {
2118 sycl::queue q{sycl::gpu_selector_v,
2219 sycl::property::queue::enable_profiling{}};
@@ -30,21 +27,21 @@ int main() {
3027 memset (data2, 0xFF , sizeof (int ) * N);
3128
3229 auto e = q.submit ([&](auto &h) {
33- h.parallel_for (
34- sycl::nd_range (sycl::range{N / 16 }, sycl::range{ 32 }),
35- [=](sycl::nd_item< 1 > it) [[ intel::reqd_sub_group_size ( 16 )]] {
36- auto sg = it. get_sub_group () ;
37- sycl::vec< int , 8 > x;
38-
39- int base = (it.get_group (0 ) * 32 +
40- sg.get_group_id ()[0 ] * sg.get_local_range ()[0 ]) *
41- 16 ;
42-
43- x = sg.load <8 >(get_multi_ptr (&(data2[base + 0 ])));
44- sg.store <8 >(get_multi_ptr (&(data[base + 0 ])), x);
45- x = sg.load <8 >(get_multi_ptr (&(data2[base + 128 ])));
46- sg.store <8 >(get_multi_ptr (&(data[base + 128 ])), x);
47- });
30+ h.parallel_for (sycl::nd_range (sycl::range{N / 16 }, sycl::range{ 32 }),
31+ [=] (sycl::nd_item< 1 > it) [[ intel::reqd_sub_group_size ( 16 )]] {
32+ auto sg = it. get_sub_group ();
33+ sycl::vec< int , 8 > x ;
34+
35+ int base =
36+ (it.get_group (0 ) * 32 +
37+ sg.get_group_id ()[0 ] * sg.get_local_range ()[0 ]) *
38+ 16 ;
39+
40+ x = sg.load <8 >(get_multi_ptr (&(data2[base + 0 ])));
41+ sg.store <8 >(get_multi_ptr (&(data[base + 0 ])), x);
42+ x = sg.load <8 >(get_multi_ptr (&(data2[base + 128 ])));
43+ sg.store <8 >(get_multi_ptr (&(data[base + 128 ])), x);
44+ });
4845 });
4946 // Snippet end
5047 q.wait ();
0 commit comments