8
8
// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
9
9
// REQUIRES-INTEL-DRIVER: lin: 31155
10
10
11
- // XFAIL: linux && gpu-intel-dg2
12
- // XFAIL-TRACKER: https://github.com/intel/llvm/issues/15812
13
-
14
11
// RUN: %{build} -o %t.out
15
12
// RUN: %{run} %t.out
16
13
21
18
static constexpr int WorkGroupSize = 16 ;
22
19
23
20
static constexpr int VL = 16 ;
24
- int main ( ) {
21
+ template < bool UseThisWorkItemAPI> bool test (sycl::queue &q ) {
25
22
bool Pass = true ;
26
- sycl::queue q;
27
- esimd_test::printTestLabel (q);
28
23
const auto MaxWGs = 8 ;
29
24
size_t WorkItemCount = MaxWGs * WorkGroupSize * VL;
30
-
25
+ std::cout << " Test case UseThisWorkItemAPI="
26
+ << std::to_string (UseThisWorkItemAPI) << std::endl;
31
27
const auto Props = sycl::ext::oneapi::experimental::properties{
32
28
sycl::ext::oneapi::experimental::use_root_sync};
33
29
sycl::buffer<int > DataBuf{sycl::range{WorkItemCount}};
@@ -40,14 +36,14 @@ int main() {
40
36
// Write data to another kernel's data to verify the barrier works.
41
37
__ESIMD_NS::block_store (
42
38
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) {
47
40
auto Root =
48
41
sycl::ext::oneapi::experimental::this_work_item::get_root_group<
49
42
1 >();
50
43
sycl::group_barrier (Root);
44
+ } else {
45
+ auto Root = it.ext_oneapi_get_root_group ();
46
+ sycl::group_barrier (Root);
51
47
}
52
48
__ESIMD_NS::simd<int , VL> VOther (ID * VL, 1 );
53
49
__ESIMD_NS::block_store (Data, ID * sizeof (int ) * VL, VOther);
@@ -63,6 +59,14 @@ int main() {
63
59
<< " ] != " << std::to_string (I) << " \n " ;
64
60
}
65
61
}
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);
66
70
if (Pass)
67
71
std::cout << " Passed\n " ;
68
72
else
0 commit comments