Skip to content

Commit 1ef407e

Browse files
AlexeySachkovAlexeySachkov
authored
[NFC][sycl-post-link] Add test with zeroinitialied spec constants (#3591)
Adjusted assertions to accept zeroinitializer for specialization_id variables, as it is a valid situation. The rest of the code wasn't changed, because it works fine with such initializers already.
1 parent c7a2004 commit 1ef407e

File tree

2 files changed

+84
-1
lines changed

2 files changed

+84
-1
lines changed

‎llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,81 @@
1+
; RUN: sycl-post-link --ir-output-only --spec-const=rt %s -S -o - | FileCheck %s
2+
;
3+
; This test is intended to check that SpecConstantsPass is able to handle the
4+
; situation where specialization constants have zeroinitializer in LLVM IR
5+
6+
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"
7+
target triple = "spir64-unknown-unknown-sycldevice"
8+
9+
%"class._ZTSN2cl4sycl17specialization_idIiEE.cl::sycl::specialization_id" = type { i32 }
10+
%"class._ZTSN2cl4sycl17specialization_idIdEE.cl::sycl::specialization_id" = type { double }
11+
%"class._ZTSN2cl4sycl17specialization_idI9compositeEE.cl::sycl::specialization_id" = type { %struct._ZTS9composite.composite }
12+
%struct._ZTS9composite.composite = type { %struct._ZTS6nested.nested, i8, i8, i64 }
13+
%struct._ZTS6nested.nested = type { float }
14+
%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
15+
%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] }
16+
%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }
17+
18+
$"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1" = comdat any
19+
20+
@__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL6int_idEiLPv0EEET0_v = private unnamed_addr addrspace(1) constant [70 x i8] c"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL6int_idEEE\00", align 1
21+
@_ZL6int_id = internal addrspace(1) constant %"class._ZTSN2cl4sycl17specialization_idIiEE.cl::sycl::specialization_id" zeroinitializer, align 4
22+
@__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL9double_idEdLPv0EEET0_v = private unnamed_addr addrspace(1) constant [73 x i8] c"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL9double_idEEE\00", align 1
23+
@_ZL9double_id = internal addrspace(1) constant %"class._ZTSN2cl4sycl17specialization_idIdEE.cl::sycl::specialization_id" zeroinitializer, align 8
24+
@__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL12composite_idE9compositeLPv0EEET0_v = private unnamed_addr addrspace(1) constant [77 x i8] c"_ZTSN2cl4sycl6detail32specialization_id_name_generatorIL_ZL12composite_idEEE\00", align 1
25+
@_ZL12composite_id = internal addrspace(1) constant %"class._ZTSN2cl4sycl17specialization_idI9compositeEE.cl::sycl::specialization_id" zeroinitializer, align 8
26+
27+
; Function Attrs: convergent norecurse
28+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE7Kernel1"() local_unnamed_addr #0 comdat {
29+
entry:
30+
%ref.tmp.i = alloca %struct._ZTS9composite.composite, align 8
31+
%ref.tmp.ascast.i = addrspacecast %struct._ZTS9composite.composite* %ref.tmp.i to %struct._ZTS9composite.composite addrspace(4)*
32+
%call.i.i.i = tail call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* getelementptr inbounds ([70 x i8], [70 x i8] addrspace(4)* addrspacecast ([70 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL6int_idEiLPv0EEET0_v to [70 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idIiEE.cl::sycl::specialization_id" addrspace(1)* @_ZL6int_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3
33+
; CHECK: call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID0:]], i32 0)
34+
35+
%call.i.i23.i = tail call spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPKvS4_(i8 addrspace(4)* getelementptr inbounds ([73 x i8], [73 x i8] addrspace(4)* addrspacecast ([73 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL9double_idEdLPv0EEET0_v to [73 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idIdEE.cl::sycl::specialization_id" addrspace(1)* @_ZL9double_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3
36+
; CHECK: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID1:]], double 0.000000e+00)
37+
38+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI9compositeET_PKcPKvS5_(%struct._ZTS9composite.composite addrspace(4)* sret(%struct._ZTS9composite.composite) align 8 %ref.tmp.ascast.i, i8 addrspace(4)* getelementptr inbounds ([77 x i8], [77 x i8] addrspace(4)* addrspacecast ([77 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL12composite_idE9compositeLPv0EEET0_v to [77 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idI9compositeEE.cl::sycl::specialization_id" addrspace(1)* @_ZL12composite_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3
39+
; CHECK: call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
40+
; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID3:]], i8 0)
41+
; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 0)
42+
; CHECK: call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID5:]], i64 0)
43+
44+
ret void
45+
}
46+
47+
; Function Attrs: convergent
48+
declare dso_local spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2
49+
50+
; Function Attrs: convergent
51+
declare dso_local spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPKvS4_(i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2
52+
53+
; Function Attrs: convergent
54+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI9compositeET_PKcPKvS5_(%struct._ZTS9composite.composite addrspace(4)* sret(%struct._ZTS9composite.composite) align 8, i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2
55+
56+
attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="t.cpp" "uniform-work-group-size"="true" }
57+
attributes #1 = { argmemonly nofree nosync nounwind willreturn }
58+
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
59+
attributes #3 = { convergent }
60+
attributes #4 = { nounwind }
61+
62+
!llvm.module.flags = !{!0}
63+
!opencl.spir.version = !{!1}
64+
!spirv.Source = !{!2}
65+
!llvm.ident = !{!3}
66+
67+
!0 = !{i32 1, !"wchar_size", i32 4}
68+
!1 = !{i32 1, i32 2}
69+
!2 = !{i32 4, i32 100000}
70+
!3 = !{!"clang version 13.0.0 (/data/github.com/intel/llvm/clang c1f0fd875de45645c66776667682af55059968b7)"}
71+
!4 = !{i32 -1, i32 -1, i32 -1, i32 -1}
72+
!5 = !{!6, !6, i64 0}
73+
!6 = !{!"int", !7, i64 0}
74+
!7 = !{!"omnipotent char", !8, i64 0}
75+
!8 = !{!"Simple C++ TBAA"}
76+
!9 = !{!10, !7, i64 5}
77+
!10 = !{!"_ZTS9composite", !11, i64 0, !13, i64 4, !7, i64 5, !14, i64 8}
78+
!11 = !{!"_ZTS6nested", !12, i64 0}
79+
!12 = !{!"float", !7, i64 0}
80+
!13 = !{!"bool", !7, i64 0}
81+
!14 = !{!"long long", !7, i64 0}

‎llvm/tools/sycl-post-link/SpecConstants.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -508,8 +508,10 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
508508
auto *GV = dyn_cast<GlobalVariable>(
509509
CI->getArgOperand(NameArgNo + 1)->stripPointerCasts());
510510
if (GV) {
511+
assert(GV->hasInitializer() && "expected initializer");
511512
auto *Initializer = GV->getInitializer();
512-
assert(isa<ConstantAggregate>(Initializer) &&
513+
assert((isa<ConstantAggregate>(Initializer) ||
514+
Initializer->isZeroValue()) &&
513515
"expected specialization_id instance");
514516
// specialization_id structure contains a single field which is the
515517
// default value of corresponding specialization constant.

0 commit comments

Comments
 (0)