[SYCL] Support online_compiler::compile compiled with pre-C++11 ABI (#16269)

This is an updated version of #16179.

Run-time ABI mismatch resulting in a segmentation fault was uncovered
why enabling extra testing in #16235,
this new version fixes that.

Author: aelovikov-intel

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

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<byte>
+  compile_impl(detail::string_view Src,
+               const std::vector<detail::string_view> &Options,
+               std::pair<int, int> OutputFormatVersion,
+               sycl::info::device_type DeviceType, device_arch DeviceArch,
+               bool Is64Bit, detail::string_view DeviceStepping,
+               void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle);
+
+  std::vector<byte> compile_impl(const std::string &Source,
+                                 const std::vector<std::string> &UserArgs) {
+    std::vector<sycl::detail::string_view> 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<byte>
-online_compiler<source_language::opencl_c>::compile(
-    const std::string &src, const std::vector<std::string> &options);
+#if !defined(__SYCL_ONLINE_COMPILER_CPP) ||                                    \
+    defined(__INTEL_PREVIEW_BREAKING_CHANGES)
+inline
+#else
+__SYCL_EXPORT
+#endif
+    std::vector<byte>
+    online_compiler<source_language::opencl_c>::compile(
+        const std::string &src, const std::vector<std::string> &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<source_language::opencl_c>::compile(const std::string &src) {
 /// @param options - compilation options (implementation defined).
 template <>
 template <>
-__SYCL_EXPORT std::vector<byte> online_compiler<source_language::cm>::compile(
-    const std::string &src, const std::vector<std::string> &options);
+#if !defined(__SYCL_ONLINE_COMPILER_CPP) ||                                    \
+    defined(__INTEL_PREVIEW_BREAKING_CHANGES)
+inline
+#else
+__SYCL_EXPORT
+#endif
+    std::vector<byte>
+    online_compiler<source_language::cm>::compile(
+        const std::string &src, const std::vector<std::string> &options) {
+  return compile_impl(src, options);
+}
 
 /// Compiles the given CM source \p src.
 template <>

sycl/source/detail/online_compiler/online_compiler.cpp

@@ -6,6 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+#define __SYCL_ONLINE_COMPILER_CPP
+
 #include <sycl/detail/os_util.hpp>
 #include <sycl/detail/ur.hpp>
 #include <sycl/ext/intel/experimental/online_compiler.hpp>
@@ -19,9 +21,11 @@ inline namespace _V1 {
 namespace ext::intel::experimental {
 namespace detail {
 
+using namespace sycl::detail;
+
 static std::vector<const char *>
 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<const char *> 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<byte>
-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<std::string> &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<byte>
-online_compiler<source_language::opencl_c>::compile(
-    const std::string &Source, const std::vector<std::string> &UserArgs) {
+template <source_language Lang>
+__SYCL_EXPORT std::vector<byte> online_compiler<Lang>::compile_impl(
+    detail::string_view Src, const std::vector<detail::string_view> &Options,
+    std::pair<int, int> OutputFormatVersion, sycl::info::device_type DeviceType,
+    device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping,
+    void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle) {
 
   if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
     std::string Version = std::to_string(OutputFormatVersion.first) + ", " +
@@ -211,29 +216,31 @@ online_compiler<source_language::opencl_c>::compile(
                                Version + ") is not supported yet");
   }
 
-  return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit,
-                                DeviceStepping, CompileToSPIRVHandle,
-                                FreeSPIRVOutputsHandle, UserArgs);
-}
-
-template <>
-template <>
-__SYCL_EXPORT std::vector<byte> online_compiler<source_language::cm>::compile(
-    const std::string &Source, const std::vector<std::string> &UserArgs) {
+  std::vector<std::string> UserArgs;
+  for (auto &&Opt : Options)
+    UserArgs.emplace_back(Opt.data());
 
-  if (OutputFormatVersion != std::pair<int, int>{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<std::string> 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<byte>
+online_compiler<source_language::opencl_c>::compile_impl(
+    detail::string_view Src, const std::vector<detail::string_view> &Options,
+    std::pair<int, int> OutputFormatVersion, sycl::info::device_type DeviceType,
+    device_arch DeviceArch, bool Is64Bit, detail::string_view DeviceStepping,
+    void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle);
+
+template __SYCL_EXPORT std::vector<byte>
+online_compiler<source_language::cm>::compile_impl(
+    detail::string_view Src, const std::vector<detail::string_view> &Options,
+    std::pair<int, int> 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 {

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

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