Skip to content

Commit 8c7bb45

Browse files
kbobrovskbobrovs
authored
[SYCL] sycl-post-link changes to support invoke_simd. (#6160)
* [SYCL] sycl-post-link changes to support invoke_simd. - Split into SYCL and ESIMD callgraph is now obligatory, as these kinds of functions undergo different set of IR t-forms. If no SYCL/ESIMD splitting was requested on the command line, then the tool clones functions shared between the two callgraphs and then links them back together. This is future default for the compiler driver, as for invoke_simd to be correctly handled by the BE, both kinds of code must be in the same module. - Logic of the output file table generation refactored to be more straightforward. - ModuleDesc extended to accomodate properties and support entry point (Function object) replacement with new entry with the same name. - Make --split-esimd a first-class action of sycl-post-link - if passed alone, the tool will no longer complain that no actions are requested. - Fixed bug in ModuleSplitter.cpp/extractSubModule where no actual replacement of entry point Functions with cloned ones happenned. Signed-off-by: Konstantin S Bobrovsky <[email protected]>
1 parent 4c619e9 commit 8c7bb45

18 files changed

+987
-458
lines changed

‎llvm/test/tools/sycl-post-link/auto-module-split-1.ll

Lines changed: 27 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -10,27 +10,27 @@ target triple = "spir64-unknown-linux"
1010

1111
$_Z3barIiET_S0_ = comdat any
1212

13-
; CHECK-TU0-NOT: @{{.*}}GV{{.*}}
14-
; CHECK-TU1: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
13+
; CHECK-TU1-NOT: @{{.*}}GV{{.*}}
14+
; CHECK-TU0: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
1515
@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4
1616

17-
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
18-
; CHECK-TU0-TXT: {{.*}}TU0_kernel0{{.*}}
19-
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
20-
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
17+
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
18+
; CHECK-TU1-TXT: {{.*}}TU0_kernel0{{.*}}
19+
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
20+
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
2121

22-
; CHECK-TU0: call spir_func void @{{.*}}foo{{.*}}()
22+
; CHECK-TU1: call spir_func void @{{.*}}foo{{.*}}()
2323

2424
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 {
2525
entry:
2626
call spir_func void @_Z3foov()
2727
ret void
2828
}
2929

30-
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo{{.*}}()
31-
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo{{.*}}()
30+
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo{{.*}}()
31+
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo{{.*}}()
3232

33-
; CHECK-TU0: call spir_func i32 @{{.*}}bar{{.*}}(i32 1)
33+
; CHECK-TU1: call spir_func i32 @{{.*}}bar{{.*}}(i32 1)
3434

3535
define dso_local spir_func void @_Z3foov() {
3636
entry:
@@ -41,8 +41,8 @@ entry:
4141
ret void
4242
}
4343

44-
; CHECK-TU0: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
45-
; CHECK-TU1-NOT: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
44+
; CHECK-TU1: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
45+
; CHECK-TU0-NOT: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
4646

4747
; Function Attrs: nounwind
4848
define linkonce_odr dso_local spir_func i32 @_Z3barIiET_S0_(i32 %arg) comdat {
@@ -53,21 +53,21 @@ entry:
5353
ret i32 %0
5454
}
5555

56-
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
57-
; CHECK-TU0-TXT: {{.*}}TU0_kernel1{{.*}}
58-
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
59-
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
56+
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
57+
; CHECK-TU1-TXT: {{.*}}TU0_kernel1{{.*}}
58+
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
59+
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
6060

61-
; CHECK-TU0: call spir_func void @{{.*}}foo1{{.*}}()
61+
; CHECK-TU1: call spir_func void @{{.*}}foo1{{.*}}()
6262

6363
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 {
6464
entry:
6565
call spir_func void @_Z4foo1v()
6666
ret void
6767
}
6868

69-
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo1{{.*}}()
70-
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo1{{.*}}()
69+
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo1{{.*}}()
70+
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo1{{.*}}()
7171

7272
; Function Attrs: nounwind
7373
define dso_local spir_func void @_Z4foo1v() {
@@ -77,27 +77,27 @@ entry:
7777
ret void
7878
}
7979

80-
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
81-
; CHECK-TU0-TXT-NOT: {{.*}}TU1_kernel{{.*}}
82-
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
83-
; CHECK-TU1-TXT: {{.*}}TU1_kernel{{.*}}
80+
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
81+
; CHECK-TU1-TXT-NOT: {{.*}}TU1_kernel{{.*}}
82+
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
83+
; CHECK-TU0-TXT: {{.*}}TU1_kernel{{.*}}
8484

85-
; CHECK-TU1: call spir_func void @{{.*}}foo2{{.*}}()
85+
; CHECK-TU0: call spir_func void @{{.*}}foo2{{.*}}()
8686

8787
define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 {
8888
entry:
8989
call spir_func void @_Z4foo2v()
9090
ret void
9191
}
9292

93-
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo2{{.*}}()
94-
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo2{{.*}}()
93+
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo2{{.*}}()
94+
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo2{{.*}}()
9595

9696
; Function Attrs: nounwind
9797
define dso_local spir_func void @_Z4foo2v() {
9898
entry:
9999
%a = alloca i32, align 4
100-
; CHECK-TU1: %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}GV{{.*}} to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
100+
; CHECK-TU0: %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}GV{{.*}} to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
101101
%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
102102
%add = add nsw i32 4, %0
103103
store i32 %add, i32* %a, align 4

‎llvm/test/tools/sycl-post-link/basic-module-split.ll

Lines changed: 27 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -10,27 +10,27 @@ target triple = "spir64-unknown-linux"
1010

1111
$_Z3barIiET_S0_ = comdat any
1212

13-
;CHECK-TU0-NOT: @{{.*}}GV{{.*}}
14-
;CHECK-TU1: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
13+
;CHECK-TU1-NOT: @{{.*}}GV{{.*}}
14+
;CHECK-TU0: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
1515
@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4
1616

17-
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
18-
; CHECK-TU0-TXT: {{.*}}TU0_kernel0{{.*}}
19-
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
20-
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
17+
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
18+
; CHECK-TU1-TXT: {{.*}}TU0_kernel0{{.*}}
19+
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
20+
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
2121

22-
; CHECK-TU0: call spir_func void @{{.*}}foo{{.*}}()
22+
; CHECK-TU1: call spir_func void @{{.*}}foo{{.*}}()
2323

2424
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 {
2525
entry:
2626
call spir_func void @_Z3foov()
2727
ret void
2828
}
2929

30-
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo{{.*}}()
31-
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo{{.*}}()
30+
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo{{.*}}()
31+
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo{{.*}}()
3232

33-
; CHECK-TU0: call spir_func i32 @{{.*}}bar{{.*}}(i32 1)
33+
; CHECK-TU1: call spir_func i32 @{{.*}}bar{{.*}}(i32 1)
3434

3535
define dso_local spir_func void @_Z3foov() {
3636
entry:
@@ -41,8 +41,8 @@ entry:
4141
ret void
4242
}
4343

44-
; CHECK-TU0: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
45-
; CHECK-TU1-NOT: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
44+
; CHECK-TU1: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
45+
; CHECK-TU0-NOT: define {{.*}} spir_func i32 @{{.*}}bar{{.*}}(i32 %arg)
4646

4747
; Function Attrs: nounwind
4848
define linkonce_odr dso_local spir_func i32 @_Z3barIiET_S0_(i32 %arg) comdat {
@@ -53,21 +53,21 @@ entry:
5353
ret i32 %0
5454
}
5555

56-
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
57-
; CHECK-TU0-TXT: {{.*}}TU0_kernel1{{.*}}
58-
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
59-
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
56+
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
57+
; CHECK-TU1-TXT: {{.*}}TU0_kernel1{{.*}}
58+
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
59+
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
6060

61-
; CHECK-TU0: call spir_func void @{{.*}}foo1{{.*}}()
61+
; CHECK-TU1: call spir_func void @{{.*}}foo1{{.*}}()
6262

6363
define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 {
6464
entry:
6565
call spir_func void @_Z4foo1v()
6666
ret void
6767
}
6868

69-
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo1{{.*}}()
70-
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo1{{.*}}()
69+
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo1{{.*}}()
70+
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo1{{.*}}()
7171

7272
; Function Attrs: nounwind
7373
define dso_local spir_func void @_Z4foo1v() {
@@ -77,27 +77,27 @@ entry:
7777
ret void
7878
}
7979

80-
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
81-
; CHECK-TU0-TXT-NOT: {{.*}}TU1_kernel{{.*}}
82-
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
83-
; CHECK-TU1-TXT: {{.*}}TU1_kernel{{.*}}
80+
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
81+
; CHECK-TU1-TXT-NOT: {{.*}}TU1_kernel{{.*}}
82+
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
83+
; CHECK-TU0-TXT: {{.*}}TU1_kernel{{.*}}
8484

85-
; CHECK-TU1: call spir_func void @{{.*}}foo2{{.*}}()
85+
; CHECK-TU0: call spir_func void @{{.*}}foo2{{.*}}()
8686

8787
define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 {
8888
entry:
8989
call spir_func void @_Z4foo2v()
9090
ret void
9191
}
9292

93-
; CHECK-TU0-NOT: define dso_local spir_func void @{{.*}}foo2{{.*}}()
94-
; CHECK-TU1: define dso_local spir_func void @{{.*}}foo2{{.*}}()
93+
; CHECK-TU1-NOT: define dso_local spir_func void @{{.*}}foo2{{.*}}()
94+
; CHECK-TU0: define dso_local spir_func void @{{.*}}foo2{{.*}}()
9595

9696
; Function Attrs: nounwind
9797
define dso_local spir_func void @_Z4foo2v() {
9898
entry:
9999
%a = alloca i32, align 4
100-
; CHECK-TU1: %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}GV{{.*}} to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
100+
; CHECK-TU0: %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @{{.*}}GV{{.*}} to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
101101
%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
102102
%add = add nsw i32 4, %0
103103
store i32 %add, i32* %a, align 4

‎llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -17,11 +17,11 @@ $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any
1717
$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any
1818

1919
@dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 #0
20-
; CHECK-MOD0: @dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4
21-
; CHECK-MOD1-NOT: @dg_int2
20+
; CHECK-MOD1: @dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4
21+
; CHECK-MOD0-NOT: @dg_int2
2222
@dg_int3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4 #1
23-
; CHECK-MOD1: @dg_int3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4
24-
; CHECK-MOD0-NOT: @dg_int3
23+
; CHECK-MOD0: @dg_int3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, align 4
24+
; CHECK-MOD1-NOT: @dg_int3
2525

2626
; Third kernel that uses no device-global variables
2727
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 {

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

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -7,8 +7,8 @@
77
; Per-module split
88
; RUN: sycl-post-link -symbols -split=source -emit-exported-symbols -S %s -o %t.per_module.files.table
99
; RUN: FileCheck %s -input-file=%t.per_module.files_0.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-0-PROP
10-
; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-1-PROP
11-
; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP
10+
; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP
11+
; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-2-PROP
1212
;
1313
; Per-kernel split
1414
; RUN: sycl-post-link -symbols -split=kernel -emit-exported-symbols -S %s -o %t.per_kernel.files.table
@@ -61,31 +61,31 @@ attributes #2 = { "sycl-module-id"="c.cpp" }
6161
; CHECK-GLOBAL-PROP-NEXT: ExportedSpirFunc3
6262

6363
; Per-module split
64-
; CHECK-PERMODULE-0-PROP: [SYCL/exported symbols]
65-
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc1
66-
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc3
67-
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc2
64+
; CHECK-PERMODULE-2-PROP: [SYCL/exported symbols]
65+
; CHECK-PERMODULE-2-PROP-NEXT: ExportedSpirFunc1
66+
; CHECK-PERMODULE-2-PROP-NEXT: ExportedSpirFunc3
67+
; CHECK-PERMODULE-2-PROP-NOT: ExportedSpirFunc2
6868

69-
; CHECK-PERMODULE-1-PROP: [SYCL/exported symbols]
70-
; CHECK-PERMODULE-1-PROP-NEXT: ExportedSpirFunc2
71-
; CHECK-PERMODULE-1-PROP-NOT: ExportedSpirFunc1
72-
; CHECK-PERMODULE-1-PROP-NOT: ExportedSpirFunc3
69+
; CHECK-PERMODULE-0-PROP: [SYCL/exported symbols]
70+
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc2
71+
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc1
72+
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc3
7373

7474
; Per-kernel split
75-
; CHECK-PERKERNEL-0-PROP: [SYCL/exported symbols]
76-
; CHECK-PERKERNEL-0-PROP-NEXT: ExportedSpirFunc1
77-
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc2
78-
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc3
75+
; CHECK-PERKERNEL-2-PROP: [SYCL/exported symbols]
76+
; CHECK-PERKERNEL-2-PROP-NEXT: ExportedSpirFunc1
77+
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc2
78+
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc3
7979

8080
; CHECK-PERKERNEL-1-PROP: [SYCL/exported symbols]
8181
; CHECK-PERKERNEL-1-PROP-NEXT: ExportedSpirFunc2
8282
; CHECK-PERKERNEL-1-PROP-NOT: ExportedSpirFunc1
8383
; CHECK-PERKERNEL-1-PROP-NOT: ExportedSpirFunc3
8484

85-
; CHECK-PERKERNEL-2-PROP: [SYCL/exported symbols]
86-
; CHECK-PERKERNEL-2-PROP-NEXT: ExportedSpirFunc3
87-
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc1
88-
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc2
85+
; CHECK-PERKERNEL-0-PROP: [SYCL/exported symbols]
86+
; CHECK-PERKERNEL-0-PROP-NEXT: ExportedSpirFunc3
87+
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc1
88+
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc2
8989

9090
; Kernel-only generated modules should have no exported Symbols
9191
; CHECK-KERNELONLY-PROP-NOT: [SYCL/exported symbols]

0 commit comments

Comments
 (0)