Skip to content

Commit 0fd711c

Browse files
jzcAlexeySachkov
authored and
Chenyang-L
committed
[SYCL] Optional kernel features: implement split based on reqd-sub-group-size (#9928)
Based off #8167 --------- Co-authored-by: Alexey Sachkov <[email protected]>
1 parent 3e72b46 commit 0fd711c

File tree

13 files changed

+531
-107
lines changed

13 files changed

+531
-107
lines changed

‎llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll

Lines changed: 133 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,133 @@
1+
; This test emulates two translation units with 3 kernels:
2+
; TU0_kernel0 - 1st translation unit, no reqd_sub_group_size attribute used
3+
; TU0_kernel1 - 1st translation unit, reqd_sub_group_size attribute is used
4+
; TU1_kernel2 - 2nd translation unit, no reqd_sub_group_size attribute used
5+
6+
; The test is intended to check that sycl-post-link correctly separates kernels
7+
; that use reqd_sub_group_size attributes from kernels which doesn't use them
8+
; regardless of device code split mode
9+
10+
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
11+
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
12+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
13+
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
14+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
15+
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
16+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
17+
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
18+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
19+
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
20+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
21+
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
22+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
23+
24+
; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table
25+
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
26+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
27+
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
28+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
29+
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
30+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
31+
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
32+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
33+
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
34+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
35+
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
36+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
37+
38+
; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table
39+
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
40+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
41+
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
42+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
43+
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
44+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
45+
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
46+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
47+
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
48+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
49+
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
50+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
51+
52+
; Regardless of device code split mode, each kernel should go into a separate
53+
; device image
54+
55+
; CHECK-M2-IR: define {{.*}} @TU0_kernel0
56+
; CHECK-M2-SYMS: TU0_kernel0
57+
58+
; CHECK-M1-IR: define {{.*}} @TU0_kernel1
59+
; CHECK-M1-SYMS: TU0_kernel1
60+
61+
; CHECK-M0-IR: define {{.*}} @TU1_kernel2
62+
; CHECK-M0-SYMS: TU1_kernel2
63+
64+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
65+
target triple = "spir64-unknown-linux"
66+
67+
; FIXME: device globals should also be properly distributed across device images
68+
; if they are of optional type
69+
@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4
70+
71+
define dso_local spir_kernel void @TU0_kernel0() #0 {
72+
entry:
73+
call spir_func void @foo()
74+
ret void
75+
}
76+
77+
define dso_local spir_func void @foo() {
78+
entry:
79+
%a = alloca i32, align 4
80+
%call = call spir_func i32 @bar(i32 1)
81+
%add = add nsw i32 2, %call
82+
store i32 %add, i32* %a, align 4
83+
ret void
84+
}
85+
86+
; Function Attrs: nounwind
87+
define linkonce_odr dso_local spir_func i32 @bar(i32 %arg) {
88+
entry:
89+
%arg.addr = alloca i32, align 4
90+
store i32 %arg, i32* %arg.addr, align 4
91+
%0 = load i32, i32* %arg.addr, align 4
92+
ret i32 %0
93+
}
94+
95+
define dso_local spir_kernel void @TU0_kernel1() #0 !intel_reqd_sub_group_size !2 {
96+
entry:
97+
call spir_func void @foo1()
98+
ret void
99+
}
100+
101+
; Function Attrs: nounwind
102+
define dso_local spir_func void @foo1() {
103+
entry:
104+
%a = alloca i32, align 4
105+
store i32 2, i32* %a, align 4
106+
ret void
107+
}
108+
109+
define dso_local spir_kernel void @TU1_kernel2() #1 {
110+
entry:
111+
call spir_func void @foo2()
112+
ret void
113+
}
114+
115+
; Function Attrs: nounwind
116+
define dso_local spir_func void @foo2() {
117+
entry:
118+
%a = alloca i32, align 4
119+
%0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @_ZL2GV to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
120+
%add = add nsw i32 4, %0
121+
store i32 %add, i32* %a, align 4
122+
ret void
123+
}
124+
125+
attributes #0 = { "sycl-module-id"="TU1.cpp" }
126+
attributes #1 = { "sycl-module-id"="TU2.cpp" }
127+
128+
!opencl.spir.version = !{!0, !0}
129+
!spirv.Source = !{!1, !1}
130+
131+
!0 = !{i32 1, i32 2}
132+
!1 = !{i32 4, i32 100000}
133+
!2 = !{i32 32}

‎llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
; The test is intended to check that sycl-post-link correctly groups kernels
2+
; by unique reqd_sub_group_size values used in them
3+
4+
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
5+
; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE
6+
;
7+
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \
8+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel3
9+
;
10+
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M1-SYMS \
11+
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 \
12+
; RUN: --implicit-check-not kernel3
13+
14+
;
15+
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \
16+
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 \
17+
; RUN: --implicit-check-not kernel2
18+
19+
; CHECK-TABLE: Code
20+
; CHECK-TABLE-NEXT: _0.sym
21+
; CHECK-TABLE-NEXT: _1.sym
22+
; CHECK-TABLE-NEXT: _2.sym
23+
; CHECK-TABLE-EMPTY:
24+
25+
; CHECK-M0-SYMS: kernel1
26+
; CHECK-M0-SYMS: kernel2
27+
28+
; CHECK-M1-SYMS: kernel0
29+
30+
; CHECK-M2-SYMS: kernel3
31+
32+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
33+
target triple = "spir64-unknown-linux"
34+
35+
define dso_local spir_kernel void @kernel0() #0 !intel_reqd_sub_group_size !1 {
36+
entry:
37+
ret void
38+
}
39+
40+
define dso_local spir_kernel void @kernel1() #0 !intel_reqd_sub_group_size !2 {
41+
entry:
42+
ret void
43+
}
44+
45+
define dso_local spir_kernel void @kernel2() #0 !intel_reqd_sub_group_size !3 {
46+
entry:
47+
ret void
48+
}
49+
50+
define dso_local spir_kernel void @kernel3() #0 !intel_reqd_sub_group_size !4 {
51+
entry:
52+
ret void
53+
}
54+
55+
attributes #0 = { "sycl-module-id"="TU1.cpp" }
56+
57+
!1 = !{i32 32}
58+
!2 = !{i32 64}
59+
!3 = !{i32 64}
60+
!4 = !{i32 16}

‎llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll

Lines changed: 125 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
1+
; Original code:
2+
; Compile with: clang++ -fsycl -fsycl-device-only -fno-sycl-instrument-device-code -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S reqd-sub-group-size.cpp
3+
; #include <sycl/sycl.hpp>
4+
5+
; int main() {
6+
; sycl::queue q;
7+
; q.submit([&](sycl::handler &h) {
8+
; h.parallel_for<class KernelA>(
9+
; sycl::range<1>(32),
10+
; [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(16)]] {});
11+
; });
12+
; q.submit([&](sycl::handler &h) {
13+
; h.parallel_for<class KernelB>(
14+
; sycl::range<1>(32),
15+
; [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(32)]] {});
16+
; });
17+
; q.submit([&](sycl::handler &h) {
18+
; h.parallel_for<class KernelC>(
19+
; sycl::range<1>(32),
20+
; [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(16)]] {});
21+
; });
22+
; return 0;
23+
; }
24+
25+
; RUN: sycl-post-link -split=auto %s -o %t.table
26+
; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT-0
27+
; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-AUTO-SPLIT-1
28+
29+
; CHECK-PROP-AUTO-SPLIT-0: [SYCL/device requirements]
30+
; CHECK-PROP-AUTO-SPLIT-0: reqd_sub_group_size=1|32
31+
32+
; CHECK-PROP-AUTO-SPLIT-1: [SYCL/device requirements]
33+
; CHECK-PROP-AUTO-SPLIT-1: reqd_sub_group_size=1|16
34+
35+
; ModuleID = 'foo.cpp'
36+
source_filename = "foo.cpp"
37+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
38+
target triple = "spir64-unknown-unknown"
39+
40+
$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E7KernelA = comdat any
41+
42+
$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E7KernelB = comdat any
43+
44+
$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_E7KernelC = comdat any
45+
46+
; Function Attrs: norecurse nounwind
47+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E7KernelA() local_unnamed_addr #0 comdat !srcloc !48 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !50 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
48+
entry:
49+
ret void
50+
}
51+
52+
; Function Attrs: norecurse nounwind
53+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E7KernelB() local_unnamed_addr #0 comdat !srcloc !51 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !52 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
54+
entry:
55+
ret void
56+
}
57+
58+
; Function Attrs: norecurse nounwind
59+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_E7KernelC() local_unnamed_addr #0 comdat !srcloc !53 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !50 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
60+
entry:
61+
ret void
62+
}
63+
64+
attributes #0 = { norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="foo.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
65+
66+
!llvm.module.flags = !{!0, !1}
67+
!opencl.spir.version = !{!2}
68+
!spirv.Source = !{!3}
69+
!sycl_aspects = !{!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}
70+
!llvm.ident = !{!47}
71+
72+
!0 = !{i32 1, !"wchar_size", i32 4}
73+
!1 = !{i32 7, !"frame-pointer", i32 2}
74+
!2 = !{i32 1, i32 2}
75+
!3 = !{i32 4, i32 100000}
76+
!4 = !{!"cpu", i32 1}
77+
!5 = !{!"gpu", i32 2}
78+
!6 = !{!"accelerator", i32 3}
79+
!7 = !{!"custom", i32 4}
80+
!8 = !{!"fp16", i32 5}
81+
!9 = !{!"fp64", i32 6}
82+
!10 = !{!"image", i32 9}
83+
!11 = !{!"online_compiler", i32 10}
84+
!12 = !{!"online_linker", i32 11}
85+
!13 = !{!"queue_profiling", i32 12}
86+
!14 = !{!"usm_device_allocations", i32 13}
87+
!15 = !{!"usm_host_allocations", i32 14}
88+
!16 = !{!"usm_shared_allocations", i32 15}
89+
!17 = !{!"usm_system_allocations", i32 17}
90+
!18 = !{!"ext_intel_pci_address", i32 18}
91+
!19 = !{!"ext_intel_gpu_eu_count", i32 19}
92+
!20 = !{!"ext_intel_gpu_eu_simd_width", i32 20}
93+
!21 = !{!"ext_intel_gpu_slices", i32 21}
94+
!22 = !{!"ext_intel_gpu_subslices_per_slice", i32 22}
95+
!23 = !{!"ext_intel_gpu_eu_count_per_subslice", i32 23}
96+
!24 = !{!"ext_intel_max_mem_bandwidth", i32 24}
97+
!25 = !{!"ext_intel_mem_channel", i32 25}
98+
!26 = !{!"usm_atomic_host_allocations", i32 26}
99+
!27 = !{!"usm_atomic_shared_allocations", i32 27}
100+
!28 = !{!"atomic64", i32 28}
101+
!29 = !{!"ext_intel_device_info_uuid", i32 29}
102+
!30 = !{!"ext_oneapi_srgb", i32 30}
103+
!31 = !{!"ext_oneapi_native_assert", i32 31}
104+
!32 = !{!"host_debuggable", i32 32}
105+
!33 = !{!"ext_intel_gpu_hw_threads_per_eu", i32 33}
106+
!34 = !{!"ext_oneapi_cuda_async_barrier", i32 34}
107+
!35 = !{!"ext_oneapi_bfloat16_math_functions", i32 35}
108+
!36 = !{!"ext_intel_free_memory", i32 36}
109+
!37 = !{!"ext_intel_device_id", i32 37}
110+
!38 = !{!"ext_intel_memory_clock_rate", i32 38}
111+
!39 = !{!"ext_intel_memory_bus_width", i32 39}
112+
!40 = !{!"emulated", i32 40}
113+
!41 = !{!"ext_intel_legacy_image", i32 41}
114+
!42 = !{!"int64_base_atomics", i32 7}
115+
!43 = !{!"int64_extended_atomics", i32 8}
116+
!44 = !{!"usm_system_allocator", i32 17}
117+
!45 = !{!"usm_restricted_shared_allocations", i32 16}
118+
!46 = !{!"host", i32 0}
119+
!47 = !{!"clang version 17.0.0 (https://github.com/jzc/llvm eed5b5576bef314433e8ae7313620dae399c9d22)"}
120+
!48 = !{i32 170}
121+
!49 = !{}
122+
!50 = !{i32 16}
123+
!51 = !{i32 351}
124+
!52 = !{i32 32}
125+
!53 = !{i32 532}

‎llvm/test/tools/sycl-post-link/registerallocmode.ll

Lines changed: 28 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -2,41 +2,49 @@
22

33
; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table
44
; RUN: FileCheck %s -input-file=%t.table
5-
; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
6-
; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
7-
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-LargeGRF-IR
8-
; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
9-
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-LargeGRF-SYM
10-
; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-PROP
11-
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
12-
; RUN: FileCheck %s -input-file=%t_esimd_1.prop --check-prefixes CHECK-ESIMD-PROP
13-
; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
145

156
; CHECK: [Code|Properties|Symbols]
16-
; CHECK: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
17-
; CHECK: {{.*}}_esimd_1.ll|{{.*}}_esimd_1.prop|{{.*}}_esimd_1.sym
18-
; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
7+
; CHECK-NEXT: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
8+
; CHECK-NEXT: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
9+
; CHECK-NEXT: {{.*}}_esimd_2.ll|{{.*}}_esimd_2.prop|{{.*}}_esimd_2.sym
10+
; CHECK-NEXT: {{.*}}_3.ll|{{.*}}_3.prop|{{.*}}_3.sym
1911

20-
; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
21-
; CHECK-ESIMD-LargeGRF-PROP: sycl-register-alloc-mode=1|2
12+
; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
13+
; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
14+
; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
2215

23-
; CHECK-SYCL-LargeGRF-PROP: sycl-register-alloc-mode=1|2
16+
; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
17+
; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
2418

25-
; CHECK-SYCL-PROP-NOT: sycl-register-alloc-mode
19+
; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
20+
; CHECK-ESIMD-LargeGRF-PROP: sycl-register-alloc-mode=1|2
2621

27-
; CHECK-SYCL-SYM: __SYCL_kernel
28-
; CHECK-SYCL-SYM-EMPTY:
22+
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-SYCL-LargeGRF-IR
23+
; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
24+
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-LargeGRF-SYM
2925

3026
; CHECK-SYCL-LargeGRF-SYM: __SYCL_kernel_large_grf
3127
; CHECK-SYCL-LargeGRF-SYM-EMPTY:
3228

29+
; CHECK-SYCL-LargeGRF-PROP: sycl-register-alloc-mode=1|2
30+
31+
; RUN: FileCheck %s -input-file=%t_esimd_2.prop --check-prefixes CHECK-ESIMD-PROP
32+
; RUN: FileCheck %s -input-file=%t_esimd_2.sym --check-prefixes CHECK-ESIMD-SYM
33+
3334
; CHECK-ESIMD-SYM: __ESIMD_kernel
3435
; CHECK-ESIMD-SYM-EMPTY:
3536

3637
; CHECK-ESIMD-PROP-NOT: sycl-register-alloc-mode
3738

38-
; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
39-
; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
39+
; RUN: FileCheck %s -input-file=%t_3.prop --check-prefixes CHECK-SYCL-PROP
40+
; RUN: FileCheck %s -input-file=%t_3.sym --check-prefixes CHECK-SYCL-SYM
41+
42+
; CHECK-SYCL-SYM: __SYCL_kernel
43+
; CHECK-SYCL-SYM-EMPTY:
44+
45+
; CHECK-SYCL-PROP-NOT: sycl-register-alloc-mode
46+
47+
4048

4149
; ModuleID = 'large_grf.bc'
4250
source_filename = "llvm-link"

0 commit comments

Comments
 (0)