Skip to content

Commit 26f1c0a

Browse files
vmaksimovmaksimo
authored
[SYCL] Replace spec constants 2020 with RT-provided values in sycl-post-link (#3430)
By doing this change, SpecConstantsPass starts to replace scalar specialization constants with a user-provided value from RT (instead of using C++ default initialization value) for AOT mode. This is done with the new intrinsic __sycl_getScalar2020SpecConstantValue: the 3rd argument RTBuffer is a pointer to a runtime buffer, which holds values of all specialization constants.
1 parent b2e4a83 commit 26f1c0a

File tree

4 files changed

+200
-9
lines changed

4 files changed

+200
-9
lines changed

‎llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll

Lines changed: 41 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \
2-
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
2+
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue --implicit-check-not __sycl_getComposite2020SpecConstantValue
33
;
44
; This test checks that composite specialization constants can be correctly
55
; initialized by sycl-post-link tool for AOT use-case (default initialization
@@ -24,8 +24,10 @@ target triple = "spir64-unknown-unknown-sycldevice"
2424
%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
2525

2626
$_ZTS4Test = comdat any
27+
$_ZTS17SpecializedKernel = comdat any
2728

2829
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1
30+
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1
2931

3032
; Function Attrs: convergent norecurse uwtable
3133
define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
@@ -45,6 +47,38 @@ entry:
4547
ret void
4648
}
4749

50+
; Function Attrs: convergent norecurse
51+
define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
52+
entry:
53+
%c.i = alloca %struct._ZTS1A.A, align 4
54+
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
55+
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
56+
%2 = load i64, i64 addrspace(4)* %1, align 8
57+
%add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %2
58+
; CHECK: %[[CAST1:[0-9a-z.]+]] = addrspacecast %struct._ZTS1A.A* %c.i
59+
%c.ascast.i = addrspacecast %struct._ZTS1A.A* %c.i to %struct._ZTS1A.A addrspace(4)*
60+
%3 = bitcast %struct._ZTS1A.A* %c.i to i8*
61+
call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %3) #3
62+
%a.i.i = alloca i32, align 4
63+
%bc = bitcast i32* %a.i.i to i8*
64+
%tmp = addrspacecast i8* %bc to i8 addrspace(4)*
65+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* %tmp) #4
66+
; CHECK: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8 addrspace(4)* %tmp, i32 0
67+
; CHECK: %[[BITCAST:[0-9a-z]+]] = bitcast i8 addrspace(4)* %[[GEP]] to %struct._ZTS1A.A addrspace(4)*
68+
; CHECK: %[[LOAD:[0-9a-z]+]] = load %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %[[BITCAST]], align 4
69+
; CHECK: store %struct._ZTS1A.A %[[LOAD]], %struct._ZTS1A.A {{.*}}* %[[CAST1]]
70+
%a.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 0
71+
%4 = load i32, i32 addrspace(4)* %a.i, align 4
72+
%conv.i = sitofp i32 %4 to float
73+
%b.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 1
74+
%5 = load float, float addrspace(4)* %b.i, align 4
75+
%add.i = fadd float %5, %conv.i
76+
%ptridx.ascast.i.i = addrspacecast float addrspace(1)* %add.ptr.i to float addrspace(4)*
77+
store float %add.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !11
78+
call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %3) #3
79+
ret void
80+
}
81+
4882
; Function Attrs: argmemonly nounwind willreturn
4983
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
5084

@@ -57,6 +91,9 @@ declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture write
5791
; Function Attrs: convergent
5892
declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 8, i8 addrspace(4)*) local_unnamed_addr #2
5993

94+
; Function Attrs: convergent
95+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4, i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2
96+
6097
attributes #0 = { convergent norecurse uwtable "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="../sycl/test/spec_const/composite.cpp" "tune-cpu"="generic" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
6198
attributes #1 = { argmemonly nounwind willreturn }
6299
attributes #2 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "tune-cpu"="generic" "unsafe-fp-math"="false" "use-soft-float"="false" }
@@ -77,3 +114,6 @@ attributes #4 = { convergent }
77114
!6 = !{!7, !7, i64 0}
78115
!7 = !{!"omnipotent char", !8, i64 0}
79116
!8 = !{!"Simple C++ TBAA"}
117+
!9 = !{!"int", !7, i64 0}
118+
!10 = !{!"float", !7, i64 0}
119+
!11 = !{!10, !10, i64 0}

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

Lines changed: 93 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,93 @@
1+
; RUN: sycl-post-link --ir-output-only -spec-const=default %s -S -o - | \
2+
; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-DEF
3+
4+
; This test checks that the post link tool is able to correctly transform
5+
; specialization constant intrinsics for different types in a device code
6+
; compiled for AOT.
7+
8+
%class.specialization_id = type { double }
9+
%class.specialization_id.0 = type { i32 }
10+
%class.specialization_id.1 = type { %struct.ComposConst }
11+
%struct.ComposConst = type { i32, double, %struct.myConst }
12+
%struct.myConst = type { i32, float }
13+
%class.specialization_id.2 = type { %struct.ComposConst2 }
14+
%struct.ComposConst2 = type { i8, %struct.myConst, double }
15+
16+
@id_double = dso_local global %class.specialization_id { double 3.140000e+00 }, align 8
17+
@id_int = dso_local global %class.specialization_id.0 { i32 42 }, align 4
18+
@id_compos = dso_local global %class.specialization_id.1 { %struct.ComposConst { i32 1, double 2.000000e+00, %struct.myConst { i32 13, float 0x4020666660000000 } } }, align 8
19+
@id_compos2 = dso_local global %class.specialization_id.2 { %struct.ComposConst2 { i8 1, %struct.myConst { i32 52, float 0x40479999A0000000 }, double 2.000000e+00 } }, align 8
20+
21+
; check that no longer used globals are removed:
22+
@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v = private unnamed_addr constant [37 x i8] c"_ZTS14name_generatorIL_Z9id_doubleEE\00", align 1
23+
; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v
24+
@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v = private unnamed_addr constant [34 x i8] c"_ZTS14name_generatorIL_Z6id_intEE\00", align 1
25+
; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v
26+
@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v = private unnamed_addr constant [37 x i8] c"_ZTS14name_generatorIL_Z9id_composEE\00", align 1
27+
; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v
28+
@__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_compos2E17specialization_idI12ComposConst2ES1_ET1_v = private unnamed_addr constant [39 x i8] c"_ZTS14name_generatorIL_Z10id_compos2EE\00", align 1
29+
; CHECK-NOT: @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_compos2E17specialization_idI12ComposConst2ES1_ET1_v
30+
31+
define dso_local void @_Z4testv() local_unnamed_addr #0 {
32+
entry:
33+
%call.i = tail call fast double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8* getelementptr inbounds ([37 x i8], [37 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id* @id_double to i8*), i8* null)
34+
; CHECK-DEF: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8* null, i32 0
35+
; CHECK-DEF: %[[BITCAST:[0-9a-z]+]] = bitcast i8* %[[GEP]] to double*
36+
; CHECK-DEF: %[[LOAD:[0-9a-z]+]] = load double, double* %[[BITCAST]], align 8
37+
%call.i3 = tail call i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPvS3_(i8* getelementptr inbounds ([34 x i8], [34 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z6id_intE17specialization_idIiEiET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.0* @id_int to i8*), i8* null)
38+
; CHECK-DEF: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 8
39+
; CHECK-DEF: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to i32*
40+
; CHECK-DEF: %[[LOAD1:[0-9a-z]+]] = load i32, i32* %[[BITCAST1]], align 4
41+
%call.i4 = tail call fast double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8* getelementptr inbounds ([37 x i8], [37 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_doubleE17specialization_idIdEdET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id* @id_double to i8*), i8* null)
42+
; CHECK-DEF: %[[GEP2:[0-9a-z]+]] = getelementptr i8, i8* null, i32 0
43+
; CHECK-DEF: %[[BITCAST2:[0-9a-z]+]] = bitcast i8* %[[GEP2]] to double*
44+
; CHECK-DEF: %[[LOAD2:[0-9a-z]+]] = load double, double* %[[BITCAST2]], align 8
45+
ret void
46+
}
47+
48+
define dso_local void @_Z5test2v() local_unnamed_addr #0 {
49+
entry:
50+
%tmp = alloca %struct.ComposConst, align 8
51+
%tmp1 = alloca %struct.ComposConst2, align 8
52+
%tmp2 = alloca %struct.ComposConst, align 8
53+
%0 = bitcast %struct.ComposConst* %tmp to i8*
54+
call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %0) #3
55+
call void @_Z40__sycl_getComposite2020SpecConstantValueI11ComposConstET_PKcPvS4_(%struct.ComposConst* nonnull sret(%struct.ComposConst) align 8 %tmp, i8* getelementptr inbounds ([37 x i8], [37 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.1* @id_compos to i8*), i8* null)
56+
; CHECK: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8* null, i32 12
57+
; CHECK: %[[BITCAST:[0-9a-z]+]] = bitcast i8* %[[GEP]] to %struct.ComposConst*
58+
; CHECK: %[[LOAD:[0-9a-z]+]] = load %struct.ComposConst, %struct.ComposConst* %[[BITCAST]], align 8
59+
; CHECK: store %struct.ComposConst %[[LOAD]], %struct.ComposConst*
60+
call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %0) #3
61+
%1 = getelementptr inbounds %struct.ComposConst2, %struct.ComposConst2* %tmp1, i64 0, i32 0
62+
call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %1) #3
63+
call void @_Z40__sycl_getComposite2020SpecConstantValueI12ComposConst2ET_PKcPvS4_(%struct.ComposConst2* nonnull sret(%struct.ComposConst2) align 8 %tmp1, i8* getelementptr inbounds ([39 x i8], [39 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z10id_compos2E17specialization_idI12ComposConst2ES1_ET1_v, i64 0, i64 0), i8* getelementptr inbounds (%class.specialization_id.2, %class.specialization_id.2* @id_compos2, i64 0, i32 0, i32 0), i8* null) call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %1) #3
64+
; CHECK: %[[GEP1:[0-9a-z]+]] = getelementptr i8, i8* null, i32 36
65+
; CHECK: %[[BITCAST1:[0-9a-z]+]] = bitcast i8* %[[GEP1]] to %struct.ComposConst2*
66+
; CHECK: %[[LOAD1:[0-9a-z]+]] = load %struct.ComposConst2, %struct.ComposConst2* %[[BITCAST1]], align 8
67+
; CHECK: store %struct.ComposConst2 %[[LOAD1]], %struct.ComposConst2*
68+
%2 = bitcast %struct.ComposConst* %tmp2 to i8*
69+
call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3
70+
call void @_Z40__sycl_getComposite2020SpecConstantValueI11ComposConstET_PKcPvS4_(%struct.ComposConst* nonnull sret(%struct.ComposConst) align 8 %tmp2, i8* getelementptr inbounds ([37 x i8], [37 x i8]* @__builtin_unique_stable_name._Z27get_specialization_constantIL_Z9id_composE17specialization_idI11ComposConstES1_ET1_v, i64 0, i64 0), i8* bitcast (%class.specialization_id.1* @id_compos to i8*), i8* null)
71+
; CHECK: %[[GEP2:[0-9a-z]+]] = getelementptr i8, i8* null, i32 12
72+
; CHECK: %[[BITCAST2:[0-9a-z]+]] = bitcast i8* %[[GEP2]] to %struct.ComposConst*
73+
; CHECK: %[[LOAD2:[0-9a-z]+]] = load %struct.ComposConst, %struct.ComposConst* %[[BITCAST2]], align 8
74+
; CHECK: store %struct.ComposConst %[[LOAD2]], %struct.ComposConst*
75+
call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3
76+
ret void
77+
}
78+
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
79+
80+
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1
81+
82+
declare dso_local double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8*, i8*, i8*) local_unnamed_addr #1
83+
84+
declare dso_local i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPvS3_(i8*, i8*, i8*) local_unnamed_addr #1
85+
86+
declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI11ComposConstET_PKcPvS4_(%struct.ComposConst* sret(%struct.ComposConst) align 8, i8*, i8*, i8*) local_unnamed_addr #2
87+
88+
declare dso_local void @_Z40__sycl_getComposite2020SpecConstantValueI12ComposConst2ET_PKcPvS4_(%struct.ComposConst2* sret(%struct.ComposConst2) align 8, i8*, i8*, i8*) local_unnamed_addr #2
89+
90+
attributes #0 = { uwtable mustprogress "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-jump-tables"="false" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" "unsafe-fp-math"="true" "use-soft-float"="false" }
91+
attributes #1 = { argmemonly nofree nosync nounwind willreturn }
92+
attributes #2 = { "denormal-fp-math"="preserve-sign,preserve-sign" "denormal-fp-math-f32"="ieee,ieee" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" "unsafe-fp-math"="true" "use-soft-float"="false" }
93+
attributes #3 = { nounwind }

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

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,10 @@ define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrsp
116116
%37 = tail call spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null, i8 addrspace(4)* null)
117117
; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 11, double 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID11:[0-9]+]]
118118
%38 = fadd double %37, %36
119-
; CHECK-DEF: %[[SUM10:[0-9]+]] = fadd double 0.000000e+00, %[[SUM9]]
119+
; CHECK-DEF: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8 addrspace(4)* null, i32 0
120+
; CHECK-DEF: %[[BITCAST:[0-9a-z]+]] = bitcast i8 addrspace(4)* %[[GEP]] to double addrspace(4)*
121+
; CHECK-DEF: %[[LOAD:[0-9a-z]+]] = load double, double addrspace(4)* %[[BITCAST]], align 8
122+
; CHECK-DEF: %[[SUM10:[0-9]+]] = fadd double %[[LOAD]], %[[SUM9]]
120123
%39 = fptrunc double %38 to float
121124
; CHECK-DEF: %[[VAL8:[0-9]+]] = fptrunc double %[[SUM10]] to float
122125
%40 = addrspacecast float addrspace(1)* %7 to float addrspace(4)*

0 commit comments

Comments
 (0)