[sycl-post-link] Add handling for zero-initialized array constants #5321

ghost · 2022-01-17T13:13:56Z

The problem is the following, when the getOperand method is invoked for a
zeroinitializer array: C->getOperand(i), the assert from User.h was triggered:

assert(i < NumUserOperands && "getOperand() out of range!");

This is because zeroinitializer has no operands. The problem occurs only when
the specialization constant was not initialized directly in the point of definition.

When a specialization constant is defined in the IR as zeroinitializer, its default
value is just an array of zeros of the arbitrary size. No recursive calls of the
collectCompositeElementsDefaultValuesRecursive function is required in the case.

Also, to make processing constants with different types more robust, dynamic
casting of the constant itself (dyn_cast(C)) should be used
instead of casting the constant's type since the type can be inconsistent with
the constant behavior, for example not every constant with the type equals to
ArrayType has any components (the zeroinitializer has not).

When a specialization constant, wich is an array, is defined in the IR as zeroinitializer, the corresponding llvm::Constant object has no operands and a zero constant of arbitrary type must be created for each array element.

ghost · 2022-01-17T13:17:55Z

/verify with intel/llvm-test-suite#747

mlychkov · 2022-01-17T13:26:28Z

Could you please also attach LIT test?

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

The test is needed to proceed with the change. Fix itself LGTM

ghost · 2022-01-17T13:42:41Z

@vmaksimo @mlychkov Thank you for the comments, I'm working on the test.

ghost · 2022-01-17T15:10:38Z

@vmaksimo @mlychkov The lit test has been added.

ghost · 2022-01-17T15:11:07Z

/verify with intel/llvm-test-suite#747

vmaksimo

The test is OK, but some improvements are welcome to simplify it.

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll

vmaksimo · 2022-01-17T16:47:13Z

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll

+attributes #0 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="spec-constant-convolution-test.cpp" "uniform-work-group-size"="true" }
+attributes #1 = { argmemonly nofree nosync nounwind willreturn }
+attributes #2 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #3 = { argmemonly nofree nounwind willreturn }
+attributes #4 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #5 = { convergent inlinehint mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #6 = { convergent norecurse nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #7 = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #8 = { argmemonly nofree nounwind willreturn writeonly }
+attributes #9 = { convergent inlinehint norecurse nounwind "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #10 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #11 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="C:/Users/psamolys/Dev/sycl/libdevice/fallback-cassert.cpp" }
+attributes #12 = { nounwind }
+attributes #13 = { convergent }
+attributes #14 = { convergent nounwind }
+
+!llvm.dependent-libraries = !{!0, !0, !0}
+!opencl.spir.version = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
+!spirv.Source = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
+!llvm.ident = !{!3, !4, !4, !4, !4, !4, !4, !4, !4, !4, !4, !4}
+!llvm.module.flags = !{!5, !6}
+
+!0 = !{!"libcpmt"}
+!1 = !{i32 1, i32 2}
+!2 = !{i32 4, i32 100000}
+!3 = !{!"clang version 14.0.0 (https://github.com/intel/llvm.git bb0055ccd006e861d781dc6332df658364590350)"}
+!4 = !{!"clang version 14.0.0 (https://github.com/intel/llvm.git 23e180bce8024817348ea07b1f5eefe15c7ace32)"}
+!5 = !{i32 1, !"wchar_size", i32 2}
+!6 = !{i32 7, !"frame-pointer", i32 2}
+!7 = !{i32 -1, i32 -1, i32 -1, i32 -1}
+!8 = !{!9, !9, i64 0}
+!9 = !{!"any pointer", !10, i64 0}
+!10 = !{!"omnipotent char", !11, i64 0}
+!11 = !{!"Simple C++ TBAA"}
+!12 = !{!10, !10, i64 0}
+!13 = !{!14, !14, i64 0}
+!14 = !{!"int", !10, i64 0}
+!15 = !{!16, !16, i64 0}
+!16 = !{!"long long", !10, i64 0}
+!17 = distinct !{!17, !18, !19}
+!18 = !{!"llvm.loop.mustprogress"}
+!19 = !{!"llvm.loop.unroll.enable"}
+!20 = !{i32 -1}
+!21 = !{i64 0, i64 8, !8}
+!22 = !{!23, !9, i64 0}
+!23 = !{!"_ZTSN2cl4sycl14kernel_handlerE", !9, i64 0}
+!24 = !{!25, !14, i64 0}
+!25 = !{!"_ZTS14AssertHappened", !14, i64 0, !10, i64 4, !10, i64 261, !10, i64 518, !14, i64 648, !16, i64 656, !16, i64 664, !16, i64 672, !16, i64 680, !16, i64 688, !16, i64 696}
+!26 = distinct !{!26, !18}
+!27 = !{i64 0, i64 4, !13, i64 4, i64 257, !12, i64 261, i64 257, !12, i64 518, i64 129, !12, i64 648, i64 4, !13, i64 656, i64 8, !15, i64 664, i64 8, !15, i64 672, i64 8, !15, i64 680, i64 8, !15, i64 688, i64 8, !15, i64 696, i64 8, !15}


I suggest simplifying the test a bit - many of these metadata and attributes are not needed to check spec constants instructions.
And it'd be nice to remove unused type definitions at the beginning of the test (e.g., %"struct.cl::sycl::detail::AssertHappened"), it'll make the test more "unit". Also you can remove some unrelated instructions from the test function itself (like @llvm.lifetime intrinsic) - they should not affect the tested change).

@vmaksimo Thank you for the suggestions. I've simplified the test a bit.

ghost · 2022-01-18T07:48:06Z

/verify with intel/llvm-test-suite#747

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll

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

Signed-off-by: Pavel Samolysov <[email protected]> Co-authored-by: Mikhail Lychkov <[email protected]>

ghost · 2022-01-18T10:25:12Z

@mlychkov I've applied your suggestions about the lit test in 9687d41, the code of the solution itself should be re-designed.

Also a check that the expected .prop file has been generated and contains the definitions of the specialization constants has been added.

To make processing constants with different types more robust, dynamic casting of the constant itself (dyn_cast<ConstantType>(C)) should be used instead of casting the constant's type since the type can be incosistent with the constant behavior, for example not every constant with the type equals to ArrayType has any components (the zeroinitializer has not).

ghost · 2022-01-18T17:03:39Z

I've pushed a commit that replaces dynamic casts for the constant's types with dynamic casts to the constant's subtypes, I believe this should make the code more robust. Also, an additional case is used to check for zeroinitializer. Please, have a look.

bader

Looks better. Thanks.

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

Also some local variables were inlined into function calls. Signed-off-by: Pavel Samolysov <[email protected]> Co-authored-by: Alexey Bader <[email protected]>

ghost · 2022-01-19T08:16:45Z

I've changed the description to make it clearer and describe all the changes in the PR.

AlexeySachkov

I think that test RUN lines could be cleaned up a bit

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll

ghost · 2022-01-19T12:17:24Z

/verify with intel/llvm-test-suite#747

ghost · 2022-01-19T14:06:53Z

/verify with intel/llvm-test-suite#747

ghost · 2022-01-19T14:08:29Z

I would like to test the vector-convolution-demo example in the llvm-test-suite project. The test fails due to errors in the test itself and it should not block merging this PR.

ghost · 2022-01-19T15:13:01Z

/verify with intel/llvm-test-suite#747

ghost · 2022-01-20T08:45:34Z

/verify with intel/llvm-test-suite#747

ghost requested review from AlexeySachkov and vmaksimo January 17, 2022 13:13

ghost self-requested a review as a code owner January 17, 2022 13:13

AlexeySachkov reviewed Jan 17, 2022

View reviewed changes

llvm/tools/sycl-post-link/SpecConstants.cpp Outdated Show resolved Hide resolved

vmaksimo previously approved these changes Jan 17, 2022

View reviewed changes

vmaksimo reviewed Jan 17, 2022

View reviewed changes

Pavel Samolysov added 2 commits January 18, 2022 10:25

[sycl-post-link] Add a lit test to check array-of-arrays handling

47e3837

[sycl-post-link] Simplify the lit test

3875a0f

AlexeySachkov previously approved these changes Jan 18, 2022

View reviewed changes

mlychkov reviewed Jan 18, 2022

View reviewed changes

llvm/tools/sycl-post-link/SpecConstants.cpp Outdated Show resolved Hide resolved

[sycl-post-link] Simplify the lit test again

9687d41

Signed-off-by: Pavel Samolysov <[email protected]> Co-authored-by: Mikhail Lychkov <[email protected]>

ghost dismissed AlexeySachkov’s stale review via 9687d41 January 18, 2022 10:22

ghost marked this pull request as draft January 18, 2022 10:27

Pavel Samolysov added 3 commits January 18, 2022 16:24

[sycl-post-link] Add a lit test to check a partially initialized array

6d80223

[sycl-post-link] Add a lit test to check a structure with an array

a84ae51

Also a check that the expected .prop file has been generated and contains the definitions of the specialization constants has been added.

[sycl-post-link][NFC] Fix formatting issue

1708dae

ghost marked this pull request as ready for review January 18, 2022 17:06

ghost requested a review from vmaksimo January 18, 2022 17:06

ghost requested review from mlychkov, AlexeySachkov and bader January 18, 2022 17:06

bader reviewed Jan 18, 2022

View reviewed changes

Pavel Samolysov and others added 2 commits January 19, 2022 10:56

[sycl-post-link] Using early exit to make code more readable

662aeba

Also some local variables were inlined into function calls. Signed-off-by: Pavel Samolysov <[email protected]> Co-authored-by: Alexey Bader <[email protected]>

[sycl-post-link][NFC] Fix formatting issues

a5a3f85

ghost requested a review from bader January 19, 2022 08:14

AlexeySachkov previously approved these changes Jan 19, 2022

View reviewed changes

llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll Outdated Show resolved Hide resolved

[sycl-post-link] Clean up a bit the RUN lines for the lit test

657aa24

ghost dismissed AlexeySachkov’s stale review via 657aa24 January 19, 2022 11:07

AlexeySachkov approved these changes Jan 19, 2022

View reviewed changes

bader approved these changes Jan 19, 2022

View reviewed changes

AlexeySachkov merged commit eec22ed into intel:sycl Jan 20, 2022

[sycl-post-link] Add handling for zero-initialized array constants #5321

[sycl-post-link] Add handling for zero-initialized array constants #5321

Uh oh!

Conversation

ghost commented Jan 17, 2022 • edited by ghost Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ghost commented Jan 17, 2022

Uh oh!

mlychkov commented Jan 17, 2022

Uh oh!

Uh oh!

ghost commented Jan 17, 2022

Uh oh!

ghost commented Jan 17, 2022

Uh oh!

ghost commented Jan 17, 2022

Uh oh!

vmaksimo left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

vmaksimo Jan 17, 2022

Choose a reason for hiding this comment

Uh oh!

ghost Jan 18, 2022

Choose a reason for hiding this comment

Uh oh!

ghost commented Jan 18, 2022

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

ghost commented Jan 18, 2022

Uh oh!

ghost commented Jan 18, 2022

Uh oh!

bader left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

ghost commented Jan 19, 2022 • edited by ghost Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

AlexeySachkov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

ghost commented Jan 19, 2022

Uh oh!

ghost commented Jan 19, 2022

Uh oh!

ghost commented Jan 19, 2022 • edited by ghost Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ghost commented Jan 19, 2022

Uh oh!

ghost commented Jan 20, 2022

Uh oh!

Uh oh!

ghost commented Jan 17, 2022 •

edited by ghost

Loading

ghost commented Jan 19, 2022 •

edited by ghost

Loading

ghost commented Jan 19, 2022 •

edited by ghost

Loading