Merge branch 'sycl' into review/yang/restructure_asan_msan

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

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-0b4b682",
-      "version": "0b4b682",
-      "updated_at": "2024-11-17T01:09:50Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2197388704/zip",
+      "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",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

llvm/include/llvm/ADT/StringRef.h

@@ -263,11 +263,6 @@ namespace llvm {
       return Length >= Prefix.Length &&
              compareMemory(Data, Prefix.Data, Prefix.Length) == 0;
     }
-    [[nodiscard]] LLVM_DEPRECATED(
-        "Use starts_with instead",
-        "starts_with") bool startswith(StringRef Prefix) const {
-      return starts_with(Prefix);
-    }
     [[nodiscard]] bool starts_with(char Prefix) const {
       return !empty() && front() == Prefix;
     }
@@ -281,11 +276,6 @@ namespace llvm {
              compareMemory(end() - Suffix.Length, Suffix.Data, Suffix.Length) ==
                  0;
     }
-    [[nodiscard]] LLVM_DEPRECATED(
-        "Use ends_with instead",
-        "ends_with") bool endswith(StringRef Suffix) const {
-      return ends_with(Suffix);
-    }
     [[nodiscard]] bool ends_with(char Suffix) const {
       return !empty() && back() == Suffix;
     }

llvm/test/CodeGen/SPIRV/transcoding/TransFNeg.ll

@@ -1,8 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
-; https://github.com/intel/llvm/issues/14372
-; UNSUPPORTED: windows
-
 ; CHECK-SPIRV: OpFNegate
 ; CHECK-SPIRV: OpFNegate
 ; CHECK-SPIRV: OpFNegate

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -3,4 +3,4 @@
 # Date:   Tue Oct 22 03:30:08 2024 -0700
 #
 #     Support UR program creation from multiple device binaries (#2147)
-set(UNIFIED_RUNTIME_TAG review/yang/restructure_asan)
+set(UNIFIED_RUNTIME_TAG review/yang/restructure_asan_msan)

sycl/include/sycl/__spirv/spirv_ops.hpp

@@ -27,155 +27,6 @@
 
 extern __DPCPP_SYCL_EXTERNAL float __spirv_RoundFToTF32INTEL(float a);
 
-#ifndef __SPIRV_USE_COOPERATIVE_MATRIX
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *
-    __spirv_JointMatrixLoadINTEL(T *Ptr, std::size_t Stride,
-                                 __spv::MatrixLayout Layout = L,
-                                 __spv::Scope::Flag Sc = S, int MemOperand = 0);
-
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL void __spirv_JointMatrixStoreINTEL(
-    T *Ptr, __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *Object,
-    std::size_t Stride, __spv::MatrixLayout Layout = L,
-    __spv::Scope::Flag Sc = S, int MemOperand = 0);
-
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *
-    __spirv_CooperativeMatrixConstructCheckedINTEL(int32_t CoordX,
-                                                   int32_t CoordY,
-                                                   uint32_t Height,
-                                                   uint32_t Width,
-                                                   const T Value);
-
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *
-    __spirv_CooperativeMatrixLoadCheckedINTEL(
-        T *Ptr, int32_t CoordX, int32_t CoordY, __spv::MatrixLayout Layout = L,
-        uint32_t Height = 0, uint32_t Width = 0, std::size_t Stride = 0,
-        int MemOperand = 0);
-
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL(
-    T *Ptr, int32_t CoordX, int32_t CoordY,
-    __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *Object,
-    __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0,
-    std::size_t Stride = 0, int MemOperand = 0);
-
-template <typename TA, typename TB, typename TC, std::size_t M, std::size_t K,
-          std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
-          __spv::MatrixUse UC,
-          __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<TC, M, N, LC, S, UC> *
-    __spirv_JointMatrixMadINTEL(
-        __spv::__spirv_JointMatrixINTEL<TA, M, K, LA, S, UA> *A,
-        __spv::__spirv_JointMatrixINTEL<TB, K, N, LB, S, UB> *B,
-        __spv::__spirv_JointMatrixINTEL<TC, M, N, LC, S, UC> *C,
-        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
-
-template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
-          std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
-          __spv::MatrixUse UC,
-          __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-    __spirv_JointMatrixUUMadINTEL(
-        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
-
-template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
-          std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
-          __spv::MatrixUse UC,
-          __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-    __spirv_JointMatrixUSMadINTEL(
-        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
-
-template <typename T1, typename T2, typename T3, std::size_t M, std::size_t K,
-          std::size_t N, __spv::MatrixUse UA, __spv::MatrixUse UB,
-          __spv::MatrixUse UC,
-          __spv::MatrixLayout LA = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LB = __spv::MatrixLayout::RowMajor,
-          __spv::MatrixLayout LC = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *
-    __spirv_JointMatrixSUMadINTEL(
-        __spv::__spirv_JointMatrixINTEL<T1, M, K, LA, S, UA> *A,
-        __spv::__spirv_JointMatrixINTEL<T2, K, N, LB, S, UB> *B,
-        __spv::__spirv_JointMatrixINTEL<T3, M, N, LC, S, UC> *C,
-        __spv::Scope::Flag Sc = __spv::Scope::Flag::Subgroup);
-
-template <typename T, typename Tp, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL
-    __spv::__spirv_JointMatrixINTEL<Tp, R, C, L, S, U> *
-    __spirv_CompositeConstruct(const T v);
-
-template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL __ocl_vec_t<uint32_t, 2>
-__spirv_JointMatrixGetElementCoordINTEL(
-    __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);
-
-template <typename T, std::size_t R, std::size_t C, __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL size_t __spirv_JointMatrixWorkItemLengthINTEL(
-    __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *);
-
-template <typename Ts, typename T, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL Ts __spirv_VectorExtractDynamic(
-    __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *, size_t i);
-
-template <typename Ts, typename T, std::size_t R, std::size_t C,
-          __spv::MatrixUse U,
-          __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
-          __spv::Scope::Flag S = __spv::Scope::Flag::Subgroup>
-extern __DPCPP_SYCL_EXTERNAL __spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *
-__spirv_VectorInsertDynamic(__spv::__spirv_JointMatrixINTEL<T, R, C, L, S, U> *,
-                            Ts val, size_t i);
-#else  // __SPIRV_USE_COOPERATIVE_MATRIX
 template <typename T, typename Tp, std::size_t R, std::size_t C,
           __spv::MatrixUse U,
           __spv::MatrixLayout L = __spv::MatrixLayout::RowMajor,
@@ -304,7 +155,6 @@ extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixStoreCheckedINTEL(
     __spv::__spirv_CooperativeMatrixKHR<Tp, S, R, C, U> *Object,
     __spv::MatrixLayout Layout = L, uint32_t Height = 0, uint32_t Width = 0,
     std::size_t Stride = 0, int MemOperand = 0);
-#endif // __SPIRV_USE_COOPERATIVE_MATRIX
 
 template <typename T>
 extern __DPCPP_SYCL_EXTERNAL void __spirv_CooperativeMatrixPrefetchINTEL(