forked from intel/llvm
-
Notifications
You must be signed in to change notification settings - Fork 0
/
group_barrier.cpp
66 lines (61 loc) · 3.54 KB
/
group_barrier.cpp
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
// RUN: %clangxx -fsycl-device-only -fsycl-unnamed-lambda -S -Xclang -emit-llvm -Xclang -no-enable-noundef-analysis %s -o - | FileCheck %s
#include <sycl/sycl.hpp>
const auto TestLambda = [](auto G) {
sycl::group_barrier(G);
sycl::group_barrier(G, sycl::memory_scope_work_item);
sycl::group_barrier(G, sycl::memory_scope_sub_group);
sycl::group_barrier(G, sycl::memory_scope_work_group);
sycl::group_barrier(G, sycl::memory_scope_device);
sycl::group_barrier(G, sycl::memory_scope_system);
};
int main() {
sycl::queue Q;
Q.submit([](sycl::handler &CGH) {
CGH.parallel_for(sycl::nd_range{sycl::range{1}, sycl::range{1}},
[](sycl::nd_item<1> item) {
auto G = item.get_group();
auto SG = item.get_sub_group();
TestLambda(G);
TestLambda(SG);
});
});
Q.submit([](sycl::handler &CGH) {
CGH.parallel_for(sycl::nd_range{sycl::range{1, 1}, sycl::range{1, 1}},
[](sycl::nd_item<2> item) {
auto G = item.get_group();
TestLambda(G);
});
});
Q.submit([](sycl::handler &CGH) {
CGH.parallel_for(sycl::nd_range{sycl::range{1, 1, 1}, sycl::range{1, 1, 1}},
[](sycl::nd_item<3> item) {
auto G = item.get_group();
TestLambda(G);
});
});
return 0;
}
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 4, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 3, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 1, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 3, i32 0, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 4, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 3, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 1, i32 912)
// CHECK: tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 0, i32 912)