Merge branch 'sycl' into programbuildrefactor

Author: sergey-semenov

.github/CODEOWNERS

@@ -169,12 +169,12 @@ sycl/test-e2e/LLVMIntrinsicLowering/ @intel/dpcpp-spirv-reviewers
 
 # Sanitizer
 clang/lib/Driver/SanitizerArgs.cpp @intel/dpcpp-sanitizers-review
-libdevice/sanitizer_utils.cpp @intel/dpcpp-sanitizers-review
-libdevice/include/asan_libdevice.hpp @intel/dpcpp-sanitizers-review
-libdevice/include/sanitizer_utils.hpp @intel/dpcpp-sanitizers-review
+libdevice/include/asan_rtl.hpp @intel/dpcpp-sanitizers-review
+libdevice/include/sanitizer_defs.hpp @intel/dpcpp-sanitizers-review
+libdevice/sanitizer/ @intel/dpcpp-sanitizers-review
+llvm/include/llvm/Transforms/Instrumentation/AddressSanitizer.h @intel/dpcpp-sanitizers-review
+llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerCommon.h @intel/dpcpp-sanitizers-review
+llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerOptions.h @intel/dpcpp-sanitizers-review
 llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @intel/dpcpp-sanitizers-review
-sycl/test-e2e/AddressSanitizer/ @intel/dpcpp-sanitizers-review
 llvm/test/Instrumentation/AddressSanitizer/ @intel/dpcpp-sanitizers-review
-llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerOptions.h @intel/dpcpp-sanitizers-review
-llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerCommon.h @intel/dpcpp-sanitizers-review
-llvm/include/llvm/Transforms/Instrumentation/AddressSanitizer.h @intel/dpcpp-sanitizers-review
+sycl/test-e2e/AddressSanitizer/ @intel/dpcpp-sanitizers-review

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -350,34 +350,6 @@ llvm::Type *CodeGenTypes::ConvertFunctionTypeInternal(QualType QFT) {
   return ResultType;
 }
 
-template <bool NeedTypeInterpret = false>
-llvm::Type *getJointMatrixINTELExtType(llvm::Type *CompTy,
-                                       ArrayRef<TemplateArgument> TemplateArgs,
-                                       const unsigned Val = 0) {
-  // TODO: we should actually have exactly 5 template parameters: 1 for
-  // type and 4 for type parameters. But in previous version of the SPIR-V
-  // spec we have Layout matrix type parameter, that was later removed.
-  // Once we update to the newest version of the spec - this should be updated.
-  assert((TemplateArgs.size() == 5 || TemplateArgs.size() == 6) &&
-         "Wrong JointMatrixINTEL template parameters number");
-  // This is required to represent optional 'Component Type Interpretation'
-  // parameter
-  std::vector<unsigned> Params;
-  for (size_t I = 1; I != TemplateArgs.size(); ++I) {
-    assert(TemplateArgs[I].getKind() == TemplateArgument::Integral &&
-           "Wrong JointMatrixINTEL template parameter");
-    Params.push_back(TemplateArgs[I].getAsIntegral().getExtValue());
-  }
-  // Don't add type interpretation for legacy matrices.
-  // Legacy matrices has 5 template parameters, while new representation
-  // has 6.
-  if (NeedTypeInterpret && TemplateArgs.size() != 5)
-    Params.push_back(Val);
-
-  return llvm::TargetExtType::get(CompTy->getContext(),
-                                  "spirv.JointMatrixINTEL", {CompTy}, Params);
-}
-
 llvm::Type *
 getCooperativeMatrixKHRExtType(llvm::Type *CompTy,
                                ArrayRef<TemplateArgument> TemplateArgs) {
@@ -394,49 +366,6 @@ getCooperativeMatrixKHRExtType(llvm::Type *CompTy,
       CompTy->getContext(), "spirv.CooperativeMatrixKHR", {CompTy}, Params);
 }
 
-/// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type
-/// which is represented as a pointer to a structure to LLVM extension type
-/// with the parameters that follow SPIR-V JointMatrixINTEL type.
-/// The expected representation is:
-/// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%,
-///        %use%, (optional) %element_type_interpretation%)
-llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) {
-  auto *TemplateDecl = cast<ClassTemplateSpecializationDecl>(RD);
-  ArrayRef<TemplateArgument> TemplateArgs =
-      TemplateDecl->getTemplateArgs().asArray();
-  assert(TemplateArgs[0].getKind() == TemplateArgument::Type &&
-         "1st JointMatrixINTEL template parameter must be type");
-  llvm::Type *CompTy = ConvertType(TemplateArgs[0].getAsType());
-
-  // Per JointMatrixINTEL spec the type can have an optional
-  // 'Component Type Interpretation' parameter. We should emit it in case
-  // if on SYCL level joint matrix accepts 'bfloat16' or 'tf32' objects as
-  // matrix's components. Yet 'bfloat16' should be represented as 'int16' and
-  // 'tf32' as 'float' types.
-  if (CompTy->isStructTy()) {
-    StringRef LlvmTyName = CompTy->getStructName();
-    // Emit half/int16/float for sycl[::*]::{half,bfloat16,tf32}
-    if (LlvmTyName.starts_with("class.sycl::") ||
-        LlvmTyName.starts_with("class.__sycl_internal::"))
-      LlvmTyName = LlvmTyName.rsplit("::").second;
-    if (LlvmTyName == "half") {
-      CompTy = llvm::Type::getHalfTy(getLLVMContext());
-      return getJointMatrixINTELExtType(CompTy, TemplateArgs);
-    } else if (LlvmTyName == "tf32") {
-      CompTy = llvm::Type::getFloatTy(getLLVMContext());
-      // 'tf32' interpretation is mapped to '0'
-      return getJointMatrixINTELExtType<true>(CompTy, TemplateArgs, 0);
-    } else if (LlvmTyName == "bfloat16") {
-      CompTy = llvm::Type::getInt16Ty(getLLVMContext());
-      // 'bfloat16' interpretation is mapped to '1'
-      return getJointMatrixINTELExtType<true>(CompTy, TemplateArgs, 1);
-    } else {
-      llvm_unreachable("Wrong matrix base type!");
-    }
-  }
-  return getJointMatrixINTELExtType(CompTy, TemplateArgs);
-}
-
 /// ConvertSPVCooperativeMatrixType - Convert SYCL joint_matrix type
 /// which is represented as a pointer to a structure to LLVM extension type
 /// with the parameters that follow SPIR-V CooperativeMatrixKHR type.
@@ -733,11 +662,7 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
       if (ClangETy && ClangETy->isStructureOrClassType()) {
         RecordDecl *RD = ClangETy->getAsCXXRecordDecl();
         if (RD && RD->getQualifiedNameAsString() ==
-                      "__spv::__spirv_JointMatrixINTEL") {
-          ResultType = ConvertSYCLJointMatrixINTELType(RD);
-          break;
-        } else if (RD && RD->getQualifiedNameAsString() ==
-                             "__spv::__spirv_CooperativeMatrixKHR") {
+                      "__spv::__spirv_CooperativeMatrixKHR") {
           ResultType = ConvertSPVCooperativeMatrixType(RD);
           break;
         } else if (RD && RD->getQualifiedNameAsString() ==

clang/lib/CodeGen/CodeGenTypes.h

@@ -145,14 +145,6 @@ class CodeGenTypes {
   /// load/store type are the same.
   llvm::Type *convertTypeForLoadStore(QualType T, llvm::Type *LLVMTy = nullptr);
 
-  /// ConvertSYCLJointMatrixINTELType - Convert SYCL joint_matrix type
-  /// which is represented as a pointer to a structure to LLVM extension type
-  /// with the parameters that follow SPIR-V JointMatrixINTEL type.
-  /// The expected representation is:
-  /// target("spirv.JointMatrixINTEL", %element_type, %rows%, %cols%, %scope%,
-  ///        %use%, (optional) %element_type_interpretation%)
-  llvm::Type *ConvertSYCLJointMatrixINTELType(RecordDecl *RD);
-
   /// ConvertSPVCooperativeMatrixType - Convert SYCL joint_matrix type
   /// which is represented as a pointer to a structure to LLVM extension type
   /// with the parameters that follow SPIR-V CooperativeMatrixKHR type.

clang/lib/Driver/OffloadBundler.cpp

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

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10674,12 +10674,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
                                          ArgStringList &TranslatorArgs) {
   bool IsCPU = Triple.isSPIR() &&
                Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64;
-  // Enable NonSemanticShaderDebugInfo.200 for CPU AOT and for non-Windows
-  const bool IsWindowsMSVC =
-      Triple.isWindowsMSVCEnvironment() ||
-      C.getDefaultToolChain().getTriple().isWindowsMSVCEnvironment();
-  const bool EnableNonSemanticDebug =
-      IsCPU || (!IsWindowsMSVC && !C.getDriver().IsFPGAHWMode());
+  // Enable NonSemanticShaderDebugInfo.200 for non-FPGA targets.
+  const bool EnableNonSemanticDebug = !C.getDriver().IsFPGAHWMode();
   if (EnableNonSemanticDebug) {
     TranslatorArgs.push_back(
         "-spirv-debug-info-version=nonsemantic-shader-200");

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -658,43 +658,44 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
     addLibraries(SYCLDeviceAnnotationLibs);
 
 #if !defined(_WIN32)
+  std::string SanitizeVal;
   size_t sanitizer_lib_idx = getSingleBuildTarget();
   if (Arg *A = Args.getLastArg(options::OPT_fsanitize_EQ,
                                options::OPT_fno_sanitize_EQ)) {
     if (A->getOption().matches(options::OPT_fsanitize_EQ) &&
-        A->getValues().size() == 1) {
-      std::string SanitizeVal = A->getValue();
-      if (SanitizeVal == "address")
-        addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
-    }
+        A->getValues().size() == 1)
+      SanitizeVal = A->getValue();
   } else {
     // User can pass -fsanitize=address to device compiler via
     // -Xsycl-target-frontend, sanitize device library must be
     // linked with user's device image if so.
-    bool IsDeviceAsanEnabled = false;
-    auto SyclFEArg = Args.getAllArgValues(options::OPT_Xsycl_frontend);
-    IsDeviceAsanEnabled = (std::count(SyclFEArg.begin(), SyclFEArg.end(),
-                                      "-fsanitize=address") > 0);
-    if (!IsDeviceAsanEnabled) {
-      auto SyclFEArgEq = Args.getAllArgValues(options::OPT_Xsycl_frontend_EQ);
-      IsDeviceAsanEnabled = (std::count(SyclFEArgEq.begin(), SyclFEArgEq.end(),
-                                        "-fsanitize=address") > 0);
-    }
-
-    // User can also enable asan for SYCL device via -Xarch_device option.
-    if (!IsDeviceAsanEnabled) {
-      auto DeviceArchVals = Args.getAllArgValues(options::OPT_Xarch_device);
-      for (auto DArchVal : DeviceArchVals) {
-        if (DArchVal.find("-fsanitize=address") != std::string::npos) {
-          IsDeviceAsanEnabled = true;
-          break;
-        }
+    std::vector<std::string> EnabledDeviceSanitizers;
+
+    // NOTE: "-fsanitize=" applies to all device targets
+    auto SyclFEArgVals = Args.getAllArgValues(options::OPT_Xsycl_frontend);
+    auto SyclFEEQArgVals = Args.getAllArgValues(options::OPT_Xsycl_frontend_EQ);
+    auto ArchDeviceVals = Args.getAllArgValues(options::OPT_Xarch_device);
+
+    std::vector<std::string> ArgVals(
+        SyclFEArgVals.size() + SyclFEEQArgVals.size() + ArchDeviceVals.size());
+    ArgVals.insert(ArgVals.end(), SyclFEArgVals.begin(), SyclFEArgVals.end());
+    ArgVals.insert(ArgVals.end(), SyclFEEQArgVals.begin(),
+                   SyclFEEQArgVals.end());
+    ArgVals.insert(ArgVals.end(), ArchDeviceVals.begin(), ArchDeviceVals.end());
+
+    // Driver will report error if address sanitizer and memory sanitizer are
+    // both enabled, so we only need to check first one here.
+    for (const std::string &Arg : ArgVals) {
+      if (Arg.find("-fsanitize=address") != std::string::npos) {
+        SanitizeVal = "address";
+        break;
       }
     }
-
-    if (IsDeviceAsanEnabled)
-      addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
   }
+
+  if (SanitizeVal == "address")
+    addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
+
 #endif
 
   if (isNativeCPU)