Merge branch 'sycl' into root_group_query

.github/CODEOWNERS

@@ -41,11 +41,11 @@ sycl/include/sycl/detail/ur.hpp @intel/unified-runtime-reviewers
 sycl/source/detail/posix_ur.cpp @intel/unified-runtime-reviewers
 sycl/source/detail/ur.cpp @intel/unified-runtime-reviewers
 sycl/source/detail/windows_ur.cpp @intel/unified-runtime-reviewers
-sycl/test-e2e/Plugin/ @intel/unified-runtime-reviewers
+sycl/test-e2e/Adapters/ @intel/unified-runtime-reviewers
 
 # Win Proxy Loader
-sycl/pi_win_proxy_loader @intel/llvm-reviewers-runtime
-sycl/test-e2e/Plugin/dll-detach-order.cpp @intel/llvm-reviewers-runtime
+sycl/ur_win_proxy_loader @intel/llvm-reviewers-runtime
+sycl/test-e2e/Adapters/dll-detach-order.cpp @intel/llvm-reviewers-runtime
 
 # CUDA specific runtime implementations
 sycl/include/sycl/ext/oneapi/experimental/cuda/ @intel/llvm-reviewers-cuda

.github/workflows/sycl-linux-precommit.yml

@@ -80,6 +80,12 @@ jobs:
             image: ghcr.io/intel/llvm/ubuntu2204_build:latest
             image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
             target_devices: ext_oneapi_cuda:gpu
+          - name: AMD/HIP
+            runner: '["Linux", "amdgpu"]'
+            image: ghcr.io/intel/llvm/ubuntu2204_build:latest
+            image_options: -u 1001 --device=/dev/dri --device=/dev/kfd
+            target_devices: ext_oneapi_hip:gpu
+            reset_intel_gpu: false
           - name: Intel
             runner: '["Linux", "gen12"]'
             image: ghcr.io/intel/llvm/ubuntu2204_intel_drivers:latest

CONTRIBUTING.md

@@ -58,7 +58,7 @@ To contribute:
   - [The seven rules of a great Git commit message](https://cbea.ms/git-commit)
     are recommended read and follow.
 - To a reasonable extent, title tags can be used to signify the component
-  changed, e.g.: `[PI]`, `[CUDA]`, `[Doc]`.
+  changed, e.g.: `[UR]`, `[CUDA]`, `[Doc]`.
 - Create a pull request (PR) for your changes following
   [Creating a pull request instructions](https://help.github.com/articles/creating-a-pull-request/).
   - Make sure PR has a good description explaining all of the changes made,

buildbot/configure.py

@@ -69,7 +69,7 @@ def do_configure(args):
     if sys.platform != "darwin":
         sycl_enabled_backends.append("level_zero")
 
-    # lld is needed on Windows or for the HIP plugin on AMD
+    # lld is needed on Windows or for the HIP adapter on AMD
     if platform.system() == "Windows" or (args.hip and args.hip_platform == "AMD"):
         llvm_enable_projects += ";lld"
 
@@ -152,8 +152,8 @@ def do_configure(args):
                 libclc_targets_to_build += libclc_nvidia_target_names
             libclc_gen_remangled_variants = "ON"
 
-    if args.enable_plugin:
-        sycl_enabled_backends += args.enable_plugin
+    if args.enable_backends:
+        sycl_enabled_backends += args.enable_backends
 
     if args.disable_preview_lib:
         sycl_preview_lib = "OFF"
@@ -374,7 +374,9 @@ def main():
     parser.add_argument(
         "--ci-defaults", action="store_true", help="Enable default CI parameters"
     )
-    parser.add_argument("--enable-plugin", action="append", help="Enable SYCL plugin")
+    parser.add_argument(
+        "--enable-backends", action="append", help="Enable SYCL backend"
+    )
     parser.add_argument(
         "--disable-preview-lib",
         action="store_true",

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-ac93a93",
-      "version": "ac93a93",
-      "updated_at": "2024-11-21T02:09:35Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2216471673/zip",
+      "github_tag": "igc-dev-7dad678",
+      "version": "7dad678",
+      "updated_at": "2024-11-24T10:48:51Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2229466354/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

sycl-jit/common/include/Kernel.h

@@ -17,6 +17,7 @@
 #include <cstdint>
 #include <cstring>
 #include <functional>
+#include <string_view>
 #include <type_traits>
 
 namespace jit_compiler {
@@ -350,11 +351,60 @@ struct SYCLKernelInfo {
       : Name{KernelName}, Args{NumArgs}, Attributes{}, NDR{}, BinaryInfo{} {}
 };
 
+// RTC-related datastructures
+// TODO: Consider moving into separate header.
+
 struct InMemoryFile {
   const char *Path;
   const char *Contents;
 };
 
+using RTCBundleBinaryInfo = SYCLKernelBinaryInfo;
+using FrozenSymbolTable = DynArray<sycl::detail::string>;
+
+// Note: `FrozenPropertyValue` and `FrozenPropertySet` constructors take
+// `std::string_view` arguments instead of `const char *` because they will be
+// created from `llvm::SmallString`s, which don't contain the trailing '\0'
+// byte. Hence obtaining a C-string would cause an additional copy.
+
+struct FrozenPropertyValue {
+  sycl::detail::string Name;
+  bool IsUIntValue;
+  uint32_t UIntValue;
+  DynArray<uint8_t> Bytes;
+
+  FrozenPropertyValue() = default;
+  FrozenPropertyValue(FrozenPropertyValue &&) = default;
+  FrozenPropertyValue &operator=(FrozenPropertyValue &&) = default;
+
+  FrozenPropertyValue(std::string_view Name, uint32_t Value)
+      : Name{Name}, IsUIntValue{true}, UIntValue{Value}, Bytes{0} {}
+  FrozenPropertyValue(std::string_view Name, const uint8_t *Ptr, size_t Size)
+      : Name{Name}, IsUIntValue{false}, Bytes{Size} {
+    std::memcpy(Bytes.begin(), Ptr, Size);
+  }
+};
+
+struct FrozenPropertySet {
+  sycl::detail::string Name;
+  DynArray<FrozenPropertyValue> Values;
+
+  FrozenPropertySet() = default;
+  FrozenPropertySet(FrozenPropertySet &&) = default;
+  FrozenPropertySet &operator=(FrozenPropertySet &&) = default;
+
+  FrozenPropertySet(std::string_view Name, size_t Size)
+      : Name{Name}, Values{Size} {}
+};
+
+using FrozenPropertyRegistry = DynArray<FrozenPropertySet>;
+
+struct RTCBundleInfo {
+  RTCBundleBinaryInfo BinaryInfo;
+  FrozenSymbolTable SymbolTable;
+  FrozenPropertyRegistry Properties;
+};
+
 } // namespace jit_compiler
 
 #endif // SYCL_FUSION_COMMON_KERNEL_H

sycl-jit/jit-compiler/CMakeLists.txt

@@ -31,6 +31,7 @@ add_llvm_library(sycl-jit
    Target
    TargetParser
    MC
+   SYCLLowerIR
    ${LLVM_TARGETS_TO_BUILD}
 
    LINK_LIBS

sycl-jit/jit-compiler/include/KernelFusion.h

@@ -56,6 +56,32 @@ class JITResult {
   sycl::detail::string ErrorMessage;
 };
 
+class RTCResult {
+public:
+  explicit RTCResult(const char *ErrorMessage)
+      : Failed{true}, BundleInfo{}, ErrorMessage{ErrorMessage} {}
+
+  explicit RTCResult(RTCBundleInfo &&BundleInfo)
+      : Failed{false}, BundleInfo{std::move(BundleInfo)}, ErrorMessage{} {}
+
+  bool failed() const { return Failed; }
+
+  const char *getErrorMessage() const {
+    assert(failed() && "No error message present");
+    return ErrorMessage.c_str();
+  }
+
+  const RTCBundleInfo &getBundleInfo() const {
+    assert(!failed() && "No bundle info");
+    return BundleInfo;
+  }
+
+private:
+  bool Failed;
+  RTCBundleInfo BundleInfo;
+  sycl::detail::string ErrorMessage;
+};
+
 extern "C" {
 
 #ifdef __clang__
@@ -77,7 +103,7 @@ KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
     const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo,
     View<unsigned char> SpecConstBlob);
 
-KF_EXPORT_SYMBOL JITResult compileSYCL(InMemoryFile SourceFile,
+KF_EXPORT_SYMBOL RTCResult compileSYCL(InMemoryFile SourceFile,
                                        View<InMemoryFile> IncludeFiles,
                                        View<const char *> UserArgs);
 

sycl-jit/jit-compiler/lib/KernelFusion.cpp

@@ -25,8 +25,8 @@ using namespace jit_compiler;
 using FusedFunction = helper::FusionHelper::FusedFunction;
 using FusedFunctionList = std::vector<FusedFunction>;
 
-static JITResult errorToFusionResult(llvm::Error &&Err,
-                                     const std::string &Msg) {
+template <typename ResultType>
+static ResultType errorTo(llvm::Error &&Err, const std::string &Msg) {
   std::stringstream ErrMsg;
   ErrMsg << Msg << "\nDetailed information:\n";
   llvm::handleAllErrors(std::move(Err),
@@ -35,7 +35,7 @@ static JITResult errorToFusionResult(llvm::Error &&Err,
                           // compiled without exception support.
                           ErrMsg << "\t" << StrErr.getMessage() << "\n";
                         });
-  return JITResult{ErrMsg.str().c_str()};
+  return ResultType{ErrMsg.str().c_str()};
 }
 
 static std::vector<jit_compiler::NDRange>
@@ -95,7 +95,7 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
       translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(),
                                                  ModuleInfo.kernels());
   if (auto Error = ModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error), "Failed to load kernels");
+    return errorTo<JITResult>(std::move(Error), "Failed to load kernels");
   }
   std::unique_ptr<llvm::Module> NewMod = std::move(*ModOrError);
   if (!fusion::FusionPipeline::runMaterializerPasses(
@@ -107,8 +107,8 @@ extern "C" KF_EXPORT_SYMBOL JITResult materializeSpecConstants(
   SYCLKernelInfo &MaterializerKernelInfo = *ModuleInfo.getKernelFor(KernelName);
   if (auto Error = translation::KernelTranslator::translateKernel(
           MaterializerKernelInfo, *NewMod, JITCtx, TargetFormat)) {
-    return errorToFusionResult(std::move(Error),
-                               "Translation to output format failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Translation to output format failed");
   }
 
   return JITResult{MaterializerKernelInfo};
@@ -133,7 +133,7 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   llvm::Expected<jit_compiler::FusedNDRange> FusedNDR =
       jit_compiler::FusedNDRange::get(NDRanges);
   if (llvm::Error Err = FusedNDR.takeError()) {
-    return errorToFusionResult(std::move(Err), "Illegal ND-range combination");
+    return errorTo<JITResult>(std::move(Err), "Illegal ND-range combination");
   }
 
   if (!isTargetFormatSupported(TargetFormat)) {
@@ -180,7 +180,7 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
       translation::KernelTranslator::loadKernels(*JITCtx.getLLVMContext(),
                                                  ModuleInfo.kernels());
   if (auto Error = ModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error), "SPIR-V translation failed");
+    return errorTo<JITResult>(std::move(Error), "SPIR-V translation failed");
   }
   std::unique_ptr<llvm::Module> LLVMMod = std::move(*ModOrError);
 
@@ -197,8 +197,8 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   llvm::Expected<std::unique_ptr<llvm::Module>> NewModOrError =
       helper::FusionHelper::addFusedKernel(LLVMMod.get(), FusedKernelList);
   if (auto Error = NewModOrError.takeError()) {
-    return errorToFusionResult(std::move(Error),
-                               "Insertion of fused kernel stub failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Insertion of fused kernel stub failed");
   }
   std::unique_ptr<llvm::Module> NewMod = std::move(*NewModOrError);
 
@@ -221,8 +221,8 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
 
   if (auto Error = translation::KernelTranslator::translateKernel(
           FusedKernelInfo, *NewMod, JITCtx, TargetFormat)) {
-    return errorToFusionResult(std::move(Error),
-                               "Translation to output format failed");
+    return errorTo<JITResult>(std::move(Error),
+                              "Translation to output format failed");
   }
 
   FusedKernelInfo.NDR = FusedNDR->getNDR();
@@ -234,37 +234,47 @@ fuseKernels(View<SYCLKernelInfo> KernelInformation, const char *FusedKernelName,
   return JITResult{FusedKernelInfo};
 }
 
-extern "C" KF_EXPORT_SYMBOL JITResult
+extern "C" KF_EXPORT_SYMBOL RTCResult
 compileSYCL(InMemoryFile SourceFile, View<InMemoryFile> IncludeFiles,
             View<const char *> UserArgs) {
   auto UserArgListOrErr = parseUserArgs(UserArgs);
   if (!UserArgListOrErr) {
-    return errorToFusionResult(UserArgListOrErr.takeError(),
-                               "Parsing of user arguments failed");
+    return errorTo<RTCResult>(UserArgListOrErr.takeError(),
+                              "Parsing of user arguments failed");
   }
   llvm::opt::InputArgList UserArgList = std::move(*UserArgListOrErr);
 
   auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgList);
   if (!ModuleOrErr) {
-    return errorToFusionResult(ModuleOrErr.takeError(),
-                               "Device compilation failed");
+    return errorTo<RTCResult>(ModuleOrErr.takeError(),
+                              "Device compilation failed");
   }
 
   std::unique_ptr<llvm::LLVMContext> Context;
   std::unique_ptr<llvm::Module> Module = std::move(*ModuleOrErr);
   Context.reset(&Module->getContext());
 
   if (auto Error = linkDeviceLibraries(*Module, UserArgList)) {
-    return errorToFusionResult(std::move(Error), "Device linking failed");
+    return errorTo<RTCResult>(std::move(Error), "Device linking failed");
   }
 
-  SYCLKernelInfo Kernel;
-  if (auto Error = translation::KernelTranslator::translateKernel(
-          Kernel, *Module, JITContext::getInstance(), BinaryFormat::SPIRV)) {
-    return errorToFusionResult(std::move(Error), "SPIR-V translation failed");
+  auto BundleInfoOrError = performPostLink(*Module, UserArgList);
+  if (!BundleInfoOrError) {
+    return errorTo<RTCResult>(BundleInfoOrError.takeError(),
+                              "Post-link phase failed");
+  }
+  auto BundleInfo = std::move(*BundleInfoOrError);
+
+  auto BinaryInfoOrError =
+      translation::KernelTranslator::translateBundleToSPIRV(
+          *Module, JITContext::getInstance());
+  if (!BinaryInfoOrError) {
+    return errorTo<RTCResult>(BinaryInfoOrError.takeError(),
+                              "SPIR-V translation failed");
   }
+  BundleInfo.BinaryInfo = std::move(*BinaryInfoOrError);
 
-  return JITResult{Kernel};
+  return RTCResult{std::move(BundleInfo)};
 }
 
 extern "C" KF_EXPORT_SYMBOL void resetJITConfiguration() {