Merge branch 'sycl' into llvmspirv_pulldown

.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",

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -1789,6 +1789,9 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn,
     if (SyclOptReport.HasOptReportInfo(FD)) {
       llvm::OptimizationRemarkEmitter ORE(Fn);
       for (auto ORI : llvm::enumerate(SyclOptReport.GetInfo(FD))) {
+        // Temporarily apply arg location to ensure SourceLocToDebugLoc
+        // picks up the expected file.
+        ApplyDebugLocation TempApplyLoc(*this, ORI.value().KernelArgLoc);
         llvm::DiagnosticLocation DL =
             SourceLocToDebugLoc(ORI.value().KernelArgLoc);
         StringRef NameInDesc = ORI.value().KernelArgDescName;

clang/lib/Driver/OffloadBundler.cpp

@@ -687,12 +687,12 @@ class ObjectFileHandler final : public FileHandler {
         if (Error Err = Symbol.printName(NameOS))
           return std::move(Err);
 
-        // If we are dealing with a bitcode file do not add special globals
-        // llvm.used and llvm.compiler.used and __AsanDeviceGlobalMetadata to
+        // If we are dealing with a bitcode file do not add special globals to
         // the list of defined symbols.
         if (SF->isIR() &&
             (Name == "llvm.used" || Name == "llvm.compiler.used" ||
-             Name == "__AsanDeviceGlobalMetadata"))
+             Name == "__AsanDeviceGlobalMetadata" ||
+             Name == "__AsanKernelMetadata"))
           continue;
 
         // Add symbol name with the target prefix to the buffer.

clang/lib/Driver/SanitizerArgs.cpp

@@ -1197,6 +1197,9 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
 
       CmdArgs.push_back("-mllvm");
       CmdArgs.push_back("-asan-mapping-scale=4");
+
+      addSpecialCaseListOpt(Args, CmdArgs,
+                            "-fsanitize-ignorelist=", UserIgnorelistFiles);
     }
     return;
   }

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1618,6 +1618,23 @@ static std::vector<OptSpecifier> getUnsupportedOpts(void) {
   return UnsupportedOpts;
 }
 
+// Currently supported options by SYCL NativeCPU device compilation
+static inline bool SupportedByNativeCPU(const SYCLToolChain &TC,
+                                        const OptSpecifier &Opt) {
+  if (!TC.IsSYCLNativeCPU)
+    return false;
+
+  switch (Opt.getID()) {
+  case options::OPT_fcoverage_mapping:
+  case options::OPT_fno_coverage_mapping:
+  case options::OPT_fprofile_instr_generate:
+  case options::OPT_fprofile_instr_generate_EQ:
+  case options::OPT_fno_profile_instr_generate:
+    return true;
+  }
+  return false;
+}
+
 SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
                              const ToolChain &HostTC, const ArgList &Args)
     : ToolChain(D, Triple, Args), HostTC(HostTC),
@@ -1629,6 +1646,9 @@ SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
   // Diagnose unsupported options only once.
   for (OptSpecifier Opt : getUnsupportedOpts()) {
     if (const Arg *A = Args.getLastArg(Opt)) {
+      // Native CPU can support options unsupported by other targets.
+      if (SupportedByNativeCPU(*this, Opt))
+        continue;
       // All sanitizer options are not currently supported, except
       // AddressSanitizer
       if (A->getOption().getID() == options::OPT_fsanitize_EQ &&
@@ -1669,6 +1689,9 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
     bool Unsupported = false;
     for (OptSpecifier UnsupportedOpt : getUnsupportedOpts()) {
       if (Opt.matches(UnsupportedOpt)) {
+        // NativeCPU should allow most normal cpu options.
+        if (SupportedByNativeCPU(*this, Opt.getID()))
+          continue;
         if (Opt.getID() == options::OPT_fsanitize_EQ &&
             A->getValues().size() == 1) {
           std::string SanitizeVal = A->getValue();

clang/lib/Sema/SemaSYCL.cpp

@@ -3625,8 +3625,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
                      FinalizeStmts.end());
 
+    SourceLocation LL = NewBody ? NewBody->getBeginLoc() : SourceLocation();
+    SourceLocation LR = NewBody ? NewBody->getEndLoc() : SourceLocation();
+
     return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts,
-                                FPOptionsOverride(), {}, {});
+                                FPOptionsOverride(), LL, LR);
   }
 
   void annotateHierarchicalParallelismAPICalls() {
@@ -6969,6 +6972,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
   Policy.adjustForCPlusPlusFwdDecl();
   Policy.SuppressTypedefs = true;
   Policy.SuppressUnwrittenScope = true;
+  Policy.PrintCanonicalTypes = true;
 
   llvm::SmallSet<const VarDecl *, 8> Visited;
   bool EmittedFirstSpecConstant = false;

clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp

@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -emit-llvm %s -o -
+// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
+
+// This test checks that integration footer is emitted correctly when a
+// device_global has an explicit template specialization in template arguments.
+
+#include "sycl.hpp"
+
+namespace sycl {
+template <typename T> struct X {};
+template <> struct X<int> {};
+namespace detail {
+struct Y {};
+} // namespace detail
+template <> struct X<detail::Y> {};
+} // namespace sycl
+
+using namespace sycl;
+template <typename T, typename = X<detail::Y>> struct Arg1 { T val; };
+
+using namespace sycl::ext::oneapi;
+template <typename properties_t>
+device_global<properties_t> dev_global;
+
+SYCL_EXTERNAL auto foo() {
+  (void)dev_global<Arg1<int>>;
+}
+
+// CHECK-FOOTER: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::dev_global<Arg1<int, sycl::X<sycl::detail::Y>>>, "_Z10dev_globalI4Arg1IiN4sycl1XINS1_6detail1YEEEEE");
+// CHECK-FOOTER-NEXT: }
+// CHECK-FOOTER-NEXT: } // namespace (unnamed)

clang/test/Driver/sycl-native-cpu.cpp

@@ -24,3 +24,9 @@
 
 // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -### 2>&1 | FileCheck -check-prefix=CHECK-NONATIVECPU %s
 // CHECK-NONATIVECPU-NOT: "-D" "__SYCL_NATIVE_CPU__"
+
+// Checking that coverage testing options are accepted by native_cpu, and that device and host compilation invocations receive the same options
+// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Werror -fno-profile-instr-generate -fprofile-instr-generate -fno-coverage-mapping -fcoverage-mapping -### %s 2>&1 | FileCheck %s --check-prefix=CHECK_COV_INVO
+// CHECK_COV_INVO:{{.*}}clang{{.*}}-fsycl-is-device{{.*}}"-fsycl-is-native-cpu" "-D" "__SYCL_NATIVE_CPU__"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
+// CHECK_COV_INVO:{{.*}}clang{{.*}}"-fsycl-is-host"{{.*}}"-fprofile-instrument=clang"{{.*}}"-fcoverage-mapping" "-fcoverage-compilation-dir={{.*}}"
+

clang/test/SemaSYCL/kernel_functor_location.cpp

@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+//
+// Checks that the compound statement of the implicitly generated kernel body
+// has a valid source location (containing "line"). Previously this location
+// was invalid containing "<<invalid sloc>>" which causes asserts in the
+// llvm profiling tools.
+
+#include "Inputs/sycl.hpp"
+
+struct Functor {
+  void operator()() const {}
+};
+
+// CHECK: FunctionDecl {{.*}} _ZTS7Functor 'void ()'
+// CHECK-NEXT: |-CompoundStmt {{.*}} <{{.*}}line{{.*}}>
+
+int main() {
+
+  sycl::queue().submit([&](sycl::handler &cgh) {
+    cgh.single_task(Functor{});
+  });
+}

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-6ee988a",
+      "version": "6ee988a",
+      "updated_at": "2024-11-26T15:44:10Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2239640503/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

libdevice/sanitizer/asan_rtl.cpp

@@ -665,11 +665,16 @@ constexpr size_t AlignMask(size_t n) { return n - 1; }
 ///
 /// ASAN Load/Store Report Built-ins
 ///
+/// NOTE:
+///   if __AsanLaunchInfo equals 0, the sanitizer is disabled for this launch
+///
 
 #define ASAN_REPORT_ERROR_BASE(type, is_write, size, as)                       \
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_as##as(                  \
       uptr addr, const char __SYCL_CONSTANT__ *file, uint32_t line,            \
       const char __SYCL_CONSTANT__ *func) {                                    \
+    if (!__AsanLaunchInfo)                                                     \
+      return;                                                                  \
     if (addr & AlignMask(size)) {                                              \
       __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \
                                    func);                                      \
@@ -682,6 +687,8 @@ constexpr size_t AlignMask(size_t n) { return n - 1; }
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_as##as##_noabort(        \
       uptr addr, const char __SYCL_CONSTANT__ *file, uint32_t line,            \
       const char __SYCL_CONSTANT__ *func) {                                    \
+    if (!__AsanLaunchInfo)                                                     \
+      return;                                                                  \
     if (addr & AlignMask(size)) {                                              \
       __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \
                                    func, true);                                \
@@ -714,6 +721,8 @@ ASAN_REPORT_ERROR(store, true, 16)
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##N_as##as(                       \
       uptr addr, size_t size, const char __SYCL_CONSTANT__ *file,              \
       uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
+    if (!__AsanLaunchInfo)                                                     \
+      return;                                                                  \
     if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) {      \
       __asan_report_access_error(addr, as, size, is_write, poisoned_addr,      \
                                  file, line, func);                            \
@@ -722,6 +731,8 @@ ASAN_REPORT_ERROR(store, true, 16)
   DEVICE_EXTERN_C_NOINLINE void __asan_##type##N_as##as##_noabort(             \
       uptr addr, size_t size, const char __SYCL_CONSTANT__ *file,              \
       uint32_t line, const char __SYCL_CONSTANT__ *func) {                     \
+    if (!__AsanLaunchInfo)                                                     \
+      return;                                                                  \
     if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) {      \
       __asan_report_access_error(addr, as, size, is_write, poisoned_addr,      \
                                  file, line, func, true);                      \
@@ -743,6 +754,9 @@ ASAN_REPORT_ERROR_N(store, true)
 ///
 
 DEVICE_EXTERN_C_NOINLINE uptr __asan_mem_to_shadow(uptr ptr, uint32_t as) {
+  if (!__AsanLaunchInfo)
+    return 0;
+
   return MemToShadow(ptr, as);
 }
 
@@ -756,6 +770,9 @@ static __SYCL_CONSTANT__ const char __mem_set_shadow_local[] =
 DEVICE_EXTERN_C_NOINLINE void
 __asan_set_shadow_static_local(uptr ptr, size_t size,
                                size_t size_with_redzone) {
+  if (!__AsanLaunchInfo)
+    return;
+
   // Since ptr is aligned to ASAN_SHADOW_GRANULARITY,
   // if size != aligned_size, then the buffer tail of ptr is not aligned
   uptr aligned_size = RoundUpTo(size, ASAN_SHADOW_GRANULARITY);
@@ -795,6 +812,9 @@ static __SYCL_CONSTANT__ const char __mem_unpoison_shadow_static_local_end[] =
 DEVICE_EXTERN_C_NOINLINE void
 __asan_unpoison_shadow_static_local(uptr ptr, size_t size,
                                     size_t size_with_redzone) {
+  if (!__AsanLaunchInfo)
+    return;
+
   ASAN_DEBUG(__spirv_ocl_printf(__mem_unpoison_shadow_static_local_begin));
 
   auto shadow_begin = MemToShadow(ptr + size, ADDRESS_SPACE_LOCAL);
@@ -828,6 +848,9 @@ static __SYCL_CONSTANT__ const char __mem_report_arg_count_incorrect[] =
 
 DEVICE_EXTERN_C_NOINLINE void
 __asan_set_shadow_dynamic_local(uptr ptr, uint32_t num_args) {
+  if (!__AsanLaunchInfo)
+    return;
+
   ASAN_DEBUG(__spirv_ocl_printf(__mem_set_shadow_dynamic_local_begin));
 
   auto *launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
@@ -859,6 +882,9 @@ static __SYCL_CONSTANT__ const char __mem_unpoison_shadow_dynamic_local_end[] =
 
 DEVICE_EXTERN_C_NOINLINE void
 __asan_unpoison_shadow_dynamic_local(uptr ptr, uint32_t num_args) {
+  if (!__AsanLaunchInfo)
+    return;
+
   ASAN_DEBUG(__spirv_ocl_printf(__mem_unpoison_shadow_dynamic_local_begin));
 
   auto *launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;
@@ -895,6 +921,9 @@ static __SYCL_CONSTANT__ const char __mem_set_shadow_private[] =
 
 DEVICE_EXTERN_C_NOINLINE void __asan_set_shadow_private(uptr begin, uptr size,
                                                         char val) {
+  if (!__AsanLaunchInfo)
+    return;
+
   ASAN_DEBUG(__spirv_ocl_printf(__mem_set_shadow_private_begin));
 
   auto *launch_info = (__SYCL_GLOBAL__ const LaunchInfo *)__AsanLaunchInfo;

llvm/include/llvm/SYCLLowerIR/AsanKernelMetadata.h

@@ -0,0 +1,23 @@
+//===-- AsanKernelMetadata.h - fix kernel medatadata for sanitizer ---===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This pass fixes attributes and metadata of the global variable
+// "__AsanKernelMetadata"
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class AsanKernelMetadataPass : public PassInfoMixin<AsanKernelMetadataPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
+};
+
+} // namespace llvm

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

@@ -25,6 +25,10 @@ constexpr char GENX_KERNEL_METADATA[] = "genx.kernels";
 // sycl/ext/oneapi/experimental/invoke_simd.hpp::__builtin_invoke_simd
 // overloads instantiations:
 constexpr char INVOKE_SIMD_PREF[] = "_Z33__regcall3____builtin_invoke_simd";
+// The regexp for ESIMD intrinsics:
+// /^_Z(\d+)__esimd_\w+/
+static constexpr char ESIMD_INTRIN_PREF0[] = "_Z";
+static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_";
 
 bool isSlmAllocatorConstructor(const Function &F);
 bool isSlmAllocatorDestructor(const Function &F);
@@ -133,5 +137,9 @@ struct UpdateUint64MetaDataToMaxValue {
 // functions has changed its attribute to alwaysinline.
 bool prepareForAlwaysInliner(Module &M);
 
+// Remove mangling from an ESIMD intrinsic function.
+// Returns empty on pattern match failure.
+StringRef stripMangling(StringRef FName);
+
 } // namespace esimd
 } // namespace llvm