Skip to content

Commit f99335f

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Clean up device_globals from llvm.compiler.used (#9907)
The implementation of device_global uses llvm.compiler.used to avoid the compiler wrongfully optimizing out internally linked device_global variables. However, since this variable has appending linkage the variables may carry over to other binaries during linking. To avoid bloating other device binaries with device_global variables they have no use for, this commit makes sycl-post-link remove all occurences of device_global variables in llvm.compiler.used, similar to how it removes llvm.used. The distinction between the logic for llvm.compiler.used and llvm.used is that sycl-post-link will preserve any value in llvm.compiler.used that is not a device_global. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 008e374 commit f99335f

File tree

5 files changed

+343
-0
lines changed

5 files changed

+343
-0
lines changed

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

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
3+
;
4+
; Test checks that llvm.compiler.used is removed when all values in it are
5+
; device_global. Likewise it checks that device_global variables that have no
6+
; uses after it are dropped too.
7+
8+
source_filename = "test_global_variable.cpp"
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* }
13+
%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 }
14+
%class.anon.0 = type { i8 }
15+
16+
; CHECK-IR-NOT: @llvm.compiler.used =
17+
@llvm.compiler.used = appending global [4 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7no_dg_int1 to i8 addrspace(4)*)]
18+
19+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0
20+
@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !4 #1
21+
@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !8 #2
22+
@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !10 #3
23+
@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !19 #4
24+
25+
; CHECK-IR: @_ZL7dg_int1 =
26+
; CHECK-IR: @_ZL7dg_int2 =
27+
; CHECK-IR: @_ZL8dg_bool3 =
28+
; CHECK-IR: @_ZL8dg_bool4 =
29+
; CHECK-IR-NOT: @_ZL7no_dg_int1 =
30+
31+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
32+
define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 {
33+
entry:
34+
%this.addr = alloca %class.anon.0 addrspace(4)*, align 8
35+
%this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)*
36+
store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
37+
%this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
38+
%call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #6
39+
%call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #6
40+
%call3 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #6
41+
%call4 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #6
42+
ret void
43+
}
44+
45+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
46+
declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8)) #5 align 2
47+
48+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
49+
declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1)) #5 align 2
50+
51+
attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
52+
attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
53+
attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
54+
attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
55+
attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
56+
attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
57+
attributes #6 = { convergent nounwind }
58+
59+
!llvm.dependent-libraries = !{!13}
60+
!llvm.module.flags = !{!14, !15}
61+
!opencl.spir.version = !{!16}
62+
!spirv.Source = !{!17}
63+
!llvm.ident = !{!18}
64+
65+
!0 = !{!1, !2, !3}
66+
!1 = !{i32 6149, i32 1}
67+
!2 = !{i32 6148, i32 0}
68+
!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
69+
!4 = !{!5, !6, !7}
70+
!5 = !{i32 6149, i32 0}
71+
!6 = !{i32 6148, i32 1}
72+
!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
73+
!8 = !{!1, !2, !9}
74+
!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
75+
!10 = !{!11}
76+
!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
77+
!12 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7no_dg_int1"}
78+
!13 = !{!"libcpmt"}
79+
!14 = !{i32 1, !"wchar_size", i32 2}
80+
!15 = !{i32 7, !"frame-pointer", i32 2}
81+
!16 = !{i32 1, i32 2}
82+
!17 = !{i32 4, i32 100000}
83+
!18 = !{!"clang version 14.0.0"}
84+
!19 = !{!1, !2, !12}
85+
86+
; CHECK-IR: !"6da74a122db9f35d____ZL7dg_int1"
87+
; CHECK-IR: !"7da74a1187b9f35d____ZL7dg_int2"
88+
; CHECK-IR: !"9d329ad59055e972____ZL8dg_bool3"
89+
; CHECK-IR: !"dda2bad52c45c432____ZL8dg_bool4"
90+
; CHECK-IR-NOT: !"6da74a122db9f35d____ZL7no_dg_int1"

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

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
3+
;
4+
; Test checks that llvm.compiler.used is removed when all values in it are
5+
; device_global. Likewise it checks that device_global variables that have no
6+
; uses after it are dropped too. This case is using opaque pointers.
7+
8+
source_filename = "test_global_variable.cpp"
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.cl::sycl::ext::oneapi::device_global.0" = type { ptr addrspace(4) }
13+
%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 }
14+
%class.anon.0 = type { i8 }
15+
16+
; CHECK-IR-NOT: @llvm.compiler.used =
17+
@llvm.compiler.used = appending global [4 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL8dg_bool4 to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @_ZL7no_dg_int1 to ptr addrspace(4))]
18+
19+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0
20+
@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !4 #1
21+
@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !8 #2
22+
@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !10 #3
23+
@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !19 #4
24+
25+
; CHECK-IR: @_ZL7dg_int1 =
26+
; CHECK-IR: @_ZL7dg_int2 =
27+
; CHECK-IR: @_ZL8dg_bool3 =
28+
; CHECK-IR: @_ZL8dg_bool4 =
29+
; CHECK-IR-NOT: @_ZL7no_dg_int1 =
30+
31+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
32+
define internal spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) %this) #5 align 2 {
33+
entry:
34+
%this.addr = alloca ptr addrspace(4), align 8
35+
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
36+
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
37+
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
38+
%call1 = call spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8) addrspacecast (ptr addrspace(1) @_ZL7dg_int1 to ptr addrspace(4))) #6
39+
%call2 = call spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8) addrspacecast (ptr addrspace(1) @_ZL7dg_int2 to ptr addrspace(4))) #6
40+
%call3 = call spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) addrspacecast (ptr addrspace(1) @_ZL8dg_bool3 to ptr addrspace(4))) #6
41+
%call4 = call spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) addrspacecast (ptr addrspace(1) @_ZL8dg_bool4 to ptr addrspace(4))) #6
42+
ret void
43+
}
44+
45+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
46+
declare spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 8 dereferenceable_or_null(8)) #5 align 2
47+
48+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
49+
declare spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1)) #5 align 2
50+
51+
attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
52+
attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
53+
attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
54+
attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
55+
attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
56+
attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
57+
attributes #6 = { convergent nounwind }
58+
59+
!llvm.dependent-libraries = !{!13}
60+
!llvm.module.flags = !{!14, !15}
61+
!opencl.spir.version = !{!16}
62+
!spirv.Source = !{!17}
63+
!llvm.ident = !{!18}
64+
65+
!0 = !{!1, !2, !3}
66+
!1 = !{i32 6149, i32 1}
67+
!2 = !{i32 6148, i32 0}
68+
!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
69+
!4 = !{!5, !6, !7}
70+
!5 = !{i32 6149, i32 0}
71+
!6 = !{i32 6148, i32 1}
72+
!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
73+
!8 = !{!1, !2, !9}
74+
!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
75+
!10 = !{!11}
76+
!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
77+
!12 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7no_dg_int1"}
78+
!13 = !{!"libcpmt"}
79+
!14 = !{i32 1, !"wchar_size", i32 2}
80+
!15 = !{i32 7, !"frame-pointer", i32 2}
81+
!16 = !{i32 1, i32 2}
82+
!17 = !{i32 4, i32 100000}
83+
!18 = !{!"clang version 14.0.0"}
84+
!19 = !{!1, !2, !12}
85+
86+
; CHECK-IR: !"6da74a122db9f35d____ZL7dg_int1"
87+
; CHECK-IR: !"7da74a1187b9f35d____ZL7dg_int2"
88+
; CHECK-IR: !"9d329ad59055e972____ZL8dg_bool3"
89+
; CHECK-IR: !"dda2bad52c45c432____ZL8dg_bool4"
90+
; CHECK-IR-NOT: !"6da74a122db9f35d____ZL7no_dg_int1"

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

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
; RUN: sycl-post-link --device-globals -S < %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
3+
;
4+
; Test checks that all device_global variables in llvm.compiler.used are removed
5+
; but any other values stay in.
6+
7+
source_filename = "test_global_variable.cpp"
8+
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"
9+
target triple = "spir64-unknown-unknown"
10+
11+
%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* }
12+
%class.anon.0 = type { i8 }
13+
14+
; CHECK-IR: @llvm.compiler.used = appending global [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @_ZL16NotADeviceGlobal to i8 addrspace(4)*)]
15+
@llvm.compiler.used = appending global [3 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @_ZL16NotADeviceGlobal to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to i8 addrspace(4)*)]
16+
17+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0
18+
@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !4 #1
19+
20+
@_ZL16NotADeviceGlobal = internal addrspace(1) constant i8 zeroinitializer
21+
22+
attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
23+
attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
24+
attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
25+
attributes #6 = { convergent nounwind }
26+
27+
!llvm.dependent-libraries = !{!8}
28+
!llvm.module.flags = !{!9, !10}
29+
!opencl.spir.version = !{!11}
30+
!spirv.Source = !{!12}
31+
!llvm.ident = !{!13}
32+
33+
!0 = !{!1, !2, !3}
34+
!1 = !{i32 6149, i32 1}
35+
!2 = !{i32 6148, i32 0}
36+
!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
37+
!4 = !{!5, !6, !7}
38+
!5 = !{i32 6149, i32 0}
39+
!6 = !{i32 6148, i32 1}
40+
!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
41+
!8 = !{!"libcpmt"}
42+
!9 = !{i32 1, !"wchar_size", i32 2}
43+
!10 = !{i32 7, !"frame-pointer", i32 2}
44+
!11 = !{i32 1, i32 2}
45+
!12 = !{i32 4, i32 100000}
46+
!13 = !{!"clang version 14.0.0"}

0 commit comments

Comments
 (0)