[SYCL] Enable basic support for 2020 specialization constants in sycl-post-link #3353
Conversation
…-post-link SYCL 2020 specialization constants implementation will use new intrinsic in DPC++ headers to provide required markup for the compiler. By doing this change, SpecConstantsPass starts to support new interface for specialization constants. It doesn't implement handling of new arguments, which would be done in further PRs (for now new arguments are just ignored).
vmaksimo
force-pushed
the
basic_spec_const_2020
branch
from
ca3eebc to
f7ebe9f
Compare
| @@ -21,4 +21,14 @@ SYCL_EXTERNAL T __sycl_getSpecConstantValue(const char *ID); | |||
| template <typename T> | |||
| SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID); | |||
|
|
|||
| template <typename T> | |||
| SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID, | |||
There was a problem hiding this comment.
please add description and doc comments for each new intrinsic
| @@ -23,7 +23,17 @@ | |||
| ; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"([2 x %struct._ZTS1A.A] %[[#NA]], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" %[[#B]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD:]] | |||
| ; CHECK: store %struct._ZTS3POD.POD %[[#POD]] | |||
| ; | |||
| ; Test checks related to 2020 API for composite specialization constants. | |||
There was a problem hiding this comment.
this does not seem to have any 2020 specifics, am I missing something?
There was a problem hiding this comment.
These checks are just related to function with new intrinsic API calls, but you're right that IR doesn't look different.
I'll remove this mark as a bit confusing one.
| store float %add.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !11 | ||
| call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %3) #3 | ||
| ret void | ||
| } |
There was a problem hiding this comment.
there no single CHECK in this function - what is the purpose of adding it?
There was a problem hiding this comment.
This function was added to test new __sycl_get*2020SpecializationConstant intrinsic (which are introduced in spec constants 2020 design).
The corresponding checks for this piece of IR are located at the beginning of the file (as it was done for this test initially).
There was a problem hiding this comment.
I suggest to move the checks next to the source IR portion which are supposed to produce the checked result. Otherwise it is really hard to follow the logic of the test.
| @@ -69,6 +57,11 @@ entry: | |||
| ret void | |||
| } | |||
|
|
|||
| ; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 6]], i32 | |||
There was a problem hiding this comment.
My idea was to actually move the checks right below the __sycl_get* intrinsics which lead to the checked IR.
| // The intrinsics below are used to enable support of specialization constants | ||
| // 2020. This feature required post-link to handle more parameters. |
There was a problem hiding this comment.
Thanks! One more nit
| // The intrinsics below are used to enable support of specialization constants | |
| // 2020. This feature required post-link to handle more parameters. | |
| // The intrinsics below are used to implement support SYCL2020 specialization constants. | |
| // SYCL2020 version requires more parameters compared to the initial version. |
I don't think that we need to duplicate the description for composite intrinsic as the comment would be a copy-paste. It relates to both of the new intrinsics (the same way that a previous couple of intrinsics was documented above). |
|
@kbobrovs, please, take a look if your concerns are resolved. |
Ok, works for me |
For met too. Added comment seems enough. |
SYCL 2020 specialization constants implementation will use new intrinsic in
DPC++ headers to provide required markup for the compiler.
By doing this change, SpecConstantsPass starts to support new interface for
specialization constants. It doesn't implement handling of new arguments, which
would be done in further PRs (for now new arguments are just ignored).
Spec is under review here, but it doesn't affect the patch promotion: #3331