Merge remote-tracking branch 'origin/sycl' into llvmspirv_pulldown

Author: jsji

.github/CODEOWNERS

@@ -165,15 +165,16 @@ sycl/test-e2e/SeparateCompile/ @intel/dpcpp-tools-reviewers
 sycl/test-e2e/Printf/ @intel/dpcpp-tools-reviewers @intel/llvm-reviewers-runtime
 sycl/test-e2e/SpecConstants/ @intel/dpcpp-tools-reviewers
 sycl/test-e2e/NewOffloadDriver/ @intel/dpcpp-tools-reviewers
+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

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

@@ -169,7 +169,6 @@ jobs:
           --cmake-opt=-DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
           --cmake-opt="-DLLVM_INSTALL_UTILS=ON" \
           --cmake-opt="-DNATIVECPU_USE_OCK=Off" \
-          --cmake-opt="-DSYCL_PI_TESTS=OFF" \
           --cmake-opt="-DLLVM_EXPERIMENTAL_TARGETS_TO_BUILD=SPIRV"
     - name: Compile
       id: build

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

@@ -162,7 +162,6 @@ jobs:
             runner: '["Linux", "amdgpu"]'
             image: ghcr.io/intel/llvm/ubuntu2204_build:latest
             image_extra_opts: --device=/dev/dri --device=/dev/kfd
-            extra_cmake_args: -DHIP_PLATFORM="AMD" -DAMD_ARCH="gfx1031"
           - name: CUDA system
             runner: '["Linux", "cuda"]'
             image: ghcr.io/intel/llvm/ubuntu2204_build:latest

.github/workflows/sycl-linux-run-tests.yml

@@ -21,9 +21,6 @@ on:
         type: string
         required: True
       extra_cmake_args:
-        description: |
-          If empty, then HIP_PLATFORM and AMD_ARCH would be automatically added
-          if inputs.target_devices contains 'ext_oneapi_hip'
         type: string
         required: False
       tests_selector:
@@ -282,12 +279,6 @@ jobs:
       run: |
         if [ -n "$CMAKE_EXTRA_ARGS" ]; then
           echo "opts=$CMAKE_EXTRA_ARGS" >> $GITHUB_OUTPUT
-        else
-          if [ "${{ contains(inputs.target_devices, 'ext_oneapi_hip')  }}" == "true" ]; then
-            echo 'opts=-DHIP_PLATFORM="AMD" -DAMD_ARCH="gfx1031"' >> $GITHUB_OUTPUT
-          else
-            echo 'opts=' >> $GITHUB_OUTPUT
-          fi
         fi
     - name: Configure E2E tests
       if: inputs.tests_selector == 'e2e'

.github/workflows/sycl-macos-build-and-test.yml

@@ -52,7 +52,6 @@ jobs:
           --ci-defaults $ARGS \
           --cmake-opt=-DCMAKE_C_COMPILER_LAUNCHER=ccache \
           --cmake-opt=-DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
-          --cmake-opt="-DLLVM_INSTALL_UTILS=ON" \
-          --cmake-opt="-DSYCL_PI_TESTS=OFF"
+          --cmake-opt="-DLLVM_INSTALL_UTILS=ON"
     - name: Compile
       run: cmake --build $GITHUB_WORKSPACE/build --target deploy-sycl-toolchain

.github/workflows/sycl-windows-run-tests.yml

@@ -88,6 +88,7 @@ jobs:
         cmake --build build-e2e --target check-sycl-e2e
     - name: Detect hung tests
       shell: powershell
+      if: always()
       run: |
         $exitCode = 0
         $hungTests = Get-Process | Where-Object { ($_.Path -match "llvm\\install") -or ($_.Path -match "llvm\\build-e2e") }

clang/include/clang/AST/PrettyPrinter.h

@@ -68,17 +68,18 @@ struct PrintingPolicy {
         SuppressStrongLifetime(false), SuppressLifetimeQualifiers(false),
         SuppressTypedefs(false), SuppressFinalSpecifier(false),
         SuppressTemplateArgsInCXXConstructors(false),
-        SuppressDefaultTemplateArgs(true), Bool(LO.Bool),
-        Nullptr(LO.CPlusPlus11 || LO.C23), NullptrTypeInNamespace(LO.CPlusPlus),
-        Restrict(LO.C99), Alignof(LO.CPlusPlus11), UnderscoreAlignof(LO.C11),
+        SuppressDefaultTemplateArgs(true), EnforceDefaultTemplateArgs(false),
+        Bool(LO.Bool), Nullptr(LO.CPlusPlus11 || LO.C23),
+        NullptrTypeInNamespace(LO.CPlusPlus), Restrict(LO.C99),
+        Alignof(LO.CPlusPlus11), UnderscoreAlignof(LO.C11),
         UseVoidForZeroParams(!LO.CPlusPlus),
         SplitTemplateClosers(!LO.CPlusPlus11), TerseOutput(false),
         PolishForDeclaration(false), Half(LO.Half),
         MSWChar(LO.MicrosoftExt && !LO.WChar), IncludeNewlines(true),
         MSVCFormatting(false), ConstantsAsWritten(false),
         SuppressImplicitBase(false), FullyQualifiedName(false),
-        SuppressDefinition(false), SuppressDefaultTemplateArguments(false),
-        PrintCanonicalTypes(false),
+        EnforceScopeForElaboratedTypes(false), SuppressDefinition(false),
+        SuppressDefaultTemplateArguments(false), PrintCanonicalTypes(false),
         SkipCanonicalizationOfTemplateTypeParms(false),
         PrintInjectedClassNameWithArguments(true), UsePreferredNames(true),
         AlwaysIncludeTypeForTemplateArgument(false),
@@ -241,6 +242,11 @@ struct PrintingPolicy {
   LLVM_PREFERRED_TYPE(bool)
   unsigned SuppressDefaultTemplateArgs : 1;
 
+  /// When true, print template arguments that match the default argument for
+  /// the parameter, even if they're not specified in the source.
+  LLVM_PREFERRED_TYPE(bool)
+  unsigned EnforceDefaultTemplateArgs : 1;
+
   /// Whether we can use 'bool' rather than '_Bool' (even if the language
   /// doesn't actually have 'bool', because, e.g., it is defined as a macro).
   LLVM_PREFERRED_TYPE(bool)
@@ -339,6 +345,10 @@ struct PrintingPolicy {
   LLVM_PREFERRED_TYPE(bool)
   unsigned FullyQualifiedName : 1;
 
+  /// Enforce fully qualified name printing for elaborated types.
+  LLVM_PREFERRED_TYPE(bool)
+  unsigned EnforceScopeForElaboratedTypes : 1;
+
   /// When true does not print definition of a type. E.g.
   ///   \code
   ///   template<typename T> class C0 : public C1 {...}

clang/lib/AST/TypePrinter.cpp

@@ -101,7 +101,7 @@ class ElaboratedTypePolicyRAII {
     SuppressTagKeyword = Policy.SuppressTagKeyword;
     SuppressScope = Policy.SuppressScope;
     Policy.SuppressTagKeyword = true;
-    Policy.SuppressScope = true;
+    Policy.SuppressScope = !Policy.EnforceScopeForElaboratedTypes;
   }
 
   ~ElaboratedTypePolicyRAII() {
@@ -1728,8 +1728,10 @@ void TypePrinter::printElaboratedBefore(const ElaboratedType *T,
       Policy.SuppressScope = OldSupressScope;
       return;
     }
-    if (Qualifier && !(Policy.SuppressTypedefs &&
-                       T->getNamedType()->getTypeClass() == Type::Typedef))
+    if (Qualifier &&
+        !(Policy.SuppressTypedefs &&
+          T->getNamedType()->getTypeClass() == Type::Typedef) &&
+        !Policy.EnforceScopeForElaboratedTypes)
       Qualifier->print(OS, Policy);
   }
 
@@ -2229,15 +2231,6 @@ static void printArgument(const TemplateArgument &A, const PrintingPolicy &PP,
   A.print(PP, OS, IncludeType);
 }
 
-static void printArgument(const TemplateArgumentLoc &A,
-                          const PrintingPolicy &PP, llvm::raw_ostream &OS,
-                          bool IncludeType) {
-  const TemplateArgument::ArgKind &Kind = A.getArgument().getKind();
-  if (Kind == TemplateArgument::ArgKind::Type)
-    return A.getTypeSourceInfo()->getType().print(OS, PP);
-  return A.getArgument().print(PP, OS, IncludeType);
-}
-
 static bool isSubstitutedTemplateArgument(ASTContext &Ctx, TemplateArgument Arg,
                                           TemplateArgument Pattern,
                                           ArrayRef<TemplateArgument> Args,
@@ -2408,15 +2401,40 @@ template <typename TA>
 static void
 printTo(raw_ostream &OS, ArrayRef<TA> Args, const PrintingPolicy &Policy,
         const TemplateParameterList *TPL, bool IsPack, unsigned ParmIndex) {
-  // Drop trailing template arguments that match default arguments.
-  if (TPL && Policy.SuppressDefaultTemplateArgs &&
-      !Policy.PrintCanonicalTypes && !Args.empty() && !IsPack &&
+  llvm::SmallVector<TemplateArgument, 8> ArgsToPrint;
+  for (const TA &A : Args)
+    ArgsToPrint.push_back(getArgument(A));
+  if (TPL && !Policy.PrintCanonicalTypes && !IsPack &&
       Args.size() <= TPL->size()) {
-    llvm::SmallVector<TemplateArgument, 8> OrigArgs;
-    for (const TA &A : Args)
-      OrigArgs.push_back(getArgument(A));
-    while (!Args.empty() && getArgument(Args.back()).getIsDefaulted())
-      Args = Args.drop_back();
+    // Drop trailing template arguments that match default arguments.
+    if (Policy.SuppressDefaultTemplateArgs) {
+      while (!ArgsToPrint.empty() &&
+             getArgument(ArgsToPrint.back()).getIsDefaulted())
+        ArgsToPrint.pop_back();
+    } else if (Policy.EnforceDefaultTemplateArgs) {
+      for (unsigned I = Args.size(); I < TPL->size(); ++I) {
+        auto Param = TPL->getParam(I);
+        if (auto *TTPD = dyn_cast<TemplateTypeParmDecl>(Param)) {
+          // If we met a non default-argument past provided list of arguments,
+          // it is either a pack which must be the last arguments, or provided
+          // argument list was problematic. Bail out either way. Do the same
+          // for each kind of template argument.
+          if (!TTPD->hasDefaultArgument())
+            break;
+          ArgsToPrint.push_back(getArgument(TTPD->getDefaultArgument()));
+        } else if (auto *TTPD = dyn_cast<TemplateTemplateParmDecl>(Param)) {
+          if (!TTPD->hasDefaultArgument())
+            break;
+          ArgsToPrint.push_back(getArgument(TTPD->getDefaultArgument()));
+        } else if (auto *NTTPD = dyn_cast<NonTypeTemplateParmDecl>(Param)) {
+          if (!NTTPD->hasDefaultArgument())
+            break;
+          ArgsToPrint.push_back(getArgument(NTTPD->getDefaultArgument()));
+        } else {
+          llvm_unreachable("unexpected template parameter");
+        }
+      }
+    }
   }
 
   const char *Comma = Policy.MSVCFormatting ? "," : ", ";
@@ -2425,7 +2443,7 @@ printTo(raw_ostream &OS, ArrayRef<TA> Args, const PrintingPolicy &Policy,
 
   bool NeedSpace = false;
   bool FirstArg = true;
-  for (const auto &Arg : Args) {
+  for (const auto &Arg : ArgsToPrint) {
     // Print the argument into a string.
     SmallString<128> Buf;
     llvm::raw_svector_ostream ArgOS(Buf);

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.
@@ -739,11 +668,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

@@ -10724,12 +10724,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");
@@ -11521,6 +11517,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA,
     if (Args.hasArg(options::OPT_fsycl_embed_ir))
       CmdArgs.push_back(Args.MakeArgString("-sycl-embed-ir"));
 
+    if (Args.hasFlag(options::OPT_fsycl_allow_device_image_dependencies,
+                     options::OPT_fno_sycl_allow_device_image_dependencies,
+                     false))
+      CmdArgs.push_back(
+          Args.MakeArgString("-sycl-allow-device-image-dependencies"));
+
     // Formulate and add any offload-wrapper and AOT specific options. These
     // are additional options passed in via -Xsycl-target-linker and
     // -Xsycl-target-backend.