[SYCL RTC] Add sycl-rtc-experimental-redist-mode option

aelovikov-intel · aelovikov-intel · commit 9d0562d3c84d · 2025-10-29T17:07:43.000-07:00
diff --git a/buildbot/configure.py b/buildbot/configure.py
@@ -179,6 +179,11 @@ def do_configure(args, passthrough_args):
 
     install_dir = os.path.join(abs_obj_dir, "install")
 
+    llvm_enable_runtimes = "libcxx"
+    if platform.system() != "Windows":
+        llvm_enable_runtimes += ";libcxxabi;libunwind"
+    llvm_enable_runtimes += ";libc"
+
     cmake_cmd = [
         "cmake",
         "-G",
@@ -195,6 +200,8 @@ def do_configure(args, passthrough_args):
         "-DLLVM_EXTERNAL_LIBDEVICE_SOURCE_DIR={}".format(libdevice_dir),
         "-DLLVM_EXTERNAL_SYCL_JIT_SOURCE_DIR={}".format(jit_dir),
         "-DLLVM_ENABLE_PROJECTS={}".format(llvm_enable_projects),
+        "-DLLVM_ENABLE_RUNTIMES={}".format(llvm_enable_runtimes),
+        "-DLLVM_LIBC_FULL_BUILD=ON",
         "-DSYCL_BUILD_PI_HIP_PLATFORM={}".format(sycl_build_pi_hip_platform),
         "-DLLVM_BUILD_TOOLS=ON",
         "-DLLVM_ENABLE_ZSTD={}".format(llvm_enable_zstd),
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
@@ -7561,6 +7561,14 @@ 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_in_memory_fs_only
+        : Flag<["--"], "sycl-rtc-in-memory-fs-only">,
+          HelpText<"Disable real filesystem access for SCYl RTC compilation, "
+                   "debug/testing only">,
+          Flags<[HelpHidden]>;
   } // let Group = sycl_rtc_only_Group
 } // let Visibility = [SYCLRTCOnlyOption]
 
diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt
@@ -45,13 +45,40 @@ add_custom_target(rtc-prepare-resources
   ${SYCL_JIT_RESOURCE_FILES}
 )
 
+set(SYCL_JIT_RUNTIME_RESOURCE_DEPS)
+set(SYCL_JIT_PREPARE_RUNTIME_RESOURCE_COMMANDS)
+
+if ("libcxx" IN_LIST LLVM_ENABLE_RUNTIMES)
+  list(APPEND SYCL_JIT_RUNTIME_RESOURCE_DEPS runtimes-configure)
+
+  list(APPEND SYCL_JIT_PREPARE_RUNTIME_RESOURCE_COMMANDS
+    COMMAND ${CMAKE_COMMAND} --build ${CMAKE_BINARY_DIR}/runtimes/runtimes-bins --target generate-cxx-headers
+    COMMAND ${CMAKE_COMMAND} --install ${CMAKE_BINARY_DIR}/runtimes/runtimes-bins --prefix ${CMAKE_CURRENT_BINARY_DIR}/rtc-resources-install --component cxx-headers
+    COMMAND sed -i 's/_LIBCPP_HAS_THREAD_API_EXTERNAL 0/_LIBCPP_HAS_THREAD_API_EXTERNAL 1/' ${CMAKE_CURRENT_BINARY_DIR}/rtc-resources-install/include/x86_64-unknown-linux-gnu/c++/v1/__config_site
+  )
+endif()
+
+if ("libc" IN_LIST LLVM_ENABLE_RUNTIMES)
+  list(APPEND SYCL_JIT_RUNTIME_RESOURCE_DEPS libc)
+  list(APPEND SYCL_JIT_PREPARE_RUNTIME_RESOURCE_COMMANDS
+    COMMAND ${CMAKE_COMMAND} -E copy_directory ${CMAKE_BINARY_DIR}/runtimes/runtimes-bins/libc/include ${SYCL_JIT_RESOURCE_INSTALL_DIR}/include/libc
+  )
+endif()
+
+add_custom_target(rtc-prepare-runtime-resources
+  DEPENDS ${SYCL_JIT_RUNTIME_RESOURCE_DEPS}
+  ${SYCL_JIT_PREPARE_RUNTIME_RESOURCE_COMMANDS}
+)
+
 add_custom_command(
   OUTPUT ${SYCL_JIT_RESOURCE_CPP}
   COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/utils/generate.py --toolchain-dir ${SYCL_JIT_RESOURCE_INSTALL_DIR} --output ${SYCL_JIT_RESOURCE_CPP} --prefix ${SYCL_JIT_VIRTUAL_TOOLCHAIN_ROOT}
   DEPENDS
   rtc-prepare-resources
   ${SYCL_JIT_RESOURCE_DEPS}
   ${SYCL_JIT_RESOURCE_FILES}
+  rtc-prepare-runtime-resources
+  ${SYCL_JIT_RUNTIME_RESOURCE_DEPS}
   ${CMAKE_CURRENT_SOURCE_DIR}/utils/generate.py
 )
 
@@ -94,6 +121,7 @@ add_custom_command(
   ${SYCL_JIT_RESOURCE_CPP}
   ${SYCL_JIT_RESOURCE_DEPS}
   ${SYCL_JIT_RESOURCE_FILES}
+  ${SYCL_JIT_RUNTIME_RESOURCE_DEPS}
   ${CMAKE_CURRENT_SOURCE_DIR}/include/Resource.h
 )
 
diff --git a/sycl-jit/jit-compiler/lib/resource-includes/include/sycl-rtc-standalone/__external_threading b/sycl-jit/jit-compiler/lib/resource-includes/include/sycl-rtc-standalone/__external_threading
@@ -0,0 +1,78 @@
+_LIBCPP_BEGIN_NAMESPACE_STD
+
+using __libcpp_timespec_t = int;
+
+//
+// Mutex
+//
+using __libcpp_mutex_t = int;
+#define _LIBCPP_MUTEX_INITIALIZER 0
+
+using __libcpp_recursive_mutex_t = int;
+
+int __libcpp_recursive_mutex_init(__libcpp_recursive_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS int __libcpp_recursive_mutex_lock(__libcpp_recursive_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS bool __libcpp_recursive_mutex_trylock(__libcpp_recursive_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS int __libcpp_recursive_mutex_unlock(__libcpp_recursive_mutex_t*);
+int __libcpp_recursive_mutex_destroy(__libcpp_recursive_mutex_t*);
+
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS int __libcpp_mutex_lock(__libcpp_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS bool __libcpp_mutex_trylock(__libcpp_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS int __libcpp_mutex_unlock(__libcpp_mutex_t*);
+int __libcpp_mutex_destroy(__libcpp_mutex_t*);
+
+//
+// Condition Variable
+//
+using __libcpp_condvar_t = int;
+#define _LIBCPP_CONDVAR_INITIALIZER 0
+
+int __libcpp_condvar_signal(__libcpp_condvar_t*);
+int __libcpp_condvar_broadcast(__libcpp_condvar_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS int __libcpp_condvar_wait(__libcpp_condvar_t*, __libcpp_mutex_t*);
+_LIBCPP_NO_THREAD_SAFETY_ANALYSIS
+int __libcpp_condvar_timedwait(__libcpp_condvar_t*, __libcpp_mutex_t*, __libcpp_timespec_t*);
+int __libcpp_condvar_destroy(__libcpp_condvar_t*);
+
+//
+// Execute once
+//
+using __libcpp_exec_once_flag = int;
+#define _LIBCPP_EXEC_ONCE_INITIALIZER 0
+
+int __libcpp_execute_once(__libcpp_exec_once_flag*, void (*__init_routine)());
+
+//
+// Thread id
+//
+using __libcpp_thread_id = int;
+
+bool __libcpp_thread_id_equal(__libcpp_thread_id, __libcpp_thread_id);
+bool __libcpp_thread_id_less(__libcpp_thread_id, __libcpp_thread_id);
+
+//
+// Thread
+//
+#define _LIBCPP_NULL_THREAD 0
+using __libcpp_thread_t = int;
+
+bool __libcpp_thread_isnull(const __libcpp_thread_t*);
+int __libcpp_thread_create(__libcpp_thread_t*, void* (*__func)(void*), void* __arg);
+__libcpp_thread_id __libcpp_thread_get_current_id();
+__libcpp_thread_id __libcpp_thread_get_id(const __libcpp_thread_t*);
+int __libcpp_thread_join(__libcpp_thread_t*);
+int __libcpp_thread_detach(__libcpp_thread_t*);
+void __libcpp_thread_yield();
+void __libcpp_thread_sleep_for(const chrono::nanoseconds&);
+
+//
+// Thread local storage
+//
+#define _LIBCPP_TLS_DESTRUCTOR_CC 0
+using __libcpp_tls_key = int;
+
+int __libcpp_tls_create(__libcpp_tls_key*, void (*__at_exit)(void*));
+void* __libcpp_tls_get(__libcpp_tls_key);
+int __libcpp_tls_set(__libcpp_tls_key, void*);
+
+_LIBCPP_END_NAMESPACE_STD
diff --git a/sycl-jit/jit-compiler/lib/resource-includes/include/sycl-rtc-standalone/linux/errno.h b/sycl-jit/jit-compiler/lib/resource-includes/include/sycl-rtc-standalone/linux/errno.h
diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
@@ -262,6 +262,26 @@ class SYCLToolchain {
       DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_offload_arch_EQ), CPU);
     }
 
+    if (UserArgList.hasArg(OPT_sycl_rtc_exp_redist_mode)) {
+      DAL.AddFlagArg(nullptr, OptTable.getOption(OPT_nostdlibinc));
+      auto AddInc = [&](auto RelPath) {
+        DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_isystem),
+                         (getPrefix() + RelPath).str());
+      };
+      AddInc("include/sycl/stl_wrappers");
+      AddInc("include/x86_64-unknown-linux-gnu/c++/v1");
+      AddInc("include/c++/v1");
+      AddInc("include/libc");
+      AddInc("include/");
+      AddInc("include/sycl-rtc-standalone/");
+      AddInc("include/lib/clang/22/include/");
+      DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_D),
+                       "_LIBCPP_REMOVE_TRANSITIVE_INCLUDES");
+      DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_include), "stdio.h");
+      DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_include), "wchar.h");
+      DAL.AddJoinedArg(nullptr, OptTable.getOption(OPT_include), "time.h");
+    }
+
     ArgStringList ASL;
     for (Arg *A : DAL)
       A->render(DAL, ASL);
@@ -543,9 +563,15 @@ class SYCLToolchain {
     std::vector<std::string> CommandLine =
         createCommandLine(UserArgList, Format, SourceFilePath);
 
-    auto FS = llvm::makeIntrusiveRefCnt<llvm::vfs::OverlayFileSystem>(
-        llvm::vfs::getRealFileSystem());
-    FS->pushOverlay(getToolchainFS());
+    llvm::IntrusiveRefCntPtr<llvm::vfs::OverlayFileSystem> FS;
+    if (UserArgList.hasArg(OPT_sycl_rtc_in_memory_fs_only)) {
+      FS = llvm::makeIntrusiveRefCnt<llvm::vfs::OverlayFileSystem>(
+          getToolchainFS());
+    } else {
+      FS = llvm::makeIntrusiveRefCnt<llvm::vfs::OverlayFileSystem>(
+          llvm::vfs::getRealFileSystem());
+      FS->pushOverlay(getToolchainFS());
+    }
     if (FSOverlay)
       FS->pushOverlay(std::move(FSOverlay));
 
diff --git a/sycl-jit/jit-compiler/utils/generate.py b/sycl-jit/jit-compiler/utils/generate.py
@@ -32,11 +32,11 @@ def main():
 const resource_file ToolchainFiles[] = {"""
         )
 
-        def process_file(file_path):
+        def process_file(file_path, relative_to):
             out.write(
                 f"""
 {{
-    {{"{args.prefix}{os.path.relpath(file_path, toolchain_dir).replace(os.sep, "/")}"}} ,
+    {{"{args.prefix}{os.path.relpath(file_path, relative_to).replace(os.sep, "/")}"}} ,
     []() {{
     static const char data[] = {{
     #embed "{file_path}" if_empty(0)
@@ -50,9 +50,17 @@ def process_dir(dir):
             for root, _, files in os.walk(dir):
                 for file in files:
                     file_path = os.path.join(root, file)
-                    process_file(file_path)
+                    process_file(file_path, dir)
 
         process_dir(args.toolchain_dir)
+        process_dir(
+            os.path.realpath(
+                os.path.join(
+                    os.path.dirname(os.path.realpath(__file__)),
+                    "../lib/resource-includes/",
+                )
+            )
+        )
 
         out.write(
             f"""
diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp
@@ -27,8 +27,13 @@ extern __DPCPP_SYCL_EXTERNAL_LIBC char *strncpy(char *dest, const char *src,
 extern __DPCPP_SYCL_EXTERNAL_LIBC int strcmp(const char *s1, const char *s2);
 extern __DPCPP_SYCL_EXTERNAL_LIBC int strncmp(const char *s1, const char *s2,
                                               size_t n);
+#ifdef __LLVM_LIBC__
+extern __DPCPP_SYCL_EXTERNAL_LIBC int rand() noexcept;
+extern __DPCPP_SYCL_EXTERNAL_LIBC void srand(unsigned int seed) noexcept;
+#else
 extern __DPCPP_SYCL_EXTERNAL_LIBC int rand();
 extern __DPCPP_SYCL_EXTERNAL_LIBC void srand(unsigned int seed);
+#endif
 extern __DPCPP_SYCL_EXTERNAL long long int __imf_llmax(long long int x,
                                                        long long int y);
 extern __DPCPP_SYCL_EXTERNAL long long int __imf_llmin(long long int x,
diff --git a/sycl/test-e2e/KernelCompiler/exp_redist_mode.cpp b/sycl/test-e2e/KernelCompiler/exp_redist_mode.cpp
@@ -0,0 +1,65 @@
+// 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
+// CHECK-ERROR-LABEL:  Device compilation failed
+// CHECK-ERROR-NEXT:   Detailed information:
+// CHECK-ERROR:        	In file included from rtc_0.cpp:2:
+// CHECK-ERROR-NEXT:   In file included from /sycl-jit-toolchain//bin/../include/sycl/sycl.hpp:38:
+// CHECK-ERROR-NEXT:   In file included from /sycl-jit-toolchain//bin/../include/sycl/detail/core.hpp:21:
+// CHECK-ERROR-NEXT:   In file included from /sycl-jit-toolchain//bin/../include/sycl/accessor.hpp:11:
+// CHECK-ERROR-NEXT:   /sycl-jit-toolchain//bin/../include/sycl/access/access.hpp:14:10: fatal error: 'type_traits' file not found
+// 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
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main(int argc, char *argv[]) {
+  sycl::queue q;
+  std::string source = R"""(
+    #include <sycl/sycl.hpp>
+    namespace syclext = sycl::ext::oneapi;
+    namespace syclexp = sycl::ext::oneapi::experimental;
+
+    extern "C"
+    SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+    void foo(int *p) {
+      *p = 42;
+    }
+)""";
+  std::vector<std::string> opts;
+  for (int i = 1; i < argc; ++i)
+    opts.emplace_back(argv[i]);
+  try {
+
+    auto kb_src = syclexp::create_kernel_bundle_from_source(
+        q.get_context(), syclexp::source_language::sycl, source);
+    auto kb_exe = syclexp::build(
+        kb_src, syclexp::properties{syclexp::build_options{opts}});
+    sycl::kernel krn = kb_exe.ext_oneapi_get_kernel("foo");
+    auto *p = sycl::malloc_shared<int>(1, q);
+    q.submit([&](sycl::handler &cgh) {
+       cgh.set_args(p);
+       cgh.single_task(krn);
+     }).wait();
+    std::cout << "Result: " << *p << std::endl;
+    assert(*p == 42);
+    sycl::free(p, q);
+  } catch (const sycl::exception &e) {
+    std::cout << e.what();
+    return 1;
+  }
+}
diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp
@@ -6,7 +6,7 @@
 // CHECK-DAG: README.md
 // CHECK-DAG: lit.cfg.py
 //
-// CHECK-NUM-MATCHES: 29
+// CHECK-NUM-MATCHES: 30
 //
 // This test verifies that `<sycl/sycl.hpp>` isn't used in E2E tests. Instead,
 // fine-grained includes should used, see
