POD Specialization constants design update #2669
Conversation
Removed `__sycl_getCompositeSpecConstantValue` intrinsic, because we can check if specialization constant is a composite type by looking at return type of `__sycl_getSpecConstantValue`. This also allows us to avoid any changes in `cl::sycl::ONEAPI::experimental::spec_constant::get` method. Renamed `__spirvCompositeSpecConstant` to `__spirv_SpecConstantComposite` to better match SPIR-V friendly IR representaiton, which is aligned with opcode names from SPIR-V spec. Changed logic of generating `__spirv_SpecConstantComposite`: instead of getting list of `SpecID` for each composite member it now accepts vector of actual members, which might be another composits, undefs or regular constants. This change was done to align this intrinsic with corresponding opcode prototype to make its handling more universal and useful for the upstream translator project. Therefore, completely removed some sections which are outdated assuming new approach.
Type information is useless in precence of constant size
sycl/doc/SpecializationConstants.md
Outdated
|
|
||
| And 1 "composite" | ||
| %gold = call %struct.POD __spirv_SpecConstantComposite([2 x %struct.A] %gold_POD_A, i32 %gold_POD_b) |
There was a problem hiding this comment.
the only reason why __spirvCompositeSpecConstant accepted IDs instead of SSA values is a concern that there could be some non-trivia data flow (i.e. via store/load to alloca) of an operand, which might confuse the translator. If we can be sure that this will never happen, then using SSA with spec constant tree construction in IR seems fine.
Renamed section describing interaction with the SPIR-V translator. Returned back section about interaction with the SPIR-V translator about composite spec constants. Returned back `__sycl_getCompositeSpecConstantValue` intrinsic. Improved an example of composite spec constant by adding vector into it.
|
@AlexeySotkin - please review/approve |
| It does not need a default value, because its default value consists of default | ||
| valued of its leaf specialization constants (see below). | ||
| `__sycl_getCompositeSpecConstantValue` is a new "intrinsic" (in addition to | ||
| `__sycl_getSpecConstantValue`) recognized by the `SpecConstants` pass, which |
There was a problem hiding this comment.
@kbobrovs, what do you think about renaming existing __sycl_getSpecConstantValue into __sycl_getScalarSpecConstantValue to unify names and add more clarity?
There was a problem hiding this comment.
@AlexeySachkov, I don't mind
| struct A { | ||
| int x; | ||
| float y; | ||
| }; | ||
|
|
||
| struct POD { | ||
| A a[2]; | ||
| int b; | ||
| cl::sycl::vec<int, 2> b; |
There was a problem hiding this comment.
This example might actually be confusing, because cl::sycl::vec is not a POD-type in our implementation, because it is not trivially copyable due to the fact that there is non-trivial copy constructor (I don't know why it is coded in this way, probably @romanovvlad can say something about this).
I guess that from user point of view it would be nice to have possibility to use vectors as specialization constants, so, what should I do here: do I put a vector into the example or do I leave it as-is?
This reverts commit 4f864cd.
| value for each leaf as for a separate specialization constant. It contains the | ||
| following data: | ||
| - ID of composite constants leaf, i.e. ID of scalar spec constant | ||
| - Offset from the beginning of composite, which points to the location of a |
There was a problem hiding this comment.
I guess the offset also points to the location of the specialization value in the byte array provided by the user. Should we mention it here?
There was a problem hiding this comment.
I've updated this item in 8f323a1
|
Vào Th 5, 19 thg 11, 2020 lúc 20:04 Alexey Sotkin <[email protected]>
đã viết:
… ***@***.**** approved this pull request.
LGTM, Thanks!
—
You are receiving this because you are subscribed to this thread.
Reply to this email directly, view it on GitHub
<#2669 (review)>, or
unsubscribe
<https://github.com/notifications/unsubscribe-auth/AQBNS72VUW7VZLDVADKBNBDSQUJVJANCNFSM4S3JWPJA>
.
|
AlexeySachkov
deleted the
private/asachkov/spec-constants-design-update
branch
Besides some tiny markup & typo fixes this PR changes interfaces for defining specialization constants between
sycl-post-linkand SPIRV-LLVM-Translator, in particular:Renamed
__spirvCompositeSpecConstantto__spirv_SpecConstantCompositeto better match SPIR-V friendly IR representaiton, which is aligned with opcode names from SPIR-V spec.Changed logic of generating
__spirv_SpecConstantComposite: instead of getting list ofSpecIDfor each composite member it now accepts vector of actual members, which might be another composits, undefs or regular constants. This change was done to align this intrinsic with corresponding opcode prototype to make its handling more universal and useful for the upstream translator project.