Merge remote-tracking branch 'upstream/sycl' into embed_fallback_svp

Author: jinge90

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

@@ -1625,6 +1625,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),
@@ -1636,6 +1653,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 &&
@@ -1676,6 +1696,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() {

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-7dad678",
-      "version": "7dad678",
-      "updated_at": "2024-11-24T10:48:51Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2229466354/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