Skip to content

[SYCL] sycl-post-link changes to support invoke_simd. #6160

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
Jun 1, 2022
Merged

[SYCL] sycl-post-link changes to support invoke_simd. #6160

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
a8a294b
[SYCL] sycl-post-link changes to support invoke_simd.
kbobrovs May 12, 2022
5759c00
Test failure fixes.
kbobrovs May 18, 2022
198c42a
More test fixes
kbobrovs May 18, 2022
a8d7b9a
Fixed help message, fix more tests
kbobrovs May 18, 2022
44cbb8e
Modify names of SYCL/ESIMD shared functions before linkage, add test.
kbobrovs May 19, 2022
d067250
added test comment
kbobrovs May 19, 2022
e207813
Merge remote-tracking branch 'intel/sycl' into invoke_simd_post_link
kbobrovs May 31, 2022
fdec8e0
address review comments
kbobrovs Jun 1, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
54 changes: 27 additions & 27 deletions llvm/test/tools/sycl-post-link/auto-module-split-1.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,27 +10,27 @@ target triple = "spir64-unknown-linux"

$_Z3barIiET_S0_ = comdat any

; CHECK-TU0-NOT: @{{.*}}GV{{.*}}
; CHECK-TU1: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
; CHECK-TU1-NOT: @{{.*}}GV{{.*}}
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

all test changes are related to re-numeration of output files - E.g. ESIMD kernels appeared in files generated earlier (with smaller numeric index), now they will appear in files generated later, etc.

; CHECK-TU0: @{{.*}}GV{{.*}} = internal addrspace(1) constant [1 x i32] [i32 42], align 4
@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4

; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-TXT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-TXT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel0{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 {
entry:
call spir_func void @_Z3foov()
ret void
}

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

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

define dso_local spir_func void @_Z3foov() {
entry:
Expand All @@ -41,8 +41,8 @@ entry:
ret void
}

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

; Function Attrs: nounwind
define linkonce_odr dso_local spir_func i32 @_Z3barIiET_S0_(i32 %arg) comdat {
Expand All @@ -53,21 +53,21 @@ entry:
ret i32 %0
}

; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU0-TXT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU1-TXT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel1{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 {
entry:
call spir_func void @_Z4foo1v()
ret void
}

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

; Function Attrs: nounwind
define dso_local spir_func void @_Z4foo1v() {
Expand All @@ -77,27 +77,27 @@ entry:
ret void
}

; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU0-TXT-NOT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU1-TXT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU1-TXT-NOT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU0-TXT: {{.*}}TU1_kernel{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 {
entry:
call spir_func void @_Z4foo2v()
ret void
}

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

; Function Attrs: nounwind
define dso_local spir_func void @_Z4foo2v() {
entry:
%a = alloca i32, align 4
; 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
; 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
%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
%add = add nsw i32 4, %0
store i32 %add, i32* %a, align 4
Expand Down
54 changes: 27 additions & 27 deletions llvm/test/tools/sycl-post-link/basic-module-split.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,27 +10,27 @@ target triple = "spir64-unknown-linux"

$_Z3barIiET_S0_ = comdat any

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

; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-TXT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU1-TXT: {{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel0{{.*}}
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel0{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel0() #0 {
entry:
call spir_func void @_Z3foov()
ret void
}

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

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

define dso_local spir_func void @_Z3foov() {
entry:
Expand All @@ -41,8 +41,8 @@ entry:
ret void
}

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

; Function Attrs: nounwind
define linkonce_odr dso_local spir_func i32 @_Z3barIiET_S0_(i32 %arg) comdat {
Expand All @@ -53,21 +53,21 @@ entry:
ret i32 %0
}

; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU0-TXT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU1-TXT-NOT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU1-TXT: {{.*}}TU0_kernel1{{.*}}
; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU0_kernel1{{.*}}()
; CHECK-TU0-TXT-NOT: {{.*}}TU0_kernel1{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE11TU0_kernel1() #0 {
entry:
call spir_func void @_Z4foo1v()
ret void
}

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

; Function Attrs: nounwind
define dso_local spir_func void @_Z4foo1v() {
Expand All @@ -77,27 +77,27 @@ entry:
ret void
}

; CHECK-TU0-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU0-TXT-NOT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU1: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU1-TXT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU1-NOT: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU1-TXT-NOT: {{.*}}TU1_kernel{{.*}}
; CHECK-TU0: define dso_local spir_kernel void @{{.*}}TU1_kernel{{.*}}()
; CHECK-TU0-TXT: {{.*}}TU1_kernel{{.*}}

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

define dso_local spir_kernel void @_ZTSZ4mainE10TU1_kernel() #1 {
entry:
call spir_func void @_Z4foo2v()
ret void
}

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

; Function Attrs: nounwind
define dso_local spir_func void @_Z4foo2v() {
entry:
%a = alloca i32, align 4
; 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
; 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
%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
%add = add nsw i32 4, %0
store i32 %add, i32* %a, align 4
Expand Down
8 changes: 4 additions & 4 deletions ...test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_two_vars_ok.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,11 +17,11 @@ $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any
$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any

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

; Third kernel that uses no device-global variables
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 {
Expand Down
36 changes: 18 additions & 18 deletions llvm/test/tools/sycl-post-link/emit_exported_symbols.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,8 +7,8 @@
; Per-module split
; RUN: sycl-post-link -symbols -split=source -emit-exported-symbols -S %s -o %t.per_module.files.table
; RUN: FileCheck %s -input-file=%t.per_module.files_0.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-0-PROP
; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-1-PROP
; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP
; RUN: FileCheck %s -input-file=%t.per_module.files_1.prop -implicit-check-not="NotExported" --check-prefix=CHECK-KERNELONLY-PROP
; RUN: FileCheck %s -input-file=%t.per_module.files_2.prop -implicit-check-not="NotExported" --check-prefix=CHECK-PERMODULE-2-PROP
;
; Per-kernel split
; RUN: sycl-post-link -symbols -split=kernel -emit-exported-symbols -S %s -o %t.per_kernel.files.table
Expand Down Expand Up @@ -61,31 +61,31 @@ attributes #2 = { "sycl-module-id"="c.cpp" }
; CHECK-GLOBAL-PROP-NEXT: ExportedSpirFunc3

; Per-module split
; CHECK-PERMODULE-0-PROP: [SYCL/exported symbols]
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc1
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc3
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc2
; CHECK-PERMODULE-2-PROP: [SYCL/exported symbols]
; CHECK-PERMODULE-2-PROP-NEXT: ExportedSpirFunc1
; CHECK-PERMODULE-2-PROP-NEXT: ExportedSpirFunc3
; CHECK-PERMODULE-2-PROP-NOT: ExportedSpirFunc2

; CHECK-PERMODULE-1-PROP: [SYCL/exported symbols]
; CHECK-PERMODULE-1-PROP-NEXT: ExportedSpirFunc2
; CHECK-PERMODULE-1-PROP-NOT: ExportedSpirFunc1
; CHECK-PERMODULE-1-PROP-NOT: ExportedSpirFunc3
; CHECK-PERMODULE-0-PROP: [SYCL/exported symbols]
; CHECK-PERMODULE-0-PROP-NEXT: ExportedSpirFunc2
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc1
; CHECK-PERMODULE-0-PROP-NOT: ExportedSpirFunc3

; Per-kernel split
; CHECK-PERKERNEL-0-PROP: [SYCL/exported symbols]
; CHECK-PERKERNEL-0-PROP-NEXT: ExportedSpirFunc1
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc2
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc3
; CHECK-PERKERNEL-2-PROP: [SYCL/exported symbols]
; CHECK-PERKERNEL-2-PROP-NEXT: ExportedSpirFunc1
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc2
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc3

; CHECK-PERKERNEL-1-PROP: [SYCL/exported symbols]
; CHECK-PERKERNEL-1-PROP-NEXT: ExportedSpirFunc2
; CHECK-PERKERNEL-1-PROP-NOT: ExportedSpirFunc1
; CHECK-PERKERNEL-1-PROP-NOT: ExportedSpirFunc3

; CHECK-PERKERNEL-2-PROP: [SYCL/exported symbols]
; CHECK-PERKERNEL-2-PROP-NEXT: ExportedSpirFunc3
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc1
; CHECK-PERKERNEL-2-PROP-NOT: ExportedSpirFunc2
; CHECK-PERKERNEL-0-PROP: [SYCL/exported symbols]
; CHECK-PERKERNEL-0-PROP-NEXT: ExportedSpirFunc3
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc1
; CHECK-PERKERNEL-0-PROP-NOT: ExportedSpirFunc2

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