diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index 6af747fcd0e96..f8ada192b30df 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -28,9 +28,9 @@ sycl/doc/extensions/ @intel/dpcpp-specification-reviewers # Sub-groups sycl/include/CL/sycl/detail/spirv.hpp @Pennycook @AlexeySachkov -sycl/include/CL/sycl/intel/group_algorithm.hpp @Pennycook @AlexeySachkov -sycl/include/CL/sycl/intel/sub_group.hpp @Pennycook @AlexeySachkov -sycl/include/CL/sycl/intel/sub_group_host.hpp @Pennycook @AlexeySachkov +sycl/include/sycl/ext/intel/group_algorithm.hpp @Pennycook @AlexeySachkov +sycl/include/sycl/ext/intel/sub_group.hpp @Pennycook @AlexeySachkov +sycl/include/sycl/ext/intel/sub_group_host.hpp @Pennycook @AlexeySachkov # PI API sycl/include/CL/sycl/detail/pi.def @smaslov-intel @@ -53,17 +53,17 @@ sycl/source/detail/stream_impl.cpp @againull sycl/source/stream.cpp @againull # FPGA extensions -sycl/include/CL/sycl/intel/fpga_device_selector.hpp @MrSidims -sycl/include/CL/sycl/intel/fpga_extensions.hpp @MrSidims -sycl/include/CL/sycl/intel/fpga_reg.hpp @MrSidims -sycl/include/CL/sycl/intel/pipes.hpp @MrSidims +sycl/include/sycl/ext/intel/fpga_device_selector.hpp @MrSidims +sycl/include/sycl/ext/intel/fpga_extensions.hpp @MrSidims +sycl/include/sycl/ext/intel/fpga_reg.hpp @MrSidims +sycl/include/sycl/ext/intel/pipes.hpp @MrSidims sycl/include/CL/sycl/pipes.hpp @MrSidims # Reduction extension -sycl/include/CL/sycl/intel/reduction.hpp @v-klochkov +sycl/include/sycl/ext/intel/reduction.hpp @v-klochkov # Function pointers -sycl/include/CL/sycl/intel/function_pointer.hpp @AlexeySachkov +sycl/include/sycl/ext/intel/function_pointer.hpp @AlexeySachkov sycl/source/function_pointer.cpp @AlexeySachkov # Half Type diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5126475f8aa4c..dccc8a096d388 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -5061,14 +5061,23 @@ bool Util::isSyclHalfType(QualType Ty) { } bool Util::isSyclSpecConstantType(QualType Ty) { - std::array Scopes = { + std::array Scopes = { + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental"), + Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, + "spec_constant")}; + std::array ScopesDeprecated = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ONEAPI"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental"), Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, "spec_constant")}; - return matchQualifiedTypeName(Ty, Scopes); + return matchQualifiedTypeName(Ty, Scopes) || + matchQualifiedTypeName(Ty, ScopesDeprecated); } bool Util::isSyclSpecIdType(QualType Ty) { @@ -5089,7 +5098,16 @@ bool Util::isSyclKernelHandlerType(QualType Ty) { } bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) { - std::array Scopes = { + std::array Scopes = { + Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "ext"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "oneapi"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, "no_alias"}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, + "instance"}}; + std::array ScopesDeprecated = { Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "ONEAPI"}, @@ -5097,11 +5115,21 @@ bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) { Util::DeclContextDesc{Decl::Kind::CXXRecord, "no_alias"}, Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, "instance"}}; - return matchQualifiedTypeName(Ty, Scopes); + return matchQualifiedTypeName(Ty, Scopes) || + matchQualifiedTypeName(Ty, ScopesDeprecated); } bool Util::isSyclBufferLocationType(QualType Ty) { - std::array Scopes = { + std::array Scopes = { + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "intel"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"), + Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "buffer_location"), + Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, + "instance")}; + std::array ScopesDeprecated = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "INTEL"), @@ -5109,7 +5137,8 @@ bool Util::isSyclBufferLocationType(QualType Ty) { Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "buffer_location"), Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, "instance")}; - return matchQualifiedTypeName(Ty, Scopes); + return matchQualifiedTypeName(Ty, Scopes) || + matchQualifiedTypeName(Ty, ScopesDeprecated); } bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) { @@ -5138,13 +5167,21 @@ bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) { } bool Util::isAccessorPropertyListType(QualType Ty) { - std::array Scopes = { + std::array Scopes = { + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"), + Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"), + Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, + "accessor_property_list")}; + std::array ScopesDeprecated = { Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"), Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ONEAPI"), Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization, "accessor_property_list")}; - return matchQualifiedTypeName(Ty, Scopes); + return matchQualifiedTypeName(Ty, Scopes) || + matchQualifiedTypeName(Ty, ScopesDeprecated); } bool Util::matchContext(const DeclContext *Ctx, diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 4ba6e672c308c..ea5d9ee6917a6 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -94,28 +94,34 @@ class property_list { bool operator!=(const property_list &rhs) const { return false; } }; -namespace INTEL { +namespace ext { +namespace intel { namespace property { // Compile time known accessor property struct buffer_location { template class instance {}; }; } // namespace property -} // namespace INTEL +} // namespace intel +} // namespace ext -namespace ONEAPI { +namespace ext { +namespace oneapi { namespace property { // Compile time known accessor property struct no_alias { template class instance {}; }; } // namespace property -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext -namespace ONEAPI { +namespace ext { +namespace oneapi { template class accessor_property_list {}; -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext template struct id { @@ -166,7 +172,7 @@ struct _ImplT { template > + typename propertyListT = ext::oneapi::accessor_property_list<>> class accessor { public: @@ -286,7 +292,8 @@ struct get_kernel_name_t { using name = Type; }; -namespace ONEAPI { +namespace ext { +namespace oneapi { namespace experimental { template class spec_constant { @@ -302,7 +309,8 @@ class spec_constant { } }; } // namespace experimental -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext class kernel_handler { void __init_specialization_constants_buffer(char *specialization_constants_buffer) {} diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index cafe1c5ad496e..3ac03aeb67cd8 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp @@ -72,12 +72,12 @@ int main() { // Check accessors initialization // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 2 // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_FIELD]]) // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to i8 addrspace(4)* // CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8 addrspace(4)* [[BITCAST1]], i64 20 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8 addrspace(4)* [[GEP1]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST2]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST2]]) // CHECK C field initialization // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured addrspace(4)* [[GEP]], i32 0, i32 2 diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp index 682480cf9569f..042edbbd33f9b 100644 --- a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -12,15 +12,15 @@ int main() { cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::oneapi::property::no_alias::instance>> accessorA; cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<1>>> accessorB; cl::sycl::kernel_single_task( diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index 930822b88922c..892f84ef0b1e0 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -10,8 +10,8 @@ struct Base { cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<2>>> AccField; }; @@ -19,8 +19,8 @@ struct Captured : Base, cl::sycl::accessor>> { + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<2>>> { int C; }; @@ -29,8 +29,8 @@ int main() { cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list< + cl::sycl::ext::intel::property::buffer_location::instance<3>>> accessorA; cl::sycl::kernel_single_task( [=]() { diff --git a/clang/test/CodeGenSYCL/int_header_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_spec_const.cpp index d745e43d1b013..a1f10b441aa36 100644 --- a/clang/test/CodeGenSYCL/int_header_spec_const.cpp +++ b/clang/test/CodeGenSYCL/int_header_spec_const.cpp @@ -24,22 +24,22 @@ class MySpecConstantWithinANamespace; int main() { // Create specialization constants. - cl::sycl::ONEAPI::experimental::spec_constant i1(false); - cl::sycl::ONEAPI::experimental::spec_constant i8(0); - cl::sycl::ONEAPI::experimental::spec_constant ui8(0); - cl::sycl::ONEAPI::experimental::spec_constant i16(0); - cl::sycl::ONEAPI::experimental::spec_constant ui16(0); - cl::sycl::ONEAPI::experimental::spec_constant i32(0); + cl::sycl::ext::oneapi::experimental::spec_constant i1(false); + cl::sycl::ext::oneapi::experimental::spec_constant i8(0); + cl::sycl::ext::oneapi::experimental::spec_constant ui8(0); + cl::sycl::ext::oneapi::experimental::spec_constant i16(0); + cl::sycl::ext::oneapi::experimental::spec_constant ui16(0); + cl::sycl::ext::oneapi::experimental::spec_constant i32(0); // Constant used twice, but there must be single entry in the int header, // otherwise compilation error would be issued. - cl::sycl::ONEAPI::experimental::spec_constant i32_1(0); - cl::sycl::ONEAPI::experimental::spec_constant ui32(0); - cl::sycl::ONEAPI::experimental::spec_constant f32(0); - cl::sycl::ONEAPI::experimental::spec_constant f64(0); + cl::sycl::ext::oneapi::experimental::spec_constant i32_1(0); + cl::sycl::ext::oneapi::experimental::spec_constant ui32(0); + cl::sycl::ext::oneapi::experimental::spec_constant f32(0); + cl::sycl::ext::oneapi::experimental::spec_constant f64(0); // Kernel name can be used as a spec constant name - cl::sycl::ONEAPI::experimental::spec_constant spec1(0); + cl::sycl::ext::oneapi::experimental::spec_constant spec1(0); // Spec constant name can be declared within a namespace - cl::sycl::ONEAPI::experimental::spec_constant spec2(0); + cl::sycl::ext::oneapi::experimental::spec_constant spec2(0); double val; double *ptr = &val; // to avoid "unused" warnings diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index ede290317db32..73eaad209889b 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -39,18 +39,22 @@ enum class address_space : int { class property_list {}; -namespace INTEL { +namespace ext { +namespace intel { namespace property { struct buffer_location { template class instance {}; }; } // namespace property -} // namespace INTEL +} // namespace intel +} // namespace ext -namespace ONEAPI { +namespace ext { +namespace oneapi { template class accessor_property_list {}; -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext namespace detail { namespace half_impl { @@ -102,7 +106,7 @@ struct DeviceValueType { template > + typename propertyListT = ext::oneapi::accessor_property_list<>> class accessor { public: @@ -312,12 +316,14 @@ class stream { int FlushBufferSize; }; -namespace ONEAPI { +namespace ext { +namespace oneapi { namespace experimental { template class spec_constant {}; } // namespace experimental -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext } // namespace sycl } // namespace cl diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index b951c59054b89..036684cbc6cb9 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -10,14 +10,14 @@ class another_property_list { }; template -using buffer_location = cl::sycl::INTEL::property::buffer_location::instance; +using buffer_location = cl::sycl::ext::intel::property::buffer_location::instance; struct Base { int A, B; cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list>> AccField; }; @@ -26,7 +26,7 @@ struct Captured cl::sycl::accessor>> { + cl::sycl::ext::oneapi::accessor_property_list>> { int C; }; @@ -37,13 +37,13 @@ int main() { cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 @@ -51,14 +51,14 @@ int main() { cl::sycl::accessor> accessorC; #else cl::sycl::accessor>> + cl::sycl::ext::oneapi::accessor_property_list>> accessorD; cl::sycl::accessor, buffer_location<2>>> accessorF; diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index f56fde742b11f..775248746a3b4 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -21,7 +21,7 @@ struct StructWithSampler { }; struct StructWithSpecConst { - sycl::ONEAPI::experimental::spec_constant SC; + sycl::ext::oneapi::experimental::spec_constant SC; }; sycl::handler H; diff --git a/clang/test/SemaSYCL/spec-const-kernel-arg.cpp b/clang/test/SemaSYCL/spec-const-kernel-arg.cpp index 34f370e6eb2b3..a9ea7ba702ec3 100644 --- a/clang/test/SemaSYCL/spec-const-kernel-arg.cpp +++ b/clang/test/SemaSYCL/spec-const-kernel-arg.cpp @@ -8,12 +8,12 @@ sycl::queue myQueue; struct SpecConstantsWrapper { - sycl::ONEAPI::experimental::spec_constant SC1; - sycl::ONEAPI::experimental::spec_constant SC2; + sycl::ext::oneapi::experimental::spec_constant SC1; + sycl::ext::oneapi::experimental::spec_constant SC2; }; int main() { - sycl::ONEAPI::experimental::spec_constant SC; + sycl::ext::oneapi::experimental::spec_constant SC; SpecConstantsWrapper SCWrapper; myQueue.submit([&](sycl::handler &h) { h.single_task( @@ -27,7 +27,7 @@ int main() { // CHECK: FunctionDecl {{.*}}kernel_sc{{.*}} 'void ()' // CHECK: VarDecl {{.*}}'(lambda at {{.*}}' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}' -// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::ONEAPI::experimental::spec_constant':'sycl::ONEAPI::experimental::spec_constant' +// CHECK-NEXT: CXXConstructExpr {{.*}}'sycl::ext::oneapi::experimental::spec_constant':'sycl::ext::oneapi::experimental::spec_constant' // CHECK-NEXT: InitListExpr {{.*}} 'SpecConstantsWrapper' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ONEAPI::experimental::spec_constant':'sycl::ONEAPI::experimental::spec_constant' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ONEAPI::experimental::spec_constant':'sycl::ONEAPI::experimental::spec_constant' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant':'sycl::ext::oneapi::experimental::spec_constant' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::ext::oneapi::experimental::spec_constant':'sycl::ext::oneapi::experimental::spec_constant' diff --git a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp index 15a4ed6d7dfc6..e3e9040bf5521 100644 --- a/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp +++ b/clang/test/SemaSYCL/spec_const_and_accesor_crash.cpp @@ -10,9 +10,9 @@ __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { } int main() { - cl::sycl::ONEAPI::experimental::spec_constant spec_const; + cl::sycl::ext::oneapi::experimental::spec_constant spec_const; cl::sycl::accessor accessor; - // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::ONEAPI::experimental::spec_constant' + // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::ext::oneapi::experimental::spec_constant' // CHECK: FieldDecl {{.*}} implicit referenced 'cl::sycl::accessor' kernel([spec_const, accessor]() {}); return 0; diff --git a/llvm/include/llvm/SYCLLowerIR/LowerESIMD.h b/llvm/include/llvm/SYCLLowerIR/LowerESIMD.h index e265dcbe366c3..9c13fef9c503f 100644 --- a/llvm/include/llvm/SYCLLowerIR/LowerESIMD.h +++ b/llvm/include/llvm/SYCLLowerIR/LowerESIMD.h @@ -10,8 +10,9 @@ // - ESIMD intrinsics, e.g.: // template -// sycl::intel::gpu::vector_type_t -// __esimd_rdregion(sycl::intel::gpu::vector_type_t Input, +// sycl::ext::intel::experimental::esimd::vector_type_t +// __esimd_rdregion(sycl::ext::intel::experimental::esimd::vector_type_t Input, // uint16_t Offset); //===----------------------------------------------------------------------===// diff --git a/llvm/lib/SYCLLowerIR/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/LowerESIMD.cpp index cdd17ed7427a1..cd619a94f71bd 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMD.cpp @@ -1104,9 +1104,11 @@ static Function *createTestESIMDDeclaration(const ESIMDIntrinDesc &Desc, // // ### Source-level intrinsic: // -// sycl::intel::gpu::__vector_type::type __esimd_flat_read( -// sycl::intel::gpu::__vector_type::type, -// sycl::intel::gpu::__vector_type::type) +// sycl::ext::intel::experimental::esimd::__vector_type::type +// __esimd_flat_read( +// sycl::ext::intel::experimental::esimd::__vector_type::type, sycl::ext::intel::experimental::esimd::__vector_type::type) // // ### Itanium-mangled name: // diff --git a/llvm/test/SYCLLowerIR/esimd_lower_intrins.ll b/llvm/test/SYCLLowerIR/esimd_lower_intrins.ll index dc424a7e4b32f..c0cca3a8b1c67 100644 --- a/llvm/test/SYCLLowerIR/esimd_lower_intrins.ll +++ b/llvm/test/SYCLLowerIR/esimd_lower_intrins.ll @@ -179,7 +179,7 @@ define dso_local spir_func <16 x i32> @FUNC_32() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z14__esimd_uudp4aIjjjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z14__esimd_uudp4aIjjjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.uudp4a.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -191,7 +191,7 @@ define dso_local spir_func <16 x i32> @FUNC_33() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z14__esimd_usdp4aIjiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z14__esimd_usdp4aIjiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.usdp4a.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -203,7 +203,7 @@ define dso_local spir_func <16 x i32> @FUNC_34() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z14__esimd_sudp4aIijjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z14__esimd_sudp4aIijjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.sudp4a.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -215,7 +215,7 @@ define dso_local spir_func <16 x i32> @FUNC_35() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z14__esimd_ssdp4aIiiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z14__esimd_ssdp4aIiiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.ssdp4a.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -227,7 +227,7 @@ define dso_local spir_func <16 x i32> @FUNC_36() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z18__esimd_uudp4a_satIjjjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z18__esimd_uudp4a_satIjjjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.uudp4a.sat.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -239,7 +239,7 @@ define dso_local spir_func <16 x i32> @FUNC_37() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z18__esimd_usdp4a_satIjiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z18__esimd_usdp4a_satIjiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.usdp4a.sat.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -251,7 +251,7 @@ define dso_local spir_func <16 x i32> @FUNC_38() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z18__esimd_sudp4a_satIijjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z18__esimd_sudp4a_satIijjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.sudp4a.sat.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } @@ -263,25 +263,25 @@ define dso_local spir_func <16 x i32> @FUNC_39() { %2 = load <16 x i32>, <16 x i32>* %a_2 %a_3 = alloca <16 x i32> %3 = load <16 x i32>, <16 x i32>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z18__esimd_ssdp4a_satIiiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) + %ret_val = call spir_func <16 x i32> @_Z18__esimd_ssdp4a_satIiiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %1, <16 x i32> %2, <16 x i32> %3) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.ssdp4a.sat.v16i32.v16i32.v16i32.v16i32(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}, <16 x i32> %{{[0-9a-zA-Z_.]+}}) ret <16 x i32> %ret_val } define dso_local spir_func <8 x i32> @FUNC_40() { - %ret_val = call spir_func <8 x i32> @_Z22__esimd_slm_block_readIiLi8EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeEj(i32 0) + %ret_val = call spir_func <8 x i32> @_Z22__esimd_slm_block_readIiLi8EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT0_EE4typeEj(i32 0) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <8 x i32> @llvm.genx.oword.ld.v8i32(i32 0, i32 254, i32 0) ret <8 x i32> %ret_val } define dso_local spir_func void @FUNC_41() { - call spir_func void @_Z16__esimd_sbarrierN2cl4sycl5INTEL3gpu17EsimdSbarrierTypeE(i8 zeroext 1) + call spir_func void @_Z16__esimd_sbarrierN2cl4sycl3ext5intel3gpu17EsimdSbarrierTypeE(i8 zeroext 1) ; CHECK: call void @llvm.genx.sbarrier(i8 1) ret void } define dso_local spir_func void @FUNC_42() { - call spir_func void @_Z16__esimd_sbarrierN2cl4sycl5INTEL3gpu17EsimdSbarrierTypeE(i8 zeroext 0) + call spir_func void @_Z16__esimd_sbarrierN2cl4sycl3ext5intel3gpu17EsimdSbarrierTypeE(i8 zeroext 0) ; CHECK: call void @llvm.genx.sbarrier(i8 0) ret void } @@ -291,7 +291,7 @@ define dso_local spir_func <8 x i32> @FUNC_43() { %1 = load <16 x i32>, <16 x i32>* %a_1 %a_2 = alloca <8 x i16> %2 = load <8 x i16>, <8 x i16>* %a_2 - %ret_val = call spir_func <8 x i32> @_Z18__esimd_rdindirectIiLi16ELi8ELi0EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT1_EE4typeENS4_IS5_XT0_EE4typeENS4_ItXT1_EE4typeE(<16 x i32> %1, <8 x i16> %2) + %ret_val = call spir_func <8 x i32> @_Z18__esimd_rdindirectIiLi16ELi8ELi0EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT1_EE4typeENS4_IS5_XT0_EE4typeENS4_ItXT1_EE4typeE(<16 x i32> %1, <8 x i16> %2) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <8 x i32> @llvm.genx.rdregioni.v8i32.v16i32.v8i16(<16 x i32> %{{[0-9a-zA-Z_.]+}}, i32 0, i32 8, i32 0, <8 x i16> %{{[0-9a-zA-Z_.]+}}, i32 0) ret <8 x i32> %ret_val } @@ -303,7 +303,7 @@ define dso_local spir_func <16 x i32> @FUNC_44() { %2 = load <8 x i32>, <8 x i32>* %a_2 %a_3 = alloca <8 x i16> %3 = load <8 x i16>, <8 x i16>* %a_3 - %ret_val = call spir_func <16 x i32> @_Z18__esimd_wrindirectIiLi16ELi8ELi0EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeES7_NS4_IS5_XT1_EE4typeENS4_ItXT1_EE4typeESB_(<16 x i32> %1, <8 x i32> %2, <8 x i16> %3, <8 x i16> ) + %ret_val = call spir_func <16 x i32> @_Z18__esimd_wrindirectIiLi16ELi8ELi0EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT0_EE4typeES7_NS4_IS5_XT1_EE4typeENS4_ItXT1_EE4typeESB_(<16 x i32> %1, <8 x i32> %2, <8 x i16> %3, <8 x i16> ) ; CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.wrregioni.v16i32.v8i32.v8i16.v8i1(<16 x i32> %{{[0-9a-zA-Z_.]+}}, <8 x i32> %{{[0-9a-zA-Z_.]+}}, i32 0, i32 8, i32 0, <8 x i16> %{{[0-9a-zA-Z_.]+}}, i32 0, <8 x i1> ) ret <16 x i32> %ret_val } @@ -409,18 +409,18 @@ declare dso_local spir_func void @_Z25__esimd_media_block_storeIiLi4ELi8E14ocl_i declare dso_local spir_func <32 x i32> @_Z13__esimd_vloadIiLi32EEN2cm3gen13__vector_typeIT_XT0_EE4typeEPKS5_(<32 x i32> addrspace(4)* %0) declare dso_local spir_func void @_Z14__esimd_vstoreIfLi16EEvPN2cm3gen13__vector_typeIT_XT0_EE4typeES5_(<16 x float> addrspace(4)* %0, <16 x float> %1) declare dso_local spir_func void @_ZN2cl4sycl3ext5intel12experimental5esimd8slm_initEj(i32) -declare dso_local spir_func <16 x i32> @_Z14__esimd_uudp4aIjjjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z14__esimd_usdp4aIjiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z14__esimd_sudp4aIijjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z14__esimd_ssdp4aIiiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z18__esimd_uudp4a_satIjjjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z18__esimd_usdp4a_satIjiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z18__esimd_sudp4a_satIijjjLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <16 x i32> @_Z18__esimd_ssdp4a_satIiiiiLi16EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) -declare dso_local spir_func <8 x i32> @_Z22__esimd_slm_block_readIiLi8EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeEj(i32 %0) -declare dso_local spir_func void @_Z16__esimd_sbarrierN2cl4sycl5INTEL3gpu17EsimdSbarrierTypeE(i8 %0) -declare dso_local spir_func <8 x i32> @_Z18__esimd_rdindirectIiLi16ELi8ELi0EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT1_EE4typeENS4_IS5_XT0_EE4typeENS4_ItXT1_EE4typeE(<16 x i32>, <8 x i16>) -declare dso_local spir_func <16 x i32> @_Z18__esimd_wrindirectIiLi16ELi8ELi0EEN2cl4sycl5INTEL3gpu11vector_typeIT_XT0_EE4typeES7_NS4_IS5_XT1_EE4typeENS4_ItXT1_EE4typeESB_(<16 x i32>, <8 x i32>, <8 x i16>, <8 x i16>) +declare dso_local spir_func <16 x i32> @_Z14__esimd_uudp4aIjjjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z14__esimd_usdp4aIjiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z14__esimd_sudp4aIijjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z14__esimd_ssdp4aIiiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z18__esimd_uudp4a_satIjjjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z18__esimd_usdp4a_satIjiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z18__esimd_sudp4a_satIijjjLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <16 x i32> @_Z18__esimd_ssdp4a_satIiiiiLi16EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT3_EE4typeENS4_IT0_XT3_EE4typeENS4_IT1_XT3_EE4typeENS4_IT2_XT3_EE4typeE(<16 x i32> %0, <16 x i32> %1, <16 x i32> %2) +declare dso_local spir_func <8 x i32> @_Z22__esimd_slm_block_readIiLi8EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT0_EE4typeEj(i32 %0) +declare dso_local spir_func void @_Z16__esimd_sbarrierN2cl4sycl3ext5intel3gpu17EsimdSbarrierTypeE(i8 %0) +declare dso_local spir_func <8 x i32> @_Z18__esimd_rdindirectIiLi16ELi8ELi0EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT1_EE4typeENS4_IS5_XT0_EE4typeENS4_ItXT1_EE4typeE(<16 x i32>, <8 x i16>) +declare dso_local spir_func <16 x i32> @_Z18__esimd_wrindirectIiLi16ELi8ELi0EEN2cl4sycl3ext5intel3gpu11vector_typeIT_XT0_EE4typeES7_NS4_IS5_XT1_EE4typeENS4_ItXT1_EE4typeESB_(<16 x i32>, <8 x i32>, <8 x i16>, <8 x i16>) declare dso_local spir_func <16 x float> @_Z12__esimd_rnddILi16EEN2cl4sycl3ext5intel12experimental5esimd6detail11vector_typeIfXT_EE4typeES9_(<16 x float>) declare dso_local spir_func <16 x float> @_Z12__esimd_rnduILi16EEN2cl4sycl3ext5intel12experimental5esimd6detail11vector_typeIfXT_EE4typeES9_(<16 x float>) declare dso_local spir_func <16 x float> @_Z12__esimd_rndzILi16EEN2cl4sycl3ext5intel12experimental5esimd6detail11vector_typeIfXT_EE4typeES9_(<16 x float>) diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll index 15bc9eb81a2ad..e2a8303d88f9e 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-O0.ll @@ -38,20 +38,20 @@ target triple = "spir64-unknown-unknown-sycldevice" %"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._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlvE_.anon" = type <{ %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant", [7 x i8] }> +%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" = type { %"class._ZTSN2cl4sycl6detail18AccessorImplDeviceILi1EEE.cl::sycl::detail::AccessorImplDevice", %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 } +%union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEUt_E.anon = type { %struct._ZTS3POD.POD addrspace(1)* } +%"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" = type { i8 } +%"class._ZTSN2cl4sycl6detail15accessor_commonI3PODLi1ELNS0_6access4modeE1025ELNS4_6targetE2014ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::detail::accessor_common" = type { i8 } $_ZTS4Test = comdat any -$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev = comdat any +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev = comdat any -$_ZN2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_EC2Ev = comdat any +$_ZN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EC2Ev = comdat any -$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE = comdat any +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE = comdat any $_ZN2cl4sycl2idILi1EEC2Ev = comdat any @@ -65,23 +65,23 @@ $_ZN2cl4sycl5rangeILi1EEC2ILi1EEENSt9enable_ifIXeqT_Li1EEmE4typeE = comdat any $_ZN2cl4sycl6detail5arrayILi1EEixEi = comdat any -$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE9getOffsetEv = comdat any +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE9getOffsetEv = comdat any -$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getAccessRangeEv = comdat any +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getAccessRangeEv = comdat any -$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE14getMemoryRangeEv = comdat any +$_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getMemoryRangeEv = comdat any -$_ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv = comdat any +$_ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv = comdat any -$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE = comdat any +$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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_3ext6oneapi22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE = comdat any -$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEE15getQualifiedPtrEv = comdat any +$_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 +@__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_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 { @@ -93,11 +93,11 @@ entry: %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 + %2 = addrspacecast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor"* %1 to %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* + call spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %4 = addrspacecast %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant"* %3 to %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)* + call spir_func void @_ZN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EC2Ev(%"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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* @@ -109,24 +109,24 @@ entry: %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 + %13 = addrspacecast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor"* %5 to %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* + call spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 { +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = bitcast %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1 to %"class._ZTSN2cl4sycl6detail15accessor_commonI3PODLi1ELNS0_6access4modeE1025ELNS4_6targetE2014ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::detail::accessor_common" addrspace(4)* + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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)* @@ -140,26 +140,26 @@ entry: } ; 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 { +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EC2Ev(%"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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 + %this.addr = alloca %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)* %this, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)*, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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 { +define linkonce_odr dso_local spir_func void @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1S2_NS0_5rangeILi1EEESE_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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)* + %1 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -175,7 +175,7 @@ for.body: ; preds = %for.cond %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 + %call2 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE9getOffsetEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -185,7 +185,7 @@ for.body: ; preds = %for.cond %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 + %call5 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getAccessRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -195,7 +195,7 @@ for.body: ; preds = %for.cond %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 + %call8 = call spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getMemoryRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -213,8 +213,8 @@ for.end: ; preds = %for.cond %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)* + %25 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData11 = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -234,11 +234,11 @@ entry: %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 + call spir_func void @_ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 4 %1, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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 + %call = call spir_func align 4 dereferenceable(20) %struct._ZTS3POD.POD addrspace(4)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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) @@ -341,65 +341,65 @@ entry: } ; 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 { +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE9getOffsetEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 { +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getAccessRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 { +define linkonce_odr dso_local spir_func align 8 dereferenceable(8) %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" addrspace(4)* @_ZN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getMemoryRangeEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %impl = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 { +define linkonce_odr dso_local spir_func void @_ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS8_EE5valueES8_E4typeEv(%struct._ZTS3POD.POD addrspace(4)* noalias sret(%struct._ZTS3POD.POD) align 4 %agg.result, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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 + %this.addr = alloca %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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 + store %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)* %this, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::oneapi::experimental::spec_constant" addrspace(4)*, %"class._ZTSN2cl4sycl3ext6oneapi12experimental13spec_constantI3PODS4_EE.cl::sycl::ext::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._ZNK2cl4sycl3ext6oneapi12experimental13spec_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 { +define linkonce_odr dso_local spir_func align 4 dereferenceable(20) %struct._ZTS3POD.POD addrspace(4)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEixILi1EvEERS2_NS0_2idILi1EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %call = call spir_func i64 @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %call2 = call spir_func %struct._ZTS3POD.POD addrspace(1)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE15getQualifiedPtrEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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)* @@ -424,12 +424,12 @@ entry: 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 { +define linkonce_odr dso_local spir_func i64 @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE14getLinearIndexILi1EEEmNS0_2idIXT_EEE(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 @@ -438,13 +438,13 @@ entry: } ; 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 { +define linkonce_odr dso_local spir_func %struct._ZTS3POD.POD addrspace(1)* @_ZNK2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE15getQualifiedPtrEv(%"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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)* + %this.addr = alloca %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, align 8 + store %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %this1 = load %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)*, %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)** %this.addr, align 8 + %0 = getelementptr inbounds %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor", %"class._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* %this1, i32 0, i32 1 + %MData = bitcast %union._ZTSN2cl4sycl8accessorI3PODLi1ELNS0_6access4modeE1025ELNS3_6targetE2014ELNS3_11placeholderE0ENS0_3ext6oneapi22accessor_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 } diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll index 8821f23d3c957..d8c50e368741d 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll @@ -55,8 +55,8 @@ target triple = "spir64-unknown-unknown-sycldevice" $_ZTS4Test = comdat any $_ZTS17SpecializedKernel = comdat any -@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1 -@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1 +@__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_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._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1 ; Function Attrs: convergent norecurse uwtable define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { @@ -68,7 +68,7 @@ entry: %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(%struct._ZTS3POD.POD) 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 + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_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 @@ -87,7 +87,7 @@ entry: %c.ascast.i = addrspacecast %struct._ZTS1A.A* %c.i to %struct._ZTS1A.A addrspace(4)* %3 = bitcast %struct._ZTS1A.A* %c.i to i8* call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %3) #3 - call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* null) #4 + call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* null) #4 %a.i = getelementptr inbounds %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %c.ascast.i, i64 0, i32 0 %4 = load i32, i32 addrspace(4)* %a.i, align 4 %conv.i = sitofp i32 %4 to float diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value.ll index ab43431b68f59..cdb610d9f5801 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-default-value.ll @@ -26,8 +26,8 @@ target triple = "spir64-unknown-unknown-sycldevice" $_ZTS4Test = comdat any $_ZTS17SpecializedKernel = comdat any -@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI3PODS4_E3getIS4_EENSt9enable_ifIXsr3std6is_podIT_EE5valueES8_E4typeEv = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTS3POD\00", align 1 -@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1 +@__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_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._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv = private unnamed_addr addrspace(1) constant [20 x i8] c"_ZTS13MyComposConst\00", align 1 ; Function Attrs: convergent norecurse uwtable define weak_odr dso_local spir_kernel void @_ZTS4Test(%struct._ZTS3POD.POD addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { @@ -39,7 +39,7 @@ entry: %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(%struct._ZTS3POD.POD) 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 + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueI3PODET_PKc(%struct._ZTS3POD.POD addrspace(4)* sret(%struct._ZTS3POD.POD) align 8 %3, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_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 @@ -62,7 +62,7 @@ entry: %a.i.i = alloca i32, align 4 %bc = bitcast i32* %a.i.i to i8* %tmp = addrspacecast i8* %bc to i8 addrspace(4)* - call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* %tmp) #4 + call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI13MyComposConstET_PKcPvS4_(%struct._ZTS1A.A addrspace(4)* sret(%struct._ZTS1A.A) align 4 %c.ascast.i, i8 addrspace(4)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(4)* addrspacecast ([20 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantI13MyComposConstE3getIS4_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podIS9_EE5valueES9_E4typeEv to [20 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* null, i8 addrspace(4)* %tmp) #4 ; CHECK: %[[GEP:[0-9a-z]+]] = getelementptr i8, i8 addrspace(4)* %tmp, i32 0 ; CHECK: %[[BITCAST:[0-9a-z]+]] = bitcast i8 addrspace(4)* %[[GEP]] to %struct._ZTS1A.A addrspace(4)* ; CHECK: %[[LOAD:[0-9a-z]+]] = load %struct._ZTS1A.A, %struct._ZTS1A.A addrspace(4)* %[[BITCAST]], align 4 diff --git a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll b/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll index 1c530dcc993b2..149d15fcdd4e5 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/multiple-composite-usages-2.ll @@ -28,7 +28,7 @@ $_ZTSN4test8kernel_tIfEE = comdat any $_ZTSN4test8kernel_tIiEE = comdat any -@__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv = private unnamed_addr addrspace(1) constant [18 x i8] c"_ZTS11sc_kernel_t\00", align 1 +@__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv = private unnamed_addr addrspace(1) constant [18 x i8] c"_ZTS11sc_kernel_t\00", align 1 ; Function Attrs: convergent norecurse define weak_odr dso_local spir_kernel void @_ZTSN4test8kernel_tIfEE() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { @@ -37,7 +37,7 @@ entry: %0 = bitcast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to i8* call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0) #3 %1 = addrspacecast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to %"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* - call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0) #3 ret void } @@ -58,7 +58,7 @@ entry: %0 = bitcast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to i8* call void @llvm.lifetime.start.p0i8(i64 8, i8* nonnull %0) #3 %1 = addrspacecast %"struct._ZTSN4test5pod_tE.test::pod_t"* %ref.tmp.i to %"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* - call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl6ONEAPI12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 + call spir_func void @_Z36__sycl_getCompositeSpecConstantValueIN4test5pod_tEET_PKc(%"struct._ZTSN4test5pod_tE.test::pod_t" addrspace(4)* sret(%"struct._ZTSN4test5pod_tE.test::pod_t") align 4 %1, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds ([18 x i8], [18 x i8] addrspace(1)* @__builtin_unique_stable_name._ZNK2cl4sycl3ext6oneapi12experimental13spec_constantIN4test5pod_tE11sc_kernel_tE3getIS5_EENSt9enable_ifIXaasr3std8is_classIT_EE5valuesr3std6is_podISA_EE5valueESA_E4typeEv, i64 0, i64 0) to i8 addrspace(4)*)) #4 call void @llvm.lifetime.end.p0i8(i64 8, i8* nonnull %0) #3 ret void } diff --git a/sycl/doc/extensions/ExplicitSIMD/README.md b/sycl/doc/extensions/ExplicitSIMD/README.md index 3c714f93b83e7..42f7d8cc5eb65 100644 --- a/sycl/doc/extensions/ExplicitSIMD/README.md +++ b/sycl/doc/extensions/ExplicitSIMD/README.md @@ -5,7 +5,7 @@ lower-level Intel GPU programming. It provides APIs closely matching Intel GPU I yet allows to write explicitly vectorized device code. This helps programmer to have more control over the generated code and depend less on compiler optimizations. The [specification](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md), -[documented ESIMD APIs headers](https://github.com/intel/llvm/tree/sycl/sycl/include/CL/sycl/INTEL/esimd) and +[documented ESIMD APIs headers](https://github.com/intel/llvm/tree/sycl/sycl/include/sycl/ext/intel/experimental/esimd) and [working code examples](https://github.com/intel/llvm-test-suite/tree/intel/SYCL/ESIMD) are available on the Intel DPC++ project's github. **_NOTE:_** _This extension is under active development and lots of APIs are diff --git a/sycl/doc/extensions/IntelFPGA/FPGALsu.md b/sycl/doc/extensions/IntelFPGA/FPGALsu.md index c80f6d81723de..5c2ef0a05d914 100644 --- a/sycl/doc/extensions/IntelFPGA/FPGALsu.md +++ b/sycl/doc/extensions/IntelFPGA/FPGALsu.md @@ -1,41 +1,41 @@ # FPGA lsu -The Intel FPGA `lsu` class is implemented in `CL/sycl/INTEL/fpga_lsu.hpp` which -is included in `CL/sycl/INTEL/fpga_extensions.hpp`. +The Intel FPGA `lsu` class is implemented in `sycl/ext/intel/fpga_lsu.hpp` which +is included in `sycl/ext/intel/fpga_extensions.hpp`. -The class `cl::sycl::INTEL::lsu` allows users to explicitly request that the +The class `cl::sycl::ext::intel::lsu` allows users to explicitly request that the implementation of a global memory access is configured in a certain way. The class has two member functions, `load()` and `store()` which allow loading from and storing to a `multi_ptr`, respectively, and is templated on the following 4 optional paremeters: -1. **`cl::sycl::INTEL::burst_coalesce`, where `B` is a boolean**: request, +1. **`cl::sycl::ext::intel::burst_coalesce`, where `B` is a boolean**: request, to the extent possible, that a dynamic burst coalescer be implemented when `load` or `store` are called. The default value of this parameter is `false`. -2. **`cl::sycl::INTEL::cache`, where `N` is an integer greater or equal to +2. **`cl::sycl::ext::intel::cache`, where `N` is an integer greater or equal to 0**: request, to the extent possible, that a read-only cache of the specified size in bytes be implemented when when `load` is called. It is not allowed to use that parameter for `store`. The default value of this parameter is `0`. -3. **`cl::sycl::INTEL::statically_coalesce`, where `B` is a boolean**: +3. **`cl::sycl::ext::intel::statically_coalesce`, where `B` is a boolean**: request, to the extent possible, that `load` or `store` accesses, is allowed to be statically coalesced with other memory accesses at compile time. The default value of this parameter is `true`. -4. **`cl::sycl::INTEL::prefetch`, where `B` is a boolean**: request, to the +4. **`cl::sycl::ext::intel::prefetch`, where `B` is a boolean**: request, to the extent possible, that a prefetcher be implemented when `load` is called. It is not allowed to use that parameter for `store`. The default value of this parameter is `false`. Currently, not every combination of parameters is allowed due to limitations in the backend. The following rules apply: -1. For `store`, `cl::sycl::INTEL::cache` must be `0` and -`cl::sycl::INTEL::prefetch` must be `false`. -2. For `load`, if `cl::sycl::INTEL::cache` is set to a value greater than `0`, -then `cl::sycl::INTEL::burst_coalesce` must be set to `true`. -3. For `load`, exactly one of `cl::sycl::INTEL::prefetch` and -`cl::sycl::INTEL::burst_coalesce` is allowed to be `true`. -4. For `load`, exactly one of `cl::sycl::INTEL::prefetch` and -`cl::sycl::INTEL::cache` is allowed to be `true`. +1. For `store`, `cl::sycl::ext::intel::cache` must be `0` and +`cl::sycl::ext::intel::prefetch` must be `false`. +2. For `load`, if `cl::sycl::ext::intel::cache` is set to a value greater than `0`, +then `cl::sycl::ext::intel::burst_coalesce` must be set to `true`. +3. For `load`, exactly one of `cl::sycl::ext::intel::prefetch` and +`cl::sycl::ext::intel::burst_coalesce` is allowed to be `true`. +4. For `load`, exactly one of `cl::sycl::ext::intel::prefetch` and +`cl::sycl::ext::intel::cache` is allowed to be `true`. ## Implementation @@ -81,7 +81,7 @@ public: ## Usage ```c++ -#include +#include ... cl::sycl::buffer output_buffer(output_data, 1); cl::sycl::buffer input_buffer(input_data, 1); @@ -95,19 +95,19 @@ Queue.submit([&](cl::sycl::handler &cgh) { auto output_ptr = output_accessor.get_pointer(); using PrefetchingLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::statically_coalesce>; + cl::sycl::ext::intel::lsu, + cl::sycl::ext::intel::statically_coalesce>; using BurstCoalescedLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::statically_coalesce>; + cl::sycl::ext::intel::lsu, + cl::sycl::ext::intel::statically_coalesce>; using CachingLSU = - cl::sycl::INTEL::lsu, - cl::sycl::INTEL::cache<1024>, - cl::sycl::INTEL::statically_coalesce>; + cl::sycl::ext::intel::lsu, + cl::sycl::ext::intel::cache<1024>, + cl::sycl::ext::intel::statically_coalesce>; - using PipelinedLSU = cl::sycl::INTEL::lsu<>; + using PipelinedLSU = cl::sycl::ext::intel::lsu<>; int X = PrefetchingLSU::load(input_ptr); // int X = input_ptr[0] int Y = CachingLSU::load(input_ptr + 1); // int Y = input_ptr[1] diff --git a/sycl/doc/extensions/IntelFPGA/FPGASelector.md b/sycl/doc/extensions/IntelFPGA/FPGASelector.md index bf6ee444d1707..208117a878cba 100644 --- a/sycl/doc/extensions/IntelFPGA/FPGASelector.md +++ b/sycl/doc/extensions/IntelFPGA/FPGASelector.md @@ -1,6 +1,6 @@ # FPGA selector -Intel FPGA users can use header file: `#include` to simplify their code +Intel FPGA users can use header file: `#include` to simplify their code when they want to specify FPGA hardware device or FPGA emulation device. ## Implementation @@ -10,18 +10,18 @@ one FPGA board installed in their system (one device per platform). ## Usage: select FPGA hardware device ```c++ -#include +#include ... // force FPGA hardware device -cl::sycl::queue deviceQueue{cl::sycl::INTEL::fpga_selector{}}; +cl::sycl::queue deviceQueue{cl::sycl::ext::intel::fpga_selector{}}; ... ``` ## Usage: select FPGA emulator device ```c++ -#include +#include ... // force FPGA emulation device -cl::sycl::queue deviceQueue{cl::sycl::INTEL::fpga_emulator_selector{}}; +cl::sycl::queue deviceQueue{cl::sycl::ext::intel::fpga_emulator_selector{}}; ... ``` diff --git a/sycl/doc/extensions/OnlineCompilation/OnlineCompilation.asciidoc b/sycl/doc/extensions/OnlineCompilation/OnlineCompilation.asciidoc index b31307a3b2ac4..1491e0e3c2a39 100644 --- a/sycl/doc/extensions/OnlineCompilation/OnlineCompilation.asciidoc +++ b/sycl/doc/extensions/OnlineCompilation/OnlineCompilation.asciidoc @@ -64,7 +64,7 @@ public: === Compiler API -To compile a source, a user program must first construct an instance of the `sycl::INTEL::online_compiler` class. Then pass the source as an `std::string` object to online compiler's `compile` function along with other relevant parameters. The `online_compiler` is templated by the source language, and the `compile` function is a variadic template function. Instantiations of the `online_compiler::compile` for different languages may have different sets of formal parameters. The `compile` function returns a binary blob - an `std::vector` - with the device code compiled according to the compilation target specification provided at online compiler construction time. +To compile a source, a user program must first construct an instance of the `sycl::ext::intel::online_compiler` class. Then pass the source as an `std::string` object to online compiler's `compile` function along with other relevant parameters. The `online_compiler` is templated by the source language, and the `compile` function is a variadic template function. Instantiations of the `online_compiler::compile` for different languages may have different sets of formal parameters. The `compile` function returns a binary blob - an `std::vector` - with the device code compiled according to the compilation target specification provided at online compiler construction time. ==== Online compiler [source,c++] @@ -97,12 +97,12 @@ template class online_compiler; |=== Online compiler construction or source compilation may be unsuccessful, in which case an instance -of `sycl::INTEL::online_compile_error` is thrown. For example, when some of the compilation +of `sycl::ext::intel::online_compile_error` is thrown. For example, when some of the compilation target specification elements are not supported by the implementation, or there is a syntax error in the source program. -==== `sycl::INTEL::online_compiler` constructors. +==== `sycl::ext::intel::online_compiler` constructors. [cols="40,60",options="header"] |=== |Constructor |Description @@ -164,7 +164,7 @@ std::vector online_compiler::compile( This example compiles an OpenCL source to a generic SPIR-V. [source,c++] ----------------- -#include "CL/sycl/INTEL/online_compiler.hpp" +#include "sycl/ext/intel/online_compiler.hpp" #include #include diff --git a/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc b/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc index c34c4581d14d3..64471166407bf 100755 --- a/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc +++ b/sycl/doc/extensions/SYCL_ONEAPI_dot_accumulate.asciidoc @@ -105,7 +105,8 @@ The value that is returned is equivalent to + ---- namespace cl { namespace sycl { -namespace ONEAPI { +namespace ext { +namespace oneapi { int32_t dot_acc(vec a, vec b, int32_t c); int32_t dot_acc(vec a, vec b, int32_t c); @@ -117,7 +118,8 @@ int32_t dot_acc(int32_t a, uint32_t b, int32_t c); int32_t dot_acc(uint32_t a, int32_t b, int32_t c); int32_t dot_acc(uint32_t a, uint32_t b, int32_t c); -} // ONEAPI +} // oneapi +} // ext } // sycl } // cl ---- diff --git a/sycl/doc/extensions/SpecConstants/README.md b/sycl/doc/extensions/SpecConstants/README.md index 0c4158c12b254..1ee9a50f2b6cf 100644 --- a/sycl/doc/extensions/SpecConstants/README.md +++ b/sycl/doc/extensions/SpecConstants/README.md @@ -24,9 +24,9 @@ kernel: for (int i = 0; i < n_sc_sets; i++) { cl::sycl::program program(q.get_context()); const int *sc_set = &sc_vals[i][0]; - cl::sycl::ONEAPI::experimental::spec_constant sc0 = + cl::sycl::ext::oneapi::experimental::spec_constant sc0 = program.set_spec_constant(sc_set[0]); - cl::sycl::ONEAPI::experimental::spec_constant sc1 = + cl::sycl::ext::oneapi::experimental::spec_constant sc1 = program.set_spec_constant(sc_set[1]); program.build_with_kernel_type(); diff --git a/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc index 8ce3858513d95..d92a228aeba7a 100644 --- a/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc +++ b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_buffer_location.asciidoc @@ -57,7 +57,8 @@ Add a new property and global variable to the listing: ```c++ namespace sycl { ... -namespace INTEL { +namespace ext { +namespace intel { namespace property { struct buffer_location { template @@ -66,7 +67,8 @@ namespace property { } // namespace property template inline constexpr /* buffer location instance class */ buffer_location; -} // namespace INTEL +} // namespace intel +} // namespace ext } // namespace sycl ``` diff --git a/sycl/doc/extensions/accessor_properties/SYCL_ONEAPI_accessor_properties.asciidoc b/sycl/doc/extensions/accessor_properties/SYCL_ONEAPI_accessor_properties.asciidoc index 5116be7fdacec..8d091223418fc 100644 --- a/sycl/doc/extensions/accessor_properties/SYCL_ONEAPI_accessor_properties.asciidoc +++ b/sycl/doc/extensions/accessor_properties/SYCL_ONEAPI_accessor_properties.asciidoc @@ -126,7 +126,7 @@ Replace the second paragraph with: A synopsis of the common properties interface, the SYCL property_list class, the SYCL accessor_property_list class template and the SYCL property classes is provided below. The member functions of the common properties interface are listed in Table 4.6. The constructors of the SYCL property_list class are listed in Table 4.7. -Free functions in the ONEAPI::property namespace are listed in Table 4.7a. +Free functions in the ext::oneapi::property namespace are listed in Table 4.7a. The constructors of the accessor_property_list class are listed in Table 4.7b. Two accessor_property_lists which were created from the same set of compile-time-constant properties must either have an identical type or be convertible regardless of their runtime properties. Accessor property lists with no compile-time-constant properties must be convertible to property_lists and vice versa. @@ -167,14 +167,15 @@ class property_list { // Available only when propertyTN contains no compile-time-constant properties template - operator ONEAPI::accessor_property_list() const; + operator ext::oneapi::accessor_property_list() const; }; ``` At the end of the listing add the new property utility struct, equality and inequality operators for compile-time-constant properties and the accessor_property_list class template as follows: ```c++ -namespace ONEAPI { +namespace ext { +namespace oneapi { // New struct for querying whether a class represents a compile-time-constant property template @@ -197,7 +198,8 @@ class accessor_property_list { // Available only when properties... contains no compile-time-constant properties operator property_list() const; }; -} // namespace ONEAPI +} // namespace oneapi +} // namespace ext ``` NOTE: Implementations will need either a conversion function or a deduction guide to satisfy the requirement that accessor_property_lists which were created from the same set of compile-time-constant properties must either have an identical type or be convertible. @@ -252,7 +254,7 @@ Available only if propertyT is a compile-time-constant property. |=== -- -Add a new table, Table 4.7a: Free functions in the ONEAPI::property namespace as follows: +Add a new table, Table 4.7a: Free functions in the ext::oneapi::property namespace as follows: -- [options="header"] @@ -342,88 +344,88 @@ that instead take an accessor_property_list: /* Available only when: (dimensions == 0) */ template > accessor(buffer &bufferRef, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions == 0) */ template accessor(buffer &bufferRef, - handler &commandGroupHandlerRef, const ONEAPI::accessor_property_list &propList = {}); + handler &commandGroupHandlerRef, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, TagT tag, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, - handler &commandGroupHandlerRef, const ONEAPI::accessor_property_list &propList = {}); + handler &commandGroupHandlerRef, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, handler &commandGroupHandlerRef, TagT tag, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, - range accessRange, const ONEAPI::accessor_property_list &propList = {}); + range accessRange, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, range accessRange, TagT tag, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, range accessRange, id accessOffset, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, range accessRange, id accessOffset, - TagT tag, const ONEAPI::accessor_property_list &propList = {}); + TagT tag, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, handler &commandGroupHandlerRef, range accessRange, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, handler &commandGroupHandlerRef, range accessRange, - TagT tag, const ONEAPI::accessor_property_list &propList = {}); + TagT tag, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, handler &commandGroupHandlerRef, range accessRange, - id accessOffset, const ONEAPI::accessor_property_list &propList = {}); + id accessOffset, const ext::oneapi::accessor_property_list &propList = {}); /* Available only when: (dimensions > 0) */ template accessor(buffer &bufferRef, handler &commandGroupHandlerRef, range accessRange, id accessOffset, TagT tag, - const ONEAPI::accessor_property_list &propList = {}); + const ext::oneapi::accessor_property_list &propList = {}); ``` Apply the same additions to the accessor constructors in Table 4.48: Constructors of the accessor class template buffer specialization. NOTE: Oddly enough, due to the rules in section 4.7.6.3 about deduction guides this extension doesn't need to explicitly list the new deduction guides that it may require. Readers may find that confusing given that deduction guides are explicitly listed for other classes, but that's how the spec is written. -The deduction guides will need to ensure that property_listT is inferred to be ONEAPI::accessor_property_list. +The deduction guides will need to ensure that property_listT is inferred to be ext::oneapi::accessor_property_list. Also add to the listing a conversion function: @@ -452,7 +454,8 @@ operator accessor -#include -#include -#include -#include -#include -#include -#include #include #include #include @@ -59,3 +51,11 @@ #include #include #include +#include +#include +#include +#include +#include +#include +#include +#include diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h b/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h index 9fcde11e6e9d4..1fb5d040f5626 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h @@ -12,35 +12,12 @@ #pragma once -// for 'uint32_t' type is included in upper-level device -// interface file ('esimdcpu_device_interface.hpp') +#include -// This file defines function interfaces for ESIMD CPU Emulation -// (ESIMD_CPU) to access LibCM CPU emulation functionalities from -// kernel applications running under emulation +__SYCL_WARNING( + "CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h usage is " + "deprecated, include " + "sycl/ext/intel/experimental/esimd/emu/detail/esimd_emu_functions_v1.h " + "instead") -// CM CPU Emulation Info : -// https://github.com/intel/cm-cpu-emulation - -// Function pointers (*_ptr) with 'cm/__cm' prefixes correspond to -// LibCM functions with same name -// e.g.: cm_fence_ptr -> cm_fence() in LibCM - -// Function pointers (*_ptr) with 'sycl_' prefix correspond to LibCM -// functions dedicated to SYCL support -// e.g.: sycl_get_surface_base_addr_ptr -// -> get_surface_base_addr(int) in LibCM - -/****** DO NOT MODIFY following function pointers ******/ -/****** No reordering, No renaming, No removal ******/ - -// Intrinsics -void (*cm_barrier_ptr)(void); -void (*cm_sbarrier_ptr)(uint32_t); -void (*cm_fence_ptr)(void); - -// libcm functionalities used for intrinsics such as -// surface/buffer/slm access -char *(*sycl_get_surface_base_addr_ptr)(int); -char *(*__cm_emu_get_slm_ptr)(void); -void (*cm_slm_init_ptr)(size_t); +#include diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp index eb249f7c61781..306a0669a2b1e 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp @@ -18,102 +18,12 @@ #pragma once -#include +#include -// cstdint-type fields such as 'uint32_t' are to be used in funtion -// pointer table file ('esimd_emu_functions_v1.h') included in 'struct -// ESIMDDeviceInterface' definition. -#include +__SYCL_WARNING( + "CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp usage is " + "deprecated, include " + "sycl/ext/intel/experimental/esimd/emu/detail/" + "esimdcpu_device_interface.hpp instead") -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace detail { - -/// This is the device interface version required (and used) by this -/// implementation of the ESIMD CPU emulator. -#define ESIMD_DEVICE_INTERFACE_VERSION 1 - -// 'ESIMDDeviceInterface' structure defines interface for ESIMD CPU -// emulation (ESIMD_CPU) to access LibCM CPU emulation functionalities -// from kernel application under emulation. - -// Header files included in the structure contains only function -// pointers to access CM functionalities. Only new function can be -// added - reordering, changing, or removing existing function pointer -// is not allowed. - -// Whenever a new function(s) is added to this interface, a new header -// file must be added following naming convention that contains -// version number such as 'v1' from 'ESIMD_DEVICE_INTERFACE_VERSION'. -struct ESIMDDeviceInterface { - uintptr_t version; - void *reserved; - - ESIMDDeviceInterface(); -#include "esimd_emu_functions_v1.h" -}; - -// Denotes the data version used by the implementation. -// Increment whenever the 'data' field interpretation within PluginOpaqueData is -// changed. -#define ESIMD_EMU_PLUGIN_OPAQUE_DATA_VERSION 0 -/// This structure denotes a ESIMD EMU plugin-specific data returned via the -/// piextPluginGetOpaqueData PI call. Depending on the \c version field, the -/// second \c data field can be interpreted differently. -struct ESIMDEmuPluginOpaqueData { - uintptr_t version; - void *data; -}; -// The table below shows the correspondence between the \c version -// and the contents of the \c data field: -// version == 0, data is ESIMDDeviceInterface* - -ESIMDDeviceInterface *getESIMDDeviceInterface() { - // TODO (performance) cache the interface pointer, can make a difference - // when calling fine-grained libCM APIs through it (like memory access in a - // tight loop) - void *PIOpaqueData = nullptr; - - PIOpaqueData = getPluginOpaqueData(nullptr); - - ESIMDEmuPluginOpaqueData *OpaqueData = - reinterpret_cast(PIOpaqueData); - - // First check if opaque data version is compatible. - if (OpaqueData->version != ESIMD_EMU_PLUGIN_OPAQUE_DATA_VERSION) { - // NOTE: the version check should always be '!=' as layouts of different - // versions of PluginOpaqueData is not backward compatible, unlike - // layout of the ESIMDDeviceInterface. - - std::cerr << __FUNCTION__ << std::endl - << "Opaque data returned by ESIMD Emu plugin is incompatible with" - << "the one used in current implementation." << std::endl - << "Returned version : " << OpaqueData->version << std::endl - << "Required version : " << ESIMD_EMU_PLUGIN_OPAQUE_DATA_VERSION - << std::endl; - throw cl::sycl::feature_not_supported(); - } - // Opaque data version is OK, can cast the 'data' field. - ESIMDDeviceInterface *Interface = - reinterpret_cast(OpaqueData->data); - - // Now check that device interface version is compatible. - if (Interface->version < ESIMD_DEVICE_INTERFACE_VERSION) { - std::cerr << __FUNCTION__ << std::endl - << "The device interface version provided from plug-in " - << "library is behind required device interface version" - << std::endl - << "Found version : " << Interface->version << std::endl - << "Required version :" << ESIMD_DEVICE_INTERFACE_VERSION - << std::endl; - throw cl::sycl::feature_not_supported(); - } - return Interface; -} - -#undef ESIMD_DEVICE_INTERFACE_VERSION -#undef ESIMD_EMU_PLUGIN_OPAQUE_DATA_VERSION - -} // namespace detail -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/INTEL/fpga_device_selector.hpp b/sycl/include/CL/sycl/INTEL/fpga_device_selector.hpp index 3218cf9c9e427..63a1e88ea5613 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_device_selector.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_device_selector.hpp @@ -8,45 +8,10 @@ #pragma once -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { +__SYCL_WARNING( + "CL/sycl/INTEL/fpga_device_selector.hpp usage is deprecated, include " + "sycl/ext/intel/fpga_device_selector.hpp instead") -class platform_selector : public device_selector { -private: - std::string device_platform_name; - -public: - platform_selector(const std::string &platform_name) - : device_platform_name(platform_name) {} - - int operator()(const device &device) const override { - const platform &pf = device.get_platform(); - const std::string &platform_name = pf.get_info(); - if (platform_name == device_platform_name) { - return 10000; - } - return -1; - } -}; - -static constexpr auto EMULATION_PLATFORM_NAME = - "Intel(R) FPGA Emulation Platform for OpenCL(TM)"; -static constexpr auto HARDWARE_PLATFORM_NAME = - "Intel(R) FPGA SDK for OpenCL(TM)"; - -class fpga_selector : public platform_selector { -public: - fpga_selector() : platform_selector(HARDWARE_PLATFORM_NAME) {} -}; - -class fpga_emulator_selector : public platform_selector { -public: - fpga_emulator_selector() : platform_selector(EMULATION_PLATFORM_NAME) {} -}; - -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/INTEL/fpga_extensions.hpp b/sycl/include/CL/sycl/INTEL/fpga_extensions.hpp index c2021fcfe7658..696dc1416872a 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_extensions.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_extensions.hpp @@ -7,7 +7,10 @@ //===----------------------------------------------------------------------===// #pragma once -#include -#include -#include -#include + +#include + +__SYCL_WARNING("CL/sycl/INTEL/fpga_extensions.hpp usage is deprecated, include " + "sycl/ext/intel/fpga_extensions.hpp instead") + +#include diff --git a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp index ab2c859603bf5..baa34398ebef2 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_lsu.hpp @@ -7,122 +7,9 @@ //===----------------------------------------------------------------------===// #pragma once -#include "fpga_utils.hpp" -#include -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { -constexpr uint8_t BURST_COALESCE = 0x1; -constexpr uint8_t CACHE = 0x2; -constexpr uint8_t STATICALLY_COALESCE = 0x4; -constexpr uint8_t PREFETCH = 0x8; +__SYCL_WARNING("CL/sycl/INTEL/fpga_lsu.hpp usage is deprecated, include " + "sycl/ext/intel/fpga_lsu.hpp instead") -struct burst_coalesce_impl_id; -template -struct burst_coalesce_impl : std::integral_constant { - using type_id = burst_coalesce_impl_id; -}; - -struct cache_id; -template struct cache : std::integral_constant { - using type_id = cache_id; -}; - -struct prefetch_impl_id; -template -struct prefetch_impl : std::integral_constant { - using type_id = prefetch_impl_id; -}; - -struct statically_coalesce_impl_id; -template -struct statically_coalesce_impl : std::integral_constant { - using type_id = statically_coalesce_impl_id; -}; - -template using burst_coalesce = burst_coalesce_impl<_B>; -template using prefetch = prefetch_impl<_B>; -template using statically_coalesce = statically_coalesce_impl<_B>; - -template class lsu final { -public: - lsu() = delete; - - template - static _T load(sycl::multi_ptr<_T, _space> Ptr) { - check_space<_space>(); - check_load(); -#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem) - return *__builtin_intel_fpga_mem((_T *)Ptr, - _burst_coalesce | _cache | - _dont_statically_coalesce | _prefetch, - _cache_val); -#else - return *Ptr; -#endif - } - - template - static void store(sycl::multi_ptr<_T, _space> Ptr, _T Val) { - check_space<_space>(); - check_store(); -#if defined(__SYCL_DEVICE_ONLY__) && __has_builtin(__builtin_intel_fpga_mem) - *__builtin_intel_fpga_mem((_T *)Ptr, - _burst_coalesce | _cache | - _dont_statically_coalesce | _prefetch, - _cache_val) = Val; -#else - *Ptr = Val; -#endif - } - -private: - static constexpr int32_t _burst_coalesce_val = - _GetValue, _mem_access_params...>::value; - static constexpr uint8_t _burst_coalesce = - _burst_coalesce_val == 1 ? BURST_COALESCE : 0; - - static constexpr int32_t _cache_val = - _GetValue, _mem_access_params...>::value; - static constexpr uint8_t _cache = (_cache_val > 0) ? CACHE : 0; - - static constexpr int32_t _statically_coalesce_val = - _GetValue, _mem_access_params...>::value; - static constexpr uint8_t _dont_statically_coalesce = - _statically_coalesce_val == 0 ? STATICALLY_COALESCE : 0; - - static constexpr int32_t _prefetch_val = - _GetValue, _mem_access_params...>::value; - static constexpr uint8_t _prefetch = _prefetch_val ? PREFETCH : 0; - - static_assert(_cache_val >= 0, "cache size parameter must be non-negative"); - - template static void check_space() { - static_assert(_space == access::address_space::global_space || - _space == access::address_space::global_device_space || - _space == access::address_space::global_host_space, - "lsu controls are only supported for global_ptr, " - "device_ptr, and host_ptr objects"); - } - - static void check_load() { - static_assert(_cache == 0 || _burst_coalesce == BURST_COALESCE, - "unable to implement a cache without a burst coalescer"); - static_assert(_prefetch == 0 || _burst_coalesce == 0, - "unable to implement a prefetcher and a burst coalescer " - "simulataneously"); - static_assert( - _prefetch == 0 || _cache == 0, - "unable to implement a prefetcher and a cache simulataneously"); - } - static void check_store() { - static_assert(_cache == 0, "unable to implement a store LSU with a cache."); - static_assert(_prefetch == 0, - "unable to implement a store LSU with a prefetcher."); - } -}; -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp index 6142d0cd1996d..69f5f3a3a1fc8 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_reg.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_reg.hpp @@ -8,49 +8,9 @@ #pragma once -#include -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { +__SYCL_WARNING("CL/sycl/INTEL/fpga_reg.hpp usage is deprecated, include " + "sycl/ext/intel/fpga_reg.hpp instead") -// Returns a registered copy of the input -// This function is intended for FPGA users to instruct the compiler to insert -// at least one register stage between the input and the return value. -template -typename std::enable_if::value, _T>::type -fpga_reg(_T t) { -#if __has_builtin(__builtin_intel_fpga_reg) - return __builtin_intel_fpga_reg(t); -#else - return t; -#endif -} - -template -[[deprecated("INTEL::fpga_reg will only support trivially_copyable types in a " - "future release. The type used here will be disallowed.")]] -typename std::enable_if::value == false, - _T>::type -fpga_reg(_T t) { -#if __has_builtin(__builtin_intel_fpga_reg) - return __builtin_intel_fpga_reg(t); -#else - return t; -#endif -} - -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) - -// Keep it consistent with FPGA attributes like intelfpga::memory() -// Currently clang does not support nested namespace for attributes -namespace intelfpga { -template -[[deprecated("intelfpga::fpga_reg will be removed in a future release.")]] _T -fpga_reg(const _T &t) { - return cl::sycl::INTEL::fpga_reg(t); -} -} // namespace intelfpga +#include diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp index a9033242e9790..b03653fdfe18f 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp @@ -8,34 +8,9 @@ #pragma once -#include -#include -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { +__SYCL_WARNING("CL/sycl/INTEL/fpga_utils.hpp usage is deprecated, include " + "sycl/ext/intel/fpga_utils.hpp instead") -template -struct _MatchType : std::is_same {}; - -template struct _GetValue; - -template -struct _GetValue<_D> : std::integral_constant { -}; - -template struct _GetValue<_D, _T1, _T...> { - template - struct impl : std::integral_constant::value> {}; - - template - struct impl<_D2, _T12, std::enable_if_t<_MatchType<_D2, _T12>::value>> - : std::integral_constant {}; - - static constexpr auto value = impl<_D, _T1>::value; -}; -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index d671c0c57f4b4..30577aba66456 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -8,212 +8,9 @@ #pragma once -#include // for __SYCL_INLINE_NAMESPACE -#include // for __SYCL_EXPORT -#include +#include -#include -#include +__SYCL_WARNING("CL/sycl/INTEL/online_compiler.hpp usage is deprecated, include " + "sycl/ext/intel/online_compiler.hpp instead") -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { - -using byte = unsigned char; - -enum class compiled_code_format { - spir_v = 0 // the only format supported for now -}; - -class device_arch { -public: - static constexpr int any = 0; - - device_arch(int Val) : Val(Val) {} - - enum gpu { - gpu_any = 1, - gpu_gen9 = 2, - gpu_skl = gpu_gen9, - gpu_gen9_5 = 3, - gpu_kbl = gpu_gen9_5, - gpu_cfl = gpu_gen9_5, - gpu_gen11 = 4, - gpu_icl = gpu_gen11, - gpu_gen12 = 5 - }; - - enum cpu { - cpu_any = 1, - }; - - enum fpga { - fpga_any = 1, - }; - - operator int() { return Val; } - -private: - int Val; -}; - -/// Represents an error happend during online compilation. -class online_compile_error : public sycl::exception { -public: - online_compile_error() = default; - online_compile_error(const std::string &Msg) : sycl::exception(Msg) {} -}; - -/// Designates a source language for the online compiler. -enum class source_language { opencl_c = 0, cm = 1 }; - -/// Represents an online compiler for the language given as template -/// parameter. -template class online_compiler { -public: - /// Constructs online compiler which can target any device and produces - /// given compiled code format. Produces 64-bit device code. - /// The created compiler is "optimistic" - it assumes all applicable SYCL - /// device capabilities are supported by the target device(s). - online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), - DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), - Is64Bit(true), DeviceStepping("") {} - - /// Constructs online compiler which targets given architecture and produces - /// given compiled code format. Produces 64-bit device code. - /// Throws online_compile_error if values of constructor arguments are - /// contradictory or not supported - e.g. if the source language is not - /// supported for given device type. - online_compiler(sycl::info::device_type dev_type, device_arch arch, - compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), - DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} - - /// Constructs online compiler for the target specified by given SYCL device. - // TODO: the initial version generates the generic code (SKL now), need - // to do additional device::info calls to determine the device by it's - // features. - online_compiler(const sycl::device &) - : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), - DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), - Is64Bit(true), DeviceStepping("") {} - - /// Compiles given in-memory \c Lang source to a binary blob. Blob format, - /// other parameters are set in the constructor by the compilation target - /// specification parameters. - /// Specialization for each language will provide exact signatures, which - /// can be different for different languages. - /// Throws online_compile_error if compilation is not successful. - template - std::vector compile(const std::string &src, const Tys &... args); - - /// Sets the compiled code format of the compilation target and returns *this. - online_compiler &setOutputFormat(compiled_code_format fmt) { - OutputFormat = fmt; - return *this; - } - - /// Sets the compiled code format version of the compilation target and - /// returns *this. - online_compiler &setOutputFormatVersion(int major, int minor) { - OutputFormatVersion = {major, minor}; - return *this; - } - - /// Sets the device type of the compilation target and returns *this. - online_compiler &setTargetDeviceType(sycl::info::device_type type) { - DeviceType = type; - return *this; - } - - /// Sets the device architecture of the compilation target and returns *this. - online_compiler &setTargetDeviceArch(device_arch arch) { - DeviceArch = arch; - return *this; - } - - /// Makes the compilation target 32-bit and returns *this. - online_compiler &set32bitTarget() { - Is64Bit = false; - return *this; - }; - - /// Makes the compilation target 64-bit and returns *this. - online_compiler &set64bitTarget() { - Is64Bit = true; - return *this; - }; - - /// Sets implementation-defined target device stepping of the compilation - /// target and returns *this. - online_compiler &setTargetDeviceStepping(const std::string &id) { - DeviceStepping = id; - return *this; - } - -private: - /// Compiled code format. - compiled_code_format OutputFormat; - - /// Compiled code format version - a pair of "major" and "minor" components - std::pair OutputFormatVersion; - - /// Target device type - sycl::info::device_type DeviceType; - - /// Target device architecture - device_arch DeviceArch; - - /// Whether the target device architecture is 64-bit - bool Is64Bit; - - /// Target device stepping (implementation defined) - std::string DeviceStepping; - - /// Handles to helper functions used by the implementation. - void *CompileToSPIRVHandle = nullptr; - void *FreeSPIRVOutputsHandle = nullptr; -}; - -// Specializations of the online_compiler class and 'compile' function for -// particular languages and parameter types. - -/// Compiles the given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source. -/// @param options - compilation options (implementation defined); standard -/// OpenCL JIT compiler options must be supported. -template <> -template <> -__SYCL_EXPORT std::vector -online_compiler::compile( - const std::string &src, const std::vector &options); - -/// Compiles the given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source. -template <> -template <> -std::vector -online_compiler::compile(const std::string &src) { - return compile(src, std::vector{}); -} - -/// Compiles the given CM source \p src. -/// @param src - contents of the source. -/// @param options - compilation options (implementation defined). -template <> -template <> -__SYCL_EXPORT std::vector online_compiler::compile( - const std::string &src, const std::vector &options); - -/// Compiles the given CM source \p src. -template <> -template <> -std::vector -online_compiler::compile(const std::string &src) { - return compile(src, std::vector{}); -} - -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/INTEL/pipes.hpp b/sycl/include/CL/sycl/INTEL/pipes.hpp index eef93efd884a3..62c56d90cf200 100644 --- a/sycl/include/CL/sycl/INTEL/pipes.hpp +++ b/sycl/include/CL/sycl/INTEL/pipes.hpp @@ -8,196 +8,9 @@ #pragma once -#include -#include -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -namespace INTEL { +__SYCL_WARNING("CL/sycl/INTEL/pipes.hpp usage is deprecated, include " + "sycl/ext/intel/pipes.hpp instead") -template class pipe { -public: - // Non-blocking pipes - // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V - // friendly LLVM IR. - static _dataT read(bool &_Success) { -#ifdef __SYCL_DEVICE_ONLY__ - RPipeTy<_dataT> _RPipe = - __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage); - _dataT TempData; - _Success = !static_cast( - __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment)); - return TempData; -#else - (void)_Success; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - - // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V - // friendly LLVM IR. - static void write(const _dataT &_Data, bool &_Success) { -#ifdef __SYCL_DEVICE_ONLY__ - WPipeTy<_dataT> _WPipe = - __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage); - _Success = !static_cast( - __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment)); -#else - (void)_Success; - (void)_Data; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - - // Blocking pipes - // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V - // friendly LLVM IR. - static _dataT read() { -#ifdef __SYCL_DEVICE_ONLY__ - RPipeTy<_dataT> _RPipe = - __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage); - _dataT TempData; - __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment); - return TempData; -#else - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - - // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V - // friendly LLVM IR. - static void write(const _dataT &_Data) { -#ifdef __SYCL_DEVICE_ONLY__ - WPipeTy<_dataT> _WPipe = - __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage); - __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment); -#else - (void)_Data; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - -private: - static constexpr int32_t m_Size = sizeof(_dataT); - static constexpr int32_t m_Alignment = alignof(_dataT); - static constexpr int32_t m_Capacity = _min_capacity; -#ifdef __SYCL_DEVICE_ONLY__ - static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment, - m_Capacity}; -#endif // __SYCL_DEVICE_ONLY__ -}; - -// IO pipes that provide interface to connect with hardware peripheral. -// Their name aliases are defined in vendor-provided header, below you can see -// an example of this header. There are defined aliases to ethernet_read_pipe -// and ethernet_write_pipe that users can use in their code to connect with -// HW peripheral. -/* namespace intelfpga { -template -struct ethernet_pipe_id { - static constexpr int32_t id = ID; -}; - -template -using ethernet_read_pipe = - kernel_readable_io_pipe, _dataT, _min_capacity>; - -template -using ethernet_write_pipe = - kernel_writeable_io_pipe, _dataT, _min_capacity>; -} // namespace intelfpga */ - -template -class kernel_readable_io_pipe { -public: - // Non-blocking pipes - // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V - // friendly LLVM IR. - static _dataT read(bool &_Success) { -#ifdef __SYCL_DEVICE_ONLY__ - RPipeTy<_dataT> _RPipe = - __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage); - _dataT TempData; - _Success = !static_cast( - __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment)); - return TempData; -#else - (void)_Success; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - - // Blocking pipes - // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V - // friendly LLVM IR. - static _dataT read() { -#ifdef __SYCL_DEVICE_ONLY__ - RPipeTy<_dataT> _RPipe = - __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage); - _dataT TempData; - __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment); - return TempData; -#else - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - -private: - static constexpr int32_t m_Size = sizeof(_dataT); - static constexpr int32_t m_Alignment = alignof(_dataT); - static constexpr int32_t m_Capacity = _min_capacity; - static constexpr int32_t ID = _name::id; -#ifdef __SYCL_DEVICE_ONLY__ - static constexpr struct ConstantPipeStorage m_Storage - __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity}; -#endif // __SYCL_DEVICE_ONLY__ -}; - -template -class kernel_writeable_io_pipe { -public: - // Non-blocking pipes - // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V - // friendly LLVM IR. - static void write(const _dataT &_Data, bool &_Success) { -#ifdef __SYCL_DEVICE_ONLY__ - WPipeTy<_dataT> _WPipe = - __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage); - _Success = !static_cast( - __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment)); -#else - (void)_Data; - (void)_Success; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - - // Blocking pipes - // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V - // friendly LLVM IR. - static void write(const _dataT &_Data) { -#ifdef __SYCL_DEVICE_ONLY__ - WPipeTy<_dataT> _WPipe = - __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage); - __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment); -#else - (void)_Data; - assert(!"Pipes are not supported on a host device!"); -#endif // __SYCL_DEVICE_ONLY__ - } - -private: - static constexpr int32_t m_Size = sizeof(_dataT); - static constexpr int32_t m_Alignment = alignof(_dataT); - static constexpr int32_t m_Capacity = _min_capacity; - static constexpr int32_t ID = _name::id; -#ifdef __SYCL_DEVICE_ONLY__ - static constexpr struct ConstantPipeStorage m_Storage - __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity}; -#endif // __SYCL_DEVICE_ONLY__ -}; - -} // namespace INTEL -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +#include diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp index 4688cb47f1d38..ff5e6b05e325e 100644 --- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp +++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp @@ -8,225 +8,10 @@ #pragma once -#include -#include -#include -#include +#include -__SYCL_INLINE_NAMESPACE(cl) { -namespace sycl { -// Forward declaration -template -class accessor; -namespace detail { -// This helper template must be specialized for nested instance template -// of each compile-time-constant property. -template struct IsCompileTimePropertyInstance : std::false_type {}; -} // namespace detail -namespace ONEAPI { +__SYCL_WARNING("CL/sycl/ONEAPI/accessor_property_list.hpp usage is " + "deprecated, include " + "sycl/ext/oneapi/accessor_property_list.hpp instead") -template struct is_compile_time_property : std::false_type {}; - -/// Objects of the accessor_property_list class are containers for the SYCL -/// properties. -/// -/// Unlike \c property_list, accessor_property_list can take -/// compile-time-constant properties. -/// -/// \sa accessor -/// \sa property_list -/// -/// \ingroup sycl_api -template -class accessor_property_list : protected sycl::detail::PropertyListBase { - // These structures check if compile-time-constant property is present in - // list. For runtime properties this check is always true. - template struct AreSameTemplate : std::is_same {}; - template