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..feeaa30d9e2ba 100644 --- a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp +++ b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp @@ -81,6 +81,21 @@ 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 { + __SYCL_EXPORT std::vector + compile_impl(sycl::detail::string_view Src, + sycl::detail::string_view DeviceStepping, + const std::vector &Options); + + 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}, + std::string_view{DeviceStepping}, Args); + } + public: /// Constructs online compiler which can target any device and produces /// given compiled code format. Produces 64-bit device code. @@ -196,9 +211,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 +237,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..344d32cd61ff5 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,10 @@ 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, detail::string_view DeviceStepping, + const std::vector &Options) { if (OutputFormatVersion != std::pair{0, 0}) { std::string Version = std::to_string(OutputFormatVersion.first) + ", " + @@ -211,29 +214,27 @@ online_compiler::compile( Version + ") is not supported yet"); } - return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + std::vector UserArgs; + for (auto &&Opt : Options) + UserArgs.emplace_back(Opt.data()); + + if constexpr (Lang == source_language::cm) + UserArgs.push_back("-cmc"); + + return detail::compileToSPIRV(Src, DeviceType, DeviceArch, Is64Bit, DeviceStepping, CompileToSPIRVHandle, FreeSPIRVOutputsHandle, UserArgs); } -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, detail::string_view DeviceStepping, + const std::vector &Options); - 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"); - } - - std::vector CMUserArgs = UserArgs; - CMUserArgs.push_back("-cmc"); - return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, - DeviceStepping, CompileToSPIRVHandle, - FreeSPIRVOutputsHandle, CMUserArgs); -} +template __SYCL_EXPORT std::vector +online_compiler::compile_impl( + detail::string_view Src, detail::string_view DeviceStepping, + const std::vector &Options); } // 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 a5134a7a524ca..853ac28bad1d5 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_viewES8_RKSt6vectorIS8_SaIS8_EE _ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISE_EEEEES8_IhSaIhEERKSE_DpRKT_ +_ZN4sycl3_V13ext5intel12experimental15online_compilerILNS3_15source_languageE1EE12compile_implENS0_6detail11string_viewES8_RKSt6vectorIS8_SaIS8_EE _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 a6e6a5e47c137..8e29aba2726c9 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3741,6 +3741,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@@AEAA?AV?$vector@EV?$allocator@E@std@@@std@@Vstring_view@detail@56@0AEBV?$vector@Vstring_view@detail@_V1@sycl@@V?$allocator@Vstring_view@detail@_V1@sycl@@@std@@@8@@Z +?compile_impl@?$online_compiler@$0A@@experimental@intel@ext@_V1@sycl@@AEAA?AV?$vector@EV?$allocator@E@std@@@std@@Vstring_view@detail@56@0AEBV?$vector@Vstring_view@detail@_V1@sycl@@V?$allocator@Vstring_view@detail@_V1@sycl@@@std@@@8@@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