From e481174da6470e055e4b401e86949ed801bb6c16 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 1 Dec 2020 18:36:43 +0300 Subject: [PATCH] [sycl-post-link] Add support for composite specialization constants (#2779) Design document can be found in sycl/doc/SpecializationConstants.md. --- llvm/include/llvm/Support/PropertySetIO.h | 6 + llvm/lib/Support/PropertySetIO.cpp | 1 + .../composite-spec-constant-O0.ll | 469 ++++++++++++++++++ .../composite-spec-constant-default-value.ll | 79 +++ .../sycl-post-link/composite-spec-constant.ll | 91 ++++ .../multiple-composite-spec-const-usages.ll | 98 ++++ .../multiple-spec-const-usages.ll | 40 ++ llvm/tools/sycl-post-link/SpecConstants.cpp | 414 ++++++++++++---- llvm/tools/sycl-post-link/SpecConstants.h | 27 +- llvm/tools/sycl-post-link/sycl-post-link.cpp | 22 +- 10 files changed, 1151 insertions(+), 96 deletions(-) create mode 100644 llvm/test/tools/sycl-post-link/composite-spec-constant-O0.ll create mode 100644 llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll create mode 100644 llvm/test/tools/sycl-post-link/composite-spec-constant.ll create mode 100644 llvm/test/tools/sycl-post-link/multiple-composite-spec-const-usages.ll create mode 100644 llvm/test/tools/sycl-post-link/multiple-spec-const-usages.ll diff --git a/llvm/include/llvm/Support/PropertySetIO.h b/llvm/include/llvm/Support/PropertySetIO.h index a396ecafd5a24..d88a7e0261bf9 100644 --- a/llvm/include/llvm/Support/PropertySetIO.h +++ b/llvm/include/llvm/Support/PropertySetIO.h @@ -80,6 +80,10 @@ class PropertyValue { PropertyValue(uint32_t Val) : Ty(UINT32), Val({Val}) {} PropertyValue(const byte *Data, SizeTy DataBitSize); + template + PropertyValue(const std::vector &Data) + : PropertyValue(reinterpret_cast(Data.data()), + Data.size() * sizeof(T) * /* bits in one byte */ 8) {} PropertyValue(const PropertyValue &P); PropertyValue(PropertyValue &&P); @@ -179,6 +183,8 @@ class PropertySetRegistry { // Specific property category names used by tools. static constexpr char SYCL_SPECIALIZATION_CONSTANTS[] = "SYCL/specialization constants"; + static constexpr char SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[] = + "SYCL/composite specialization constants"; static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask"; static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt"; diff --git a/llvm/lib/Support/PropertySetIO.cpp b/llvm/lib/Support/PropertySetIO.cpp index 7c5b028d767c5..f85a46af56a9d 100644 --- a/llvm/lib/Support/PropertySetIO.cpp +++ b/llvm/lib/Support/PropertySetIO.cpp @@ -197,6 +197,7 @@ void PropertyValue::copy(const PropertyValue &P) { constexpr char PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS[]; constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[]; constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[]; +constexpr char PropertySetRegistry::SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS[]; } // namespace util } // namespace llvm diff --git a/llvm/test/tools/sycl-post-link/composite-spec-constant-O0.ll b/llvm/test/tools/sycl-post-link/composite-spec-constant-O0.ll new file mode 100644 index 0000000000000..94dfeb5abf7e0 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/composite-spec-constant-O0.ll @@ -0,0 +1,469 @@ +; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue +; +; 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 +; friendly IR operations representing those constants. +; This particular LLVM IR is generated from the same source as for +; composite-spec-constant.ll test, but -O0 optimization level was used to check +; that sycl-post-link is capable to handle this form of LLVM IR as well. +; +; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 +; CHECK: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float +; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS0]], float %[[#NS1]]) +; +; CHECK: %[[#NS2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 2]], i32 +; CHECK: %[[#NS3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 3]], float +; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS2]], float %[[#NS3]]) +; +; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) +; +; CHECK: %[[#B:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 4]], i32{{.*}}) +; +; CHECK: %[[#POD:]] = call %struct._ZTS3POD.POD @_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Ai([2 x %struct._ZTS1A.A] %[[#NA]], i32 %[[#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]]} + +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" + +%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], i32 } +%struct._ZTS1A.A = type { i32, float } +%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } +%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" = type <{ %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant", [7 x i8] }> +%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" = type { %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEUt_E.anon } +%"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" = type { %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" } +%union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEUt_E.anon = type { %struct._ZTS3POD.POD addrspace(1)* } +%"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" = type { i8 } +%"class._ZTSN2cl4sycl6detail15accessor_commonI3PODLi1ELNS0_6access4modeE1025ELNS4_6targetE2014ELNS4_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::detail::accessor_common" = type { i8 } + +$_ZTS4Test = comdat any + +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev = comdat any + +$_ZN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EC2Ev = comdat any + +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE = comdat any + +$_ZN2cl4sycl2idILi1EEC2Ev = comdat any + +$_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv = comdat any + +$_ZN2cl4sycl6detail18AccessorImplDeviceILi1EEC2ENS0_2idILi1EEENS0_5rangeILi1EEES7_ = comdat any + +$_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE = comdat any + +$_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE = comdat any + +$_ZN2cl4sycl6detail5arrayILi1EEixEi = comdat any + +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE9getOffsetEv = comdat any + +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getAccessRangeEv = comdat any + +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getMemoryRangeEv = comdat any + +$_ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv = comdat any + +$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE = comdat any + +$_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE = comdat any + +$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE = comdat any + +$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE15getQualifiedPtrEv = comdat any + +@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1 + +; Function Attrs: convergent noinline norecurse optnone mustprogress +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) #0 comdat !kernel_arg_buffer_location !4 { +entry: + %_arg_.addr = alloca %struct._ZTS3POD.POD addrspace(1)*, align 8 + %0 = alloca %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", align 8 + %agg.tmp = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + %agg.tmp4 = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + %agg.tmp5 = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + store %struct._ZTS3POD.POD addrspace(1)* %_arg_, %struct._ZTS3POD.POD addrspace(1)** %_arg_.addr, align 8 + %1 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon"* %0, i32 0, i32 0 + %2 = addrspacecast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor"* %1 to %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* + call spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %2) #8 + %3 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon"* %0, i32 0, i32 1 + %4 = addrspacecast %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant"* %3 to %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* + call spir_func void @_ZN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EC2Ev(%"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %4) #8 + %5 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon"* %0, i32 0, i32 0 + %6 = load %struct._ZTS3POD.POD addrspace(1)*, %struct._ZTS3POD.POD addrspace(1)** %_arg_.addr, align 8 + %7 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp to i8* + %8 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %_arg_1 to i8* + call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %7, i8* align 8 %8, i64 8, i1 false) + %9 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp4 to i8* + %10 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %_arg_2 to i8* + call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %9, i8* align 8 %10, i64 8, i1 false) + %11 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp5 to i8* + %12 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3 to i8* + call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %11, i8* align 8 %12, i64 8, i1 false) + %13 = addrspacecast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor"* %5 to %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* + call spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %13, %struct._ZTS3POD.POD addrspace(1)* %6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %agg.tmp, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %agg.tmp4, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %agg.tmp5) #8 + %14 = addrspacecast %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon"* %0 to %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* + call spir_func void @"_ZZZ4mainENK3$_1clERN2cl4sycl7handlerEENKUlvE_clEv"(%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* %14) #8 + ret void +} + +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + %agg.tmp = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + %agg.tmp2 = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + %agg.tmp3 = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1 to %"class._ZTSN2cl4sycl6detail15accessor_commonI3PODLi1ELNS0_6access4modeE1025ELNS4_6targetE2014ELNS4_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::detail::accessor_common" addrspace(4)* + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 0 + %1 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp to i8* + call void @llvm.memset.p0i8.i64(i8* align 8 %1, i8 0, i64 8, i1 false) + %2 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* + call spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %2) #8 + %3 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp2 to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* sret(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %3) #8 + %4 = addrspacecast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %agg.tmp3 to %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* sret(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %4) #8 + call spir_func void @_ZN2cl4sycl6detail18AccessorImplDeviceILi1EEC2ENS0_2idILi1EEENS0_5rangeILi1EEES7_(%"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %impl, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %agg.tmp, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %agg.tmp2, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %agg.tmp3) #8 + ret void +} + +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EC2Ev(%"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %this) unnamed_addr #2 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %this, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)*, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + ret void +} + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %struct._ZTS3POD.POD addrspace(1)* %Ptr, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %AccessRange, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %MemRange, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %Offset) #3 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + %Ptr.addr = alloca %struct._ZTS3POD.POD addrspace(1)*, align 8 + %I = alloca i32, align 4 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + store %struct._ZTS3POD.POD addrspace(1)* %Ptr, %struct._ZTS3POD.POD addrspace(1)** %Ptr.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = load %struct._ZTS3POD.POD addrspace(1)*, %struct._ZTS3POD.POD addrspace(1)** %Ptr.addr, align 8 + %1 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEUt_E.anon addrspace(4)* %1 to %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* + store %struct._ZTS3POD.POD addrspace(1)* %0, %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* %MData, align 8 + store i32 0, i32* %I, align 4 + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %2 = load i32, i32* %I, align 4 + %cmp = icmp slt i32 %2, 1 + br i1 %cmp, label %for.body, label %for.end + +for.body: ; preds = %for.cond + %3 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %Offset to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* + %4 = load i32, i32* %I, align 4 + %5 = addrspacecast %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* %3 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %call = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %5, i32 %4) #8 + %6 = load i64, i64 addrspace(4)* %call, align 8 + %call2 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE9getOffsetEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1) #8 + %7 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %call2 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %8 = load i32, i32* %I, align 4 + %call3 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %7, i32 %8) #8 + store i64 %6, i64 addrspace(4)* %call3, align 8 + %9 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %AccessRange to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* + %10 = load i32, i32* %I, align 4 + %11 = addrspacecast %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* %9 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %call4 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %11, i32 %10) #8 + %12 = load i64, i64 addrspace(4)* %call4, align 8 + %call5 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getAccessRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1) #8 + %13 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %call5 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %14 = load i32, i32* %I, align 4 + %call6 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %13, i32 %14) #8 + store i64 %12, i64 addrspace(4)* %call6, align 8 + %15 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %MemRange to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* + %16 = load i32, i32* %I, align 4 + %17 = addrspacecast %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* %15 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %call7 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %17, i32 %16) #8 + %18 = load i64, i64 addrspace(4)* %call7, align 8 + %call8 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getMemoryRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1) #8 + %19 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %call8 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %20 = load i32, i32* %I, align 4 + %call9 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %19, i32 %20) #8 + store i64 %18, i64 addrspace(4)* %call9, align 8 + br label %for.inc + +for.inc: ; preds = %for.body + %21 = load i32, i32* %I, align 4 + %inc = add nsw i32 %21, 1 + store i32 %inc, i32* %I, align 4 + br label %for.cond, !llvm.loop !5 + +for.end: ; preds = %for.cond + %22 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %Offset to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* + %23 = addrspacecast %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* %22 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %call10 = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %23, i32 0) #8 + %24 = load i64, i64 addrspace(4)* %call10, align 8 + %25 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData11 = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEUt_E.anon addrspace(4)* %25 to %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* + %26 = load %struct._ZTS3POD.POD addrspace(1)*, %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* %MData11, align 8 + %add.ptr = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %26, i64 %24 + store %struct._ZTS3POD.POD addrspace(1)* %add.ptr, %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* %MData11, align 8 + ret void +} + +; Function Attrs: argmemonly nofree nosync nounwind willreturn +declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #4 + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define internal spir_func void @"_ZZZ4mainENK3$_1clERN2cl4sycl7handlerEENKUlvE_clEv"(%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* %this) #3 align 2 { +entry: + %this.addr = alloca %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)*, align 8 + %ref.tmp = alloca %struct._ZTS3POD.POD, align 4 + %agg.tmp = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + store %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* %this, %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)*, %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)** %this.addr, align 8 + %0 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* %this1, i32 0, i32 1 + %1 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 4 %1, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %0) #8 + %2 = getelementptr inbounds %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon", %"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" addrspace(4)* %this1, i32 0, i32 0 + %3 = addrspacecast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp to %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* + call spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %3, i64 0) #8 + %call = call spir_func align 4 dereferenceable(20) %struct._ZTS3POD.POD addrspace(4)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %agg.tmp) #8 + %4 = bitcast %struct._ZTS3POD.POD addrspace(4)* %call to i8 addrspace(4)* + %5 = bitcast %struct._ZTS3POD.POD* %ref.tmp to i8* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 4 %4, i8* align 4 %5, i64 20, i1 false) + ret void +} + +; Function Attrs: argmemonly nofree nosync nounwind willreturn writeonly +declare void @llvm.memset.p0i8.i64(i8* nocapture writeonly, i8, i64, i1 immarg) #5 + +; Function Attrs: convergent noinline norecurse optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2Ev(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this) unnamed_addr #1 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this1 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %0, i64 0) #8 + ret void +} + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail14InitializedValILi1ENS0_5rangeEE3getILi0EEENS3_ILi1EEEv(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* noalias sret(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %agg.result) #3 comdat align 2 { +entry: + call spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %agg.result, i64 0) #8 + ret void +} + +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail18AccessorImplDeviceILi1EEC2ENS0_2idILi1EEENS0_5rangeILi1EEES7_(%"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %this, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %Offset, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %AccessRange, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %MemoryRange) unnamed_addr #2 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %this, %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)*, %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)** %this.addr, align 8 + %Offset2 = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %this1, i32 0, i32 0 + %0 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %Offset2 to i8 addrspace(4)* + %1 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %Offset to i8* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 %0, i8* align 8 %1, i64 8, i1 false) + %AccessRange3 = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %this1, i32 0, i32 1 + %2 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %AccessRange3 to i8 addrspace(4)* + %3 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %AccessRange to i8* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 %2, i8* align 8 %3, i64 8, i1 false) + %MemRange = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %this1, i32 0, i32 2 + %4 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %MemRange to i8 addrspace(4)* + %5 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %MemoryRange to i8* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 %4, i8* align 8 %5, i64 8, i1 false) + ret void +} + +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this, i64 %dim0) unnamed_addr #2 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, align 8 + %dim0.addr = alloca i64, align 8 + store %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr, align 8 + store i64 %dim0, i64* %dim0.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr, align 8 + %common_array = getelementptr inbounds %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array", %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this1, i32 0, i32 0 + %arrayinit.begin = getelementptr inbounds [1 x i64], [1 x i64] addrspace(4)* %common_array, i64 0, i64 0 + %0 = load i64, i64* %dim0.addr, align 8 + store i64 %0, i64 addrspace(4)* %arrayinit.begin, align 8 + ret void +} + +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %this, i64 %dim0) unnamed_addr #2 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)*, align 8 + %dim0.addr = alloca i64, align 8 + store %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %this, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)** %this.addr, align 8 + store i64 %dim0, i64* %dim0.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)*, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %this1 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %1 = load i64, i64* %dim0.addr, align 8 + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %0, i64 %1) #8 + ret void +} + +; Function Attrs: argmemonly nofree nosync nounwind willreturn +declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #4 + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this, i32 %dimension) #3 comdat align 2 { +entry: + %this.addr.i = alloca %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, align 8 + %dimension.addr.i = alloca i32, align 4 + %this.addr = alloca %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, align 8 + %dimension.addr = alloca i32, align 4 + store %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr, align 8 + store i32 %dimension, i32* %dimension.addr, align 4 + %this1 = load %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr, align 8 + %0 = load i32, i32* %dimension.addr, align 4 + store %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this1, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr.i, align 8 + store i32 %0, i32* %dimension.addr.i, align 4 + %this1.i = load %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)*, %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)** %this.addr.i, align 8 + %common_array = getelementptr inbounds %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array", %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %this1, i32 0, i32 0 + %1 = load i32, i32* %dimension.addr, align 4 + %idxprom = sext i32 %1 to i64 + %arrayidx = getelementptr inbounds [1 x i64], [1 x i64] addrspace(4)* %common_array, i64 0, i64 %idxprom + ret i64 addrspace(4)* %arrayidx +} + +; Function Attrs: convergent noinline norecurse nounwind optnone mustprogress +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE9getOffsetEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this) #6 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 0 + %Offset = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %impl, i32 0, i32 0 + ret %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %Offset +} + +; Function Attrs: convergent noinline norecurse nounwind optnone mustprogress +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getAccessRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this) #6 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 0 + %AccessRange = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %impl, i32 0, i32 1 + ret %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %AccessRange +} + +; Function Attrs: convergent noinline norecurse nounwind optnone mustprogress +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getMemoryRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this) #6 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 0 + %MemRange = getelementptr inbounds %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice" addrspace(4)* %impl, i32 0, i32 2 + ret %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* %MemRange +} + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func void @_ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv(%struct._ZTS3POD.POD addrspace(4)* noalias sret(%struct._ZTS3POD.POD) align 4 %agg.result, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %this) #3 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)*, align 8 + %TName = alloca i8 addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)* %this, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)*, %"class._ZTSN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EE.cl::sycl::ONEAPI::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + store i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)** %TName, align 8 + %0 = load i8 addrspace(4)*, i8 addrspace(4)** %TName, align 8 + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 4 %agg.result, i8 addrspace(4)* %0) #8 + ret void +} + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func align 4 dereferenceable(20) %struct._ZTS3POD.POD addrspace(4)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %Index) #3 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + %LinearIndex = alloca i64, align 8 + %agg.tmp = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %agg.tmp to i8* + %1 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %Index to i8* + call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %0, i8* align 8 %1, i64 8, i1 false) + %call = call spir_func i64 @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %agg.tmp) #8 + store i64 %call, i64* %LinearIndex, align 8 + %call2 = call spir_func %struct._ZTS3POD.POD addrspace(1)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE15getQualifiedPtrEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1) #8 + %2 = load i64, i64* %LinearIndex, align 8 + %ptridx = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %call2, i64 %2 + %ptridx.ascast = addrspacecast %struct._ZTS3POD.POD addrspace(1)* %ptridx to %struct._ZTS3POD.POD addrspace(4)* + ret %struct._ZTS3POD.POD addrspace(4)* %ptridx.ascast +} + +; Function Attrs: convergent noinline norecurse nounwind optnone +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl2idILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this, i64 %dim0) unnamed_addr #2 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*, align 8 + %dim0.addr = alloca i64, align 8 + store %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)** %this.addr, align 8 + store i64 %dim0, i64* %dim0.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)*, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* %this1 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %1 = load i64, i64* %dim0.addr, align 8 + call spir_func void @_ZN2cl4sycl6detail5arrayILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %0, i64 %1) #8 + ret void +} + +; Function Attrs: convergent +declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 4, i8 addrspace(4)*) #7 + +; Function Attrs: convergent noinline norecurse optnone mustprogress +define linkonce_odr dso_local spir_func i64 @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %Id) #3 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + %Result = alloca i64, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %Id to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* + %1 = addrspacecast %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array"* %0 to %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* + %call = call spir_func align 8 dereferenceable(8) i64 addrspace(4)* @_ZN2cl4sycl6detail5arrayILi1EEixEi(%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" addrspace(4)* %1, i32 0) #8 + %2 = load i64, i64 addrspace(4)* %call, align 8 + ret i64 %2 +} + +; Function Attrs: convergent noinline norecurse nounwind optnone mustprogress +define linkonce_odr dso_local spir_func %struct._ZTS3POD.POD addrspace(1)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE15getQualifiedPtrEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this) #6 comdat align 2 { +entry: + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEUt_E.anon addrspace(4)* %0 to %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* + %1 = load %struct._ZTS3POD.POD addrspace(1)*, %struct._ZTS3POD.POD addrspace(1)* addrspace(4)* %MData, align 8 + ret %struct._ZTS3POD.POD addrspace(1)* %1 +} + +attributes #0 = { convergent noinline norecurse optnone mustprogress "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"="./composite.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { convergent noinline norecurse optnone "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" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent noinline norecurse nounwind optnone "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" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #3 = { convergent noinline norecurse optnone mustprogress "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" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #4 = { argmemonly nofree nosync nounwind willreturn } +attributes #5 = { argmemonly nofree nosync nounwind willreturn writeonly } +attributes #6 = { convergent noinline norecurse nounwind optnone mustprogress "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" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #7 = { 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" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #8 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 12.0.0 "} +!4 = !{i32 -1, i32 -1, i32 -1, i32 -1} +!5 = distinct !{!5, !6, !7} +!6 = !{!"llvm.loop.mustprogress"} +!7 = !{!"llvm.loop.unroll.enable"} diff --git a/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll b/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll new file mode 100644 index 0000000000000..f73c5ffeb7242 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/composite-spec-constant-default-value.ll @@ -0,0 +1,79 @@ +; RUN: sycl-post-link -spec-const=default --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue +; +; This test checks that composite specialization constants can be correctly +; initialized by sycl-post-link tool for AOT use-case (default initialization +; should be used according to the type of constant) +; +; TODO: consider adding a test case with vector type: the pass itself already +; supports this, but at the moment, sycl::vec type is not a POD type, which +; means we can't have it within a spec constant, i.e. we can't generate LLVM IR +; from a real-life application to use it as a test here. +; +; CHECK: %[[#CAST:]] = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i +; CHECK: store %struct._ZTS3POD.POD zeroinitializer, %struct._ZTS3POD.POD {{.*}}* %[[#CAST]] + +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" + +%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" } +%struct._ZTS1A.A = type { i32, float } +%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> } +%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } +%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } + +$_ZTS4Test = 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 + +; 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 { +entry: + %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8 + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 + %1 = load i64, i64* %0, align 8 + %add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1 + %2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8* + call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3 + %3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 + %4 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)* + %5 = addrspacecast i8 addrspace(1)* %4 to i8 addrspace(4)* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %5, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false), !tbaa.struct !5 + call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3 + ret void +} + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1 + +; Function Attrs: convergent +declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8, 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" } +attributes #3 = { nounwind } +attributes #4 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 12.0.0 (/data/github.com/intel/llvm/clang 56ee5b054b5a1f2f703722fc414fcb05af18b40a)"} +!4 = !{i32 -1, i32 -1, i32 -1, i32 -1} +!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6} +!6 = !{!7, !7, i64 0} +!7 = !{!"omnipotent char", !8, i64 0} +!8 = !{!"Simple C++ TBAA"} diff --git a/llvm/test/tools/sycl-post-link/composite-spec-constant.ll b/llvm/test/tools/sycl-post-link/composite-spec-constant.ll new file mode 100644 index 0000000000000..68ba71994473b --- /dev/null +++ b/llvm/test/tools/sycl-post-link/composite-spec-constant.ll @@ -0,0 +1,91 @@ +; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue +; +; 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 +; friendly IR operations representing those constants. +; +; CHECK: %[[#NS0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 +; CHECK: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float +; CHECK: %[[#NA0:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS0]], float %[[#NS1]]) +; +; CHECK: %[[#NS2:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 2]], i32 +; CHECK: %[[#NS3:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 3]], float +; CHECK: %[[#NA1:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif(i32 %[[#NS2]], float %[[#NS3]]) +; +; CHECK: %[[#NA:]] = call [2 x %struct._ZTS1A.A] @_Z29__spirv_SpecConstantCompositestruct._ZTS1A.Astruct._ZTS1A.A(%struct._ZTS1A.A %[[#NA0]], %struct._ZTS1A.A %[[#NA1]]) +; +; CHECK: %[[#B0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 4]], i32{{.*}}) +; CHECK: %[[#B1:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID + 5]], i32{{.*}}) +; CHECK: %[[#BV:]] = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 %[[#B0]], i32 %[[#B1]]) +; CHECK: %[[#B:]] = call %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" @_Z29__spirv_SpecConstantCompositeDv2_i(<2 x i32> %[[#BV]]) +; +; 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" + +%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" } +%struct._ZTS1A.A = type { i32, float } +%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> } +%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } +%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } + +$_ZTS4Test = 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 + +; 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 { +entry: + %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8 + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 + %1 = load i64, i64* %0, align 8 + %add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1 + %2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8* + call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3 + %3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 + %4 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)* + %5 = addrspacecast i8 addrspace(1)* %4 to i8 addrspace(4)* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %5, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false), !tbaa.struct !5 + call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3 + ret void +} + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1 + +; Function Attrs: convergent +declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8, 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" } +attributes #3 = { nounwind } +attributes #4 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 12.0.0 "} +!4 = !{i32 -1, i32 -1, i32 -1, i32 -1} +!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6} +!6 = !{!7, !7, i64 0} +!7 = !{!"omnipotent char", !8, i64 0} +!8 = !{!"Simple C++ TBAA"} diff --git a/llvm/test/tools/sycl-post-link/multiple-composite-spec-const-usages.ll b/llvm/test/tools/sycl-post-link/multiple-composite-spec-const-usages.ll new file mode 100644 index 0000000000000..c50f29d30fc22 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/multiple-composite-spec-const-usages.ll @@ -0,0 +1,98 @@ +; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s --implicit-check-not __sycl_getCompositeSpecConstantValue +; +; This test is intended to check that sycl-post-link tool is capable of handling +; situations when the same composite specialization constants is used more than +; once +; +; CHECK-LABEL: @foo1 +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD0:]] +; CHECK-LABEL: @_ZTS4Test +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD1:]] +; CHECK-LABEL: @foo2 +; CHECK: call %struct._ZTS3POD.POD @"_Z29__spirv_SpecConstantCompositeAstruct._ZTS1A.Aclass._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec"({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[#MD0:]] +; +; CHECK-DAG: ![[#MD0]] = !{!"_ZTS3PO2", i32 [[#ID:]], +; CHECK-SAME: i32 [[#ID + 1]], i32 [[#ID + 2]], i32 [[#ID + 3]], i32 [[#ID + 4]], i32 [[#ID + 5]]} +; CHECK-DAG: ![[#MD1]] = !{!"_ZTS3POD", i32 [[#ID1:]], +; CHECK-SAME: i32 [[#ID1 + 1]], i32 [[#ID1 + 2]], i32 [[#ID1 + 3]], i32 [[#ID1 + 4]], i32 [[#ID1 + 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" + +%struct._ZTS3POD.POD = type { [2 x %struct._ZTS1A.A], %"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" } +%struct._ZTS1A.A = type { i32, float } +%"class._ZTSN2cl4sycl3vecIiLi2EEE.cl::sycl::vec" = type { <2 x i32> } +%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } +%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } +%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } + +$_ZTS4Test = 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.2 = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3PO2\00", align 1 + +define spir_func void @foo1() { + %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8 + %1 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name.2, i64 0, i64 0) to i8 addrspace(4)*)) #4 + ret void +} + +; 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 { +entry: + %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8 + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 + %1 = load i64, i64* %0, align 8 + %add.ptr.i = getelementptr inbounds %struct._ZTS3POD.POD, %struct._ZTS3POD.POD addrspace(1)* %_arg_, i64 %1 + %2 = bitcast %struct._ZTS3POD.POD* %ref.tmp.i to i8* + call void @llvm.lifetime.start.p0i8(i64 24, i8* nonnull %2) #3 + %3 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 + %4 = bitcast %struct._ZTS3POD.POD addrspace(1)* %add.ptr.i to i8 addrspace(1)* + %5 = addrspacecast i8 addrspace(1)* %4 to i8 addrspace(4)* + call void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* align 8 dereferenceable(24) %5, i8* nonnull align 8 dereferenceable(24) %2, i64 24, i1 false), !tbaa.struct !5 + call void @llvm.lifetime.end.p0i8(i64 24, i8* nonnull %2) #3 + ret void +} + +define spir_func void @foo2() { + %ref.tmp.i = alloca %struct._ZTS3POD.POD, align 8 + %1 = addrspacecast %struct._ZTS3POD.POD* %ref.tmp.i to %struct._ZTS3POD.POD addrspace(4)* + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name.2, i64 0, i64 0) to i8 addrspace(4)*)) #4 + ret void +} + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 + +; Function Attrs: argmemonly nounwind willreturn +declare void @llvm.memcpy.p4i8.p0i8.i64(i8 addrspace(4)* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #1 + +; Function Attrs: convergent +declare dso_local spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret align 8, 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" } +attributes #3 = { nounwind } +attributes #4 = { convergent } + +!llvm.module.flags = !{!0} +!opencl.spir.version = !{!1} +!spirv.Source = !{!2} +!llvm.ident = !{!3} + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 1, i32 2} +!2 = !{i32 4, i32 100000} +!3 = !{!"clang version 12.0.0 "} +!4 = !{i32 -1, i32 -1, i32 -1, i32 -1} +!5 = !{i64 0, i64 16, !6, i64 16, i64 8, !6} +!6 = !{!7, !7, i64 0} +!7 = !{!"omnipotent char", !8, i64 0} +!8 = !{!"Simple C++ TBAA"} diff --git a/llvm/test/tools/sycl-post-link/multiple-spec-const-usages.ll b/llvm/test/tools/sycl-post-link/multiple-spec-const-usages.ll new file mode 100644 index 0000000000000..e7ef11cff4653 --- /dev/null +++ b/llvm/test/tools/sycl-post-link/multiple-spec-const-usages.ll @@ -0,0 +1,40 @@ +; This test checks that IDs assigned to spec constants are correct, i.e. if some +; spec constant is accessed twice, then metadata for both accesses should point +; to the same ID + +; RUN: sycl-post-link -spec-const=rt --ir-output-only %s -S -o - \ +; RUN: | FileCheck %s + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" +target triple = "spir64-unknown-unknown-sycldevice" + +%"spec_constant" = type { i8 } + +@SCSymID = private unnamed_addr constant [10 x i8] c"SpecConst\00", align 1 +@SCSymID1 = private unnamed_addr constant [11 x i8] c"SpecConst1\00", align 1 + +declare dso_local spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)*) + +; Function Attrs: norecurse +define weak_odr dso_local spir_kernel void @Kernel() { + %1 = call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*)) +; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0:[0-9]+]] + ret void +} + +; Function Attrs: norecurse +define dso_local spir_func float @foo_float(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 { + %2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([11 x i8], [11 x i8]* @SCSymID1, i64 0, i64 0) to i8 addrspace(4)*)) +; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID1:[0-9]+]] + ret float %2 +} + +; Function Attrs: norecurse +define dso_local spir_func float @foo_float2(%"spec_constant" addrspace(4)* nocapture readnone dereferenceable(1) %0) local_unnamed_addr #3 { + %2 = tail call spir_func float @_Z27__sycl_getSpecConstantValueIfET_PKc(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([10 x i8], [10 x i8]* @SCSymID, i64 0, i64 0) to i8 addrspace(4)*)) +; CHECK: call float @_Z20__spirv_SpecConstantif({{.*}}), !SYCL_SPEC_CONST_SYM_ID ![[ID0]] + ret float %2 +} + +; CHECK: ![[ID0]] = !{!"SpecConst", i32 0} +; CHECK: ![[ID1]] = !{!"SpecConst1", i32 1} diff --git a/llvm/tools/sycl-post-link/SpecConstants.cpp b/llvm/tools/sycl-post-link/SpecConstants.cpp index 35475e02d750c..6bdd4997fe241 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.cpp +++ b/llvm/tools/sycl-post-link/SpecConstants.cpp @@ -26,21 +26,29 @@ namespace { // all mangled names of __sycl_getSpecConstantValue intrinsics, which differ by // the template type parameter and the specialization constant value type. constexpr char SYCL_GET_SPEC_CONST_VAL[] = "_Z27__sycl_getSpecConstantValue"; +constexpr char SYCL_GET_COMPOSITE_SPEC_CONST_VAL[] = + "_Z36__sycl_getCompositeSpecConstantValue"; + // Unmangled base name of all __spirv_SpecConstant intrinsics which differ by // the value type. constexpr char SPIRV_GET_SPEC_CONST_VAL[] = "__spirv_SpecConstant"; +// Unmangled base name of all __spirv_SpecConstantComposite intrinsics which +// differ by the value type. +constexpr char SPIRV_GET_SPEC_CONST_COMPOSITE[] = + "__spirv_SpecConstantComposite"; + // Metadata ID string added to calls to __spirv_SpecConstant to record the -// original symbolic spec constant ID. +// original symbolic spec constant ID. For composite spec constants it contains +// IDs of all scalar spec constants included into a composite constexpr char SPEC_CONST_SYM_ID_MD_STRING[] = "SYCL_SPEC_CONST_SYM_ID"; -static void AssertRelease(bool Cond, const char *Msg) { +void AssertRelease(bool Cond, const char *Msg) { if (!Cond) report_fatal_error((Twine("SpecConstants.cpp: ") + Msg).str().c_str()); } StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, - SmallVectorImpl &DelInsts, - GlobalVariable *&SymGlob) { + SmallVectorImpl &DelInsts) { Value *V = CI->getArgOperand(ArgNo)->stripPointerCasts(); if (auto *L = dyn_cast(V)) { @@ -95,51 +103,82 @@ StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo, V = Store->getValueOperand()->stripPointerCasts(); } const Constant *Init = cast(V)->getInitializer(); - SymGlob = cast(V); StringRef Res = cast(Init)->getAsString(); if (Res.size() > 0 && Res[Res.size() - 1] == '\0') Res = Res.substr(0, Res.size() - 1); return Res; } -// TODO support spec constant types other than integer or -// floating-point. -Value *genDefaultValue(Type *T, Instruction *At) { +Value *getDefaultCPPValue(Type *T) { if (T->isIntegerTy()) - return ConstantInt::get(T, 0); + return Constant::getIntegerValue(T, APInt(T->getScalarSizeInBits(), 0)); if (T->isFloatingPointTy()) return ConstantFP::get(T, 0.0); - llvm_unreachable("non-numeric specialization constants are NYI"); + if (auto *VecTy = dyn_cast(T)) + return ConstantVector::getSplat( + VecTy->getElementCount(), + cast(getDefaultCPPValue(VecTy->getElementType()))); + if (auto *ArrTy = dyn_cast(T)) { + SmallVector Elements( + ArrTy->getNumElements(), + cast(getDefaultCPPValue(ArrTy->getElementType()))); + return ConstantArray::get(ArrTy, Elements); + } + if (auto *StructTy = dyn_cast(T)) { + SmallVector Elements; + for (Type *ElTy : StructTy->elements()) { + Elements.push_back(cast(getDefaultCPPValue(ElTy))); + } + return ConstantStruct::get(StructTy, Elements); + } + llvm_unreachable( + "non-numeric (or composites consisting of non-numeric types) " + "specialization constants are NYI"); return nullptr; } -std::string manglePrimitiveType(Type *T) { +std::string manglePrimitiveType(const Type *T) { if (T->isFloatTy()) return "f"; if (T->isDoubleTy()) return "d"; - assert(T->isIntegerTy() && - "unsupported spec const type, must've been guarded in headers"); - switch (T->getIntegerBitWidth()) { - case 1: - return "b"; - case 8: - return "a"; - case 16: - return "s"; - case 32: - return "i"; - case 64: - return "x"; - default: - llvm_unreachable("unsupported spec const integer type"); + if (T->isIntegerTy()) { + switch (T->getIntegerBitWidth()) { + case 1: + return "b"; + case 8: + return "a"; + case 16: + return "s"; + case 32: + return "i"; + case 64: + return "x"; + default: + llvm_unreachable("unsupported spec const integer type"); + } } + // Mangling, which is generated below is not conformant with C++ ABI rules + // (https://itanium-cxx-abi.github.io/cxx-abi/abi.html#mangle.unqualified-name) + // But it should be more or less okay, because these declarations only + // exists in the module between invocations of sycl-post-link and llvm-spirv, + // llvm-spirv doesn't care about the mangling and the only intent here is to + // make sure that we won't encounter redefinition error when we proceed two + // spec constants with different types. + if (T->isStructTy()) + return T->getStructName().str(); + if (T->isArrayTy()) + return "A" + manglePrimitiveType(T->getArrayElementType()); + if (auto *VecTy = dyn_cast(T)) + return "Dv" + std::to_string(VecTy->getNumElements()) + "_" + + manglePrimitiveType(VecTy->getElementType()); + llvm_unreachable("unsupported spec const type"); return ""; } // This is a very basic mangler which can mangle non-templated and non-member // functions with primitive types in the signature. -std::string mangleFuncItanium(StringRef BaseName, FunctionType *FT) { +std::string mangleFuncItanium(StringRef BaseName, const FunctionType *FT) { std::string Res = (Twine("_Z") + Twine(BaseName.size()) + Twine(BaseName)).str(); for (unsigned I = 0; I < FT->getNumParams(); ++I) @@ -147,16 +186,20 @@ std::string mangleFuncItanium(StringRef BaseName, FunctionType *FT) { return Res; } -void setSpecConstMetadata(Instruction *I, StringRef SymID, int IntID) { +void setSpecConstSymIDMetadata(Instruction *I, StringRef SymID, + ArrayRef IntIDs) { LLVMContext &Ctx = I->getContext(); - MDString *SymV = MDString::get(Ctx, SymID); - ConstantAsMetadata *IntV = - ConstantAsMetadata::get(ConstantInt::get(Ctx, APInt(32, IntID))); - MDNode *Entry = MDNode::get(Ctx, {SymV, IntV}); + SmallVector MDOperands; + MDOperands.push_back(MDString::get(Ctx, SymID)); + for (unsigned ID : IntIDs) + MDOperands.push_back( + ConstantAsMetadata::get(ConstantInt::get(Ctx, APInt(32, ID)))); + MDNode *Entry = MDNode::get(Ctx, MDOperands); I->setMetadata(SPEC_CONST_SYM_ID_MD_STRING, Entry); } -std::pair getSpecConstMetadata(Instruction *I) { +std::pair +getScalarSpecConstMetadata(const Instruction *I) { const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING); if (!N) return std::make_pair("", 0); @@ -167,21 +210,183 @@ std::pair getSpecConstMetadata(Instruction *I) { return std::make_pair(MDSym->getString(), ID); } -static Value *getDefaultCPPValue(Type *T) { - if (T->isIntegerTy()) - return Constant::getIntegerValue(T, APInt(T->getScalarSizeInBits(), 0)); - if (T->isFloatingPointTy()) - return ConstantFP::get(T, 0); - llvm_unreachable("unsupported spec const type"); - return nullptr; +/// Recursively iterates over a composite type in order to collect information +/// about its scalar elements. +void collectCompositeElementsInfoRecursive( + const Type *Ty, unsigned &Index, unsigned &Offset, + std::vector &Result) { + if (auto *ArrTy = dyn_cast(Ty)) { + for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { + // TODO: this is a spot for potential optimization: for arrays we could + // just make a single recursive call here and use it to populate Result + // in a loop. + collectCompositeElementsInfoRecursive(ArrTy->getElementType(), Index, + Offset, Result); + } + } else if (auto *StructTy = dyn_cast(Ty)) { + for (Type *ElTy : StructTy->elements()) { + collectCompositeElementsInfoRecursive(ElTy, Index, Offset, Result); + } + } else if (auto *VecTy = dyn_cast(Ty)) { + for (size_t I = 0; I < VecTy->getNumElements(); ++I) { + // TODO: this is a spot for potential optimization: for vectors we could + // just make a single recursive call here and use it to populate Result + // in a loop. + collectCompositeElementsInfoRecursive(VecTy->getElementType(), Index, + Offset, Result); + } + } else { // Assume that we encountered some scalar element + CompositeSpecConstElementDescriptor Desc; + Desc.ID = 0; // To be filled later + Desc.Offset = Offset; + Desc.Size = Ty->getPrimitiveSizeInBits() / 8; + Result[Index++] = Desc; + Offset += Desc.Size; + } +} + +std::pair> +getCompositeSpecConstMetadata(const Instruction *I) { + const MDNode *N = I->getMetadata(SPEC_CONST_SYM_ID_MD_STRING); + if (!N) + return std::make_pair("", + std::vector{}); + const auto *MDSym = cast(N->getOperand(0)); + + std::vector Result(N->getNumOperands() - + 1); + unsigned Index = 0, Offset = 0; + collectCompositeElementsInfoRecursive(I->getType(), Index, Offset, Result); + + for (unsigned I = 1; I < N->getNumOperands(); ++I) { + const auto *MDInt = cast(N->getOperand(I)); + unsigned ID = static_cast( + cast(MDInt->getValue())->getValue().getZExtValue()); + Result[I - 1].ID = ID; + } + return std::make_pair(MDSym->getString(), Result); +} + +Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, + ArrayRef Args, Instruction *InsertBefore) { + SmallVector ArgTys(Args.size()); + for (unsigned I = 0; I < Args.size(); ++I) { + ArgTys[I] = Args[I]->getType(); + } + auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/); + std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT); + Module *M = InsertBefore->getFunction()->getParent(); + FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT); + assert(FC.getCallee() && "SPIRV intrinsic creation failed"); + auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); + return Call; +} + +Instruction *emitSpecConstant(unsigned NumericID, Type *Ty, + Instruction *InsertBefore) { + Function *F = InsertBefore->getFunction(); + // Generate arguments needed by the SPIRV version of the intrinsic + // - integer constant ID: + Value *ID = ConstantInt::get(Type::getInt32Ty(F->getContext()), NumericID); + // - default value: + Value *Def = getDefaultCPPValue(Ty); + // ... Now replace the call with SPIRV intrinsic version. + Value *Args[] = {ID, Def}; + return emitCall(Ty, SPIRV_GET_SPEC_CONST_VAL, Args, InsertBefore); +} + +Instruction *emitSpecConstantComposite(Type *Ty, + ArrayRef Elements, + Instruction *InsertBefore) { + SmallVector Args(Elements.size()); + for (unsigned I = 0; I < Elements.size(); ++I) { + Args[I] = cast(Elements[I]); + } + return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Args, InsertBefore); +} + +/// For specified specialization constant type emits LLVM IR which is required +/// in order to correctly handle it later during LLVM IR -> SPIR-V translation. +/// +/// @param Ty [in] Specialization constant type to handle. +/// @param InsertBefore [in] Location in the module where new instructions +/// should be inserted. +/// @param IDs [in,out] List of IDs which are assigned for scalar specialization +/// constants. If \c IsNewSpecConstant is true, this vector is expected to +/// contain a single element with ID of the first spec constant - the rest of +/// generated spec constants will have their IDs generated by incrementing that +/// first ID. If \c IsNewSpecConstant is false, this vector is expected to +/// contain enough elements to assign ID to each scalar element encountered in +/// the specified composite type. +/// @param IsNewSpecConstant [in] Flag to specify whether \c IDs vector should +/// be filled with new IDs or it should be used as-is to replicate an existing +/// spec constant +/// @param [in,out] IsFirstElement Flag indicating whether this function is +/// handling the first scalar element encountered in the specified composite +/// type \c Ty or not. +/// +/// @returns Instruction* representing specialization constant in LLVM IR, which +/// is in SPIR-V friendly LLVM IR form. +/// For scalar types it results in a single __spirv_SpecConstant call. +/// For composite types it results in a number of __spirv_SpecConstant calls +/// for each scalar member of the composite plus in a number of +/// __spirvSpecConstantComposite calls for each composite member of the +/// composite (plus for the top-level composite). Also enumerates all +/// encountered scalars and assigns them IDs (or re-uses existing ones). +Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore, + SmallVectorImpl &IDs, + bool IsNewSpecConstant, + bool &IsFirstElement) { + if (!Ty->isArrayTy() && !Ty->isStructTy() && !Ty->isVectorTy()) { // Scalar + if (IsNewSpecConstant && !IsFirstElement) { + // If it is a new specialization constant, we need to generate IDs for + // scalar elements, starting with the second one. + IDs.push_back(IDs.back() + 1); + } + IsFirstElement = false; + return emitSpecConstant(IDs.back(), Ty, InsertBefore); + } + + SmallVector Elements; + auto LoopIteration = [&](Type *Ty) { + Elements.push_back(emitSpecConstantRecursiveImpl( + Ty, InsertBefore, IDs, IsNewSpecConstant, IsFirstElement)); + }; + + if (auto *ArrTy = dyn_cast(Ty)) { + for (size_t I = 0; I < ArrTy->getNumElements(); ++I) { + LoopIteration(ArrTy->getElementType()); + } + } else if (auto *StructTy = dyn_cast(Ty)) { + for (Type *ElTy : StructTy->elements()) { + LoopIteration(ElTy); + } + } else if (auto *VecTy = dyn_cast(Ty)) { + for (size_t I = 0; I < VecTy->getNumElements(); ++I) { + LoopIteration(VecTy->getElementType()); + } + } else { + llvm_unreachable("Unexpected spec constant type"); + } + + return emitSpecConstantComposite(Ty, Elements, InsertBefore); +} + +/// Wrapper intended to hide IsFirstElement argument from the caller +Instruction *emitSpecConstantRecursive(Type *Ty, Instruction *InsertBefore, + SmallVectorImpl &IDs, + bool IsNewSpecConstant) { + bool IsFirstElement = true; + return emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, IsNewSpecConstant, + IsFirstElement); } } // namespace PreservedAnalyses SpecConstantsPass::run(Module &M, ModuleAnalysisManager &MAM) { - int NextID = 0; - StringMap IDMap; + unsigned NextID = 0; + StringMap> IDMap; // Iterate through all declarations of instances of function template // template T __sycl_getSpecConstantValue(const char *ID) @@ -193,15 +398,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, if (!F.isDeclaration()) continue; - if (!F.getName().startswith(SYCL_GET_SPEC_CONST_VAL)) + if (!F.getName().startswith(SYCL_GET_SPEC_CONST_VAL) && + !F.getName().startswith(SYCL_GET_COMPOSITE_SPEC_CONST_VAL)) continue; SmallVector SCIntrCalls; for (auto *U : F.users()) { - auto *CI = dyn_cast(U); - if (!CI) - continue; - SCIntrCalls.push_back(CI); + if (auto *CI = dyn_cast(U)) + SCIntrCalls.push_back(CI); } IRModified = IRModified || (SCIntrCalls.size() > 0); @@ -211,65 +415,96 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, // to the intrinsic - this should always be possible, as only string // 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); + SmallVector DelInsts; DelInsts.push_back(CI); - GlobalVariable *SymGlob = nullptr; - StringRef SymID = getStringLiteralArg(CI, 0, DelInsts, SymGlob); Type *SCTy = CI->getType(); + unsigned NameArgNo = 0; + if (IsComposite) { // structs are returned via sret arguments. + NameArgNo = 1; + auto *PtrTy = cast(CI->getArgOperand(0)->getType()); + SCTy = PtrTy->getElementType(); + } + StringRef SymID = getStringLiteralArg(CI, NameArgNo, DelInsts); if (SetValAtRT) { // 2. Spec constant value will be set at run time - then add the literal - // to a "spec const string literal ID" -> "integer ID" map, uniquing - // the integer ID if this is new literal - auto Ins = IDMap.insert(std::make_pair(SymID, 0)); - if (Ins.second) - Ins.first->second = NextID++; - // 3. Transform to spirv intrinsic _Z*__spirv_SpecConstant*. - LLVMContext &Ctx = F.getContext(); - // Generate arguments needed by the SPIRV version of the intrinsic - // - integer constant ID: - Value *ID = ConstantInt::get(Type::getInt32Ty(Ctx), NextID - 1); - // - default value: - Value *Def = genDefaultValue(SCTy, CI); - // ... Now replace the call with SPIRV intrinsic version. - Value *Args[] = {ID, Def}; - constexpr size_t NArgs = sizeof(Args) / sizeof(Args[0]); - Type *ArgTys[NArgs] = {nullptr}; - for (unsigned int I = 0; I < NArgs; ++I) - ArgTys[I] = Args[I]->getType(); - FunctionType *FT = FunctionType::get(SCTy, ArgTys, false /*isVarArg*/); - Module &M = *F.getParent(); - std::string SPIRVName = mangleFuncItanium(SPIRV_GET_SPEC_CONST_VAL, FT); - FunctionCallee FC = M.getOrInsertFunction(SPIRVName, FT); - assert(FC.getCallee() && "SPIRV intrinsic creation failed"); - CallInst *SPIRVCall = - CallInst::Create(FT, FC.getCallee(), Args, "", CI); - CI->replaceAllUsesWith(SPIRVCall); - // Mark the instruction with pair for later + // to a "spec const string literal ID" -> "integer ID" map or + // "composite spec const string literal ID" -> "vector of integer IDs" + // map, uniquing the integer IDs if this is new literal + auto Ins = + IDMap.insert(std::make_pair(SymID, SmallVector{})); + bool IsNewSpecConstant = Ins.second; + auto &IDs = Ins.first->second; + if (IsNewSpecConstant) { + // For any spec constant type there will be always at least one ID + // generatedA. + IDs.push_back(NextID); + } + + // 3. Transform to spirv intrinsic _Z*__spirv_SpecConstant* or + // _Z*__spirv_SpecConstantComposite + auto *SPIRVCall = + emitSpecConstantRecursive(SCTy, CI, IDs, IsNewSpecConstant); + if (IsNewSpecConstant) { + // emitSpecConstantRecursive might emit more than one spec constant + // (because of composite types) and therefore, we need to ajudst + // NextID according to the actual amount of emitted spec constants. + NextID += IDs.size(); + } + + if (IsComposite) { + // __sycl_getCompositeSpecConstant returns through argument, so, the + // only thing we need to do here is to store into a memory pointed by + // that argument + new StoreInst(SPIRVCall, CI->getArgOperand(0), CI); + } else { + CI->replaceAllUsesWith(SPIRVCall); + } + + // Mark the instruction with list for later // recollection by collectSpecConstantMetadata method. - setSpecConstMetadata(SPIRVCall, SymID, NextID - 1); + setSpecConstSymIDMetadata(SPIRVCall, SymID, IDs); // Example of the emitted call when spec constant is integer: // %6 = call i32 @_Z20__spirv_SpecConstantii(i32 0, i32 0), \ // !SYCL_SPEC_CONST_SYM_ID !22 + // !22 = {!"string-id", i32 0} + // Example of the emitted call when spec constant is vector consisting + // of two integers: + // %1 = call i32 @_Z20__spirv_SpecConstantii(i32 3, i32 0) + // %2 = call i32 @_Z20__spirv_SpecConstantii(i32 4, i32 0) + // %3 = call <2 x i32> @_Z29__spirv_SpecConstantCompositeii(i32 \ + // %1, i32 %2), !SYCL_SPEC_CONST_SYM_ID !23 + // !23 = {!"string-id-2", i32 3, i32 4} } else { // 2a. Spec constant must be resolved at compile time - just replace // the intrinsic with default C++ value for the spec constant type. - CI->replaceAllUsesWith(getDefaultCPPValue(SCTy)); + Value *Default = getDefaultCPPValue(SCTy); + if (IsComposite) { + // __sycl_getCompositeSpecConstant returns through argument, so, the + // only thing we need to do here is to store into a memory pointed by + // that argument + new StoreInst(Default, CI->getArgOperand(0), CI); + } else { + CI->replaceAllUsesWith(Default); + } } + for (auto *I : DelInsts) { assert(I->getNumUses() == 0 && "removing live instruction"); I->removeFromParent(); I->deleteValue(); } - // Don't delete SymGlob here, as it may be referenced from multiple - // functions if __sycl_getSpecConstantValue is inlined. } } return IRModified ? PreservedAnalyses::none() : PreservedAnalyses::all(); } bool SpecConstantsPass::collectSpecConstantMetadata( - Module &M, std::map &IDMap) { + Module &M, ScalarSpecIDMapTy &ScalarIDMap, + CompositeSpecIDMapTy &CompositeIDMap) { bool Met = false; for (Function &F : M) { @@ -282,13 +517,22 @@ bool SpecConstantsPass::collectSpecConstantMetadata( Function *Callee = nullptr; if (!CI || CI->isIndirectCall() || !(Callee = CI->getCalledFunction())) continue; - std::pair Res = getSpecConstMetadata(CI); - if (!Res.first.empty()) { - IDMap[Res.first] = Res.second; - Met = true; + if (Callee->getName().contains(SPIRV_GET_SPEC_CONST_COMPOSITE)) { + auto Res = getCompositeSpecConstMetadata(CI); + if (!Res.first.empty()) { + CompositeIDMap[Res.first] = Res.second; + Met = true; + } + } else if (Callee->getName().contains(SPIRV_GET_SPEC_CONST_VAL)) { + auto Res = getScalarSpecConstMetadata(CI); + if (!Res.first.empty()) { + ScalarIDMap[Res.first] = Res.second; + Met = true; + } } } } + return Met; } diff --git a/llvm/tools/sycl-post-link/SpecConstants.h b/llvm/tools/sycl-post-link/SpecConstants.h index 98dc5e2f73edc..dbb071b2cbd4f 100644 --- a/llvm/tools/sycl-post-link/SpecConstants.h +++ b/llvm/tools/sycl-post-link/SpecConstants.h @@ -24,6 +24,25 @@ using namespace llvm; +using ScalarSpecIDMapTy = std::map; +// Represents an element of a composite speciailization constant - at SYCL RT +// level composite specialization constants are being represented as a single +// byte-array, while at SPIR-V level they are represented by a number of scalar +// specialization constants. +struct CompositeSpecConstElementDescriptor { + // Encodes ID of a scalar specialization constants which is a leaf of some + // composite specialization constant. + unsigned ID; + // Encodes offset from the beginning of composite, where scalar resides, i.e. + // location of the scalar value within a byte-array containing the whole + // composite specialization constant. + unsigned Offset; + // Encodes size of scalar specialization constant. + unsigned Size; +}; +using CompositeSpecIDMapTy = + std::map>; + class SpecConstantsPass : public PassInfoMixin { public: // SetValAtRT parameter controls spec constant lowering mode: @@ -34,9 +53,13 @@ class SpecConstantsPass : public PassInfoMixin { // Searches given module for occurences of specialization constant-specific // metadata at call instructions and builds a - // "spec constant name" -> "spec constant int ID" map from this information. + // "spec constant name" -> "spec constant int ID" map for scalar spec + // constants and + // "spec constant name" -> vector<"spec constant int ID"> map for composite + // spec constants static bool collectSpecConstantMetadata(Module &M, - std::map &IDMap); + ScalarSpecIDMapTy &ScalarIDMap, + CompositeSpecIDMapTy &CompositeIDMap); private: bool SetValAtRT; diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index e359ba494ddc0..0f09f6d5c9329 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -42,7 +42,6 @@ using namespace llvm; using string_vector = std::vector; -using SpecIDMapTy = std::map; cl::OptionCategory PostLinkCat{"sycl-post-link options"}; @@ -509,14 +508,19 @@ static string_vector saveDeviceImageProperty( RMEntry); } if (ImgPSInfo.DoSpecConst && ImgPSInfo.SetSpecConstAtRT) { - // extract spec constant maps per each module - SpecIDMapTy TmpSpecIDMap; - if (ImgPSInfo.SpecConstsMet) - SpecConstantsPass::collectSpecConstantMetadata(*ResultModules[I].get(), - TmpSpecIDMap); - PropSet.add( - llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS, - TmpSpecIDMap); + if (ImgPSInfo.SpecConstsMet) { + // extract spec constant maps per each module + ScalarSpecIDMapTy TmpScalarSpecIDMap; + CompositeSpecIDMapTy TmpCompositeSpecIDMap; + SpecConstantsPass::collectSpecConstantMetadata( + *ResultModules[I].get(), TmpScalarSpecIDMap, TmpCompositeSpecIDMap); + PropSet.add( + llvm::util::PropertySetRegistry::SYCL_SPECIALIZATION_CONSTANTS, + TmpScalarSpecIDMap); + PropSet.add(llvm::util::PropertySetRegistry:: + SYCL_COMPOSITE_SPECIALIZATION_CONSTANTS, + TmpCompositeSpecIDMap); + } } if (ImgPSInfo.EmitKernelParamInfo) { // extract kernel parameter optimization info per module