18
18
19
19
namespace syclex = sycl::ext::oneapi::experimental;
20
20
21
- static constexpr int WorkGroupSize = 16 ;
21
+ static constexpr int WorkGroupSize = 32 ;
22
22
23
23
static constexpr int VL = 16 ;
24
24
@@ -33,20 +33,18 @@ template <bool UseThisWorkItemAPI> bool test(sycl::queue &q) {
33
33
auto Bundle =
34
34
sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context ());
35
35
auto Kernel = Bundle.template get_kernel <MyKernel<UseThisWorkItemAPI>>();
36
- sycl::range<1 > LocalRange{WorkGroupSize};
36
+ sycl::range<3 > LocalRange{WorkGroupSize, 1 , 1 };
37
37
auto MaxWGs = Kernel.template ext_oneapi_get_info <
38
38
syclex::info::kernel_queue_specific::max_num_work_groups>(q, LocalRange,
39
39
0 );
40
40
auto GlobalRange = LocalRange;
41
- GlobalRange[0 ] *= MaxWGs / VL;
42
41
size_t WorkItemCount = GlobalRange.size () * VL;
43
42
sycl::buffer<int > DataBuf{WorkItemCount};
44
- const auto Range = sycl::nd_range<1 >{GlobalRange, LocalRange};
45
-
43
+ const auto Range = sycl::nd_range<3 >{GlobalRange, LocalRange};
46
44
q.submit ([&](sycl::handler &h) {
47
45
sycl::accessor Data{DataBuf, h};
48
46
h.parallel_for <MyKernel<UseThisWorkItemAPI>>(
49
- Range, Props, [=](sycl::nd_item<1 > it) SYCL_ESIMD_KERNEL {
47
+ Range, Props, [=](sycl::nd_item<3 > it) SYCL_ESIMD_KERNEL {
50
48
int ID = it.get_global_linear_id ();
51
49
__ESIMD_NS::simd<int , VL> V (ID, 1 );
52
50
// Write data to another kernel's data to verify the barrier works.
0 commit comments