88// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
99// REQUIRES-INTEL-DRIVER: lin: 31155
1010
11- // XFAIL: linux && gpu-intel-dg2
12- // XFAIL-TRACKER: https://github.com/intel/llvm/issues/15812
13-
1411// RUN: %{build} -o %t.out
1512// RUN: %{run} %t.out
1613
2118static constexpr int WorkGroupSize = 16 ;
2219
2320static constexpr int VL = 16 ;
24- int main ( ) {
21+ template < bool UseThisWorkItemAPI> bool test (sycl::queue &q ) {
2522 bool Pass = true ;
26- sycl::queue q;
27- esimd_test::printTestLabel (q);
2823 const auto MaxWGs = 8 ;
2924 size_t WorkItemCount = MaxWGs * WorkGroupSize * VL;
30-
25+ std::cout << " Test case UseThisWorkItemAPI="
26+ << std::to_string (UseThisWorkItemAPI) << std::endl;
3127 const auto Props = sycl::ext::oneapi::experimental::properties{
3228 sycl::ext::oneapi::experimental::use_root_sync};
3329 sycl::buffer<int > DataBuf{sycl::range{WorkItemCount}};
@@ -40,14 +36,14 @@ int main() {
4036 // Write data to another kernel's data to verify the barrier works.
4137 __ESIMD_NS::block_store (
4238 Data, (WorkItemCount * sizeof (int )) - (ID * sizeof (int ) * VL), V);
43- if (ID % 2 == 1 ) {
44- auto Root = it.ext_oneapi_get_root_group ();
45- sycl::group_barrier (Root);
46- } else {
39+ if constexpr (UseThisWorkItemAPI) {
4740 auto Root =
4841 sycl::ext::oneapi::experimental::this_work_item::get_root_group<
4942 1 >();
5043 sycl::group_barrier (Root);
44+ } else {
45+ auto Root = it.ext_oneapi_get_root_group ();
46+ sycl::group_barrier (Root);
5147 }
5248 __ESIMD_NS::simd<int , VL> VOther (ID * VL, 1 );
5349 __ESIMD_NS::block_store (Data, ID * sizeof (int ) * VL, VOther);
@@ -63,6 +59,14 @@ int main() {
6359 << " ] != " << std::to_string (I) << " \n " ;
6460 }
6561 }
62+ return Pass;
63+ }
64+ int main () {
65+ sycl::queue q;
66+ esimd_test::printTestLabel (q);
67+ bool Pass = true ;
68+ Pass &= test<true >(q);
69+ Pass &= test<false >(q);
6670 if (Pass)
6771 std::cout << " Passed\n " ;
6872 else
0 commit comments