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

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

Author: vmaksimo

llvm/test/tools/sycl-post-link/composite-spec-constant.ll

@@ -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
@@ -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"
@@ -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 {
@@ -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
+}
+
 ; Function Attrs: argmemonly nounwind willreturn
 declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1
 
@@ -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" }
@@ -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}
@@ -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}

llvm/test/tools/sycl-post-link/spec_const_O2.ll

@@ -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 {
@@ -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
 }
 
@@ -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" }
 
@@ -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}

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

@@ -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.
@@ -426,7 +430,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;
@@ -443,7 +449,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);

sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp

@@ -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,
+                                                      void *DefaultValue,
+                                                      void *RTBuffer);
+
+template <typename T>
+SYCL_EXTERNAL T __sycl_getComposite2020SpecConstantValue(const char *SymbolicID,
+                                                         void *DefaultValue,
+                                                         void *RTBuffer);
+
 #endif