Skip to content

[sycl-post-link] Add a check for device globals with device_image_scope #5517

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
bader merged 17 commits into from Feb 14, 2022
Merged

[sycl-post-link] Add a check for device globals with device_image_scope #5517

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
a2f51f5
[sycl-post-link] Add the DeviceGlobalsMet property to GlobalBinImageP…
Jan 20, 2022
fbeb394
Revert "[sycl-post-link] Add the DeviceGlobalsMet property to GlobalB…
Feb 7, 2022
bcf8950
[sycl-post-link] Extract the hasDeviceImageScopeProperty function
Feb 8, 2022
66ab81c
[sycl-post-link] Add a check for device globals with device_image_scope
Feb 8, 2022
19fbd7d
[sycl-post-link] Add lit tests for a multi-module application
Feb 9, 2022
2cf8211
[sycl-post-link][NFC] Fix formatting issues
Feb 9, 2022
2fb04a6
Use a clearer wording in a comment
Feb 11, 2022
625166c
[sycl-post-link] Use Optional<StringRef> instead of StringRef*
Feb 11, 2022
d8e9ea0
Revert "Use a clearer wording in a comment"
Feb 11, 2022
98bbe22
[sycl-post-link] Use SmallSetVector for graph traversing
Feb 11, 2022
f278edd
[sycl-post-link] Use continue for early exit from graph traversing
Feb 11, 2022
5d09f19
[sycl-post-link] Remove unused functions from the lit tests
Feb 11, 2022
29a842d
[sycl-post-link] Fix Optional<StringRef> comparison
Feb 11, 2022
7d40f39
[NFC] Change error message according to the design document
Feb 11, 2022
1b508b3
Remove redundant braces
Feb 11, 2022
74caa29
Remove additional responsibilities from the ModuleSplitter class
Feb 14, 2022
6b9655e
[NFC] Fix formatting inssues
Feb 14, 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
189 changes: 189 additions & 0 deletions ...st/tools/sycl-post-link/device-globals/test_global_variable_many_kernels_in_one_module.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,189 @@
; RUN: sycl-post-link --device-globals --split=source -S %s -o %t.files.table
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-MOD0
; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1

; This test is intended to check that sycl-post-link generates no errors
; when a device global variable with the 'device_image_scope' property
; is used in many kernels grouped to one module.

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"
target triple = "spir64-unknown-unknown"

%"class.cl::sycl::ext::oneapi::device_global" = type { i32 }
%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" }
%"class.cl::sycl::detail::array" = type { [1 x i64] }
%"class.cl::sycl::detail::accessor_common" = type { i8 }

$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any
$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any
$_ZN2cl4sycl2idILi1EEC2Ev = comdat any
$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any
$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE = comdat any
$_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any
$_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any
$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any
$_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_ = comdat any

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

; Third kernel that uses no device-global variables
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #6 comdat !kernel_arg_buffer_location !6 {
entry:
ret void
}

; Function is required to increase the function list only
Comment thread
This conversation was marked as resolved.
Outdated
; Function Attrs: convergent noinline norecurse nounwind optnone
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 {
entry:
ret void
}

; Function is required to increase the function list only
; Function Attrs: convergent noinline norecurse nounwind optnone
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 {
entry:
ret void
}

; Function is required to increase the function list only
; Function Attrs: convergent noinline norecurse optnone
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this) unnamed_addr #1 comdat align 2 {
entry:
%this.addr = alloca %"class.cl::sycl::range" addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %"class.cl::sycl::range" addrspace(4)** %this.addr to %"class.cl::sycl::range" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::range" addrspace(4)* %this, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %"class.cl::sycl::range" addrspace(4)*, %"class.cl::sycl::range" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%0 = bitcast %"class.cl::sycl::range" addrspace(4)* %this1 to %"class.cl::sycl::detail::array" addrspace(4)*
call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::detail::array" addrspace(4)* align 8 dereferenceable_or_null(8) %0, i64 0) #7
ret void
}

; Function is required to increase the function list only
; Function Attrs: convergent mustprogress noinline norecurse optnone
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class.cl::sycl::range" addrspace(4)* noalias sret(%"class.cl::sycl::range") align 8 %agg.result) #2 comdat align 2 {
entry:
call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %agg.result, i64 0) #7
ret void
}

; Function is required to increase the function list only
; Function Attrs: convergent noinline norecurse nounwind optnone
define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEyE4typeE(%"class.cl::sycl::range" addrspace(4)* align 8 dereferenceable_or_null(8) %this, i64 %dim0) unnamed_addr #3 comdat align 2 {
entry:
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse optnone
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel4RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 {
entry:
%0 = alloca %"class.cl::sycl::detail::accessor_common", align 1
%1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)*
call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse optnone
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 {
entry:
%0 = alloca %"class.cl::sycl::detail::accessor_common", align 1
%1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)*
call spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 {
entry:
%this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
call spir_func void @_Z14kernel1_level1v() #7
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define dso_local spir_func void @_Z14kernel1_level1v() #4 {
entry:
%dg_int_ptr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8
%dg_int_ptr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %dg_int_ptr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*), %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8
%0 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %dg_int_ptr.ascast, align 8
%call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %0) #9
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 {
entry:
%retval = alloca i32 addrspace(4)*, align 8
%this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8
%retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)*
%this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0
ret i32 addrspace(4)* %val
}

; Function Attrs: convergent mustprogress noinline norecurse optnone
define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #5 comdat !kernel_arg_buffer_location !6 {
entry:
%0 = alloca %"class.cl::sycl::detail::accessor_common", align 1
%1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)*
call spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %1) #7
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 {
entry:
%this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8
%this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::detail::accessor_common" addrspace(4)* %this, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %"class.cl::sycl::detail::accessor_common" addrspace(4)*, %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #8
ret void
}

; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #4 align 2 {
entry:
%retval = alloca i32 addrspace(4)*, align 8
%this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8
%retval.ascast = addrspacecast i32 addrspace(4)** %retval to i32 addrspace(4)* addrspace(4)*
%this.addr.ascast = addrspacecast %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)** %this.addr to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)*
store %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%this1 = load %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
%val = getelementptr inbounds %"class.cl::sycl::ext::oneapi::device_global", %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* %this1, i32 0, i32 0
ret i32 addrspace(4)* %val
}

attributes #0 = { "sycl-unique-id"="dg_int2" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" }
attributes #1 = { convergent noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #3 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #5 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" }
attributes #6 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" }
attributes #7 = { convergent }
attributes #8 = { convergent nounwind }
attributes #9 = { nobuiltin allocsize(0) "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }

!llvm.dependent-libraries = !{!0}
!opencl.spir.version = !{!1}
!spirv.Source = !{!2}
!llvm.ident = !{!3}
!llvm.module.flags = !{!4, !5}

!0 = !{!"libcpmt"}
!1 = !{i32 1, i32 2}
!2 = !{i32 4, i32 100000}
!3 = !{!"clang version 14.0.0"}
!4 = !{i32 1, !"wchar_size", i32 2}
!5 = !{i32 7, !"frame-pointer", i32 2}
!6 = !{}
Loading