[SYCL] Support online_compiler::compile compiled with pre-C++11 ABI

This is an updated version of intel#16179.

Run-time ABI mismatch resulting in a segmentation fault was uncovered
why enabling extra testing in intel#16235.

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,112 +84,139 @@ 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 {
-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 <typename... Tys>
-  std::vector<byte> compile(const std::string &src, const Tys &...args);
-
-  /// Sets the compiled code format of the compilation target and returns *this.
-  online_compiler<Lang> &setOutputFormat(compiled_code_format fmt) {
-    OutputFormat = fmt;
-    return *this;
-  }
-
-  /// Sets the compiled code format version of the compilation target and
-  /// returns *this.
-  online_compiler<Lang> &setOutputFormatVersion(int major, int minor) {
-    OutputFormatVersion = {major, minor};
-    return *this;
-  }
-
-  /// Sets the device type of the compilation target and returns *this.
-  online_compiler<Lang> &setTargetDeviceType(sycl::info::device_type type) {
-    DeviceType = type;
-    return *this;
+#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);
   }
 
-  /// Sets the device architecture of the compilation target and returns *this.
-  online_compiler<Lang> &setTargetDeviceArch(device_arch arch) {
-    DeviceArch = arch;
-    return *this;
-  }
-
-  /// Makes the compilation target 32-bit and returns *this.
-  online_compiler<Lang> &set32bitTarget() {
-    Is64Bit = false;
-    return *this;
-  };
-
-  /// Makes the compilation target 64-bit and returns *this.
-  online_compiler<Lang> &set64bitTarget() {
-    Is64Bit = true;
-    return *this;
+  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 <typename... Tys>
+    std::vector<byte> compile(const std::string &src, const Tys &...args);
+
+    /// Sets the compiled code format of the compilation target and returns
+    /// *this.
+    online_compiler<Lang> &setOutputFormat(compiled_code_format fmt) {
+      OutputFormat = fmt;
+      return *this;
+    }
+
+    /// Sets the compiled code format version of the compilation target and
+    /// returns *this.
+    online_compiler<Lang> &setOutputFormatVersion(int major, int minor) {
+      OutputFormatVersion = {major, minor};
+      return *this;
+    }
+
+    /// Sets the device type of the compilation target and returns *this.
+    online_compiler<Lang> &setTargetDeviceType(sycl::info::device_type type) {
+      DeviceType = type;
+      return *this;
+    }
+
+    /// Sets the device architecture of the compilation target and returns
+    /// *this.
+    online_compiler<Lang> &setTargetDeviceArch(device_arch arch) {
+      DeviceArch = arch;
+      return *this;
+    }
+
+    /// Makes the compilation target 32-bit and returns *this.
+    online_compiler<Lang> &set32bitTarget() {
+      Is64Bit = false;
+      return *this;
+    };
+
+    /// Makes the compilation target 64-bit and returns *this.
+    online_compiler<Lang> &set64bitTarget() {
+      Is64Bit = true;
+      return *this;
+    };
+
+    /// Sets implementation-defined target device stepping of the compilation
+    /// target and returns *this.
+    online_compiler<Lang> &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<int, int> 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;
   };
 
-  /// Sets implementation-defined target device stepping of the compilation
-  /// target and returns *this.
-  online_compiler<Lang> &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<int, int> 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.
 
@@ -196,9 +226,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 +252,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 <>