Skip to content

[SYCL] Enable basic support for 2020 specialization constants in sycl-post-link #3353

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bader merged 3 commits into from
Mar 24, 2021
Merged

[SYCL] Enable basic support for 2020 specialization constants in sycl-post-link #3353

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
f7ebe9f
[SYCL] Enable basic support for 2020 specialization constants in sycl…
vmaksimo Mar 12, 2021
5ba3353
Apply code review comments
vmaksimo Mar 17, 2021
9a782af
Apply more comments
vmaksimo Mar 19, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 7 additions & 12 deletions llvm/test/tools/sycl-post-link/composite-spec-constant.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,18 +22,6 @@
;
; 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.
;
; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 6]], i32
; CHECK: %[[#N1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 7]], float
; CHECK: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#N0]], float %[[#N1]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD1:]]
; CHECK: %struct._ZTS1A.A %[[#CONST]]
;
; Common metadata checks
;
; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]}
; CHECK: ![[#MD1]] = !{!"_ZTS13MyComposConst", i32 [[#ID + 6]], i32 [[#ID + 7]]}

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"
target triple = "spir64-unknown-unknown-sycldevice"
Expand Down Expand Up @@ -69,6 +57,11 @@ entry:
ret void
}

; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 6]], i32
Copy link
Contributor

@kbobrovs kbobrovs Mar 18, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My idea was to actually move the checks right below the __sycl_get* intrinsics which lead to the checked IR.

; CHECK: %[[#N1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 7]], float
; CHECK: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#N0]], float %[[#N1]]), !SYCL_SPEC_CONST_SYM_ID ![[#MD1:]]
; CHECK: %struct._ZTS1A.A %[[#CONST]]

; Function Attrs: convergent norecurse
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 {
entry:
Expand Down Expand Up @@ -119,6 +112,8 @@ attributes #4 = { convergent }
!spirv.Source = !{!2}
!llvm.ident = !{!3}

; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]}
; CHECK: ![[#MD1]] = !{!"_ZTS13MyComposConst", i32 [[#ID + 6]], i32 [[#ID + 7]]}
!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 2}
!2 = !{i32 4, i32 100000}
Expand Down
1 change: 1 addition & 0 deletions llvm/test/tools/sycl-post-link/spec_const_O2.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,3 +180,4 @@ attributes #1 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-
; CHECK-RT: ![[ID8]] = !{!"_ZTS13MyUInt64Const", i32 8}
; CHECK-RT: ![[ID9]] = !{!"_ZTS12MyFloatConst", i32 9}
; CHECK-RT: ![[ID10]] = !{!"_ZTS13MyDoubleConst", i32 10}
; CHECK-RT: ![[ID11]] = !{!"_ZTS14MyDoubleConst2", i32 11}
10 changes: 10 additions & 0 deletions sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,16 @@ SYCL_EXTERNAL T __sycl_getSpecConstantValue(const char *ID);
template <typename T>
SYCL_EXTERNAL T __sycl_getCompositeSpecConstantValue(const char *ID);

// The intrinsics below are used to enable support of specialization constants
// 2020. This feature required post-link to handle more parameters.
Copy link
Contributor

@kbobrovs kbobrovs Mar 18, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks! One more nit

Suggested change
// 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.


// Get the value of the specialization constant with given symbolic ID.
// `SymbolicID` is a unique string ID of a specialization constant.
// `DefaultValue` contains a pointer to a global variable with the initializer,
// which should be used as the default value of the specialization constants.
// `RTBuffer` is a pointer to a runtime buffer, which holds values of all
// specialization constant and should be used if native specialization constants
// are not available.
template <typename T>
SYCL_EXTERNAL T __sycl_getScalar2020SpecConstantValue(const char *SymbolicID,
Copy link
Contributor

@kbobrovs kbobrovs Mar 17, 2021

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

please add description and doc comments for each new intrinsic

void *DefaultValue,
Expand Down