[WIP] Flip the default for libcxx/libc usage

Author: aelovikov-intel

clang/include/clang/Driver/Options.td

@@ -7588,9 +7588,9 @@ let Visibility = [SYCLRTCOnlyOption] in {
         : Joined<["--"], "persistent-auto-pch=">,
           HelpText<"Use Persistent Auto-PCH cache located at <dir> for SYCL "
                    "RTC Compilation">;
-    def sycl_rtc_exp_redist_mode
-        : Flag<["--"], "sycl-rtc-experimental-redist-mode">,
-          HelpText<"Use in-memory system includes">;
+    def sycl_rtc_use_system_includes
+        : Flag<["--"], "sycl-rtc-use-system-includes">,
+          HelpText<"Use system includes instead of in-memory libcxx/libc">;
     def sycl_rtc_in_memory_fs_only
         : Flag<["--"], "sycl-rtc-in-memory-fs-only">,
           HelpText<"Disable real filesystem access for SYCL RTC compilation, "

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -261,7 +261,7 @@ class SYCLToolchain {
       DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_offload_arch_EQ), CPU);
     }
 
-    // Reasons why this is done here and not in the clang driver:
+    // Reasons why the following is done here and not in the clang driver:
     //
     //  1) Unlike libcxx, upstream libc is installed directly into
     //     `<toolchain>/include` or `<toolchain>/<target>/include` together with
@@ -288,7 +288,22 @@ class SYCLToolchain {
     //      headers installation to a separate directory (similar to libcxx), at
     //      that time we might have support for this in the clang driver
     //      directly and would be able to avoid doing that here.
-    if (UserArgList.hasArg(OPT_sycl_rtc_exp_redist_mode)) {
+
+    // Prefer using in-memory as that's friendlier for the end users of SYCL
+    // applications as that mode doesn't require any C/C++ toolchain to be
+    // installed on the system.
+    bool UseInMemoryCxxCHeaders = true;
+
+    // Unless explicitly told not to:
+    if (UserArgList.hasArg(OPT_sycl_rtc_use_system_includes))
+      UseInMemoryCxxCHeaders = false;
+
+    // CUDA/HIP need SDK headers that we can't distribute ourselves, so we have
+    // to use system includes as well:
+    if (Format == BinaryFormat::PTX || Format == BinaryFormat::AMDGCN)
+      UseInMemoryCxxCHeaders = false;
+
+    if (UseInMemoryCxxCHeaders) {
       DAL.AddFlagArg(nullptr, OptTable.getOption(OPT_nostdlibinc));
       auto AddInc = [&](auto RelPath) {
         DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_isystem),
@@ -310,16 +325,16 @@ class SYCLToolchain {
 #endif
       // libcxx headers, must come before libc headers:
       AddInc("include/c++/v1");
-       // libc headers, our (SYCL RTC) custom non-standard location:
+      // libc headers, our (SYCL RTC) custom non-standard location:
       AddInc("include/libc");
       // SYCL/SYCL-related headers actually, because `<sycl/sycl.hpp>` and not
       // just `<sycl.hpp>`. Can be argued that actual installation layout should
       // actually be `include/sycl/ur_api.h` and `include/sycl/sycl/sycl.hpp`
       // but that's outside the SYCL RTC scope. I think any relative order in
       // relation to libcxx/libc is allowed.
       AddInc("include/");
-      // NOTE: `include/lib/clang/<version>/include/` is added automatically (we use
-      // `--nostdlibinc` and not `--nostdinc`).
+      // NOTE: `include/lib/clang/<version>/include/` is added automatically (we
+      // use `--nostdlibinc` and not `--nostdinc`).
 
       DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_D),
                        "_LIBCPP_REMOVE_TRANSITIVE_INCLUDES");

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -1046,6 +1046,28 @@ sycl::kernel k_float = kb.ext_oneapi_get_kernel("bartmpl<float>");
 sycl::kernel k_int = kb.ext_oneapi_get_kernel("bartmpl<int>");
 ----
 
+=== Restrictions on the source code when language is `sycl`
+
+Currently, SYCL applications are linked with MSVC/GNU C/C\++ runtimes on Win/Lin
+respectively. Unfortunately, we cannot distribute those headers together with
+SYCL applications using kernel compiler due to license restrictions. Instead, we
+distribute LLVM's libc\++/libc headers as part of the kernel compiler and default
+to using them for SPIR-V-based targets. That results in a restriction not to
+pass any data that transitively uses any types defined in C/C++ headers (with
+the exception of `[u]intN_t` type aliases) as those might be ABI-incompatible
+between headers used to compiler device code and those used to compile the host
+part of the SYCL application.
+
+Technical implementation of that disables system include paths entirely (as
+having multiple C library implementation in include search paths would break
+things) and uses virtual file system containing in-memory copy of LLVM's
+libc\++/libc headers. That implicitly removes CUDA/HIP SKDs and would break
+online compilation for those targets, thus we do **not** employ that mechanism
+for those targets. Instead, end user's system is required to have a C/C++
+toolchain installation.
+
+SYCL application can also pass `--sycl-rtc-use-system-includes` option to change
+the behavior for SPIR-V targets and force usage of the system toolchain.
 
 == Examples
 
@@ -1351,56 +1373,11 @@ Some notes about the current behavior:
   using the conflicting hash would proceed without pre-compiled preamble support
   as if this option wasn't enabled.
 
-==== `--sycl-rtc-experimental-redist-mode`
+==== `--sycl-rtc-use-system-includes`
 
-Highly experimental option to facilitate distribution of applications that use
-SYCL RTC.
-
-While SYCL-specific header/bitcode files are embedded/distributed inside
-`sycl-jit` DSO we still have a dependency on the C/C++ compiler toolchain on the
-end-user system (header files only actually, but that's not too important). That
-requirement is troublesome for many of the SYCL RTC customers and this option
-tries to address that limitation.
-
-The main problem we face is that default host toolchain (MSVC/GCC on Win/Lin,
-respectively) has license requirements that prohibit distribution of necessary
-headers as part of SYCL redistributables. Instead, we embed LLVM's libc++ and
-libc headers into the `sycl-jit` DSO and provide this option to use those
-headers instead. This obviously comes with many limitations/obstacles. Using
-this options results in roughly following options being passed implicitly:
-
-* `--nostdlibinc`, effectively excluding CUDA/HIP SDKs, unless the application
-  deals with that somehow. The reason we need to do that is because having
-  different C library implementation in includes search path doesn't work.
-
-* Setup virtual file system include paths to enable the use of LLVM's
-  `libc++`/`libc` headers embedded into `sycl-jit` DSO.
-
-* Define/undefine some macros found by trial-and-error to make that fragile
-  configuration work, at least for some examples.
-
-* Add implicit `-include` directive to several C headers. Again, found by
-  trial-and-error to make things work for some of the examples we tried.
-
-This also comes with lots of limitations:
-
-* First, huge potential for ABI mismatch between host (SYCL app) ABI and the
-  device code (JIT-compiled via SYCL RTC). As such, it is recommended to limit
-  types passing the host/device boundary to:
-
-  - Fixed-width integer types
-  - `float`/`double`
-  - Aggregate types of the above.
-  - Anything else needs **very thorough** testing at the SYCL application side.
-
-  Using SYCL/STL types in the device code only, without passing host-device
-  boundary is expected to work.
-
-* CUDA/HIP aren't expected to work out-of-the-box, see above (`--nostdlibinc`).
-
-* Option is **highly** experimentall, **no** support is guaranteed and is
-  subject to change at **any** time, including the possibility of a complete
-  removal.
+Force usage of system C/C++ toolchain headers instead of the in-memory
+distribution of LLVM's libc\+\+/libc. Option has no effect if the target
+defaults to using system toolchain by default.
 
 === Known issues and limitations when the language is `sycl`
 

…t-e2e/KernelCompiler/exp_redist_mode.cpp …st-e2e/KernelCompiler/in_memory_only.cppsycl/test-e2e/KernelCompiler/exp_redist_mode.cpp renamed to sycl/test-e2e/KernelCompiler/in_memory_only.cpp sycl/test-e2e/KernelCompiler/exp_redist_mode.cpp renamed to sycl/test-e2e/KernelCompiler/in_memory_only.cpp

@@ -1,9 +1,7 @@
 // RUN: %{build} -o %t.out
 
-// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out
-
 // Make sure that debug/test-only option `--sycl-rtc-in-memory-fs-only` works
-// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} not %t.out --sycl-rtc-in-memory-fs-only | FileCheck %s --check-prefix CHECK-ERROR
+// RUN: %{run} not %t.out --sycl-rtc-in-memory-fs-only --sycl-rtc-use-system-includes | FileCheck %s --check-prefix CHECK-ERROR
 // CHECK-ERROR-LABEL:  Device compilation failed
 // CHECK-ERROR-NEXT:   Detailed information:
 // CHECK-ERROR:        	In file included from rtc_0.cpp:2:
@@ -14,12 +12,9 @@
 // CHECK-ERROR-NEXT:      14 | #include <type_traits>
 // CHECK-ERROR-NEXT:         |          ^~~~~~~~~~~~~
 
-// Now actually test the `--sycl-rtc-experimental-redist-mode` option:
-// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out --sycl-rtc-experimental-redist-mode --sycl-rtc-in-memory-fs-only
-// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out --sycl-rtc-experimental-redist-mode
-
-// XFAIL: target-native_cpu
-// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142
+// Extra check that our in-memory libcxx/libc headers can really work on a
+// system with no C/C++ toolchain:
+// RUN: %{run} %t.out --sycl-rtc-in-memory-fs-only
 
 // CUDA/HIP have SDK dependencies but exclude system includes so those aren't
 // satisfied.