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
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 all commits
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
42 changes: 39 additions & 3 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
@@ -1,5 +1,5 @@
; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue
; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue --implicit-check-not __sycl_getComposite2020SpecConstantValue
;
; This test is intended to check that sycl-post-link tool is capable of handling
; composite specialization constants by lowering them into a set of SPIR-V
Expand All @@ -22,8 +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]]
;
; CHECK: ![[#MD]] = !{!"_ZTS3POD", i32 [[#ID]], i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]}

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 All @@ -36,8 +34,10 @@ target triple = "spir64-unknown-unknown-sycldevice"
%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" }

$_ZTS4Test = comdat any
$_ZTS17SpecializedKernel = comdat any

@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1
@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1

; Function Attrs: convergent norecurse uwtable
define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD 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 {
Expand All @@ -57,6 +57,34 @@ entry:
ret void
}

; 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:
%c.i = alloca %struct._ZTS1A.A, align 4
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
%2 = load i64, i64 addrspace(4)* %1, align 8
%add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %2
%c.ascast.i = addrspacecast %struct._ZTS1A.A* %c.i to %struct._ZTS1A.A addrspace(4)*
%3 = bitcast %struct._ZTS1A.A* %c.i to i8*
call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %3) #3
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* null) #4
; 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]]
%a.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 0
%4 = load i32, i32 addrspace(4)* %a.i, align 4
%conv.i = sitofp i32 %4 to float
%b.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 1
%5 = load float, float addrspace(4)* %b.i, align 4
%add.i = fadd float %5, %conv.i
%ptridx.ascast.i.i = addrspacecast float addrspace(1)* %add.ptr.i to float addrspace(4)*
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
}
Copy link
Contributor

Choose a reason for hiding this comment

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

there no single CHECK in this function - what is the purpose of adding it?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

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).

Copy link
Contributor

Choose a reason for hiding this comment

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

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.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Agree, done


; Function Attrs: argmemonly nounwind willreturn
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1

Expand All @@ -69,6 +97,9 @@ declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture write
; Function Attrs: convergent
declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 8, i8 addrspace(4)*) local_unnamed_addr #2

; Function Attrs: convergent
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4, i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2

attributes #0 = { convergent norecurse uwtable "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="../sycl/test/spec_const/composite.cpp" "tune-cpu"="generic" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { argmemonly nounwind willreturn }
attributes #2 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "tune-cpu"="generic" "unsafe-fp-math"="false" "use-soft-float"="false" }
Expand All @@ -80,6 +111,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 All @@ -89,3 +122,6 @@ attributes #4 = { convergent }
!6 = !{!7, !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!"int", !7, i64 0}
!10 = !{!"float", !7, i64 0}
!11 = !{!10, !10, i64 0}
17 changes: 13 additions & 4 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 @@ -39,6 +39,8 @@ $_ZTS17SpecializedKernel = comdat any
; CHECK-NOT: @__unique_stable_name.SC_If12MyFloatConstE3getEv
@__unique_stable_name.SC_Id13MyDoubleConstE3getEv = private unnamed_addr constant [20 x i8] c"_ZTS13MyDoubleConst\00", align 1
; CHECK-NOT: @__unique_stable_name.SC_Id13MyDoubleConstE3getEv
@__unique_stable_name.SC_Id14MyDoubleConst2E3getEv = private unnamed_addr constant [21 x i8] c"_ZTS14MyDoubleConst2\00", align 1
; CHECK-NOT: @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv

; Function Attrs: norecurse
define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrspace(1)* %0, %"cl::sycl::range"* byval(%"cl::sycl::range") align 8 %1, %"cl::sycl::range"* byval(%"cl::sycl::range") align 8 %2, %"cl::sycl::id"* byval(%"cl::sycl::id") align 8 %3) local_unnamed_addr #0 comdat !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
Expand Down Expand Up @@ -111,10 +113,14 @@ define weak_odr dso_local spir_kernel void @_ZTS17SpecializedKernel(float addrsp
; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 10, double 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID10:[0-9]+]]
%36 = fadd double %35, %34
; CHECK-DEF: %[[SUM9:[0-9]+]] = fadd double 0.000000e+00, %[[VAL7]]
%37 = fptrunc double %36 to float
; CHECK-DEF: %[[VAL8:[0-9]+]] = fptrunc double %[[SUM9]] to float
%38 = addrspacecast float addrspace(1)* %7 to float addrspace(4)*
store float %37, float addrspace(4)* %38, align 4, !tbaa !8
%37 = tail call spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([21 x i8], [21 x i8]* @__unique_stable_name.SC_Id14MyDoubleConst2E3getEv, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* null, i8 addrspace(4)* null)
; CHECK-RT: %{{[0-9]+}} = call double @_Z20__spirv_SpecConstantid(i32 11, double 0.000000e+00), !SYCL_SPEC_CONST_SYM_ID ![[ID11:[0-9]+]]
%38 = fadd double %37, %36
; CHECK-DEF: %[[SUM10:[0-9]+]] = fadd double 0.000000e+00, %[[SUM9]]
%39 = fptrunc double %38 to float
; CHECK-DEF: %[[VAL8:[0-9]+]] = fptrunc double %[[SUM10]] to float
%40 = addrspacecast float addrspace(1)* %7 to float addrspace(4)*
store float %39, float addrspace(4)* %40, align 4, !tbaa !8
ret void
}

Expand All @@ -140,6 +146,8 @@ declare dso_local spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 ad

declare dso_local spir_func double @_Z27__sycl_getSpecConstantValueIdET_PKc(i8 addrspace(4)*) local_unnamed_addr #1

declare dso_local spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPvS3_(i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #1

attributes #0 = { norecurse "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "sycl-module-id"="/iusers/kbobrovs/ws/kbobrovs_llvm/sycl/test/spec_const/spec_const_types.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }

Expand Down Expand Up @@ -172,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}
11 changes: 9 additions & 2 deletions llvm/tools/sycl-post-link/SpecConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,10 @@ namespace {
constexpr char SYCL_GET_SPEC_CONST_VAL[] = "_Z27__sycl_getSpecConstantValue";
constexpr char SYCL_GET_COMPOSITE_SPEC_CONST_VAL[] =
"_Z36__sycl_getCompositeSpecConstantValue";
constexpr char SYCL_GET_SCALAR_2020_SPEC_CONST_VAL[] =
"_Z37__sycl_getScalar2020SpecConstantValue";
constexpr char SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL[] =
"_Z40__sycl_getComposite2020SpecConstantValue";

// Unmangled base name of all __spirv_SpecConstant intrinsics which differ by
// the value type.
Expand Down Expand Up @@ -403,7 +407,9 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
continue;

if (!F.getName().startswith(SYCL_GET_SPEC_CONST_VAL) &&
!F.getName().startswith(SYCL_GET_COMPOSITE_SPEC_CONST_VAL))
!F.getName().startswith(SYCL_GET_COMPOSITE_SPEC_CONST_VAL) &&
!F.getName().startswith(SYCL_GET_SCALAR_2020_SPEC_CONST_VAL) &&
!F.getName().startswith(SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL))
continue;

SmallVector<CallInst *, 32> SCIntrCalls;
Expand All @@ -420,7 +426,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
// literals are passed to it in the SYCL RT source code, and application
// code can't use this intrinsic directly.
bool IsComposite =
F.getName().startswith(SYCL_GET_COMPOSITE_SPEC_CONST_VAL);
F.getName().startswith(SYCL_GET_COMPOSITE_SPEC_CONST_VAL) ||
F.getName().startswith(SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL);

SmallVector<Instruction *, 3> DelInsts;
DelInsts.push_back(CI);
Expand Down
21 changes: 21 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,4 +21,25 @@ 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 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

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,
void *RTBuffer);

template <typename T>
SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(const char *SymbolicID,
void *DefaultValue,
void *RTBuffer);

#endif