diff --git a/sycl/include/sycl/detail/string_view.hpp b/sycl/include/sycl/detail/string_view.hpp index 7815ecc4ce676..d36301efbfdac 100644 --- a/sycl/include/sycl/detail/string_view.hpp +++ b/sycl/include/sycl/detail/string_view.hpp @@ -37,14 +37,19 @@ class string_view { const char *data() const noexcept { return str; } - friend bool operator==(const string_view &lhs, - std::string_view rhs) noexcept { + friend bool operator==(string_view lhs, std::string_view rhs) noexcept { return rhs == lhs.data(); } - friend bool operator==(std::string_view lhs, - const string_view &rhs) noexcept { + friend bool operator==(std::string_view lhs, string_view rhs) noexcept { return lhs == rhs.data(); } + + friend bool operator!=(string_view lhs, std::string_view rhs) noexcept { + return rhs != lhs.data(); + } + friend bool operator!=(std::string_view lhs, string_view rhs) noexcept { + return lhs != rhs.data(); + } }; } // namespace detail diff --git a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp index 769dbe45050fc..365265e762c86 100644 --- a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp +++ b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp @@ -18,6 +18,9 @@ namespace sycl { inline namespace _V1 { namespace ext::intel::experimental { +namespace detail { +using namespace sycl::detail; +} using byte = unsigned char; @@ -81,6 +84,30 @@ class __SYCL2020_DEPRECATED( "experimental online_compiler is being deprecated. See " "'sycl_ext_oneapi_kernel_compiler.asciidoc' instead for new kernel " "compiler extension to kernel_bundle implementation.") online_compiler { +#if __INTEL_PREVIEW_BREAKING_CHANGES + // Refactor this during next ABI Breaking window. We have an `std::string` + // data member so cannot be accessing `this` when crossing ABI boundary. +#endif + __SYCL_EXPORT static std::vector + compile_impl(detail::string_view Src, + const std::vector &Options, + std::pair OutputFormatVersion, + sycl::info::device_type DeviceType, device_arch DeviceArch, + bool Is64Bit, detail::string_view DeviceStepping, + void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle); + + std::vector compile_impl(const std::string &Source, + const std::vector &UserArgs) { + std::vector Args; + for (auto &&Arg : UserArgs) + Args.emplace_back(Arg); + + return compile_impl(std::string_view{Source}, Args, OutputFormatVersion, + DeviceType, DeviceArch, Is64Bit, + std::string_view{DeviceStepping}, CompileToSPIRVHandle, + FreeSPIRVOutputsHandle); + } + public: /// Constructs online compiler which can target any device and produces /// given compiled code format. Produces 64-bit device code. @@ -196,9 +223,17 @@ class __SYCL2020_DEPRECATED( /// OpenCL JIT compiler options must be supported. template <> template <> -__SYCL_EXPORT std::vector -online_compiler::compile( - const std::string &src, const std::vector &options); +#if !defined(__SYCL_ONLINE_COMPILER_CPP) || \ + defined(__INTEL_PREVIEW_BREAKING_CHANGES) +inline +#else +__SYCL_EXPORT +#endif + std::vector + online_compiler::compile( + const std::string &src, const std::vector &options) { + return compile_impl(src, options); +} /// Compiles the given OpenCL source. May throw \c online_compile_error. /// @param src - contents of the source. @@ -214,8 +249,17 @@ online_compiler::compile(const std::string &src) { /// @param options - compilation options (implementation defined). template <> template <> -__SYCL_EXPORT std::vector online_compiler::compile( - const std::string &src, const std::vector &options); +#if !defined(__SYCL_ONLINE_COMPILER_CPP) || \ + defined(__INTEL_PREVIEW_BREAKING_CHANGES) +inline +#else +__SYCL_EXPORT +#endif + std::vector + online_compiler::compile( + const std::string &src, const std::vector &options) { + return compile_impl(src, options); +} /// Compiles the given CM source \p src. template <> diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp index 5d3c3a381607b..f5e62f8a1d4d8 100644 --- a/sycl/source/detail/online_compiler/online_compiler.cpp +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -6,6 +6,8 @@ // //===----------------------------------------------------------------------===// +#define __SYCL_ONLINE_COMPILER_CPP + #include #include #include @@ -19,9 +21,11 @@ inline namespace _V1 { namespace ext::intel::experimental { namespace detail { +using namespace sycl::detail; + static std::vector prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, - bool Is64Bit, const std::string &DeviceStepping, + bool Is64Bit, string_view DeviceStepping, const std::string &UserArgs) { std::vector Args = {"ocloc", "-q", "-spv_only", "-device"}; @@ -54,7 +58,7 @@ prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, if (DeviceStepping != "") { Args.push_back("-revision_id"); - Args.push_back(DeviceStepping.c_str()); + Args.push_back(DeviceStepping.data()); } Args.push_back(Is64Bit ? "-64" : "-32"); @@ -82,11 +86,11 @@ prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, /// allocated during the compilation. /// @param UserArgs - User's options to ocloc compiler. static std::vector -compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, - device_arch DeviceArch, bool Is64Bit, - const std::string &DeviceStepping, void *&CompileToSPIRVHandle, - void *&FreeSPIRVOutputsHandle, +compileToSPIRV(string_view Src, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, string_view DeviceStepping, + void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle, const std::vector &UserArgs) { + std::string Source{Src.data()}; if (!CompileToSPIRVHandle) { #ifdef __SYCL_RT_OS_WINDOWS @@ -198,11 +202,12 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, } } // namespace detail -template <> -template <> -__SYCL_EXPORT std::vector -online_compiler::compile( - const std::string &Source, const std::vector &UserArgs) { +template +__SYCL_EXPORT std::vector online_compiler::compile_impl( + detail::string_view Src, const std::vector &Options, + std::pair OutputFormatVersion, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping, + void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle) { if (OutputFormatVersion != std::pair{0, 0}) { std::string Version = std::to_string(OutputFormatVersion.first) + ", " + @@ -211,29 +216,31 @@ online_compiler::compile( Version + ") is not supported yet"); } - return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, - DeviceStepping, CompileToSPIRVHandle, - FreeSPIRVOutputsHandle, UserArgs); -} - -template <> -template <> -__SYCL_EXPORT std::vector online_compiler::compile( - const std::string &Source, const std::vector &UserArgs) { + std::vector UserArgs; + for (auto &&Opt : Options) + UserArgs.emplace_back(Opt.data()); - if (OutputFormatVersion != std::pair{0, 0}) { - std::string Version = std::to_string(OutputFormatVersion.first) + ", " + - std::to_string(OutputFormatVersion.second); - throw online_compile_error(std::string("The output format version (") + - Version + ") is not supported yet"); - } + if constexpr (Lang == source_language::cm) + UserArgs.push_back("-cmc"); - std::vector CMUserArgs = UserArgs; - CMUserArgs.push_back("-cmc"); - return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + return detail::compileToSPIRV(Src, DeviceType, DeviceArch, Is64Bit, DeviceStepping, CompileToSPIRVHandle, - FreeSPIRVOutputsHandle, CMUserArgs); + FreeSPIRVOutputsHandle, UserArgs); } + +template __SYCL_EXPORT std::vector +online_compiler::compile_impl( + detail::string_view Src, const std::vector &Options, + std::pair OutputFormatVersion, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping, + void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle); + +template __SYCL_EXPORT std::vector +online_compiler::compile_impl( + detail::string_view Src, const std::vector &Options, + std::pair OutputFormatVersion, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping, + void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle); } // namespace ext::intel::experimental namespace ext { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 88ba9663dcfc5..e8632d585cbd6 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -2985,7 +2985,9 @@ _ZN4sycl3_V121__isgreaterequal_implEdd _ZN4sycl3_V121__isgreaterequal_implEff _ZN4sycl3_V122accelerator_selector_vERKNS0_6deviceE _ZN4sycl3_V128verifyUSMAllocatorPropertiesERKNS0_13property_listE +_ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE0EE12compile_implENS0_6detail11string_viewERKSt6vectorIS8_SaIS8_EESt4pairIiiENS0_4info11device_typeENS3_11device_archEbS8_RPvSK_ _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ +_ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE1EE12compile_implENS0_6detail11string_viewERKSt6vectorIS8_SaIS8_EESt4pairIiiENS0_4info11device_typeENS3_11device_archEbS8_RPvSK_ _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ _ZN4sycl3_V13ext5intel12experimental9pipe_base13get_pipe_nameB5cxx11EPKv _ZN4sycl3_V13ext5intel12experimental9pipe_base17wait_non_blockingERKNS0_5eventE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 03ce01bbfe456..4cd02d0bc61c3 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3743,6 +3743,8 @@ ?category@exception@_V1@sycl@@QEBAAEBVerror_category@std@@XZ ?clearArgs@handler@_V1@sycl@@AEAAXXZ ?code@exception@_V1@sycl@@QEBAAEBVerror_code@std@@XZ +?compile_impl@?$online_compiler@$00@experimental@intel@ext@_V1@sycl@@CA?AV?$vector@EV?$allocator@E@std@@@std@@Vstring_view@detail@56@AEBV?$vector@Vstring_view@detail@_V1@sycl@@V?$allocator@Vstring_view@detail@_V1@sycl@@@std@@@8@U?$pair@HH@8@W4device_type@info@56@Vdevice_arch@23456@_N0AEAPEAX6@Z +?compile_impl@?$online_compiler@$0A@@experimental@intel@ext@_V1@sycl@@CA?AV?$vector@EV?$allocator@E@std@@@std@@Vstring_view@detail@56@AEBV?$vector@Vstring_view@detail@_V1@sycl@@V?$allocator@Vstring_view@detail@_V1@sycl@@@std@@@8@U?$pair@HH@8@W4device_type@info@56@Vdevice_arch@23456@_N0AEAPEAX6@Z ?compile_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$kernel_bundle@$0A@@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z ?complete_fusion@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAA?AVevent@56@AEBVproperty_list@56@@Z ?computeFallbackKernelBounds@handler@_V1@sycl@@AEAA?AV?$id@$01@23@_K0@Z