|
14 | 14 | #include "esimd_test_utils.hpp" |
15 | 15 | #include <sycl/ext/oneapi/experimental/root_group.hpp> |
16 | 16 | #include <sycl/group_barrier.hpp> |
| 17 | +#include <sycl/kernel_bundle.hpp> |
| 18 | + |
| 19 | +namespace syclex = sycl::ext::oneapi::experimental; |
17 | 20 |
|
18 | 21 | static constexpr int WorkGroupSize = 16; |
19 | 22 |
|
20 | 23 | static constexpr int VL = 16; |
| 24 | + |
| 25 | +template <int Val> class MyKernel; |
| 26 | + |
21 | 27 | template <bool UseThisWorkItemAPI> bool test(sycl::queue &q) { |
22 | 28 | bool Pass = true; |
23 | | - const auto MaxWGs = 8; |
24 | | - size_t WorkItemCount = MaxWGs * WorkGroupSize * VL; |
25 | 29 | std::cout << "Test case UseThisWorkItemAPI=" |
26 | 30 | << std::to_string(UseThisWorkItemAPI) << std::endl; |
27 | 31 | const auto Props = sycl::ext::oneapi::experimental::properties{ |
28 | 32 | sycl::ext::oneapi::experimental::use_root_sync}; |
29 | | - sycl::buffer<int> DataBuf{sycl::range{WorkItemCount}}; |
30 | | - const auto Range = sycl::nd_range<1>{MaxWGs * WorkGroupSize, WorkGroupSize}; |
| 33 | + auto Bundle = |
| 34 | + sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context()); |
| 35 | + auto Kernel = Bundle.template get_kernel<MyKernel<UseThisWorkItemAPI>>(); |
| 36 | + sycl::range<1> LocalRange{WorkGroupSize}; |
| 37 | + auto MaxWGs = Kernel.template ext_oneapi_get_info< |
| 38 | + syclex::info::kernel_queue_specific::max_num_work_groups>(q, LocalRange, |
| 39 | + 0); |
| 40 | + auto GlobalRange = LocalRange; |
| 41 | + GlobalRange[0] *= MaxWGs / VL; |
| 42 | + size_t WorkItemCount = GlobalRange.size() * VL; |
| 43 | + sycl::buffer<int> DataBuf{WorkItemCount}; |
| 44 | + const auto Range = sycl::nd_range<1>{GlobalRange, LocalRange}; |
| 45 | + |
31 | 46 | q.submit([&](sycl::handler &h) { |
32 | 47 | sycl::accessor Data{DataBuf, h}; |
33 | | - h.parallel_for(Range, Props, [=](sycl::nd_item<1> it) SYCL_ESIMD_KERNEL { |
34 | | - int ID = it.get_global_linear_id(); |
35 | | - __ESIMD_NS::simd<int, VL> V(ID, 1); |
36 | | - // Write data to another kernel's data to verify the barrier works. |
37 | | - __ESIMD_NS::block_store( |
38 | | - Data, (WorkItemCount * sizeof(int)) - (ID * sizeof(int) * VL), V); |
39 | | - if constexpr (UseThisWorkItemAPI) { |
40 | | - auto Root = |
41 | | - sycl::ext::oneapi::experimental::this_work_item::get_root_group< |
42 | | - 1>(); |
43 | | - sycl::group_barrier(Root); |
44 | | - } else { |
45 | | - auto Root = it.ext_oneapi_get_root_group(); |
46 | | - sycl::group_barrier(Root); |
47 | | - } |
48 | | - __ESIMD_NS::simd<int, VL> VOther(ID * VL, 1); |
49 | | - __ESIMD_NS::block_store(Data, ID * sizeof(int) * VL, VOther); |
50 | | - }); |
| 48 | + h.parallel_for<MyKernel<UseThisWorkItemAPI>>( |
| 49 | + Range, Props, [=](sycl::nd_item<1> it) SYCL_ESIMD_KERNEL { |
| 50 | + int ID = it.get_global_linear_id(); |
| 51 | + __ESIMD_NS::simd<int, VL> V(ID, 1); |
| 52 | + // Write data to another kernel's data to verify the barrier works. |
| 53 | + __ESIMD_NS::block_store( |
| 54 | + Data, (WorkItemCount * sizeof(int)) - (ID * sizeof(int) * VL), |
| 55 | + V); |
| 56 | + if constexpr (UseThisWorkItemAPI) { |
| 57 | + auto Root = sycl::ext::oneapi::experimental::this_work_item:: |
| 58 | + get_root_group<1>(); |
| 59 | + sycl::group_barrier(Root); |
| 60 | + } else { |
| 61 | + auto Root = it.ext_oneapi_get_root_group(); |
| 62 | + sycl::group_barrier(Root); |
| 63 | + } |
| 64 | + __ESIMD_NS::simd<int, VL> VOther(ID * VL, 1); |
| 65 | + __ESIMD_NS::block_store(Data, ID * sizeof(int) * VL, VOther); |
| 66 | + }); |
51 | 67 | }).wait(); |
52 | 68 | sycl::host_accessor Data{DataBuf}; |
53 | 69 | int ErrCnt = 0; |
|
0 commit comments