diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index 9dd094e85d253..129e73345b2a2 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -1794,7 +1794,7 @@ correctly in any circumstances. It can be used if:
   metaprogramming algorithms to be able to specify/detect types generically.
 
 - the generated kernel binary does not contain indirect calls because they
-  are eliminated using compiler optimizations e.g. devirtualization. 
+  are eliminated using compiler optimizations e.g. devirtualization.
 
 - the selected target supports the function pointer like functionality e.g.
   most CPU targets.
@@ -2404,29 +2404,39 @@ argument.
   int *pb =__builtin_preserve_access_index(&v->c[3].b);
   __builtin_preserve_access_index(v->j);
 
-``__builtin_unique_stable_name``
---------------------------------
+``__builtin_sycl_unique_stable_name``
+-------------------------------------
 
-``__builtin_unique_stable_name()`` is a builtin that takes a type or expression and
-produces a string literal containing a unique name for the type (or type of the
-expression) that is stable across split compilations.
+``__builtin_sycl_unique_stable_name()`` is a builtin that takes a type and
+produces a string literal containing a unique name for the type that is stable
+across split compilations, mainly to support SYCL/Data Parallel C++ language.
 
 In cases where the split compilation needs to share a unique token for a type
 across the boundary (such as in an offloading situation), this name can be used
-for lookup purposes.
+for lookup purposes, such as in the SYCL Integration Header.
+
+The value of this builtin is computed entirely at compile time, so it can be
+used in constant expressions. This value encodes lambda functions based on a
+stable numbering order in which they appear in their local declaration contexts.
+Once this builtin is evaluated in a constexpr context, it is erroneous to use
+it in an instantiation which changes its value.
+
+In order to produce the unique name, the current implementation of the bultin
+uses Itanium mangling even if the host compilation uses a different name
+mangling scheme at runtime. The mangler marks all the lambdas required to name
+the SYCL kernel and emits a stable local ordering of the respective lambdas,
+starting from ``10000``. The initial value of ``10000`` serves as an obvious
+differentiator from ordinary lambda mangling numbers but does not serve any
+other purpose and may change in the future. The resulting pattern is
+demanglable. When non-lambda types are passed to the builtin, the mangler emits
+their usual pattern without any special treatment.
+
+**Syntax**:
 
-This builtin is superior to RTTI for this purpose for two reasons.  First, this
-value is computed entirely at compile time, so it can be used in constant
-expressions. Second, this value encodes lambda functions based on line-number
-rather than the order in which it appears in a function. This is valuable
-because it is stable in cases where an unrelated lambda is introduced
-conditionally in the same function.
+.. code-block:: c
 
-The current implementation of this builtin uses a slightly modified Itanium
-Mangler to produce the unique name. The lambda ordinal is replaced with one or
-more line/column pairs in the format ``LINE->COL``, separated with a ``~``
-character. Typically, only one pair will be included, however in the case of
-macro expansions the entire macro expansion stack is expressed.
+  // Computes a unique stable name for the given type.
+  constexpr const char * __builtin_sycl_unique_stable_name( type-id );
 
 Multiprecision Arithmetic Builtins
 ----------------------------------
@@ -2622,7 +2632,7 @@ Guaranteed inlined copy
 ``__builtin_memcpy_inline`` has been designed as a building block for efficient
 ``memcpy`` implementations. It is identical to ``__builtin_memcpy`` but also
 guarantees not to call any external functions. See LLVM IR `llvm.memcpy.inline
-<https://llvm.org/docs/LangRef.html#llvm-memcpy-inline-intrinsic>`_ intrinsic 
+<https://llvm.org/docs/LangRef.html#llvm-memcpy-inline-intrinsic>`_ intrinsic
 for more information.
 
 This is useful to implement a custom version of ``memcpy``, implement a
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index 55f15003acab8..25d3008ff4346 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -103,6 +103,7 @@ class DynTypedNode;
 class DynTypedNodeList;
 class Expr;
 class GlobalDecl;
+class ItaniumMangleContext;
 class MangleContext;
 class MangleNumberingContext;
 class MaterializeTemporaryExpr;
@@ -2360,6 +2361,12 @@ class ASTContext : public RefCountedBase<ASTContext> {
   /// If \p T is null pointer, assume the target in ASTContext.
   MangleContext *createMangleContext(const TargetInfo *T = nullptr);
 
+  /// Creates a device mangle context to correctly mangle lambdas in a mixed
+  /// architecture compile by setting the lambda mangling number source to the
+  /// DeviceLambdaManglingNumber. Currently this asserts that the TargetInfo
+  /// (from the AuxTargetInfo) is a an itanium target.
+  MangleContext *createDeviceMangleContext(const TargetInfo &T);
+
   void DeepCollectObjCIvars(const ObjCInterfaceDecl *OI, bool leafClass,
                             SmallVectorImpl<const ObjCIvarDecl*> &Ivars) const;
 
@@ -3163,10 +3170,31 @@ OPT_LIST(V)
 
   StringRef getCUIDHash() const;
 
+  void AddSYCLKernelNamingDecl(const CXXRecordDecl *RD);
+  bool IsSYCLKernelNamingDecl(const NamedDecl *RD) const;
+  unsigned GetSYCLKernelNamingIndex(const NamedDecl *RD);
+  /// A SourceLocation to store whether we have evaluated a kernel name already,
+  /// and where it happened.  If so, we need to diagnose an illegal use of the
+  /// builtin.
+  llvm::MapVector<const SYCLUniqueStableNameExpr *, std::string>
+      SYCLUniqueStableNameEvaluatedValues;
+
 private:
   /// All OMPTraitInfo objects live in this collection, one per
   /// `pragma omp [begin] declare variant` directive.
   SmallVector<std::unique_ptr<OMPTraitInfo>, 4> OMPTraitInfoVector;
+
+  /// A list of the (right now just lambda decls) declarations required to
+  /// name all the SYCL kernels in the translation unit, so that we can get the
+  /// correct kernel name, as well as implement
+  /// __builtin_sycl_unique_stable_name.
+  llvm::DenseMap<const DeclContext *,
+                 llvm::SmallPtrSet<const CXXRecordDecl *, 4>>
+      SYCLKernelNamingTypes;
+  std::unique_ptr<ItaniumMangleContext> SYCLKernelFilterContext;
+  void FilterSYCLKernelNamingDecls(
+      const CXXRecordDecl *RD,
+      llvm::SmallVectorImpl<const CXXRecordDecl *> &Decls);
 };
 
 /// Insertion operator for diagnostics.
diff --git a/clang/include/clang/AST/ComputeDependence.h b/clang/include/clang/AST/ComputeDependence.h
index 04e8e2c7d2ccb..7dde42ee71ba0 100644
--- a/clang/include/clang/AST/ComputeDependence.h
+++ b/clang/include/clang/AST/ComputeDependence.h
@@ -78,6 +78,7 @@ class MaterializeTemporaryExpr;
 class CXXFoldExpr;
 class TypeTraitExpr;
 class ConceptSpecializationExpr;
+class SYCLUniqueStableNameExpr;
 class PredefinedExpr;
 class CallExpr;
 class OffsetOfExpr;
@@ -165,6 +166,7 @@ ExprDependence computeDependence(TypeTraitExpr *E);
 ExprDependence computeDependence(ConceptSpecializationExpr *E,
                                  bool ValueDependent);
 
+ExprDependence computeDependence(SYCLUniqueStableNameExpr *E);
 ExprDependence computeDependence(PredefinedExpr *E);
 ExprDependence computeDependence(CallExpr *E, llvm::ArrayRef<Expr *> PreArgs);
 ExprDependence computeDependence(OffsetOfExpr *E);
diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h
index 9453b447270bb..55ee96feae98c 100644
--- a/clang/include/clang/AST/Expr.h
+++ b/clang/include/clang/AST/Expr.h
@@ -1948,17 +1948,13 @@ class StringLiteral final
 /// [C99 6.4.2.2] - A predefined identifier such as __func__.
 class PredefinedExpr final
     : public Expr,
-      private llvm::TrailingObjects<PredefinedExpr, Stmt *, Expr *,
-                                    TypeSourceInfo *> {
+      private llvm::TrailingObjects<PredefinedExpr, Stmt *> {
   friend class ASTStmtReader;
   friend TrailingObjects;
 
   // PredefinedExpr is optionally followed by a single trailing
   // "Stmt *" for the predefined identifier. It is present if and only if
   // hasFunctionName() is true and is always a "StringLiteral *".
-  // It can also be followed by a Expr* in the case of a
-  // __builtin_unique_stable_name with an expression, or TypeSourceInfo * if
-  // __builtin_unique_stable_name with a type.
 
 public:
   enum IdentKind {
@@ -1971,18 +1967,12 @@ class PredefinedExpr final
     PrettyFunction,
     /// The same as PrettyFunction, except that the
     /// 'virtual' keyword is omitted for virtual member functions.
-    PrettyFunctionNoVirtual,
-    UniqueStableNameType,
-    UniqueStableNameExpr,
+    PrettyFunctionNoVirtual
   };
 
 private:
   PredefinedExpr(SourceLocation L, QualType FNTy, IdentKind IK,
                  StringLiteral *SL);
-  PredefinedExpr(SourceLocation L, QualType FNTy, IdentKind IK,
-                 TypeSourceInfo *Info);
-  PredefinedExpr(SourceLocation L, QualType FNTy, IdentKind IK,
-                 Expr *E);
 
   explicit PredefinedExpr(EmptyShell Empty, bool HasFunctionName);
 
@@ -1995,39 +1985,10 @@ class PredefinedExpr final
     *getTrailingObjects<Stmt *>() = SL;
   }
 
-  void setTypeSourceInfo(TypeSourceInfo *Info) {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameType &&
-           "TypeSourceInfo only valid for UniqueStableName of a Type");
-    *getTrailingObjects<TypeSourceInfo *>() = Info;
-  }
-
-  void setExpr(Expr *E) {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameExpr &&
-           "TypeSourceInfo only valid for UniqueStableName of n Expression.");
-    *getTrailingObjects<Expr *>() = E;
-  }
-
-  size_t numTrailingObjects(OverloadToken<Stmt *>) const {
-    return hasFunctionName();
-  }
-
-  size_t numTrailingObjects(OverloadToken<TypeSourceInfo *>) const {
-    return getIdentKind() == UniqueStableNameType && !hasFunctionName();
-  }
-  size_t numTrailingObjects(OverloadToken<Expr *>) const {
-    return getIdentKind() == UniqueStableNameExpr && !hasFunctionName();
-  }
-
 public:
   /// Create a PredefinedExpr.
   static PredefinedExpr *Create(const ASTContext &Ctx, SourceLocation L,
                                 QualType FNTy, IdentKind IK, StringLiteral *SL);
-  static PredefinedExpr *Create(const ASTContext &Ctx, SourceLocation L,
-                                QualType FNTy, IdentKind IK, StringLiteral *SL,
-                                TypeSourceInfo *Info);
-  static PredefinedExpr *Create(const ASTContext &Ctx, SourceLocation L,
-                                QualType FNTy, IdentKind IK, StringLiteral *SL,
-                                Expr *E);
 
   /// Create an empty PredefinedExpr.
   static PredefinedExpr *CreateEmpty(const ASTContext &Ctx,
@@ -2052,38 +2013,12 @@ class PredefinedExpr final
                : nullptr;
   }
 
-  TypeSourceInfo *getTypeSourceInfo() {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameType &&
-           "TypeSourceInfo only valid for UniqueStableName of a Type");
-    return *getTrailingObjects<TypeSourceInfo *>();
-  }
-
-  const TypeSourceInfo *getTypeSourceInfo() const {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameType &&
-           "TypeSourceInfo only valid for UniqueStableName of a Type");
-    return *getTrailingObjects<TypeSourceInfo *>();
-  }
-
-  Expr *getExpr() {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameExpr &&
-           "TypeSourceInfo only valid for UniqueStableName of n Expression.");
-    return *getTrailingObjects<Expr *>();
-  }
-
-  const Expr *getExpr() const {
-    assert(!hasFunctionName() && getIdentKind() == UniqueStableNameExpr &&
-           "TypeSourceInfo only valid for UniqueStableName of n Expression.");
-    return *getTrailingObjects<Expr *>();
-  }
-
   static StringRef getIdentKindName(IdentKind IK);
   StringRef getIdentKindName() const {
     return getIdentKindName(getIdentKind());
   }
 
   static std::string ComputeName(IdentKind IK, const Decl *CurrentDecl);
-  static std::string ComputeName(ASTContext &Context, IdentKind IK,
-                                 const QualType Ty);
 
   SourceLocation getBeginLoc() const { return getLocation(); }
   SourceLocation getEndLoc() const { return getLocation(); }
@@ -2104,6 +2039,64 @@ class PredefinedExpr final
   }
 };
 
+// This represents a use of the __builtin_sycl_unique_stable_name, which takes a
+// type-id, and at CodeGen time emits a unique string representation of the
+// type in a way that permits us to properly encode information about the SYCL
+// kernels.
+class SYCLUniqueStableNameExpr final : public Expr {
+  friend class ASTStmtReader;
+  SourceLocation OpLoc, LParen, RParen;
+  TypeSourceInfo *TypeInfo;
+
+  SYCLUniqueStableNameExpr(EmptyShell Empty, QualType ResultTy);
+  SYCLUniqueStableNameExpr(SourceLocation OpLoc, SourceLocation LParen,
+                           SourceLocation RParen, QualType ResultTy,
+                           TypeSourceInfo *TSI);
+
+  void setTypeSourceInfo(TypeSourceInfo *Ty) { TypeInfo = Ty; }
+
+  void setLocation(SourceLocation L) { OpLoc = L; }
+  void setLParenLocation(SourceLocation L) { LParen = L; }
+  void setRParenLocation(SourceLocation L) { RParen = L; }
+
+public:
+  TypeSourceInfo *getTypeSourceInfo() { return TypeInfo; }
+
+  const TypeSourceInfo *getTypeSourceInfo() const { return TypeInfo; }
+
+  static SYCLUniqueStableNameExpr *
+  Create(const ASTContext &Ctx, SourceLocation OpLoc, SourceLocation LParen,
+         SourceLocation RParen, TypeSourceInfo *TSI);
+
+  static SYCLUniqueStableNameExpr *CreateEmpty(const ASTContext &Ctx);
+
+  SourceLocation getBeginLoc() const { return getLocation(); }
+  SourceLocation getEndLoc() const { return RParen; }
+  SourceLocation getLocation() const { return OpLoc; }
+  SourceLocation getLParenLocation() const { return LParen; }
+  SourceLocation getRParenLocation() const { return RParen; }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == SYCLUniqueStableNameExprClass;
+  }
+
+  // Iterators
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+
+  const_child_range children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  // Convenience function to generate the name of the currently stored type.
+  std::string ComputeName(ASTContext &Context) const;
+
+  // Get the generated name of the type.  Note that this only works after all
+  // kernels have been instantiated.
+  static std::string ComputeName(ASTContext &Context, QualType Ty);
+};
+
 /// ParenExpr - This represents a parethesized expression, e.g. "(1)".  This
 /// AST node is only formed if full location information is requested.
 class ParenExpr : public Expr {
diff --git a/clang/include/clang/AST/JSONNodeDumper.h b/clang/include/clang/AST/JSONNodeDumper.h
index 4e7162992418a..a8b731a7cd31d 100644
--- a/clang/include/clang/AST/JSONNodeDumper.h
+++ b/clang/include/clang/AST/JSONNodeDumper.h
@@ -263,6 +263,7 @@ class JSONNodeDumper
   void VisitBlockDecl(const BlockDecl *D);
 
   void VisitDeclRefExpr(const DeclRefExpr *DRE);
+  void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E);
   void VisitPredefinedExpr(const PredefinedExpr *PE);
   void VisitUnaryOperator(const UnaryOperator *UO);
   void VisitBinaryOperator(const BinaryOperator *BO);
diff --git a/clang/include/clang/AST/Mangle.h b/clang/include/clang/AST/Mangle.h
index be0e89730ab1f..7d02f08e0120c 100644
--- a/clang/include/clang/AST/Mangle.h
+++ b/clang/include/clang/AST/Mangle.h
@@ -107,9 +107,6 @@ class MangleContext {
   virtual bool shouldMangleCXXName(const NamedDecl *D) = 0;
   virtual bool shouldMangleStringLiteral(const StringLiteral *SL) = 0;
 
-  virtual bool isDeviceMangleContext() const { return false; }
-  virtual void setDeviceMangleContext(bool) {}
-
   virtual bool isUniqueInternalLinkageDecl(const NamedDecl *ND) {
     return false;
   }
@@ -172,14 +169,11 @@ class MangleContext {
 };
 
 class ItaniumMangleContext : public MangleContext {
-  bool IsUniqueNameMangler = false;
 public:
+  using DiscriminatorOverrideTy =
+      llvm::Optional<unsigned> (*)(ASTContext &, const NamedDecl *);
   explicit ItaniumMangleContext(ASTContext &C, DiagnosticsEngine &D)
       : MangleContext(C, D, MK_Itanium) {}
-  explicit ItaniumMangleContext(ASTContext &C, DiagnosticsEngine &D,
-                                bool IsUniqueNameMangler)
-      : MangleContext(C, D, MK_Itanium),
-        IsUniqueNameMangler(IsUniqueNameMangler) {}
 
   virtual void mangleCXXVTable(const CXXRecordDecl *RD, raw_ostream &) = 0;
   virtual void mangleCXXVTT(const CXXRecordDecl *RD, raw_ostream &) = 0;
@@ -200,15 +194,18 @@ class ItaniumMangleContext : public MangleContext {
 
   virtual void mangleDynamicStermFinalizer(const VarDecl *D, raw_ostream &) = 0;
 
-  bool isUniqueNameMangler() { return IsUniqueNameMangler; }
-
+  // This has to live here, otherwise the CXXNameMangler won't have access to
+  // it.
+  virtual DiscriminatorOverrideTy getDiscriminatorOverride() const = 0;
   static bool classof(const MangleContext *C) {
     return C->getKind() == MK_Itanium;
   }
 
+  static ItaniumMangleContext *create(ASTContext &Context,
+                                      DiagnosticsEngine &Diags);
   static ItaniumMangleContext *create(ASTContext &Context,
                                       DiagnosticsEngine &Diags,
-                                      bool IsUniqueNameMangler = false);
+                                      DiscriminatorOverrideTy Discriminator);
 };
 
 class MicrosoftMangleContext : public MangleContext {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 4770a2fd42a2a..b3ea6f7817ef9 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -2651,6 +2651,9 @@ DEF_TRAVERSE_STMT(ObjCBridgedCastExpr, {
 DEF_TRAVERSE_STMT(ObjCAvailabilityCheckExpr, {})
 DEF_TRAVERSE_STMT(ParenExpr, {})
 DEF_TRAVERSE_STMT(ParenListExpr, {})
+DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
+  TRY_TO(TraverseTypeLoc(S->getTypeSourceInfo()->getTypeLoc()));
+})
 DEF_TRAVERSE_STMT(PredefinedExpr, {})
 DEF_TRAVERSE_STMT(ShuffleVectorExpr, {})
 DEF_TRAVERSE_STMT(ConvertVectorExpr, {})
diff --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 0955c494a74ad..e3aa94d5a1e83 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -249,6 +249,7 @@ class TextNodeDumper
   void VisitCastExpr(const CastExpr *Node);
   void VisitImplicitCastExpr(const ImplicitCastExpr *Node);
   void VisitDeclRefExpr(const DeclRefExpr *Node);
+  void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *Node);
   void VisitPredefinedExpr(const PredefinedExpr *Node);
   void VisitCharacterLiteral(const CharacterLiteral *Node);
   void VisitIntegerLiteral(const IntegerLiteral *Node);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ae636aa9db5c3..ea3e3a0583310 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6394,6 +6394,11 @@ def warn_pointer_arith_null_ptr : Warning<
 def warn_gnu_null_ptr_arith : Warning<
   "arithmetic on a null pointer treated as a cast from integer to pointer is a GNU extension">,
   InGroup<NullPointerArithmetic>, DefaultIgnore;
+def err_kernel_invalidates_sycl_unique_stable_name
+    : Error<"kernel instantiation changes the result of an evaluated "
+            "'__builtin_sycl_unique_stable_name'">;
+def note_sycl_unique_stable_name_evaluated_here
+    : Note<"'__builtin_sycl_unique_stable_name' evaluated here">;
 
 def warn_floatingpoint_eq : Warning<
   "comparing floating point with == or != is unsafe">,
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 21003b6838574..cb8b02eeadac7 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -462,6 +462,8 @@ class LangOptions : public LangOptionsBase {
   bool hasWasmExceptions() const {
     return getExceptionHandling() == ExceptionHandlingKind::Wasm;
   }
+
+  bool isSYCL() const { return SYCLIsDevice || SYCLIsHost; }
 };
 
 /// Floating point control options
diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index ecaf7b027e778..6134e60f35414 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -57,6 +57,7 @@ def CoreturnStmt : StmtNode<Stmt>;
 // Expressions
 def Expr : StmtNode<ValueStmt, 1>;
 def PredefinedExpr : StmtNode<Expr>;
+def SYCLUniqueStableNameExpr : StmtNode<Expr>;
 def DeclRefExpr : StmtNode<Expr>;
 def IntegerLiteral : StmtNode<Expr>;
 def FixedPointLiteral : StmtNode<Expr>;
diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def
index 09207a9a3d6e8..9113f0c9c8651 100644
--- a/clang/include/clang/Basic/TokenKinds.def
+++ b/clang/include/clang/Basic/TokenKinds.def
@@ -695,12 +695,12 @@ ALIAS("_declspec"        , __declspec , KEYMS)
 ALIAS("_pascal"      , __pascal   , KEYBORLAND)
 
 // Clang Extensions.
-KEYWORD(__builtin_convertvector     , KEYALL)
-ALIAS("__char16_t"   , char16_t     , KEYCXX)
-ALIAS("__char32_t"   , char32_t     , KEYCXX)
-KEYWORD(__builtin_bit_cast          , KEYALL)
-KEYWORD(__builtin_available         , KEYALL)
-KEYWORD(__builtin_unique_stable_name, KEYALL)
+KEYWORD(__builtin_convertvector          , KEYALL)
+ALIAS("__char16_t"   , char16_t          , KEYCXX)
+ALIAS("__char32_t"   , char32_t          , KEYCXX)
+KEYWORD(__builtin_bit_cast               , KEYALL)
+KEYWORD(__builtin_available              , KEYALL)
+KEYWORD(__builtin_sycl_unique_stable_name, KEYSYCL)
 
 // Clang-specific keywords enabled only in testing.
 TESTING_KEYWORD(__unknown_anytype , KEYALL)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 215e0e36509da..bfc2016df645e 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -1800,7 +1800,7 @@ class Parser : public CodeCompletionHandler {
   ExprResult ParsePostfixExpressionSuffix(ExprResult LHS);
   ExprResult ParseUnaryExprOrTypeTraitExpression();
   ExprResult ParseBuiltinPrimaryExpression();
-  ExprResult ParseUniqueStableNameExpression();
+  ExprResult ParseSYCLUniqueStableNameExpression();
 
   ExprResult ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
                                                      bool &isCastExpr,
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index d754e47a738fb..7113f3a998450 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -1069,6 +1069,10 @@ class Sema final {
     OpaqueParser = P;
   }
 
+  // Does the work necessary to deal with a SYCL kernel lambda. At the moment,
+  // this just marks the list of lambdas required to name the kernel.
+  void AddSYCLKernelLambda(const FunctionDecl *FD);
+
   class DelayedDiagnostics;
 
   class DelayedDiagnosticsState {
@@ -5421,14 +5425,14 @@ class Sema final {
   ExprResult ActOnPredefinedExpr(SourceLocation Loc, tok::TokenKind Kind);
   ExprResult ActOnIntegerConstant(SourceLocation Loc, uint64_t Val);
 
-  ExprResult BuildUniqueStableName(SourceLocation Loc, TypeSourceInfo *Operand);
-  ExprResult BuildUniqueStableName(SourceLocation Loc, Expr *E);
-  ExprResult ActOnUniqueStableNameExpr(SourceLocation OpLoc,
-                                       SourceLocation LParen,
-                                       SourceLocation RParen, ParsedType Ty);
-  ExprResult ActOnUniqueStableNameExpr(SourceLocation OpLoc,
-                                       SourceLocation LParen,
-                                       SourceLocation RParen, Expr *E);
+  ExprResult BuildSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                           SourceLocation LParen,
+                                           SourceLocation RParen,
+                                           TypeSourceInfo *TSI);
+  ExprResult ActOnSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                           SourceLocation LParen,
+                                           SourceLocation RParen,
+                                           ParsedType ParsedTy);
 
   bool CheckLoopHintExpr(Expr *E, SourceLocation Loc);
 
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index 6d0fde4418841..eb1fff1cca6ad 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1965,6 +1965,9 @@ enum StmtCode {
 
   // FixedPointLiteral
   EXPR_FIXEDPOINT_LITERAL,
+
+  // SYCLUniqueStableNameExpr
+  EXPR_SYCL_UNIQUE_STABLE_NAME,
 };
 
 /// The kinds of designators that can occur in a
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 4f07c5a180800..84c9627987de9 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -2470,7 +2470,7 @@ unsigned ASTContext::getPreferredTypeAlign(const Type *T) const {
   // The preferred alignment of member pointers is that of a pointer.
   if (T->isMemberPointerType())
     return getPreferredTypeAlign(getPointerDiffType().getTypePtr());
- 
+
   if (!Target->allowsLargerPreferedTypeAlignment())
     return ABIAlign;
 
@@ -11116,6 +11116,33 @@ MangleContext *ASTContext::createMangleContext(const TargetInfo *T) {
   llvm_unreachable("Unsupported ABI");
 }
 
+MangleContext *ASTContext::createDeviceMangleContext(const TargetInfo &T) {
+  assert(T.getCXXABI().getKind() != TargetCXXABI::Microsoft &&
+         "Device mangle context does not support Microsoft mangling.");
+  switch (T.getCXXABI().getKind()) {
+  case TargetCXXABI::AppleARM64:
+  case TargetCXXABI::Fuchsia:
+  case TargetCXXABI::GenericAArch64:
+  case TargetCXXABI::GenericItanium:
+  case TargetCXXABI::GenericARM:
+  case TargetCXXABI::GenericMIPS:
+  case TargetCXXABI::iOS:
+  case TargetCXXABI::WebAssembly:
+  case TargetCXXABI::WatchOS:
+  case TargetCXXABI::XL:
+    return ItaniumMangleContext::create(
+        *this, getDiagnostics(),
+        [](ASTContext &, const NamedDecl *ND) -> llvm::Optional<unsigned> {
+          if (const auto *RD = dyn_cast<CXXRecordDecl>(ND))
+            return RD->getDeviceLambdaManglingNumber();
+          return llvm::None;
+        });
+  case TargetCXXABI::Microsoft:
+    return MicrosoftMangleContext::create(*this, getDiagnostics());
+  }
+  llvm_unreachable("Unsupported ABI");
+}
+
 CXXABI::~CXXABI() = default;
 
 size_t ASTContext::getSideTableAllocatedMemory() const {
@@ -11689,3 +11716,86 @@ StringRef ASTContext::getCUIDHash() const {
   CUIDHash = llvm::utohexstr(llvm::MD5Hash(LangOpts.CUID), /*LowerCase=*/true);
   return CUIDHash;
 }
+
+// Get the closest named parent, so we can order the sycl naming decls somewhere
+// that mangling is meaningful.
+static const DeclContext *GetNamedParent(const CXXRecordDecl *RD) {
+  const DeclContext *DC = RD->getDeclContext();
+
+  while (!isa<NamedDecl, TranslationUnitDecl>(DC))
+    DC = DC->getParent();
+  return DC;
+}
+
+void ASTContext::AddSYCLKernelNamingDecl(const CXXRecordDecl *RD) {
+  assert(getLangOpts().isSYCL() && "Only valid for SYCL programs");
+  RD = RD->getCanonicalDecl();
+  const DeclContext *DC = GetNamedParent(RD);
+
+  assert(RD->getLocation().isValid() &&
+         "Invalid location on kernel naming decl");
+
+  (void)SYCLKernelNamingTypes[DC].insert(RD);
+}
+
+bool ASTContext::IsSYCLKernelNamingDecl(const NamedDecl *ND) const {
+  assert(getLangOpts().isSYCL() && "Only valid for SYCL programs");
+  const auto *RD = dyn_cast<CXXRecordDecl>(ND);
+  if (!RD)
+    return false;
+  RD = RD->getCanonicalDecl();
+  const DeclContext *DC = GetNamedParent(RD);
+
+  auto Itr = SYCLKernelNamingTypes.find(DC);
+
+  if (Itr == SYCLKernelNamingTypes.end())
+    return false;
+
+  return Itr->getSecond().count(RD);
+}
+
+// Filters the Decls list to those that share the lambda mangling with the
+// passed RD.
+void ASTContext::FilterSYCLKernelNamingDecls(
+    const CXXRecordDecl *RD,
+    llvm::SmallVectorImpl<const CXXRecordDecl *> &Decls) {
+
+  if (!SYCLKernelFilterContext)
+    SYCLKernelFilterContext.reset(
+        ItaniumMangleContext::create(*this, getDiagnostics()));
+
+  llvm::SmallString<128> LambdaSig;
+  llvm::raw_svector_ostream Out(LambdaSig);
+  SYCLKernelFilterContext->mangleLambdaSig(RD, Out);
+
+  llvm::erase_if(Decls, [this, &LambdaSig](const CXXRecordDecl *LocalRD) {
+    llvm::SmallString<128> LocalLambdaSig;
+    llvm::raw_svector_ostream LocalOut(LocalLambdaSig);
+    SYCLKernelFilterContext->mangleLambdaSig(LocalRD, LocalOut);
+    return LambdaSig != LocalLambdaSig;
+  });
+}
+
+unsigned ASTContext::GetSYCLKernelNamingIndex(const NamedDecl *ND) {
+  assert(getLangOpts().isSYCL() && "Only valid for SYCL programs");
+  assert(IsSYCLKernelNamingDecl(ND) &&
+         "Lambda not involved in mangling asked for a naming index?");
+
+  const CXXRecordDecl *RD = cast<CXXRecordDecl>(ND)->getCanonicalDecl();
+  const DeclContext *DC = GetNamedParent(RD);
+
+  auto Itr = SYCLKernelNamingTypes.find(DC);
+  assert(Itr != SYCLKernelNamingTypes.end() && "Not a valid DeclContext?");
+
+  const llvm::SmallPtrSet<const CXXRecordDecl *, 4> &Set = Itr->getSecond();
+
+  llvm::SmallVector<const CXXRecordDecl *> Decls{Set.begin(), Set.end()};
+
+  FilterSYCLKernelNamingDecls(RD, Decls);
+
+  llvm::sort(Decls, [](const CXXRecordDecl *LHS, const CXXRecordDecl *RHS) {
+    return LHS->getLambdaManglingNumber() < RHS->getLambdaManglingNumber();
+  });
+
+  return llvm::find(Decls, RD) - Decls.begin();
+}
diff --git a/clang/lib/AST/ComputeDependence.cpp b/clang/lib/AST/ComputeDependence.cpp
index 4026fdc76fd6f..1a5d2f7075fb5 100644
--- a/clang/lib/AST/ComputeDependence.cpp
+++ b/clang/lib/AST/ComputeDependence.cpp
@@ -556,6 +556,10 @@ ExprDependence clang::computeDependence(RecoveryExpr *E) {
   return D;
 }
 
+ExprDependence clang::computeDependence(SYCLUniqueStableNameExpr *E) {
+  return toExprDependence(E->getTypeSourceInfo()->getType()->getDependence());
+}
+
 ExprDependence clang::computeDependence(PredefinedExpr *E) {
   return toExprDependence(E->getType()->getDependence()) &
          ~ExprDependence::UnexpandedPack;
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index be0f2855f2980..49e1305beea97 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -504,6 +504,70 @@ SourceLocation DeclRefExpr::getEndLoc() const {
   return getNameInfo().getEndLoc();
 }
 
+SYCLUniqueStableNameExpr::SYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                                   SourceLocation LParen,
+                                                   SourceLocation RParen,
+                                                   QualType ResultTy,
+                                                   TypeSourceInfo *TSI)
+    : Expr(SYCLUniqueStableNameExprClass, ResultTy, VK_RValue, OK_Ordinary),
+      OpLoc(OpLoc), LParen(LParen), RParen(RParen) {
+  setTypeSourceInfo(TSI);
+  setDependence(computeDependence(this));
+}
+
+SYCLUniqueStableNameExpr::SYCLUniqueStableNameExpr(EmptyShell Empty,
+                                                   QualType ResultTy)
+    : Expr(SYCLUniqueStableNameExprClass, ResultTy, VK_RValue, OK_Ordinary) {}
+
+SYCLUniqueStableNameExpr *
+SYCLUniqueStableNameExpr::Create(const ASTContext &Ctx, SourceLocation OpLoc,
+                                 SourceLocation LParen, SourceLocation RParen,
+                                 TypeSourceInfo *TSI) {
+  QualType ResultTy = Ctx.getPointerType(Ctx.CharTy.withConst());
+  return new (Ctx)
+      SYCLUniqueStableNameExpr(OpLoc, LParen, RParen, ResultTy, TSI);
+}
+
+SYCLUniqueStableNameExpr *
+SYCLUniqueStableNameExpr::CreateEmpty(const ASTContext &Ctx) {
+  QualType ResultTy = Ctx.getPointerType(Ctx.CharTy.withConst());
+  return new (Ctx) SYCLUniqueStableNameExpr(EmptyShell(), ResultTy);
+}
+
+std::string SYCLUniqueStableNameExpr::ComputeName(ASTContext &Context) const {
+  return SYCLUniqueStableNameExpr::ComputeName(Context,
+                                               getTypeSourceInfo()->getType());
+}
+
+std::string SYCLUniqueStableNameExpr::ComputeName(ASTContext &Context,
+                                                  QualType Ty) {
+  auto MangleCallback = [](ASTContext &Ctx,
+                           const NamedDecl *ND) -> llvm::Optional<unsigned> {
+    // This replaces the 'lambda number' in the mangling with a unique number
+    // based on its order in the declaration.  To provide some level of visual
+    // notability (actual uniqueness from normal lambdas isn't necessary, as
+    // these are used differently), we add 10,000 to the number.
+    // For example:
+    // _ZTSZ3foovEUlvE10005_
+    // Demangles to: typeinfo name for foo()::'lambda10005'()
+    // Note that the mangler subtracts 2, since with normal lambdas the lambda
+    // mangling number '0' is an anonymous struct mangle, and '1' is omitted.
+    // So 10,002 results in the first number being 10,000.
+    if (Ctx.IsSYCLKernelNamingDecl(ND))
+      return 10'002 + Ctx.GetSYCLKernelNamingIndex(ND);
+    return llvm::None;
+  };
+  std::unique_ptr<MangleContext> Ctx{ItaniumMangleContext::create(
+      Context, Context.getDiagnostics(), MangleCallback)};
+
+  std::string Buffer;
+  Buffer.reserve(128);
+  llvm::raw_string_ostream Out(Buffer);
+  Ctx->mangleTypeName(Ty, Out);
+
+  return Out.str();
+}
+
 PredefinedExpr::PredefinedExpr(SourceLocation L, QualType FNTy, IdentKind IK,
                                StringLiteral *SL)
     : Expr(PredefinedExprClass, FNTy, VK_LValue, OK_Ordinary) {
@@ -518,34 +582,6 @@ PredefinedExpr::PredefinedExpr(SourceLocation L, QualType FNTy, IdentKind IK,
   setDependence(computeDependence(this));
 }
 
-PredefinedExpr::PredefinedExpr(SourceLocation L, QualType FnTy, IdentKind IK,
-                               TypeSourceInfo *Info)
-    : Expr(PredefinedExprClass, FnTy, VK_LValue, OK_Ordinary) {
-  PredefinedExprBits.Kind = IK;
-  assert((getIdentKind() == IK) &&
-         "IdentKind do not fit in PredefinedExprBitFields!");
-  assert(IK == UniqueStableNameType &&
-         "Constructor only valid with UniqueStableNameType");
-  PredefinedExprBits.HasFunctionName = false;
-  PredefinedExprBits.Loc = L;
-  setTypeSourceInfo(Info);
-  setDependence(computeDependence(this));
-}
-
-PredefinedExpr::PredefinedExpr(SourceLocation L, QualType FnTy, IdentKind IK,
-                               Expr *E)
-    : Expr(PredefinedExprClass, FnTy, VK_LValue, OK_Ordinary) {
-  PredefinedExprBits.Kind = IK;
-  assert((getIdentKind() == IK) &&
-         "IdentKind do not fit in PredefinedExprBitFields!");
-  assert(IK == UniqueStableNameExpr &&
-         "Constructor only valid with UniqueStableNameExpr");
-  PredefinedExprBits.HasFunctionName = false;
-  PredefinedExprBits.Loc = L;
-  setExpr(E);
-  setDependence(computeDependence(this));
-}
-
 PredefinedExpr::PredefinedExpr(EmptyShell Empty, bool HasFunctionName)
     : Expr(PredefinedExprClass, Empty) {
   PredefinedExprBits.HasFunctionName = HasFunctionName;
@@ -555,44 +591,15 @@ PredefinedExpr *PredefinedExpr::Create(const ASTContext &Ctx, SourceLocation L,
                                        QualType FNTy, IdentKind IK,
                                        StringLiteral *SL) {
   bool HasFunctionName = SL != nullptr;
-  void *Mem = Ctx.Allocate(
-      totalSizeToAlloc<Stmt *, Expr *, TypeSourceInfo *>(HasFunctionName, 0, 0),
-      alignof(PredefinedExpr));
-  return new (Mem) PredefinedExpr(L, FNTy, IK, SL);
-}
-
-PredefinedExpr *PredefinedExpr::Create(const ASTContext &Ctx, SourceLocation L,
-                                       QualType FNTy, IdentKind IK,
-                                       StringLiteral *SL,
-                                       TypeSourceInfo *Info) {
-  assert(IK == UniqueStableNameType && "Only valid with UniqueStableNameType");
-  bool HasFunctionName = SL != nullptr;
-  void *Mem = Ctx.Allocate(totalSizeToAlloc<Stmt *, Expr *, TypeSourceInfo *>(
-                               HasFunctionName, 0, !HasFunctionName),
-                           alignof(PredefinedExpr));
-  if (HasFunctionName)
-    return new (Mem) PredefinedExpr(L, FNTy, IK, SL);
-  return new (Mem) PredefinedExpr(L, FNTy, IK, Info);
-}
-
-PredefinedExpr *PredefinedExpr::Create(const ASTContext &Ctx, SourceLocation L,
-                                       QualType FNTy, IdentKind IK,
-                                       StringLiteral *SL, Expr *E) {
-  assert(IK == UniqueStableNameExpr && "Only valid with UniqueStableNameExpr");
-  bool HasFunctionName = SL != nullptr;
-  void *Mem = Ctx.Allocate(totalSizeToAlloc<Stmt *, Expr *, TypeSourceInfo *>(
-                               HasFunctionName, !HasFunctionName, 0),
+  void *Mem = Ctx.Allocate(totalSizeToAlloc<Stmt *>(HasFunctionName),
                            alignof(PredefinedExpr));
-  if (HasFunctionName)
-    return new (Mem) PredefinedExpr(L, FNTy, IK, SL);
-  return new (Mem) PredefinedExpr(L, FNTy, IK, E);
+  return new (Mem) PredefinedExpr(L, FNTy, IK, SL);
 }
 
 PredefinedExpr *PredefinedExpr::CreateEmpty(const ASTContext &Ctx,
                                             bool HasFunctionName) {
-  void *Mem = Ctx.Allocate(
-      totalSizeToAlloc<Stmt *, Expr *, TypeSourceInfo *>(HasFunctionName, 0, 0),
-      alignof(PredefinedExpr));
+  void *Mem = Ctx.Allocate(totalSizeToAlloc<Stmt *>(HasFunctionName),
+                           alignof(PredefinedExpr));
   return new (Mem) PredefinedExpr(EmptyShell(), HasFunctionName);
 }
 
@@ -612,28 +619,12 @@ StringRef PredefinedExpr::getIdentKindName(PredefinedExpr::IdentKind IK) {
     return "__FUNCSIG__";
   case LFuncSig:
     return "L__FUNCSIG__";
-  case UniqueStableNameType:
-  case UniqueStableNameExpr:
-    return "__builtin_unique_stable_name";
   case PrettyFunctionNoVirtual:
     break;
   }
   llvm_unreachable("Unknown ident kind for PredefinedExpr");
 }
 
-std::string PredefinedExpr::ComputeName(ASTContext &Context, IdentKind IK,
-                                        QualType Ty) {
-  std::unique_ptr<MangleContext> Ctx{ItaniumMangleContext::create(
-      Context, Context.getDiagnostics(), /*IsUniqueNameMangler*/ true)};
-
-  Ty = Ty.getCanonicalType();
-
-  SmallString<256> Buffer;
-  llvm::raw_svector_ostream Out(Buffer);
-  Ctx->mangleTypeName(Ty, Out);
-  return std::string(Buffer.str());
-}
-
 // FIXME: Maybe this should use DeclPrinter with a special "print predefined
 // expr" policy instead.
 std::string PredefinedExpr::ComputeName(IdentKind IK, const Decl *CurrentDecl) {
@@ -3454,6 +3445,7 @@ bool Expr::HasSideEffects(const ASTContext &Ctx,
   case SourceLocExprClass:
   case ConceptSpecializationExprClass:
   case RequiresExprClass:
+  case SYCLUniqueStableNameExprClass:
     // These never have a side-effect.
     return false;
 
diff --git a/clang/lib/AST/ExprClassification.cpp b/clang/lib/AST/ExprClassification.cpp
index 0286c730ce4e3..dd7c4a4cbe00b 100644
--- a/clang/lib/AST/ExprClassification.cpp
+++ b/clang/lib/AST/ExprClassification.cpp
@@ -433,6 +433,9 @@ static Cl::Kinds ClassifyInternal(ASTContext &Ctx, const Expr *E) {
   case Expr::CoawaitExprClass:
   case Expr::CoyieldExprClass:
     return ClassifyInternal(Ctx, cast<CoroutineSuspendExpr>(E)->getResumeExpr());
+  case Expr::SYCLUniqueStableNameExprClass:
+    return Cl::CL_PRValue;
+    break;
   }
 
   llvm_unreachable("unhandled expression kind in classification");
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index 5508a71508604..22cb01f946c4f 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -8669,6 +8669,26 @@ class PointerExprEvaluator
     return true;
   }
 
+  bool VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E) {
+    std::string ResultStr = E->ComputeName(Info.Ctx);
+
+    Info.Ctx.SYCLUniqueStableNameEvaluatedValues[E] = ResultStr;
+
+    QualType CharTy = Info.Ctx.CharTy.withConst();
+    APInt Size(Info.Ctx.getTypeSize(Info.Ctx.getSizeType()),
+               ResultStr.size() + 1);
+    QualType ArrayTy = Info.Ctx.getConstantArrayType(CharTy, Size, nullptr,
+                                                     ArrayType::Normal, 0);
+
+    StringLiteral *SL =
+        StringLiteral::Create(Info.Ctx, ResultStr, StringLiteral::Ascii,
+                              /*Pascal*/ false, ArrayTy, E->getLocation());
+
+    evaluateLValue(SL, Result);
+    Result.addArray(Info, E, cast<ConstantArrayType>(ArrayTy));
+    return true;
+  }
+
   // FIXME: Missing: @protocol, @selector
 };
 } // end anonymous namespace
@@ -10367,7 +10387,8 @@ namespace {
 
       Result = APValue(APValue::UninitArray(), 0,
                        CAT->getSize().getZExtValue());
-      if (!Result.hasArrayFiller()) return true;
+      if (!Result.hasArrayFiller())
+        return true;
 
       // Zero-initialize all elements.
       LValue Subobject = This;
@@ -15159,6 +15180,7 @@ static ICEDiag CheckICE(const Expr* E, const ASTContext &Ctx) {
   case Expr::CoawaitExprClass:
   case Expr::DependentCoawaitExprClass:
   case Expr::CoyieldExprClass:
+  case Expr::SYCLUniqueStableNameExprClass:
     return ICEDiag(IK_NotICE, E->getBeginLoc());
 
   case Expr::InitListExprClass: {
diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp
index 3ad3ad01122ba..6369d23630314 100644
--- a/clang/lib/AST/ItaniumMangle.cpp
+++ b/clang/lib/AST/ItaniumMangle.cpp
@@ -125,15 +125,16 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
   typedef std::pair<const DeclContext*, IdentifierInfo*> DiscriminatorKeyTy;
   llvm::DenseMap<DiscriminatorKeyTy, unsigned> Discriminator;
   llvm::DenseMap<const NamedDecl*, unsigned> Uniquifier;
+  const DiscriminatorOverrideTy DiscriminatorOverride = nullptr;
 
-  bool IsDevCtx = false;
   bool NeedsUniqueInternalLinkageNames = false;
 
 public:
-  explicit ItaniumMangleContextImpl(ASTContext &Context,
-                                    DiagnosticsEngine &Diags,
-                                    bool IsUniqueNameMangler)
-      : ItaniumMangleContext(Context, Diags, IsUniqueNameMangler) {}
+  explicit ItaniumMangleContextImpl(
+      ASTContext &Context, DiagnosticsEngine &Diags,
+      DiscriminatorOverrideTy DiscriminatorOverride)
+      : ItaniumMangleContext(Context, Diags),
+        DiscriminatorOverride(DiscriminatorOverride) {}
 
   /// @name Mangler Entry Points
   /// @{
@@ -148,9 +149,6 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
     NeedsUniqueInternalLinkageNames = true;
   }
 
-  bool isDeviceMangleContext() const override { return IsDevCtx; }
-  void setDeviceMangleContext(bool IsDev) override { IsDevCtx = IsDev; }
-
   void mangleCXXName(GlobalDecl GD, raw_ostream &) override;
   void mangleThunk(const CXXMethodDecl *MD, const ThunkInfo &Thunk,
                    raw_ostream &) override;
@@ -247,6 +245,10 @@ class ItaniumMangleContextImpl : public ItaniumMangleContext {
     return Name;
   }
 
+  DiscriminatorOverrideTy getDiscriminatorOverride() const override {
+    return DiscriminatorOverride;
+  }
+
   /// @}
 };
 
@@ -1517,7 +1519,8 @@ void CXXNameMangler::mangleUnqualifiedName(GlobalDecl GD,
     //     # Parameter types or 'v' for 'void'.
     if (const CXXRecordDecl *Record = dyn_cast<CXXRecordDecl>(TD)) {
       if (Record->isLambda() && (Record->getLambdaManglingNumber() ||
-                                 Context.isUniqueNameMangler())) {
+                                 Context.getDiscriminatorOverride()(
+                                     Context.getASTContext(), Record))) {
         assert(!AdditionalAbiTags &&
                "Lambda type cannot have additional abi tags");
         mangleLambda(Record);
@@ -1921,37 +1924,6 @@ void CXXNameMangler::mangleTemplateParamDecl(const NamedDecl *Decl) {
   }
 }
 
-// Handles the __builtin_unique_stable_name feature for lambdas.  Instead of the
-// ordinal of the lambda in its mangling, this does line/column to uniquely and
-// reliably identify the lambda.  Additionally, macro expansions are expressed
-// as well to prevent macros causing duplicates.
-static void mangleUniqueNameLambda(CXXNameMangler &Mangler, SourceManager &SM,
-                                   raw_ostream &Out,
-                                   const CXXRecordDecl *Lambda) {
-  SourceLocation Loc = Lambda->getLocation();
-
-  PresumedLoc PLoc = SM.getPresumedLoc(Loc);
-  Mangler.mangleNumber(PLoc.getLine());
-  Out << "_";
-  Mangler.mangleNumber(PLoc.getColumn());
-
-  while(Loc.isMacroID()) {
-    SourceLocation SLToPrint = Loc;
-    if (SM.isMacroArgExpansion(Loc))
-      SLToPrint = SM.getImmediateExpansionRange(Loc).getBegin();
-
-    PLoc = SM.getPresumedLoc(SM.getSpellingLoc(SLToPrint));
-    Out << "m";
-    Mangler.mangleNumber(PLoc.getLine());
-    Out << "_";
-    Mangler.mangleNumber(PLoc.getColumn());
-
-    Loc = SM.getImmediateMacroCallerLoc(Loc);
-    if (Loc.isFileID())
-      Loc = SM.getImmediateMacroCallerLoc(SLToPrint);
-  }
-}
-
 void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
   // When trying to be ABI-compatibility with clang 12 and before, mangle a
   // <data-member-prefix> now, with no substitutions.
@@ -1975,12 +1947,6 @@ void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
   mangleLambdaSig(Lambda);
   Out << "E";
 
-  if (Context.isUniqueNameMangler()) {
-    mangleUniqueNameLambda(
-        *this, Context.getASTContext().getSourceManager(), Out, Lambda);
-    return;
-  }
-
   // The number is omitted for the first closure type with a given
   // <lambda-sig> in a given context; it is n-2 for the nth closure type
   // (in lexical order) with that same <lambda-sig> and context.
@@ -1992,9 +1958,11 @@ void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
   // if the host-side CXX ABI has different numbering for lambda. In such case,
   // if the mangle context is that device-side one, use the device-side lambda
   // mangling number for this lambda.
-  unsigned Number = Context.isDeviceMangleContext()
-                        ? Lambda->getDeviceLambdaManglingNumber()
-                        : Lambda->getLambdaManglingNumber();
+  llvm::Optional<unsigned> DeviceNumber =
+      Context.getDiscriminatorOverride()(Context.getASTContext(), Lambda);
+  unsigned Number = DeviceNumber.hasValue() ? *DeviceNumber
+                                            : Lambda->getLambdaManglingNumber();
+
   assert(Number > 0 && "Lambda should be mangled as an unnamed class");
   if (Number > 1)
     mangleNumber(Number - 2);
@@ -5073,6 +5041,16 @@ void CXXNameMangler::mangleExpression(const Expr *E, unsigned Arity,
     Out << "v18co_yield";
     mangleExpression(cast<CoawaitExpr>(E)->getOperand());
     break;
+  case Expr::SYCLUniqueStableNameExprClass: {
+    const auto *USN = cast<SYCLUniqueStableNameExpr>(E);
+    NotPrimaryExpr();
+
+    Out << "u33__builtin_sycl_unique_stable_name";
+    mangleType(USN->getTypeSourceInfo()->getType());
+
+    Out << "E";
+    break;
+  }
   }
 
   if (AsTemplateArg && !IsPrimaryExpr)
@@ -6426,7 +6404,16 @@ void ItaniumMangleContextImpl::mangleLambdaSig(const CXXRecordDecl *Lambda,
 }
 
 ItaniumMangleContext *ItaniumMangleContext::create(ASTContext &Context,
-                                                   DiagnosticsEngine &Diags,
-                                                   bool IsUniqueNameMangler) {
-  return new ItaniumMangleContextImpl(Context, Diags, IsUniqueNameMangler);
+                                                   DiagnosticsEngine &Diags) {
+  return new ItaniumMangleContextImpl(
+      Context, Diags,
+      [](ASTContext &, const NamedDecl *) -> llvm::Optional<unsigned> {
+        return llvm::None;
+      });
+}
+
+ItaniumMangleContext *
+ItaniumMangleContext::create(ASTContext &Context, DiagnosticsEngine &Diags,
+                             DiscriminatorOverrideTy DiscriminatorOverride) {
+  return new ItaniumMangleContextImpl(Context, Diags, DiscriminatorOverride);
 }
diff --git a/clang/lib/AST/JSONNodeDumper.cpp b/clang/lib/AST/JSONNodeDumper.cpp
index 95e60f8b19a6a..5d5388880e675 100644
--- a/clang/lib/AST/JSONNodeDumper.cpp
+++ b/clang/lib/AST/JSONNodeDumper.cpp
@@ -1164,6 +1164,12 @@ void JSONNodeDumper::VisitDeclRefExpr(const DeclRefExpr *DRE) {
   }
 }
 
+void JSONNodeDumper::VisitSYCLUniqueStableNameExpr(
+    const SYCLUniqueStableNameExpr *E) {
+  JOS.attribute("typeSourceInfo",
+                createQualType(E->getTypeSourceInfo()->getType()));
+}
+
 void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) {
   JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind()));
 }
diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 58afc020f457e..3e275744931a5 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1081,6 +1081,13 @@ void StmtPrinter::VisitObjCSubscriptRefExpr(ObjCSubscriptRefExpr *Node) {
   OS << "]";
 }
 
+void StmtPrinter::VisitSYCLUniqueStableNameExpr(
+    SYCLUniqueStableNameExpr *Node) {
+  OS << "__builtin_sycl_unique_stable_name(";
+  Node->getTypeSourceInfo()->getType().print(OS, Policy);
+  OS << ")";
+}
+
 void StmtPrinter::VisitPredefinedExpr(PredefinedExpr *Node) {
   OS << PredefinedExpr::getIdentKindName(Node->getIdentKind());
 }
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index fc1f8a5327244..dc91043b44a67 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -1190,6 +1190,12 @@ void StmtProfiler::VisitDeclRefExpr(const DeclRefExpr *S) {
   }
 }
 
+void StmtProfiler::VisitSYCLUniqueStableNameExpr(
+    const SYCLUniqueStableNameExpr *S) {
+  VisitExpr(S);
+  VisitType(S->getTypeSourceInfo()->getType());
+}
+
 void StmtProfiler::VisitPredefinedExpr(const PredefinedExpr *S) {
   VisitExpr(S);
   ID.AddInteger(S->getIdentKind());
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index ce5089ff5ecee..cb5b4230a5ab9 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -1018,6 +1018,11 @@ void TextNodeDumper::VisitObjCIvarRefExpr(const ObjCIvarRefExpr *Node) {
     OS << " isFreeIvar";
 }
 
+void TextNodeDumper::VisitSYCLUniqueStableNameExpr(
+    const SYCLUniqueStableNameExpr *Node) {
+  dumpType(Node->getTypeSourceInfo()->getType());
+}
+
 void TextNodeDumper::VisitPredefinedExpr(const PredefinedExpr *Node) {
   OS << " " << PredefinedExpr::getIdentKindName(Node->getIdentKind());
 }
diff --git a/clang/lib/Basic/IdentifierTable.cpp b/clang/lib/Basic/IdentifierTable.cpp
index 0d68a2b86c387..5bee625b6b5a8 100644
--- a/clang/lib/Basic/IdentifierTable.cpp
+++ b/clang/lib/Basic/IdentifierTable.cpp
@@ -157,6 +157,8 @@ static KeywordStatus getKeywordStatus(const LangOptions &LangOpts,
   if (LangOpts.CPlusPlus && (Flags & KEYALLCXX)) return KS_Future;
   if (LangOpts.CPlusPlus && !LangOpts.CPlusPlus20 && (Flags & CHAR8SUPPORT))
     return KS_Future;
+  if (LangOpts.isSYCL() && (Flags & KEYSYCL))
+    return KS_Enabled;
   return KS_Disabled;
 }
 
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 995b6a0b5fec6..bfd2c16c12827 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -191,12 +191,27 @@ CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
   return ((Twine("__cuda") + Twine(FuncName)).str());
 }
 
+static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) {
+  // If the host and device have different C++ ABIs, mark it as the device
+  // mangle context so that the mangling needs to retrieve the additional
+  // device lambda mangling number instead of the regular host one.
+  if (CGM.getContext().getAuxTargetInfo() &&
+      CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
+      CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily()) {
+    return std::unique_ptr<MangleContext>(
+        CGM.getContext().createDeviceMangleContext(
+            *CGM.getContext().getAuxTargetInfo()));
+  }
+
+  return std::unique_ptr<MangleContext>(CGM.getContext().createMangleContext(
+      CGM.getContext().getAuxTargetInfo()));
+}
+
 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
     : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
       TheModule(CGM.getModule()),
       RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode),
-      DeviceMC(CGM.getContext().createMangleContext(
-          CGM.getContext().getAuxTargetInfo())) {
+      DeviceMC(InitDeviceMC(CGM)) {
   CodeGen::CodeGenTypes &Types = CGM.getTypes();
   ASTContext &Ctx = CGM.getContext();
 
@@ -207,14 +222,6 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
   CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
   VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
   VoidPtrPtrTy = VoidPtrTy->getPointerTo();
-  if (CGM.getContext().getAuxTargetInfo()) {
-    // If the host and device have different C++ ABIs, mark it as the device
-    // mangle context so that the mangling needs to retrieve the additonal
-    // device lambda mangling number instead of the regular host one.
-    DeviceMC->setDeviceMangleContext(
-        CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
-        CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily());
-  }
 }
 
 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
diff --git a/clang/lib/CodeGen/CGExprScalar.cpp b/clang/lib/CodeGen/CGExprScalar.cpp
index 89ab9abc7dadb..d67d6abe93c99 100644
--- a/clang/lib/CodeGen/CGExprScalar.cpp
+++ b/clang/lib/CodeGen/CGExprScalar.cpp
@@ -486,6 +486,8 @@ class ScalarExprEmitter
     return CGF.EmitPseudoObjectRValue(E).getScalarVal();
   }
 
+  Value *VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E);
+
   Value *VisitOpaqueValueExpr(OpaqueValueExpr *E) {
     if (E->isGLValue())
       return EmitLoadOfLValue(CGF.getOrCreateOpaqueLValueMapping(E),
@@ -1581,6 +1583,25 @@ Value *ScalarExprEmitter::VisitExpr(Expr *E) {
   return llvm::UndefValue::get(CGF.ConvertType(E->getType()));
 }
 
+Value *
+ScalarExprEmitter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
+  ASTContext &Context = CGF.getContext();
+  llvm::Optional<LangAS> GlobalAS =
+      Context.getTargetInfo().getConstantAddressSpace();
+  llvm::Constant *GlobalConstStr = Builder.CreateGlobalStringPtr(
+      E->ComputeName(Context), "__usn_str",
+      static_cast<unsigned>(GlobalAS.getValueOr(LangAS::Default)));
+
+  unsigned ExprAS = Context.getTargetAddressSpace(E->getType());
+
+  if (GlobalConstStr->getType()->getPointerAddressSpace() == ExprAS)
+    return GlobalConstStr;
+
+  llvm::Type *EltTy = GlobalConstStr->getType()->getPointerElementType();
+  llvm::PointerType *NewPtrTy = llvm::PointerType::get(EltTy, ExprAS);
+  return Builder.CreateAddrSpaceCast(GlobalConstStr, NewPtrTy, "usn_addr_cast");
+}
+
 Value *ScalarExprEmitter::VisitShuffleVectorExpr(ShuffleVectorExpr *E) {
   // Vector Mask Case
   if (E->getNumSubExprs() == 2) {
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index 8be38faee923d..22f3b7624c45b 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -1469,9 +1469,10 @@ ExprResult Parser::ParseCastExpression(CastParseKind ParseKind,
   case tok::kw_this:
     Res = ParseCXXThis();
     break;
-  case tok::kw___builtin_unique_stable_name:
-    Res = ParseUniqueStableNameExpression();
+  case tok::kw___builtin_sycl_unique_stable_name:
+    Res = ParseSYCLUniqueStableNameExpression();
     break;
+
   case tok::annot_typename:
     if (isStartOfObjCClassMessageMissingOpenBracket()) {
       TypeResult Type = getTypeAnnotation(Tok);
@@ -2326,42 +2327,32 @@ Parser::ParseExprAfterUnaryExprOrTypeTrait(const Token &OpTok,
   return Operand;
 }
 
-
-ExprResult Parser::ParseUniqueStableNameExpression() {
-  assert(Tok.is(tok::kw___builtin_unique_stable_name) &&
-         "Not __bulitin_unique_stable_name");
+/// Parse a __builtin_sycl_unique_stable_name expression.  Accepts a type-id as
+/// a parameter.
+ExprResult Parser::ParseSYCLUniqueStableNameExpression() {
+  assert(Tok.is(tok::kw___builtin_sycl_unique_stable_name) &&
+         "Not __bulitin_sycl_unique_stable_name");
 
   SourceLocation OpLoc = ConsumeToken();
   BalancedDelimiterTracker T(*this, tok::l_paren);
 
-  // typeid expressions are always parenthesized.
+  // __builtin_sycl_unique_stable_name expressions are always parenthesized.
   if (T.expectAndConsume(diag::err_expected_lparen_after,
-                         "__builtin_unique_stable_name"))
+                         "__builtin_sycl_unique_stable_name"))
     return ExprError();
 
-  if (isTypeIdInParens()) {
-    TypeResult Ty = ParseTypeName();
-    T.consumeClose();
+  TypeResult Ty = ParseTypeName();
 
-    if (Ty.isInvalid())
-      return ExprError();
-
-    return Actions.ActOnUniqueStableNameExpr(OpLoc, T.getOpenLocation(),
-                                             T.getCloseLocation(), Ty.get());
+  if (Ty.isInvalid()) {
+    T.skipToEnd();
+    return ExprError();
   }
 
-  EnterExpressionEvaluationContext Unevaluated(
-      Actions, Sema::ExpressionEvaluationContext::Unevaluated);
-  ExprResult Result = ParseExpression();
-
-  if (Result.isInvalid()) {
-    SkipUntil(tok::r_paren, StopAtSemi);
-    return Result;
-  }
+  if (T.consumeClose())
+    return ExprError();
 
-  T.consumeClose();
-  return Actions.ActOnUniqueStableNameExpr(OpLoc, T.getOpenLocation(),
-                                           T.getCloseLocation(), Result.get());
+  return Actions.ActOnSYCLUniqueStableNameExpr(OpLoc, T.getOpenLocation(),
+                                               T.getCloseLocation(), Ty.get());
 }
 
 /// Parse a sizeof or alignof expression.
diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index f0b9e6a5bc3ba..fd6a4940c174f 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1575,6 +1575,8 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
     return mergeCanThrow(CT, canThrow(TS->getTryBody()));
   }
 
+  case Stmt::SYCLUniqueStableNameExprClass:
+    return CT_Cannot;
   case Stmt::NoStmtClass:
     llvm_unreachable("Invalid class for statement");
   }
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 02042b05c724e..03da561f2061a 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -3549,68 +3549,26 @@ ExprResult Sema::BuildPredefinedExpr(SourceLocation Loc,
   return PredefinedExpr::Create(Context, Loc, ResTy, IK, SL);
 }
 
-static std::pair<QualType, StringLiteral *>
-GetUniqueStableNameInfo(ASTContext &Context, QualType OpType,
-                        SourceLocation OpLoc, PredefinedExpr::IdentKind K) {
-  std::pair<QualType, StringLiteral*> Result{{}, nullptr};
-
-  if (OpType->isDependentType()) {
-      Result.first = Context.DependentTy;
-      return Result;
-  }
-
-  std::string Str = PredefinedExpr::ComputeName(Context, K, OpType);
-  llvm::APInt Length(32, Str.length() + 1);
-  Result.first =
-      Context.adjustStringLiteralBaseType(Context.CharTy.withConst());
-  Result.first = Context.getConstantArrayType(
-      Result.first, Length, nullptr, ArrayType::Normal, /*IndexTypeQuals*/ 0);
-  Result.second = StringLiteral::Create(Context, Str, StringLiteral::Ascii,
-                                        /*Pascal*/ false, Result.first, OpLoc);
-  return Result;
-}
-
-ExprResult Sema::BuildUniqueStableName(SourceLocation OpLoc,
-                                       TypeSourceInfo *Operand) {
-  QualType ResultTy;
-  StringLiteral *SL;
-  std::tie(ResultTy, SL) = GetUniqueStableNameInfo(
-      Context, Operand->getType(), OpLoc, PredefinedExpr::UniqueStableNameType);
-
-  return PredefinedExpr::Create(Context, OpLoc, ResultTy,
-                                PredefinedExpr::UniqueStableNameType, SL,
-                                Operand);
-}
-
-ExprResult Sema::BuildUniqueStableName(SourceLocation OpLoc,
-                                       Expr *E) {
-  QualType ResultTy;
-  StringLiteral *SL;
-  std::tie(ResultTy, SL) = GetUniqueStableNameInfo(
-      Context, E->getType(), OpLoc, PredefinedExpr::UniqueStableNameExpr);
-
-  return PredefinedExpr::Create(Context, OpLoc, ResultTy,
-                                PredefinedExpr::UniqueStableNameExpr, SL, E);
+ExprResult Sema::BuildSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                               SourceLocation LParen,
+                                               SourceLocation RParen,
+                                               TypeSourceInfo *TSI) {
+  return SYCLUniqueStableNameExpr::Create(Context, OpLoc, LParen, RParen, TSI);
 }
 
-ExprResult Sema::ActOnUniqueStableNameExpr(SourceLocation OpLoc,
-                                           SourceLocation L, SourceLocation R,
-                                           ParsedType Ty) {
-  TypeSourceInfo *TInfo = nullptr;
-  QualType T = GetTypeFromParser(Ty, &TInfo);
+ExprResult Sema::ActOnSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                               SourceLocation LParen,
+                                               SourceLocation RParen,
+                                               ParsedType ParsedTy) {
+  TypeSourceInfo *TSI = nullptr;
+  QualType Ty = GetTypeFromParser(ParsedTy, &TSI);
 
-  if (T.isNull())
+  if (Ty.isNull())
     return ExprError();
-  if (!TInfo)
-    TInfo = Context.getTrivialTypeSourceInfo(T, OpLoc);
-
-  return BuildUniqueStableName(OpLoc, TInfo);
-}
+  if (!TSI)
+    TSI = Context.getTrivialTypeSourceInfo(Ty, LParen);
 
-ExprResult Sema::ActOnUniqueStableNameExpr(SourceLocation OpLoc,
-                                           SourceLocation L, SourceLocation R,
-                                           Expr *E) {
-  return BuildUniqueStableName(OpLoc, E);
+  return BuildSYCLUniqueStableNameExpr(OpLoc, LParen, RParen, TSI);
 }
 
 ExprResult Sema::ActOnPredefinedExpr(SourceLocation Loc, tok::TokenKind Kind) {
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index 1c07732fe8aa4..cf5e20f7b6c1e 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -461,11 +461,15 @@ void Sema::handleLambdaNumbering(
   std::tie(MCtx, ManglingContextDecl) =
       getCurrentMangleNumberContext(Class->getDeclContext());
   bool HasKnownInternalLinkage = false;
-  if (!MCtx && getLangOpts().CUDA) {
+  if (!MCtx && (getLangOpts().CUDA || getLangOpts().SYCLIsDevice ||
+                getLangOpts().SYCLIsHost)) {
     // Force lambda numbering in CUDA/HIP as we need to name lambdas following
     // ODR. Both device- and host-compilation need to have a consistent naming
     // on kernel functions. As lambdas are potential part of these `__global__`
     // function names, they needs numbering following ODR.
+    // Also force for SYCL, since we need this for the
+    // __builtin_sycl_unique_stable_name implementation, which depends on lambda
+    // mangling.
     MCtx = getMangleNumberingContext(Class, ManglingContextDecl);
     assert(MCtx && "Retrieving mangle numbering context failed!");
     HasKnownInternalLinkage = true;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8ce5ee4d51ec8..3ce184fb1a3ff 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -905,10 +905,8 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc,
 
   MC.mangleTypeName(KernelNameType, Out);
 
-  return {std::string(Out.str()),
-          PredefinedExpr::ComputeName(S.getASTContext(),
-                                      PredefinedExpr::UniqueStableNameType,
-                                      KernelNameType)};
+  return {std::string(Out.str()), SYCLUniqueStableNameExpr::ComputeName(
+                                      S.getASTContext(), KernelNameType)};
 }
 
 static bool isDefaultSPIRArch(ASTContext &Context) {
@@ -3212,9 +3210,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     // Get specialization constant ID type, which is the second template
     // argument.
     QualType SpecConstIDTy = TemplateArgs.get(1).getAsType().getCanonicalType();
-    const std::string SpecConstName = PredefinedExpr::ComputeName(
-        SemaRef.getASTContext(), PredefinedExpr::UniqueStableNameType,
-        SpecConstIDTy);
+    const std::string SpecConstName = SYCLUniqueStableNameExpr::ComputeName(
+        SemaRef.getASTContext(), SpecConstIDTy);
     Header.addSpecConstant(SpecConstName, SpecConstIDTy);
     return true;
   }
@@ -3519,6 +3516,10 @@ class SYCLKernelNameTypeVisitor
 
 void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
                                ArrayRef<const Expr *> Args) {
+  // FIXME: In place until the library works around its 'host' invocation
+  // issues.
+  if (!LangOpts.SYCLIsDevice)
+    return;
   const CXXRecordDecl *KernelObj = getKernelObjectType(KernelFunc);
   QualType KernelNameType =
       calculateKernelNameType(getASTContext(), KernelFunc);
@@ -4687,14 +4688,15 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) {
 void SYCLIntegrationFooter::emitSpecIDName(raw_ostream &O, const VarDecl *VD) {
   // FIXME: Figure out the spec-constant unique name here.
   // Note that this changes based on the linkage of the variable.
-  // We typically want to use the __builtin_unique_stable_name for the variable
-  // (or the newer-equivilent for values, see the JIRA), but we also have to
-  // figure out if this has internal or external linkage.  In external-case this
-  // should be the same as the the unique-name.  However, this isn't the case
-  // with local-linkage, where we want to put the driver-provided random-value
-  // ahead of it, so that we make sure it is unique across translation units.
-  // This name should come from the yet implemented__builtin_unique_stable_name
-  // feature that accepts variables and gives the mangling for that.
+  // We typically want to use the __builtin_sycl_unique_stable_name for the
+  // variable (or the newer-equivilent for values, see the JIRA), but we also
+  // have to figure out if this has internal or external linkage.  In
+  // external-case this should be the same as the the unique-name.  However,
+  // this isn't the case with local-linkage, where we want to put the
+  // driver-provided random-value ahead of it, so that we make sure it is unique
+  // across translation units. This name should come from the yet
+  // implemented __builtin_sycl_unique_stable_name feature that accepts
+  // variables and gives the mangling for that.
   O << "";
 }
 
@@ -5006,3 +5008,35 @@ bool Util::matchQualifiedTypeName(QualType Ty,
   const auto *Ctx = cast<DeclContext>(RecTy);
   return Util::matchContext(Ctx, Scopes);
 }
+
+// The SYCL kernel's 'object type' used for diagnostics and naming/mangling is
+// the first parameter to a sycl_kernel labeled function template. In SYCL1.2.1,
+// this was passed by value, and in SYCL2020, it is passed by reference.
+static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
+  assert(KernelCaller->getNumParams() > 0 && "Insufficient kernel parameters");
+  QualType KernelParamTy = KernelCaller->getParamDecl(0)->getType();
+
+  // SYCL 2020 kernels are passed by reference.
+  if (KernelParamTy->isReferenceType())
+    return KernelParamTy->getPointeeType();
+
+  // SYCL 1.2.1
+  return KernelParamTy;
+}
+
+void Sema::AddSYCLKernelLambda(const FunctionDecl *FD) {
+  auto MangleCallback = [](ASTContext &Ctx,
+                           const NamedDecl *ND) -> llvm::Optional<unsigned> {
+    if (const auto *RD = dyn_cast<CXXRecordDecl>(ND))
+      Ctx.AddSYCLKernelNamingDecl(RD);
+    // We always want to go into the lambda mangling (skipping the unnamed
+    // struct version), so make sure we return a value here.
+    return 1;
+  };
+
+  QualType Ty = GetSYCLKernelObjectType(FD);
+  std::unique_ptr<MangleContext> Ctx{ItaniumMangleContext::create(
+      Context, Context.getDiagnostics(), MangleCallback)};
+  llvm::raw_null_ostream Out;
+  Ctx->mangleTypeName(Ty, Out);
+}
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 37684b68cb43e..5ba50747d84ca 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -1440,52 +1440,11 @@ TemplateName TemplateInstantiator::TransformTemplateName(
                                           AllowInjectedClassName);
 }
 
-static ExprResult TransformUniqueStableName(TemplateInstantiator &TI,
-                                            PredefinedExpr *E) {
-  if (E->getIdentKind() == PredefinedExpr::UniqueStableNameType) {
-    TypeSourceInfo *Info =
-        TI.getDerived().TransformType(E->getTypeSourceInfo());
-
-    if (!Info)
-      return ExprError();
-
-    if (!TI.getDerived().AlwaysRebuild() && Info == E->getTypeSourceInfo())
-      return E;
-
-    return TI.getSema().BuildUniqueStableName(E->getLocation(), Info);
-  }
-
-  if (E->getIdentKind() == PredefinedExpr::UniqueStableNameExpr) {
-    EnterExpressionEvaluationContext Unevaluated(
-        TI.getSema(), Sema::ExpressionEvaluationContext::Unevaluated);
-    ExprResult SubExpr = TI.getDerived().TransformExpr(E->getExpr());
-
-    if (SubExpr.isInvalid())
-      return ExprError();
-
-    SubExpr = TI.getSema().CheckPlaceholderExpr(SubExpr.get());
-
-    if (SubExpr.isInvalid())
-      return ExprError();
-
-    if (!TI.getDerived().AlwaysRebuild() && SubExpr.get() == E->getExpr())
-      return E;
-
-    return TI.getSema().BuildUniqueStableName(E->getLocation(), SubExpr.get());
-  }
-
-  llvm_unreachable("Only valid for UniqueStableNameType/Expr");
-}
-
 ExprResult
 TemplateInstantiator::TransformPredefinedExpr(PredefinedExpr *E) {
   if (!E->isTypeDependent())
     return E;
 
-  if (E->getIdentKind() == PredefinedExpr::UniqueStableNameType ||
-      E->getIdentKind() == PredefinedExpr::UniqueStableNameExpr)
-    return TransformUniqueStableName(*this, E);
-
   return getSema().BuildPredefinedExpr(E->getLocation(), E->getIdentKind());
 }
 
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 1d19f81c50bcd..fdbdfaff2aa65 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -761,6 +761,35 @@ static void instantiateWorkGroupSizeHintAttr(
                              ZResult.get());
 }
 
+// This doesn't take any template parameters, but we have a custom action that
+// needs to happen when the kernel itself is instantiated. We need to run the
+// ItaniumMangler to mark the names required to name this kernel.
+static void instantiateDependentSYCLKernelAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const SYCLKernelAttr &Attr, Decl *New) {
+  // Functions cannot be partially specialized, so if we are being instantiated,
+  // we are obviously a complete specialization. Since this attribute is only
+  // valid on function template declarations, we know that this is a full
+  // instantiation of a kernel.
+  S.AddSYCLKernelLambda(cast<FunctionDecl>(New));
+
+  // Evaluate whether this would change any of the already evaluated
+  // __builtin_sycl_unique_stable_name values.
+  for (auto &Itr : S.Context.SYCLUniqueStableNameEvaluatedValues) {
+    const std::string &CurName = Itr.first->ComputeName(S.Context);
+    if (Itr.second != CurName) {
+      S.Diag(New->getLocation(),
+             diag::err_kernel_invalidates_sycl_unique_stable_name);
+      S.Diag(Itr.first->getLocation(),
+             diag::note_sycl_unique_stable_name_evaluated_here);
+      // Update this so future diagnostics work correctly.
+      Itr.second = CurName;
+    }
+  }
+
+  New->addAttr(Attr.clone(S.getASTContext()));
+}
+
 /// Determine whether the attribute A might be relevent to the declaration D.
 /// If not, we can skip instantiating it. The attribute may or may not have
 /// been instantiated yet.
@@ -1040,6 +1069,11 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
       continue;
     }
 
+    if (auto *A = dyn_cast<SYCLKernelAttr>(TmplAttr)) {
+      instantiateDependentSYCLKernelAttr(*this, TemplateArgs, *A, New);
+      continue;
+    }
+
     assert(!TmplAttr->isPackExpansion());
     if (TmplAttr->isLateParsed() && LateAttrs) {
       // Late parsed attributes must be instantiated and attached after the
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 0ae76326717b8..60f1eb3716def 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -2403,6 +2403,19 @@ class TreeTransform {
     return SEHFinallyStmt::Create(getSema().getASTContext(), Loc, Block);
   }
 
+  ExprResult RebuildSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                             SourceLocation LParen,
+                                             SourceLocation RParen, Expr *E) {
+    return getSema().BuildSYCLUniqueStableNameExpr(OpLoc, LParen, RParen, E);
+  }
+
+  ExprResult RebuildSYCLUniqueStableNameExpr(SourceLocation OpLoc,
+                                             SourceLocation LParen,
+                                             SourceLocation RParen,
+                                             TypeSourceInfo *TSI) {
+    return getSema().BuildSYCLUniqueStableNameExpr(OpLoc, LParen, RParen, TSI);
+  }
+
   /// Build a new predefined expression.
   ///
   /// By default, performs semantic analysis to build the new expression.
@@ -10179,6 +10192,24 @@ TreeTransform<Derived>::TransformConstantExpr(ConstantExpr *E) {
   return TransformExpr(E->getSubExpr());
 }
 
+template <typename Derived>
+ExprResult TreeTransform<Derived>::TransformSYCLUniqueStableNameExpr(
+    SYCLUniqueStableNameExpr *E) {
+  if (!E->isTypeDependent())
+    return E;
+
+  TypeSourceInfo *NewT = getDerived().TransformType(E->getTypeSourceInfo());
+
+  if (!NewT)
+    return ExprError();
+
+  if (!getDerived().AlwaysRebuild() && E->getTypeSourceInfo() == NewT)
+    return E;
+
+  return getDerived().RebuildSYCLUniqueStableNameExpr(
+      E->getLocation(), E->getLParenLocation(), E->getRParenLocation(), NewT);
+}
+
 template<typename Derived>
 ExprResult
 TreeTransform<Derived>::TransformPredefinedExpr(PredefinedExpr *E) {
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 04adbabc9b89b..c1f66709c8515 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -581,6 +581,16 @@ void ASTStmtReader::VisitConstantExpr(ConstantExpr *E) {
   E->setSubExpr(Record.readSubExpr());
 }
 
+void ASTStmtReader::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
+  VisitExpr(E);
+
+  E->setLocation(readSourceLocation());
+  E->setLParenLocation(readSourceLocation());
+  E->setRParenLocation(readSourceLocation());
+
+  E->setTypeSourceInfo(Record.readTypeSourceInfo());
+}
+
 void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) {
   VisitExpr(E);
   bool HasFunctionName = Record.readInt();
@@ -2802,6 +2812,10 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
                        /*StorageKind=*/Record[ASTStmtReader::NumExprFields]));
       break;
 
+    case EXPR_SYCL_UNIQUE_STABLE_NAME:
+      S = SYCLUniqueStableNameExpr::CreateEmpty(Context);
+      break;
+
     case EXPR_PREDEFINED:
       S = PredefinedExpr::CreateEmpty(
           Context,
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 017de0db357b5..10be9b894c277 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -580,6 +580,17 @@ void ASTStmtWriter::VisitConstantExpr(ConstantExpr *E) {
   Code = serialization::EXPR_CONSTANT;
 }
 
+void ASTStmtWriter::VisitSYCLUniqueStableNameExpr(SYCLUniqueStableNameExpr *E) {
+  VisitExpr(E);
+
+  Record.AddSourceLocation(E->getLocation());
+  Record.AddSourceLocation(E->getLParenLocation());
+  Record.AddSourceLocation(E->getRParenLocation());
+  Record.AddTypeSourceInfo(E->getTypeSourceInfo());
+
+  Code = serialization::EXPR_SYCL_UNIQUE_STABLE_NAME;
+}
+
 void ASTStmtWriter::VisitPredefinedExpr(PredefinedExpr *E) {
   VisitExpr(E);
 
diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index ccfc84bd7a58a..d40161f289f88 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1419,6 +1419,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
     case Stmt::OMPArraySectionExprClass:
     case Stmt::OMPArrayShapingExprClass:
     case Stmt::OMPIteratorExprClass:
+    case Stmt::SYCLUniqueStableNameExprClass:
     case Stmt::TypeTraitExprClass: {
       Bldr.takeNodes(Pred);
       ExplodedNodeSet preVisit;
diff --git a/clang/test/AST/ast-print-sycl-unique-stable-name.cpp b/clang/test/AST/ast-print-sycl-unique-stable-name.cpp
new file mode 100644
index 0000000000000..3f49ea9ee733c
--- /dev/null
+++ b/clang/test/AST/ast-print-sycl-unique-stable-name.cpp
@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -ast-print -fsycl-is-device %s -o - -triple spir64-sycldevice | FileCheck %s
+
+template <typename T>
+void WrappedInTemplate(T t) {
+  (void)__builtin_sycl_unique_stable_name(T);
+  (void)__builtin_sycl_unique_stable_name(typename T::type);
+  (void)__builtin_sycl_unique_stable_name(decltype(t.foo()));
+}
+
+struct Type {
+  using type = int;
+
+  double foo();
+};
+
+void use() {
+  WrappedInTemplate(Type{});
+}
+
+// CHECK: template <typename T> void WrappedInTemplate(T t)
+// CHECK: __builtin_sycl_unique_stable_name(T);
+// CHECK: __builtin_sycl_unique_stable_name(typename T::type);
+// CHECK: __builtin_sycl_unique_stable_name(decltype(t.foo()));
+
+// CHECK: template<> void WrappedInTemplate<Type>(Type t)
+// CHECK: __builtin_sycl_unique_stable_name(Type);
+// CHECK: __builtin_sycl_unique_stable_name(typename Type::type);
+// CHECK: __builtin_sycl_unique_stable_name(decltype(t.foo()));
diff --git a/clang/test/Analysis/eval-predefined-exprs.cpp b/clang/test/Analysis/eval-predefined-exprs.cpp
index 947ef148eeef0..1eec4476a065f 100644
--- a/clang/test/Analysis/eval-predefined-exprs.cpp
+++ b/clang/test/Analysis/eval-predefined-exprs.cpp
@@ -7,12 +7,6 @@ template <typename T>
 void clang_analyzer_dump(const T *);
 void clang_analyzer_warnIfReached();
 
-void builtin_unique_stable_name_of_lambda() {
-  auto y = [] {};
-  clang_analyzer_dump(__builtin_unique_stable_name(y));
-  // expected-warning@-1 {{&Element{"_ZTSZ36builtin_unique_stable_name_of_lambdavEUlvE11_12",0 S64b,char}}}
-}
-
 template <typename T, auto Value, typename U>
 void func(U param) {
   clang_analyzer_dump(__func__);
@@ -56,11 +50,6 @@ void foo() {
   func<struct Class, 42ull>('b'); // instantiate template
 }
 
-void test_builtin_unique_stable_name(int a) {
-  clang_analyzer_dump(__builtin_unique_stable_name(a));
-  // expected-warning@-1 {{&Element{"_ZTSi",0 S64b,char}}}
-}
-
 struct A {
   A() {
     clang_analyzer_dump(__func__);
diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
index 4946d1583036b..cafe1c5ad496e 100644
--- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp
+++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp
@@ -49,8 +49,8 @@ int main() {
 // CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
 //
 // Lambda object alloca
-// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %"class{{.*}}.anon"
-// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}.anon"* [[KERNEL]] to %"class{{.*}}.anon" addrspace(4)*
+// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon
+// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
 //
 // Kernel argument stores
 // CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast
@@ -60,7 +60,7 @@ int main() {
 // CHECK: store i32 [[ARG_C]], i32 addrspace(4)* [[ARG_C]].addr.ascast
 //
 // Check A and B scalar fields initialization
-// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
 // CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
 // CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
 // CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
@@ -85,13 +85,13 @@ int main() {
 // CHECK: store i32 [[ARG_C_LOAD]], i32 addrspace(4)* [[FIELD_C]]
 //
 // Check __init method calls
-// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
 // CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
 // CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
 // CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
 //
-// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
+// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
 // CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
 // CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]
diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
index c5cfc1cd1c14b..ea4e366de739e 100644
--- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
@@ -27,8 +27,8 @@ int main() {
 // Check alloca for pointer argument
 // CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
 // Check lambda object alloca
-// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %"class.{{.*}}.anon"
-// CHECK: [[ANON:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[ANONALLOCA]] to %"class.{{.*}}.anon" addrspace(4)*
+// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.{{.*}}.anon
+// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[ANONALLOCA]] to %class.{{.*}}.anon addrspace(4)*
 // Check allocas for ranges
 // CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range"
 // CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ARANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)*
@@ -44,7 +44,7 @@ int main() {
 // CHECK: call spir_func {{.*}}accessor
 
 // Check accessor GEP
-// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[ANON]], i32 0, i32 0
+// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[ANON]], i32 0, i32 0
 
 // Check load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast
@@ -56,4 +56,4 @@ int main() {
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
 
 // Check lambda "()" operator call
-// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}})
+// CHECK: call spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* {{[^,]*}})
diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
index 9c2300075f152..273707ed6c021 100644
--- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
+++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
@@ -25,9 +25,9 @@ int main() {
 }
 
 // CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
-// CHECK: getelementptr inbounds %"class.{{.*}}.anon"{{.*}} !dbg [[LINE_A0:![0-9]+]]
+// CHECK: getelementptr inbounds %class.{{.*}}.anon{{.*}} !dbg [[LINE_A0:![0-9]+]]
 // CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
-// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"{{.*}} !dbg [[LINE_B0:![0-9]+]]
+// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]
 // CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]
 // CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}})
 // CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "{{.*}}19use_kernel_for_test"
diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp
index 07da10edd5ba1..8bb7189f56dcf 100644
--- a/clang/test/CodeGenSYCL/device-functions.cpp
+++ b/clang/test/CodeGenSYCL/device-functions.cpp
@@ -22,6 +22,6 @@ int main() {
   return 0;
 }
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE11fake_kernel()
-// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %this)
+// CHECK: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %this)
 // CHECK: define {{.*}}spir_func void @_Z3foov()
 // CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)
diff --git a/clang/test/CodeGenSYCL/device-variables.cpp b/clang/test/CodeGenSYCL/device-variables.cpp
index b0964bf6b4f28..4c8d6844738c9 100644
--- a/clang/test/CodeGenSYCL/device-variables.cpp
+++ b/clang/test/CodeGenSYCL/device-variables.cpp
@@ -32,10 +32,10 @@ int main() {
     // CHECK: store i32 1, i32 addrspace(4)* %b
     foo(local_value);
     // Local variables and constexprs captured by lambda
-    // CHECK:  [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 0
+    // CHECK:  [[GEP:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 0
     // CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) [[GEP]])
     int some_device_local_var = some_local_var;
-    // CHECK:  [[GEP1:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 1
+    // CHECK:  [[GEP1:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 1
     // CHECK:  [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]]
     // CHECK:  store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var
   });
diff --git a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
index 97c4c754896b5..4f44b8c4607d0 100644
--- a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
+++ b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp
@@ -29,8 +29,8 @@ int main() {
   return 0;
 }
 
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} ![[NUM4:[0-9]+]]
-// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
+// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel1() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
+// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel2() #0 {{.*}} ![[NUM4:[0-9]+]]
+// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel3() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
 // CHECK: ![[NUM4]] = !{}
 // CHECK: ![[NUM5]] = !{i32 1}
diff --git a/clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp b/clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp
index 97f66783e22bd..c8a2e3363d195 100644
--- a/clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp
+++ b/clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp
@@ -26,7 +26,7 @@ void test(int val) {
   });
 
   // --- Name
-  // CHECK-LABEL: define {{.*}}spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE12esimd_kernel"(
+  // CHECK-LABEL: define {{.*}}spir_kernel void @_ZTSZZ4testiENKUlRN2cl4sycl7handlerEE_clES2_E12esimd_kernel(
   // --- Attributes
   // CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
   // --- init_esimd call is expected instead of __init:
diff --git a/clang/test/CodeGenSYCL/initiation_interval.cpp b/clang/test/CodeGenSYCL/initiation_interval.cpp
index 0fa5699b32e99..93ba8190b358e 100644
--- a/clang/test/CodeGenSYCL/initiation_interval.cpp
+++ b/clang/test/CodeGenSYCL/initiation_interval.cpp
@@ -39,10 +39,10 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} ![[NUM0:[0-9]+]]
 // CHECK: ![[NUM0]] = !{}
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM42]] = !{i32 42}
diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
index f8edf692f66e2..d85295bc53178 100644
--- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
+++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
@@ -40,10 +40,10 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !no_global_work_offset ![[NUM5]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} ![[NUM4:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !no_global_work_offset ![[NUM5]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !no_global_work_offset ![[NUM5]]
 // CHECK-NOT: ![[NUM4]]  = !{i32 0}
 // CHECK: ![[NUM5]] = !{}
diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
index c127e98280114..53ab36ec9ab92 100644
--- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
@@ -37,9 +37,9 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_global_work_dim ![[NUM2]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_global_work_dim ![[NUM2]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_global_work_dim ![[NUM2]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_global_work_dim ![[NUM2]]
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM2]] = !{i32 2}
diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
index f59c8e91f9191..28c0bae228871 100644
--- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
@@ -45,11 +45,11 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !max_work_group_size ![[NUM1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_work_group_size ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]]
 // CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1}
 // CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8}
 // CHECK: ![[NUM6]] = !{i32 6, i32 3, i32 1}
diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp
index e66040e1a3747..1da504f291b14 100644
--- a/clang/test/CodeGenSYCL/kernel-handler.cpp
+++ b/clang/test/CodeGenSYCL/kernel-handler.cpp
@@ -22,16 +22,16 @@ void test(int val) {
   });
 }
 
-// NONATIVESUPPORT: define dso_local void @"{{.*}}test_kernel_handler{{.*}}"
+// NONATIVESUPPORT: define dso_local void @{{.*}}test_kernel_handler{{[^(]*}}
 // NONATIVESUPPORT-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer)
 // NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1
 // NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8
 // NONATIVESUPPORT: %[[ADDRSPACECAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[KH]] to i8*
 // NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]])
-// NONATIVESUPPORT: void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]"
+// NONATIVESUPPORT: void @[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]
 // NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler")
 
-// NATIVESUPPORT: define dso_local spir_kernel void @"{{.*}}test_kernel_handler{{.*}}"
+// NATIVESUPPORT: define dso_local spir_kernel void @{{.*}}test_kernel_handler{{[^(]*}}
 // NATIVESUPPORT-SAME: (i32 %_arg_)
 // NATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler"
 // NATIVESUPPORT-NOT: __init_specialization_constants_buffer
diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
index f4746b8024ceb..8d268ad5b0802 100644
--- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
@@ -39,8 +39,8 @@ int main() {
 // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
 
 // CHECK lambda object alloca
-// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %"class.{{.*}}.anon", align 4
-// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECTA]] to %"class.{{.*}}.anon" addrspace(4)*
+// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.{{.*}}.anon, align 4
+// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class.{{.*}}.anon* [[LOCAL_OBJECTA]] to %class.{{.*}}.anon addrspace(4)*
 
 // CHECK allocas for ranges
 // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
@@ -57,7 +57,7 @@ int main() {
 // CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)*
 
 // CHECK accessor array default inits
-// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0
 // Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
 // CTOR Call #1
@@ -67,7 +67,7 @@ int main() {
 // CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]])
 
 // CHECK acc[0] __init method call
-// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0
 // CHECK load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]]
@@ -77,7 +77,7 @@ int main() {
 // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
 
 // CHECK acc[1] __init method call
-// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY2]], i64 0, i64 1
 // CHECK load from kernel pointer argument alloca
 // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG2]]
diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
index 614259d7b7ed7..7e6d1c49a0fa5 100644
--- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
@@ -42,8 +42,8 @@ int main() {
 // CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
 
 // Check lambda object alloca
-// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %"class{{.*}}.anon", align 4
-// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %"class{{.*}}.anon"* %0 to %"class{{.*}}.anon" addrspace(4)*
+// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class{{.*}}.anon, align 4
+// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class{{.*}}.anon* %0 to %class{{.*}}.anon addrspace(4)*
 
 // Check allocas for ranges
 // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
@@ -60,7 +60,7 @@ int main() {
 // CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)*
 
 // CHECK accessor array default inits
-// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[ACCESSOR_WRAPPER]], i32 0, i32 0
 // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0
 // CTOR Call #1
@@ -70,7 +70,7 @@ int main() {
 // CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]])
 
 // Check acc[0] __init method call
-// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA1]], i32 0, i32 0
 // CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC1]], i64 0, i64 0
 // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr
@@ -80,7 +80,7 @@ int main() {
 // CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ARRAY_IDX1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
 
 // Check acc[1] __init method call
-// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}.struct_acc_t, %struct{{.*}}.struct_acc_t addrspace(4)* [[GEP_LAMBDA2]], i32 0, i32 0
 // CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[GEP_MEMBER_ACC2]], i64 0, i64 1
 // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]].addr
diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
index 9e1ca7c028c9d..27806edaa29e1 100644
--- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
+++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
@@ -49,11 +49,11 @@ int main() {
 // CHECK-SAME:(%struct{{.*}}.__wrapper_class* byval(%struct{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon", align 4
-// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon" addrspace(4)*
+// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon, align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon addrspace(4)*
 
 // Check for Array init loop
-// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class, %struct{{.*}}.__wrapper_class addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0
 // CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x i32], [2 x i32] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0
 // CHECK: br label %[[ARRAYINITBODY:.+]]
@@ -74,11 +74,11 @@ int main() {
 // CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon{{.*}}", align 4
-// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon{{.*}}" addrspace(4)*
+// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
 
 // Check for Array init loop
-// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon{{.*}}", %"class{{.*}}.anon{{.*}}" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, %class{{.*}}.anon{{.*}} addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class{{.*}}, %struct{{.*}}.__wrapper_class{{.*}} addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0
 // CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x %struct{{.*}}.foo], [2 x %struct{{.*}}.foo] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0
 // CHECK: br label %[[ARRAYINITBODY:.+]]
@@ -100,11 +100,11 @@ int main() {
 // CHECK-SAME:(%struct{{.*}}.__wrapper_class{{.*}}* byval(%struct{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]])
 
 // Check local lambda object alloca
-// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %"class{{.*}}.anon{{.*}}", align 4
-// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %"class{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECTA]] to %"class{{.*}}.anon{{.*}}" addrspace(4)*
+// CHECK: %[[LOCAL_OBJECTA:[0-9]+]] = alloca %class{{.*}}.anon{{.*}}, align 4
+// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = addrspacecast %class{{.*}}.anon{{.*}}* %[[LOCAL_OBJECTA]] to %class{{.*}}.anon{{.*}} addrspace(4)*
 
 // Check for Array init loop
-// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class{{.*}}.anon{{.*}}", %"class{{.*}}.anon{{.*}}" addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
+// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %class{{.*}}.anon{{.*}}, %class{{.*}}.anon{{.*}} addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 0
 // CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct{{.*}}.__wrapper_class{{.*}}, %struct{{.*}}.__wrapper_class{{.*}} addrspace(4)* %[[ARR_ARG]].ascast, i32 0, i32 0
 // CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]] addrspace(4)* %[[LAMBDA_PTR]], i64 0, i64 0
 // CHECK: br label %[[ARRAYINITBODY:.+]]
diff --git a/clang/test/CodeGenSYCL/loop_fuse_device.cpp b/clang/test/CodeGenSYCL/loop_fuse_device.cpp
index 645bb908d24c1..7a37d101b58e0 100644
--- a/clang/test/CodeGenSYCL/loop_fuse_device.cpp
+++ b/clang/test/CodeGenSYCL/loop_fuse_device.cpp
@@ -37,12 +37,12 @@ void bar() {
   });
 }
 
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_1"() {{.*}} !loop_fuse ![[LF5:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_2"() {{.*}} !loop_fuse ![[LF1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_3"() {{.*}} !loop_fuse ![[LF0:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_4"() {{.*}} !loop_fuse ![[LF1]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_5"() {{.*}} !loop_fuse ![[LF10:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_6"()
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_1() {{.*}} !loop_fuse ![[LF5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_2() {{.*}} !loop_fuse ![[LF1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_3() {{.*}} !loop_fuse ![[LF0:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_4() {{.*}} !loop_fuse ![[LF1]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_5() {{.*}} !loop_fuse ![[LF10:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_6()
 // CHECK-NOT: !loop_fuse
 // CHECK-SAME: {
 // CHECK: define {{.*}}spir_func void @{{.*}}foo{{.*}} !loop_fuse ![[LF5]]
diff --git a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp
index 49cc7640e1de7..c28c37443e471 100644
--- a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp
+++ b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp
@@ -37,12 +37,12 @@ void bar() {
   });
 }
 
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_1"() {{.*}} !loop_fuse ![[LFI5:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_2"() {{.*}} !loop_fuse ![[LFI1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_3"() {{.*}} !loop_fuse ![[LFI0:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_4"() {{.*}} !loop_fuse ![[LFI1]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_5"() {{.*}} !loop_fuse ![[LFI10:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name_6"()
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_1() {{.*}} !loop_fuse ![[LFI5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_2() {{.*}} !loop_fuse ![[LFI1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_3() {{.*}} !loop_fuse ![[LFI0:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_4() {{.*}} !loop_fuse ![[LFI1]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_5() {{.*}} !loop_fuse ![[LFI10:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name_6()
 // CHECK-NOT: !loop_fuse
 // CHECK-SAME: {
 // CHECK: define {{.*}}spir_func void @{{.*}}foo{{.*}} !loop_fuse ![[LFI5]]
diff --git a/clang/test/CodeGenSYCL/max-concurrency.cpp b/clang/test/CodeGenSYCL/max-concurrency.cpp
index 5bafcd12ab89e..4b92d9b6183da 100644
--- a/clang/test/CodeGenSYCL/max-concurrency.cpp
+++ b/clang/test/CodeGenSYCL/max-concurrency.cpp
@@ -18,7 +18,7 @@
 // CHECK: br label %for.cond2,  !llvm.loop ![[MD_MC_1:[0-9]+]]
 // CHECK: ret void
 
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name1"() [[ATTR0:#.*]] {{.*}} !max_concurrency ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() [[ATTR0:#.*]] {{.*}} !max_concurrency ![[NUM1:[0-9]+]]
 // CHECK: entry:
 // CHECK: [[F1:%.*]] = alloca [[CLASS_F1:%.*]], align 1
 // CHECK: [[F1_ASCAST:%.*]] = addrspacecast [[CLASS_F1]]* [[F1]] to [[CLASS_F1]] addrspace(4)*
@@ -29,7 +29,7 @@
 // CHECK: call void @llvm.lifetime.end.p0i8(i64 1, i8* [[TMP1]])
 // CHECK: ret void
 
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}kernel_name4"() [[ATTR0]] {{.*}} !max_concurrency ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() [[ATTR0]] {{.*}} !max_concurrency ![[NUM1:[0-9]+]]
 // CHECK: entry
 // CHECK: [[F3:%.*]] = alloca [[CLASS_F3:%.*]], align 1
 // CHECK: [[F3_ASCAST:%.*]] = addrspacecast [[CLASS_F3]]* [[F3]] to [[CLASS_F3]] addrspace(4)*
@@ -48,18 +48,18 @@
 // CHECK: %this1 = load [[CLASS_F3]] addrspace(4)*, [[CLASS_F3]] addrspace(4)* addrspace(4)* [[ADDR1_CAST]], align 8
 // CHECK: ret void
 
-// CHECK: define dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE12kernel_name5"()
+// CHECK: define dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E12kernel_name5()
 // CHECK: entry:
 // CHECK: [[H1:%.*]] = alloca [[H:%.*]], align 1
 // CHECK: [[H2:%.*]] = addrspacecast [[H]]* [[H1]] to [[H]] addrspace(4)*
 // CHECK: [[H3:%.*]] = bitcast [[H]]* [[H1]] to i8*
 // CHECK: call void @llvm.lifetime.start.p0i8(i64 1, i8* [[H3]])
-// CHECK: call spir_func void @"_ZZZ4mainENK3$_1clERN2cl4sycl7handlerEENKUlvE_clEv"([[H]] addrspace(4)* align 1 dereferenceable_or_null(1) [[H2]])
+// CHECK: call spir_func void @_ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlvE_clEv([[H]] addrspace(4)* align 1 dereferenceable_or_null(1) [[H2]])
 // CHECK: [[TMP4:%.*]] = bitcast [[H]]* [[H1]] to i8*
 // CHECK: call void @llvm.lifetime.end.p0i8(i64 1, i8* [[TMP4]])
 // CHECK: ret void
 
-// CHECK: define {{.*}}spir_func void @"_ZZZ4mainENK3$_1clERN2cl4sycl7handlerEENKUlvE_clEv"
+// CHECK: define {{.*}}spir_func void @_ZZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_ENKUlvE_clEv
 // CHECK: entry:
 // CHECK: [[ADDR_1:%.*]] = alloca [[HH:%.*]] addrspace(4)*, align 8
 // CHECK: [[ADDR1_CAST:%.*]] = addrspacecast [[HH]] addrspace(4)** [[ADDR_1]] to [[HH]] addrspace(4)* addrspace(4)*
diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
index a5c28285f0ef0..eebce1f408de4 100644
--- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp
+++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp
@@ -37,10 +37,10 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !num_simd_work_items ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !num_simd_work_items ![[NUM4:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !num_simd_work_items ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !num_simd_work_items ![[NUM4:[0-9]+]]
 // CHECK: ![[NUM1]] = !{i32 1}
 // CHECK: ![[NUM42]] = !{i32 42}
 // CHECK: ![[NUM2]] = !{i32 2}
diff --git a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
index 25269187e9346..020807bcb130b 100755
--- a/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
+++ b/clang/test/CodeGenSYCL/parallel_for_this_item.cpp
@@ -10,17 +10,17 @@
 
 // CHECK: static constexpr
 // CHECK-NEXT: const char* const kernel_names[] = {
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3GNU",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX",
-// CHECK-NEXT:   "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3GNU",
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3EMU",
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL",
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3RAT",
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX",
+// CHECK-NEXT:   "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"
 // CHECK-NEXT: };
 
 // CHECK:template <> struct KernelInfo<class GNU> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3GNU"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3GNU"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -36,7 +36,7 @@
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class EMU> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3EMU"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3EMU"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -52,7 +52,7 @@
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class OWL> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3OWL"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3OWL"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -68,7 +68,7 @@
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class RAT> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3RAT"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3RAT"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -84,7 +84,7 @@
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class FOX> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3FOX"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3FOX"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
@@ -100,7 +100,7 @@
 // CHECK-NEXT:};
 // CHECK-NEXT:template <> struct KernelInfo<class BEE> {
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
-// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3BEE"; }
+// CHECK-NEXT:    static constexpr const char* getName() { return "_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E3BEE"; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
 // CHECK-NEXT:    static constexpr unsigned getNumParams() { return 0; }
 // CHECK-NEXT:  __SYCL_DLL_LOCAL
diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
index aa9c104cfc61f..128a69052118e 100644
--- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
@@ -49,11 +49,11 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]]
 // CHECK: ![[SGSIZE16]] = !{i32 16}
 // CHECK: ![[SGSIZE8]] = !{i32 8}
 // CHECK: ![[SGSIZE4]] = !{i32 4}
diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
index 11047dbff3168..0986e89738514 100644
--- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
+++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp
@@ -52,12 +52,12 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE32:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE8:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE88:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE22:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE44:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !reqd_work_group_size ![[WGSIZE32:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !reqd_work_group_size ![[WGSIZE8:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !reqd_work_group_size ![[WGSIZE88:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !reqd_work_group_size ![[WGSIZE22:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !reqd_work_group_size ![[WGSIZE44:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]]
 // CHECK: ![[WGSIZE32]] = !{i32 16, i32 16, i32 32}
 // CHECK: ![[WGSIZE8]] = !{i32 1, i32 1, i32 8}
 // CHECK: ![[WGSIZE88]] = !{i32 8, i32 8, i32 8}
diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp
index b771290db7929..eadd5f7dec2ba 100644
--- a/clang/test/CodeGenSYCL/sampler.cpp
+++ b/clang/test/CodeGenSYCL/sampler.cpp
@@ -2,12 +2,12 @@
 // CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
 // CHECK-NEXT: entry:
 // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
-// CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8
-// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[ANON]] to %"class.{{.*}}.anon" addrspace(4)*
+// CHECK: [[ANON:%[0-9]+]] = alloca %class.{{.*}}.anon, align 8
+// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[ANON]] to %class.{{.*}}.anon addrspace(4)*
 // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
-// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8*
+// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.{{.*}}.anon* [[ANON]] to i8*
 // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
-// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[ANONCAST]], i32 0, i32 0
+// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0
 // CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8
 // CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
 //
@@ -17,21 +17,21 @@
 // Check alloca
 // CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
 // CHECK: [[ARG_A]].addr = alloca i32, align 4
-// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 8
-// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon.0"* [[LAMBDAA]] to %"class.{{.*}}.anon.0" addrspace(4)*
+// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.{{.*}}.anon, align 8
+// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[LAMBDAA]] to %class.{{.*}}.anon addrspace(4)*
 
 // Check argument store
 // CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8
 // CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4
 
 // Initialize 'a'
-// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0" addrspace(4)* [[LAMBDA]], i32 0, i32 0
+// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LAMBDA]], i32 0, i32 0
 // CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1
 // CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4
 // CHECK: store i32 [[LOAD_A]], i32 addrspace(4)* [[GEP_A]], align 8
 
 // Initialize wrapped sampler 'smpl'
-// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0" addrspace(4)* [[LAMBDA]], i32 0, i32 0
+// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LAMBDA]], i32 0, i32 0
 // CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.sampler_wrapper, %struct.{{.*}}.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0
 // CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8
 // CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]])
diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
index 715e2dd204459..55cf1f176e9fe 100644
--- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
+++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp
@@ -46,11 +46,11 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]]
 // CHECK: ![[NUM5]] = !{i32 5}
 // CHECK: ![[NUM42]] = !{i32 42}
 // CHECK: ![[NUM7]] = !{i32 7}
diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
index d810ba897f4e5..afd04ac4ae5ee 100644
--- a/clang/test/CodeGenSYCL/spir-calling-conv.cpp
+++ b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
@@ -9,9 +9,9 @@ int main() {
 
   // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function()
 
-  // CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %{{.+}})
+  // CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %{{.+}})
 
-  // CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* {{[^,]*}} %this)
+  // CHECK: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.{{.*}}anon addrspace(4)* {{[^,]*}} %this)
 
   kernel_single_task<class kernel_function>([]() {});
   return 0;
diff --git a/clang/test/CodeGenSYCL/spir-enum.cpp b/clang/test/CodeGenSYCL/spir-enum.cpp
index 09d12be1d190a..bb71291353cb4 100644
--- a/clang/test/CodeGenSYCL/spir-enum.cpp
+++ b/clang/test/CodeGenSYCL/spir-enum.cpp
@@ -22,8 +22,8 @@ int main() {
 
   // CHECK: define {{.*}}spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_)
 
-  // CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)*
-  // CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %{{.+}})
+  // CHECK: getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)*
+  // CHECK: call spir_func void @_ZZ4test9enum_typeENKUlvE_clEv(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} %{{.+}})
 
   test( enum_type::B );
   return 0;
diff --git a/clang/test/CodeGenSYCL/stall_enable.cpp b/clang/test/CodeGenSYCL/stall_enable.cpp
index 1cf6036afb50d..0c9a0ba7653b6 100644
--- a/clang/test/CodeGenSYCL/stall_enable.cpp
+++ b/clang/test/CodeGenSYCL/stall_enable.cpp
@@ -21,6 +21,6 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !stall_enable ![[NUM5:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} !stall_enable ![[NUM5]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel1() #0 {{.*}}!stall_enable ![[NUM5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}test_kernel2() #0 {{.*}}!stall_enable ![[NUM5]]
 // CHECK: ![[NUM5]] = !{i32 1}
diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
index 31a9de56d2056..73f1fe27df465 100644
--- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp
+++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp
@@ -2,7 +2,7 @@
 // RUN: FileCheck -input-file=%t.h %s
 
 // CHECK:     const kernel_param_desc_t kernel_signatures[] = {
-// CHECK-NEXT:  //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel
+// CHECK-NEXT:  //--- _ZTSZZ5test0vENKUlRN2cl4sycl7handlerEE_clES2_E8MyKernel
 // Accessor
 // CHECK-NEXT:  { kernel_param_kind_t::kind_accessor, 4062, 0 },
 // FldInt, offset to 16 because the float* causes the alignment of the structs
diff --git a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
index 61bab3383c0bb..e1ac2cf63a035 100644
--- a/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
+++ b/clang/test/CodeGenSYCL/sycl-device-static-init.cpp
@@ -7,7 +7,7 @@
 // CHECK-NOT: @_ZN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE = weak_odr addrspace(1) global %struct._ZTS16RegisterBaseInit.RegisterBaseInit zeroinitializer, comdat, align 1
 // CHECK-NOT: @_ZGVN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE = weak_odr global i64 0, comdat($_ZN8BaseInitI12TestBaseTypeE9s_regbaseE), align 8
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE11fake_kernel()
-// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv
+// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv
 
 struct TestBaseType {};
 struct RegisterBaseInit {
diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
index f2cecd45fcf8d..3c7c80016cc23 100644
--- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
+++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
@@ -42,10 +42,10 @@ int main() {
   return 0;
 }
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE1:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE3:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !reqd_work_group_size ![[WGSIZE1:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !reqd_work_group_size ![[WGSIZE3:[0-9]+]]
 // CHECK: ![[WGSIZE]] = !{i32 16, i32 16, i32 32}
 // CHECK: ![[SGSIZE]] = !{i32 4}
 // CHECK: ![[WGSIZE1]] = !{i32 32, i32 32, i32 64}
diff --git a/clang/test/CodeGenSYCL/sycl_kernel-host.cpp b/clang/test/CodeGenSYCL/sycl_kernel-host.cpp
index 5bd3cdde3345c..89d70698e5d59 100644
--- a/clang/test/CodeGenSYCL/sycl_kernel-host.cpp
+++ b/clang/test/CodeGenSYCL/sycl_kernel-host.cpp
@@ -10,7 +10,7 @@
 // CHECK: define internal spir_func void @{{.*}}Kernel{{.*}}Bar{{.*}}({{.*}}) #[[SKA]] {
 // CHECK: call spir_func void @{{.*}}KernelImpl{{.*}}({{.*}}, i32 1, double 2.000000e+00)
 // CHECK: define internal spir_func void @{{.*}}KernelImpl{{.*}}({{.*}} %f, i32 %i, double %d) #[[SKA]] {
-// CHECK: call spir_func void @"{{.*}}func{{.*}}"(%class
+// CHECK: call spir_func void @{{.*}}func{{.*}}(%class
 // CHECK: define internal spir_func void @{{.*}}func{{.*}}(%class.anon* {{[^,]*}} %this, i32 %i, double %d) #[[ALWAYSINLINE:[0-9]+]]
 // CHECK: define linkonce_odr spir_func void @{{.*}}KernelImpl{{.*}}Functor{{.*}}({{.*}}, i32 %i, double %d) #[[SKA]] comdat {
 // CHECK: call spir_func void @{{.*}}Functor{{.*}}(%struct
diff --git a/clang/test/CodeGenSYCL/union-kernel-param.cpp b/clang/test/CodeGenSYCL/union-kernel-param.cpp
index c0a07ae26a79b..19d07cb0f0729 100644
--- a/clang/test/CodeGenSYCL/union-kernel-param.cpp
+++ b/clang/test/CodeGenSYCL/union-kernel-param.cpp
@@ -31,12 +31,12 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A(%union.{{.*}}.MyUnion* byval(%union.{{.*}}.MyUnion) align 4 [[MEM_ARG:%[a-zA-Z0-9_]+]])
 
 // Check lambda object alloca
-// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
+// CHECK: [[LOCAL_OBJECT:%0]] = alloca %class.{{.*}}.anon, align 4
 
-// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECT]] to %"class.{{.*}}.anon" addrspace(4)*
+// CHECK: [[LOCAL_OBJECTAS:%.*]] = addrspacecast %class.{{.*}}.anon* [[LOCAL_OBJECT]] to %class.{{.*}}.anon addrspace(4)*
 // CHECK: [[MEM_ARGAS:%.*]] = addrspacecast %union.{{.*}}.MyUnion* [[MEM_ARG]] to %union.{{.*}}.MyUnion addrspace(4)*
-// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0
+// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECTAS]], i32 0, i32 0
 // CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[L_STRUCT_ADDR]] to i8 addrspace(4)*
 // CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion addrspace(4)* [[MEM_ARGAS]] to i8 addrspace(4)*
 // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 4 [[MEMCPY_DST]], i8 addrspace(4)* align 4 [[MEMCPY_SRC]], i64 12, i1 false)
-// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]])
+// CHECK: call spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* {{[^,]*}} [[LOCAL_OBJECTAS]])
diff --git a/clang/test/CodeGenSYCL/unique-stable-name-placeholder-crash.cpp b/clang/test/CodeGenSYCL/unique-stable-name-placeholder-crash.cpp
deleted file mode 100644
index b6da2c9a9d070..0000000000000
--- a/clang/test/CodeGenSYCL/unique-stable-name-placeholder-crash.cpp
+++ /dev/null
@@ -1,28 +0,0 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -internal-isystem %S/Inputs -std=c++17 -sycl-std=2020 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
-
-#include <sycl.hpp>
-
-using namespace cl::sycl;
-
-struct A {
-  int a = 0;
-  A() = default;
-};
-constexpr A THE_NAME;
-
-template <auto &R> void temp() {}
-template <auto &R> void foo(const char *out) {
-  out = __builtin_unique_stable_name(temp<R>);
-}
-
-int main() {
-  kernel_single_task<class kernel>(
-      []() {
-        const char *c;
-        foo<THE_NAME>(c);
-      });
-}
-
-// Note: the mangling here is actually the 'typeinfo name for void ()'.  That is
-// because the type of temp<R> is actually the function type (which is void()).
-// CHECK: @__builtin_unique_stable_name._Z3fooIL_ZL8THE_NAMEEEvPKc = private unnamed_addr addrspace(1) constant [9 x i8] c"_ZTSFvvE\00", align 1
diff --git a/clang/test/CodeGenSYCL/unique-stable-name.cpp b/clang/test/CodeGenSYCL/unique-stable-name.cpp
deleted file mode 100644
index 2259548c1a29e..0000000000000
--- a/clang/test/CodeGenSYCL/unique-stable-name.cpp
+++ /dev/null
@@ -1,77 +0,0 @@
-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
-// CHECK: @[[INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
-// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE46_16\00"
-// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_18\00"
-// CHECK: @[[MACRO_Y:[^\w]+]] =  private unnamed_addr addrspace(1) constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE52_7m28_41\00"
-// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_18m33_4\00"
-// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE42_5clEvEUlvE55_7m28_41m33_4\00"
-// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE23_12\00",
-// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE42_5clEvEUlvE46_16EvvEUlvE23_12\00",
-
-extern "C" void printf(const char *) {}
-
-template <typename T>
-void template_param() {
-  printf(__builtin_unique_stable_name(T));
-}
-
-template <typename T>
-T getT() { return T{}; }
-
-template <typename T>
-void lambda_in_dependent_function() {
-  auto y = [] {};
-  printf(__builtin_unique_stable_name(y));
-}
-
-#define DEF_IN_MACRO()                                  \
-  auto MACRO_X = []() {};auto MACRO_Y = []() {};        \
-  printf(__builtin_unique_stable_name(MACRO_X));        \
-  printf(__builtin_unique_stable_name(MACRO_Y));
-
-#define MACRO_CALLS_MACRO()                             \
-  {DEF_IN_MACRO();}{DEF_IN_MACRO();}
-
-template <typename KernelName, typename KernelType>
-[[clang::sycl_kernel]] void kernel_single_task(const KernelType &kernelFunc) {
-  kernelFunc();
-}
-
-int main() {
-  kernel_single_task<class kernel>(
-    []() {
-      printf(__builtin_unique_stable_name(int));
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(4)* addrspacecast ([[INT_SIZE]] addrspace(1)* @[[INT]] to [[INT_SIZE]] addrspace(4)*
-
-      auto x = [](){};
-      printf(__builtin_unique_stable_name(x));
-      printf(__builtin_unique_stable_name(decltype(x)));
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)*
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)*
-
-      DEF_IN_MACRO();
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_SIZE]] addrspace(1)* @[[MACRO_X]] to [[MACRO_SIZE]] addrspace(4)*
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_SIZE]] addrspace(1)* @[[MACRO_Y]] to [[MACRO_SIZE]] addrspace(4)*
-      MACRO_CALLS_MACRO();
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_X]] to [[MACRO_MACRO_SIZE]] addrspace(4)*
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]] addrspace(4)* addrspacecast ([[MACRO_MACRO_SIZE]] addrspace(1)* @[[MACRO_MACRO_Y]] to [[MACRO_MACRO_SIZE]] addrspace(4)*
-
-      template_param<int>();
-      // CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]] addrspace(4)* addrspacecast ([[INT_SIZE]] addrspace(1)* @[[INT]] to [[INT_SIZE]] addrspace(4)*
-
-      template_param<decltype(x)>();
-      // CHECK: define internal spir_func void @"_Z14template_paramIZZ4mainENK3
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]] addrspace(4)* addrspacecast ([[LAMBDA_X_SIZE]] addrspace(1)* @[[LAMBDA_X]] to [[LAMBDA_X_SIZE]] addrspace(4)*
-
-      lambda_in_dependent_function<int>();
-      // CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]] addrspace(4)* addrspacecast ([[DEP_INT_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_INT]] to [[DEP_INT_SIZE]] addrspace(4)*
-
-      lambda_in_dependent_function<decltype(x)>();
-      // CHECK: define internal spir_func void @"_Z28lambda_in_dependent_functionIZZ4mainENK3$_0clEvEUlvE_Evv
-      // CHECK: call spir_func void @printf(i8 addrspace(4)* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]] addrspace(4)* addrspacecast ([[DEP_LAMBDA_SIZE]] addrspace(1)* @[[LAMBDA_IN_DEP_X]] to [[DEP_LAMBDA_SIZE]] addrspace(4)*
-
-    });
-}
-
diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp
new file mode 100644
index 0000000000000..e084cc64eaea2
--- /dev/null
+++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp
@@ -0,0 +1,164 @@
+// RUN: %clang_cc1 -triple x86_64-linux-pc  -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE10000_E10000_\00"
+// CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
+// CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00",
+// CHECK: @[[INT2:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
+// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE_\00"
+// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE0_\00"
+// CHECK: @[[MACRO_Y:[^\w]+]] =  private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE1_\00"
+// CHECK: @{{.*}} = private unnamed_addr constant [36 x i8] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE2_\00", align 1
+// CHECK: @{{.*}} = private unnamed_addr constant [36 x i8] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE3_\00", align 1
+// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE4_\00"
+// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE5_\00"
+// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
+// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE10001_clEvEUlvE_\00"
+// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00",
+// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE10001_clEvEUlvE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE10001_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00",
+// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE10001_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00",
+
+extern "C" void puts(const char *) {}
+
+template <typename T>
+void template_param() {
+  puts(__builtin_sycl_unique_stable_name(T));
+}
+
+template <typename T>
+void lambda_in_dependent_function() {
+  auto y = [] {};
+  puts(__builtin_sycl_unique_stable_name(decltype(y)));
+}
+
+template <typename Tw, typename Tz>
+void lambda_two_dep() {
+  auto z = [] {};
+  puts(__builtin_sycl_unique_stable_name(decltype(z)));
+}
+
+template <typename Tw, typename Tz>
+void lambda_no_dep(Tw a, Tz b) {
+  auto p = [](Tw a, Tz b) { return ((Tz)a + b); };
+  puts(__builtin_sycl_unique_stable_name(decltype(p)));
+}
+
+#define DEF_IN_MACRO()                                        \
+  auto MACRO_X = []() {};                                     \
+  auto MACRO_Y = []() {};                                     \
+  puts(__builtin_sycl_unique_stable_name(decltype(MACRO_X))); \
+  puts(__builtin_sycl_unique_stable_name(decltype(MACRO_Y)));
+
+#define MACRO_CALLS_MACRO() \
+  { DEF_IN_MACRO(); }       \
+  { DEF_IN_MACRO(); }
+
+template <typename Ty>
+auto func() -> decltype(__builtin_sycl_unique_stable_name(decltype(Ty::str)));
+
+struct Derp {
+  static constexpr const char str[] = "derp derp derp";
+};
+
+template <typename KernelName, typename KernelType>
+[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
+template <typename KernelName, typename KernelType>
+void not_kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  not_kernel_single_task<class kernel2>(func<Derp>);
+  // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8* ()* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
+
+  auto l1 = []() { return 1; };
+  auto l2 = [](decltype(l1) *l = nullptr) { return 2; };
+  kernel_single_task<class kernel3>(l2);
+  puts(__builtin_sycl_unique_stable_name(decltype(l2)));
+  // CHECK: call void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]]* @[[LAMBDA_KERNEL3]], i32 0, i32 0))
+
+  constexpr const char str[] = "lalala";
+  static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling");
+
+  int i = 0;
+  puts(__builtin_sycl_unique_stable_name(decltype(i++)));
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT1]], i32 0, i32 0))
+
+  // FIXME: Ensure that j is incremented because VLAs are terrible.
+  int j = 55;
+  puts(__builtin_sycl_unique_stable_name(int[++j]));
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]]* @[[STRING]], i32 0, i32 0))
+
+  // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
+  // CHECK: declare i8* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_
+
+  kernel_single_task<class kernel>(
+      []() {
+        puts(__builtin_sycl_unique_stable_name(int));
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT2]], i32 0, i32 0))
+
+        auto x = []() {};
+        puts(__builtin_sycl_unique_stable_name(decltype(x)));
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]], i32 0, i32 0))
+
+        DEF_IN_MACRO();
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]], i32 0, i32 0))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]], i32 0, i32 0))
+
+        MACRO_CALLS_MACRO();
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]], i32 0, i32 0))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]], i32 0, i32 0))
+
+        template_param<int>();
+        // CHECK: call void @_Z14template_paramIiEvv
+
+        template_param<decltype(x)>();
+        // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+
+        lambda_in_dependent_function<int>();
+        // CHECK: call void @_Z28lambda_in_dependent_functionIiEvv
+
+        lambda_in_dependent_function<decltype(x)>();
+        // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+
+        lambda_no_dep<int, double>(3, 5.5);
+        // CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 3, double 5.500000e+00)
+
+        int a = 5;
+        double b = 10.7;
+        auto y = [](int a) { return a; };
+        auto z = [](double b) { return b; };
+        lambda_two_dep<decltype(y), decltype(z)>();
+        // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+
+        lambda_two_dep<decltype(z), decltype(y)>();
+        // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+      });
+}
+
+// CHECK: define linkonce_odr void @_Z14template_paramIiEvv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT3]], i32 0, i32 0))
+
+// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]]* @[[LAMBDA]], i32 0, i32 0))
+
+// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0))
+
+// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0))
+
+// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 %a, double %b)
+// CHECK: call void @puts(i8* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]]* @[[LAMBDA_NO_DEP]], i32 0, i32 0))
+
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]]* @[[LAMBDA_TWO_DEP]], i32 0, i32 0))
+
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]]* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0))
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
new file mode 100644
index 0000000000000..88e9174d32cb1
--- /dev/null
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+
+
+template<typename KN, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func F){
+  F();
+}
+
+template<typename KN, typename Func>
+__attribute__((sycl_kernel)) void kernel2(Func F){
+  F(1);
+}
+
+template<typename KN, typename Func>
+__attribute__((sycl_kernel)) void kernel3(Func F){
+  F(1.1);
+}
+
+int main() {
+  int i;
+  double d;
+  float f;
+  auto lambda1 = [](){};
+  auto lambda2 = [](int){};
+  auto lambda3 = [](double){};
+
+  kernel<class K1>(lambda1);
+  kernel2<class K2>(lambda2);
+  kernel3<class K3>(lambda3);
+
+  // Ensure the kernels are named the same between the device and host
+  // invocations.
+  kernel<class USN_Instances>([](){
+  (void)__builtin_sycl_unique_stable_name(decltype(lambda1));
+  (void)__builtin_sycl_unique_stable_name(decltype(lambda2));
+  (void)__builtin_sycl_unique_stable_name(decltype(lambda3));
+  });
+
+  // Make sure the following 3 are the same between the host and device compile.
+  // Note that these are NOT the same value as eachother, they differ by the
+  // signature.
+  // CHECK: private unnamed_addr constant [22 x i8] c"_ZTSZ4mainEUlvE10000_\00"
+  // CHECK: private unnamed_addr constant [22 x i8] c"_ZTSZ4mainEUliE10000_\00"
+  // CHECK: private unnamed_addr constant [22 x i8] c"_ZTSZ4mainEUldE10000_\00"
+}
diff --git a/clang/test/ParserSYCL/unique-stable-name.cpp b/clang/test/ParserSYCL/unique-stable-name.cpp
deleted file mode 100644
index d1f1304cf8b45..0000000000000
--- a/clang/test/ParserSYCL/unique-stable-name.cpp
+++ /dev/null
@@ -1,33 +0,0 @@
-// RUN: %clang_cc1 -fsyntax-only -verify -Wno-unused %s
-
-namespace NS{};
-
-void f(int var) {
-  // expected-error@+1{{expected '(' after '__builtin_unique_stable_name'}}
-  __builtin_unique_stable_name int;
-  // expected-error@+1{{expected '(' after '__builtin_unique_stable_name'}}
-  __builtin_unique_stable_name {int};
-
-  __builtin_unique_stable_name(var);
-  // expected-error@+1{{use of undeclared identifier 'bad_var'}}
-  __builtin_unique_stable_name(bad_var);
-  // expected-error@+1{{use of undeclared identifier 'bad'}}
-  __builtin_unique_stable_name(bad::type);
-  // expected-error@+1{{no member named 'still_bad' in namespace 'NS'}}
-  __builtin_unique_stable_name(NS::still_bad);
-}
-
-template <typename T>
-void f2() {
-  // expected-error@+1{{no member named 'bad_val' in 'S'}}
-  __builtin_unique_stable_name(T::bad_val);
-  // expected-error@+1{{no type named 'bad_type' in 'S'}}
-  __builtin_unique_stable_name(typename T::bad_type);
-}
-
-struct S{};
-
-void use() {
-  // expected-note@+1{{in instantiation of}}
-  f2<S>();
-}
diff --git a/clang/test/ParserSYCL/unique_stable_name.cpp b/clang/test/ParserSYCL/unique_stable_name.cpp
new file mode 100644
index 0000000000000..aa50f54ab2137
--- /dev/null
+++ b/clang/test/ParserSYCL/unique_stable_name.cpp
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -Wno-unused %s
+
+namespace NS {
+using good = double;
+}
+
+void f(int var) {
+  // expected-error@+1{{expected '(' after '__builtin_sycl_unique_stable_name'}}
+  __builtin_sycl_unique_stable_name int; // Correct usage is __builtin_sycl_unique_stable_name(int);
+
+  // expected-error@+1{{expected '(' after '__builtin_sycl_unique_stable_name'}}
+  __builtin_sycl_unique_stable_name{int}; // Correct usage is __builtin_sycl_unique_stable_name(int);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+  __builtin_sycl_unique_stable_name(int; // Missing paren before semicolon
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+  __builtin_sycl_unique_stable_name(int, float); // Missing paren before comma
+
+  // expected-error@+1{{unknown type name 'var'}}
+  __builtin_sycl_unique_stable_name(var);
+  __builtin_sycl_unique_stable_name(NS::good);
+
+  // expected-error@+1{{expected a type}}
+  __builtin_sycl_unique_stable_name(for (int i = 0; i < 10; ++i) {})
+  __builtin_sycl_unique_stable_name({
+    (for (int i = 0; i < 10; ++i){})})
+}
+
+template <typename T>
+void f2() {
+  __builtin_sycl_unique_stable_name(typename T::good_type);
+}
+
+struct S {
+  class good_type {};
+};
+
+void use() {
+  f2<S>();
+}
diff --git a/clang/test/ParserSYCL/unique_stable_name_sycl_only.cpp b/clang/test/ParserSYCL/unique_stable_name_sycl_only.cpp
new file mode 100644
index 0000000000000..5f2a2e3313df5
--- /dev/null
+++ b/clang/test/ParserSYCL/unique_stable_name_sycl_only.cpp
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fsyntax-only -verify=notsycl -Wno-unused %s
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-host -verify=sycl -Wno-unused %s
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -verify=sycl -Wno-unused %s
+
+// sycl-no-diagnostics
+void foo() {
+  // notsycl-error@+1{{expected '(' for function-style cast or type construction}}
+  __builtin_sycl_unique_stable_name(int);
+}
diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
index cf6b28cb4f9ac..46e780f2379c6 100644
--- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
+++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
@@ -402,7 +402,7 @@ int main() {
 // SPIR: Name:{{.*}}Region
 // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // SPIR-NEXT: Line: 53, Column: 9 }
-// SPIR-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// SPIR-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // SPIR-NEXT: Args:
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '0'
@@ -423,7 +423,7 @@ int main() {
 // SPIR: Name:{{.*}}Region
 // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // SPIR-NEXT: Line: 53, Column: 9 }
-// SPIR-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// SPIR-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // SPIR-NEXT: Args:
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '1'
@@ -444,7 +444,7 @@ int main() {
 // SPIR: Name:{{.*}}Region
 // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // SPIR-NEXT: Line: 53, Column: 9 }
-// SPIR-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// SPIR-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // SPIR-NEXT: Args:
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '2'
@@ -465,7 +465,7 @@ int main() {
 // SPIR: Name:{{.*}}Region
 // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // SPIR-NEXT: Line: 53, Column: 9 }
-// SPIR-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// SPIR-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // SPIR-NEXT: Args:
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '3'
@@ -486,7 +486,7 @@ int main() {
 // SPIR: Name:{{.*}}Region
 // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // SPIR-NEXT: Line: 53, Column: 9 }
-// SPIR-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// SPIR-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // SPIR-NEXT: Args:
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '4'
@@ -507,7 +507,7 @@ int main() {
 // NVPTX: Name:{{.*}}Region
 // NVPTX: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp',
 // NVPTX: Line: 53, Column: 9 }
-// NVPTX-NEXT: Function:        '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ'
+// NVPTX-NEXT: Function:        _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_E3XYZ
 // NVPTX-NEXT: Args:
 // NVPTX-NEXT: String:          'Arg '
 // NVPTX: Argument:        '5'
diff --git a/clang/test/SemaSYCL/kernel-not-functor.cpp b/clang/test/SemaSYCL/kernel-not-functor.cpp
index 5a1dbf00448a9..67f50d2679a89 100644
--- a/clang/test/SemaSYCL/kernel-not-functor.cpp
+++ b/clang/test/SemaSYCL/kernel-not-functor.cpp
@@ -1,5 +1,6 @@
 // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
+// Disabled while we are no longer checking in host mode.
+// RUNX: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
 
 template <typename Name, typename F>
 __attribute__((sycl_kernel)) void kernel(F kernelFunc) {
diff --git a/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp b/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp
index 3d682754532fd..eca7e9937d47c 100644
--- a/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp
+++ b/clang/test/SemaSYCL/mangle-unnamed-kernel.cpp
@@ -8,5 +8,5 @@ int main() {
   return 0;
 }
 
-// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE6_12clES2_EUlvE6_54{{.*}}
-// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE7_12clES2_EUlvE7_54{{.*}}
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE10000_clES2_EUlvE10000_
+// CHECK: _ZTSZZ4mainENKUlRN2cl4sycl7handlerEE10001_clES2_EUlvE10000_
diff --git a/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp b/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp
new file mode 100644
index 0000000000000..c169010be537a
--- /dev/null
+++ b/clang/test/SemaSYCL/unique-stable-name-multiple-target-crash.cpp
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 %s %s -std=c++17 -triple x86_64-linux-gnu -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused
+
+// This would crash due to the double-inputs, since the 'magic static' use in
+// the AST Context SCYL Filtering would end up caching an old version of the
+// ASTContext object, which no longer exists in the second file's invocation.
+//
+// expected-no-diagnostics
+class Empty {};
+template <typename, typename F> __attribute__((sycl_kernel)) void kernel(F) {
+    __builtin_sycl_unique_stable_name(F);
+}
+
+void use() {
+  [](Empty) {
+    auto lambda = []{};
+    kernel<class i>(lambda);
+  };
+}
diff --git a/clang/test/SemaSYCL/unique_stable_name.cpp b/clang/test/SemaSYCL/unique_stable_name.cpp
new file mode 100644
index 0000000000000..7497aa1194285
--- /dev/null
+++ b/clang/test/SemaSYCL/unique_stable_name.cpp
@@ -0,0 +1,215 @@
+// RUN: %clang_cc1 %s -std=c++17 -triple x86_64-pc-windows-msvc -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused
+// RUN: %clang_cc1 %s -std=c++17 -triple x86_64-linux-gnu -Wno-sycl-2020-compat -fsycl-is-device -verify -fsyntax-only -Wno-unused
+
+template <typename KernelName, typename KernelType>
+[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { // #kernelSingleTask
+  kernelFunc();
+}
+
+// kernel1 - expect error
+// The current function is named with a lambda (i.e., takes a lambda as a
+// template parameter. Call the builtin on the current function then it is
+// passed to a kernel. Test that passing the given function to the unique
+// stable name builtin and then to the kernel throws an error because the
+// latter causes its name mangling to change.
+template <typename Func>
+void kernel1func(const Func &F1) {
+  constexpr const char *F1_output = __builtin_sycl_unique_stable_name(Func); // #USN_F1
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#kernel1func_call{{in instantiation of function template specialization}}
+  // expected-note@#USN_F1{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel1>(F1); // #kernel1_call
+}
+
+void callkernel1() {
+  kernel1func([]() {}); // #kernel1func_call
+}
+
+// kernel2 - expect error
+// The current function is named with a lambda (i.e., takes a lambda as a
+// template parameter). Call the builtin on the given function,
+// then an empty lambda is passed to kernel.
+// Test that passing the given function to the unique stable name builtin and
+// then passing a different lambda to the kernel still throws an error because
+// the calling context is part of naming the kernel. Even though the given
+// function (F2) is not passed to the kernel, its mangling changes due to
+// kernel call with the unrelated lambda.
+template <typename Func>
+void kernel2func(const Func &F2) {
+  constexpr const char *F2_output = __builtin_sycl_unique_stable_name(Func); // #USN_F2
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#kernel2func_call{{in instantiation of function template specialization}}
+  // expected-note@#USN_F2{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel2>([]() {});
+}
+
+void callkernel2() {
+  kernel2func([]() {}); // #kernel2func_call
+}
+
+template <template <typename> typename Outer, typename Inner>
+struct S {
+  void operator()() const {}
+};
+
+template <typename Ty>
+struct Tangerine {};
+
+template <typename Func>
+void kernel3_4func(const Func &F) {
+  // Test that passing the same lambda to two kernels does not cause an error
+  // because the kernel uses do not interfere with each other or invalidate
+  // the stable name in any way.
+  kernel_single_task<class kernel3>(F);
+  kernel_single_task<class kernel4>(F);
+  // Using the same functor twice should be fine
+}
+
+// kernel3 and kernel4 - expect no errors
+void callkernel3_4() {
+  kernel3_4func([]() {});
+}
+
+template <typename T>
+static constexpr const char *output1 = __builtin_sycl_unique_stable_name(T);
+
+#define MACRO()                      \
+  auto l14 = []() { return 1; };     \
+  constexpr const char *l14_output = \
+      __builtin_sycl_unique_stable_name(decltype(l14));
+
+int main() {
+
+  // kernel5 - expect no error
+  // Test that passing the lambda to the unique stable name builtin and then
+  // using the lambda in a way that does not  contribute to the kernel name
+  // does not cause an error because the  stable name is not invalidated in
+  // this situation.
+  auto l5 = []() {};
+  constexpr const char *l5_output =
+      __builtin_sycl_unique_stable_name(decltype(l5));
+  kernel_single_task<class kernel5>(
+      [=]() { l5(); }); // Used in the kernel, but not the kernel name itself
+
+  // kernel6 - expect error
+  // Test that passing the lambda to the unique stable name builtin and then
+  // using the same lambda in the naming of a kernel causes a diagnostic on the
+  // kernel use due to the change in results to the stable name.
+  auto l6 = []() { return 1; };
+  constexpr const char *l6_output =
+      __builtin_sycl_unique_stable_name(decltype(l6)); // #USN_l6
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#USN_l6{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel6>(l6); // Used in the kernel name after builtin
+
+  // kernel7 - expect error
+  // Same as kernel11 (below) except make the lambda part of naming the kernel.
+  // Test that passing a lambda to the unique stable name builtin and then
+  // passing a second lambda to the kernel throws an error because the first
+  // lambda is included in the signature of the second lambda, hence it changes
+  // the mangling of the kernel.
+  auto l7 = []() { return 1; };
+  auto l8 = [](decltype(l7) *derp = nullptr) { return 2; };
+  constexpr const char *l7_output =
+      __builtin_sycl_unique_stable_name(decltype(l7)); // #USN_l7
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#USN_l7{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel7>(l8);
+
+  // kernel8 and kernel9 - expect error
+  // Tests that passing a lambda to the unique stable name builtin and passing it
+  // to a kernel called with an if constexpr branch causes a diagnostic on the
+  // kernel9 use due to the change in the results to the stable name. This happens
+  // even though the use of kernel9 happens in the false branch of a constexpr if
+  // because both the true and the false branches cause the instantiation of
+  // kernel_single_task.
+  auto l9 = []() { return 1; };
+  auto l10 = []() { return 2; };
+  constexpr const char *l10_output =
+      __builtin_sycl_unique_stable_name(decltype(l10)); // #USN_l10
+  if constexpr (1) {
+    kernel_single_task<class kernel8>(l9);
+  } else {
+    // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+    // expected-note@#USN_l10{{'__builtin_sycl_unique_stable_name' evaluated here}}
+    // expected-note@+1{{in instantiation of function template specialization}}
+    kernel_single_task<class kernel9>(l10);
+  }
+
+  // kernel11 - expect no error
+  // Test that passing a lambda to the unique stable name builtin and then
+  // passing a second lambda capturing the first one to the kernel does not
+  // throw an error because the first lambda is not involved in naming the
+  // kernel i.e., the mangling does not change.
+  auto l11 = []() { return 1; };
+  auto l12 = [l11]() { return 2; };
+  constexpr const char *l11_output =
+      __builtin_sycl_unique_stable_name(decltype(l11));
+  kernel_single_task<class kernel11>(l12);
+
+  // kernel12 - expect an error
+  // Test that passing a lambda to the unique stable name builtin and then
+  // passing it to the kernel as a template template parameter causes a
+  // diagnostic on the kernel use due to template template parameter being
+  // involved in the mangling of the kernel name.
+  auto l13 = []() { return 1; };
+  constexpr const char *l13_output =
+      __builtin_sycl_unique_stable_name(decltype(l13)); // #USN_l13
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#USN_l13{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel12>(S<Tangerine, decltype(l13)>{});
+
+  // kernel13 - expect an error
+  // Test that passing a lambda to the unique stable name builtin within a macro
+  // and then calling the macro within the kernel causes an error on the kernel
+  // and diagnoses in all the expected places despite the use of a macro.
+  // expected-error@#kernelSingleTask{{kernel instantiation changes the result of an evaluated '__builtin_sycl_unique_stable_name'}}
+  // expected-note@#USN_MACRO{{'__builtin_sycl_unique_stable_name' evaluated here}}
+  // expected-note@+1{{in instantiation of function template specialization}}
+  kernel_single_task<class kernel13>(
+      []() {
+        MACRO(); // #USN_MACRO
+      });
+}
+
+namespace NS {}
+
+void f() {
+  // expected-error@+1{{unknown type name 'bad_var'}}
+  __builtin_sycl_unique_stable_name(bad_var);
+  // expected-error@+1{{use of undeclared identifier 'bad'}}
+  __builtin_sycl_unique_stable_name(bad::type);
+  // expected-error@+1{{no type named 'still_bad' in namespace 'NS'}}
+  __builtin_sycl_unique_stable_name(NS::still_bad);
+
+  // FIXME: warning about side-effects in an unevaluated context expected, but
+  // none currently emitted.
+  int i = 0;
+  __builtin_sycl_unique_stable_name(decltype(i++));
+
+  // Tests that use within a VLA does not diagnose as a side-effecting use in
+  // an unevaluated context because the use within a VLA extent forces
+  // evaluation.
+  int j = 55;
+  __builtin_sycl_unique_stable_name(int[++j]); // no warning expected
+}
+
+template <typename T>
+void f2() {
+  // expected-error@+1{{no type named 'bad_val' in 'St'}}
+  __builtin_sycl_unique_stable_name(typename T::bad_val);
+  // expected-error@+1{{no type named 'bad_type' in 'St'}}
+  __builtin_sycl_unique_stable_name(typename T::bad_type);
+}
+
+struct St {};
+
+void use() {
+  // expected-note@+1{{in instantiation of}}
+  f2<St>();
+}
diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index b4822d3b1e934..7c6b7bcd05937 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -336,6 +336,7 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
   case Stmt::ObjCBoxedExprClass:
   case Stmt::ObjCSubscriptRefExprClass:
   case Stmt::RecoveryExprClass:
+  case Stmt::SYCLUniqueStableNameExprClass:
     K = CXCursor_UnexposedExpr;
     break;
 
diff --git a/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp b/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp
index 7fd47d36e7db2..e89d56ba9ce9c 100644
--- a/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp
@@ -50,7 +50,7 @@ template <typename T, typename ID = T> class spec_constant {
   typename sycl::detail::enable_if_t<std::is_arithmetic<V>::value, V>
   get() const { // explicit access.
 #ifdef __SYCL_DEVICE_ONLY__
-    const char *TName = __builtin_unique_stable_name(ID);
+    const char *TName = __builtin_sycl_unique_stable_name(ID);
     return __sycl_getScalarSpecConstantValue<T>(TName);
 #else
     return Val;
@@ -63,7 +63,7 @@ template <typename T, typename ID = T> class spec_constant {
                                      V>
   get() const { // explicit access.
 #ifdef __SYCL_DEVICE_ONLY__
-    const char *TName = __builtin_unique_stable_name(ID);
+    const char *TName = __builtin_sycl_unique_stable_name(ID);
     return __sycl_getCompositeSpecConstantValue<T>(TName);
 #else
     return Val;
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 5faa502102ba3..0aa80c63bbef8 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -184,34 +184,24 @@ constexpr bool isKernelLambdaCallableWithKernelHandler() {
 
 // Helpers for running kernel lambda on the host device
 
-template <typename KernelType,
-          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
-              KernelType>()> * = nullptr>
-constexpr void runKernelWithoutArg(KernelType KernelName) {
-  kernel_handler KH;
-  KernelName(KH);
-}
-
-template <typename KernelType,
-          typename std::enable_if_t<!isKernelLambdaCallableWithKernelHandler<
-              KernelType>()> * = nullptr>
-constexpr void runKernelWithoutArg(KernelType KernelName) {
-  KernelName();
-}
-
-template <typename ArgType, typename KernelType,
-          typename std::enable_if_t<isKernelLambdaCallableWithKernelHandler<
-              KernelType, ArgType>()> * = nullptr>
-constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
-  kernel_handler KH;
-  KernelName(Arg, KH);
+template <typename KernelType> void runKernelWithoutArg(KernelType KernelName) {
+  if constexpr (isKernelLambdaCallableWithKernelHandler<KernelType>()) {
+    kernel_handler KH;
+    KernelName(KH);
+  } else {
+    KernelName();
+  }
 }
 
-template <typename ArgType, typename KernelType,
-          typename std::enable_if_t<!isKernelLambdaCallableWithKernelHandler<
-              KernelType, ArgType>()> * = nullptr>
+template <typename ArgType, typename KernelType>
 constexpr void runKernelWithArg(KernelType KernelName, ArgType Arg) {
-  KernelName(Arg);
+  if constexpr (isKernelLambdaCallableWithKernelHandler<KernelType,
+                                                        ArgType>()) {
+    kernel_handler KH;
+    KernelName(Arg, KH);
+  } else {
+    KernelName(Arg);
+  }
 }
 
 // The pure virtual class aimed to store lambda/functors of any type.
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index d3a4de78657b3..046c494064a36 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -61,7 +61,7 @@ template <class Name> struct SpecConstantInfo {
 // Translates SYCL 2020 specialization constant type to its name.
 template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
 #ifdef SYCL_LANGUAGE_VERSION
-  return __builtin_unique_stable_name(
+  return __builtin_sycl_unique_stable_name(
       specialization_id_name_generator<SpecName>);
 #else
   return "";
@@ -105,7 +105,7 @@ using make_index_sequence =
 
 template <typename T> struct KernelInfoImpl {
 private:
-  static constexpr auto n = __builtin_unique_stable_name(T);
+  static constexpr auto n = __builtin_sycl_unique_stable_name(T);
   template <unsigned long long... I>
   static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
     return {};
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 059a8fb4344e0..7ec129a2c4463 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -825,9 +825,8 @@ class __SYCL_EXPORT handler {
 
       range<Dims> AdjustedRange = NumWorkItems;
       AdjustedRange.set_range_dim0(NewValX);
-#ifdef __SYCL_DEVICE_ONLY__
       kernel_parallel_for_wrapper<NameWT, TransformedArgType>(Wrapper);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
       detail::checkValueRange<Dims>(AdjustedRange);
       MNDRDesc.set(std::move(AdjustedRange));
       StoreLambda<NameWT, decltype(Wrapper), Dims, TransformedArgType>(
@@ -839,10 +838,9 @@ class __SYCL_EXPORT handler {
        // !DPCPP_HOST_DEVICE_OPENMP && !DPCPP_HOST_DEVICE_PERF_NATIVE &&       \
        // SYCL_LANGUAGE_VERSION >= 202001
     {
-#ifdef __SYCL_DEVICE_ONLY__
       (void)NumWorkItems;
       kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
       detail::checkValueRange<Dims>(NumWorkItems);
       MNDRDesc.set(std::move(NumWorkItems));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
@@ -871,79 +869,94 @@ class __SYCL_EXPORT handler {
     MKernelName = getKernelName();
   }
 
-#ifdef __SYCL_DEVICE_ONLY__
-
+#ifdef SYCL_LANGUAGE_VERSION
+#define __SYCL_KERNEL_ATTR__ __attribute__((sycl_kernel))
+#else
+#define __SYCL_KERNEL_ATTR__
+#endif
   // NOTE: the name of this function - "kernel_single_task" - is used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task(KernelType KernelFunc) {
 #else
   kernel_single_task(const KernelType &KernelFunc) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc();
+#endif
   }
 
   // NOTE: the name of this function - "kernel_single_task" - is used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task(KernelType KernelFunc, kernel_handler KH) {
 #else
   kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc(KH);
+#endif
   }
 
   // NOTE: the name of these functions - "kernel_parallel_for" - are used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for(KernelType KernelFunc) {
 #else
   kernel_parallel_for(const KernelType &KernelFunc) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
+#endif
   }
 
   // NOTE: the name of these functions - "kernel_parallel_for" - are used by the
   // Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for(KernelType KernelFunc, kernel_handler KH) {
 #else
   kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
+#endif
   }
 
   // NOTE: the name of this function - "kernel_parallel_for_work_group" - is
   // used by the Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group(KernelType KernelFunc) {
 #else
   kernel_parallel_for_work_group(const KernelType &KernelFunc) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
+#endif
   }
 
   // NOTE: the name of this function - "kernel_parallel_for_work_group" - is
   // used by the Front End to determine kernel invocation kind.
   template <typename KernelName, typename ElementType, typename KernelType>
-  __attribute__((sycl_kernel)) void
+  __SYCL_KERNEL_ATTR__ void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group(KernelType KernelFunc, kernel_handler KH) {
 #else
   kernel_parallel_for_work_group(const KernelType &KernelFunc,
                                  kernel_handler KH) {
 #endif
+#ifdef __SYCL_DEVICE_ONLY__
     KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
+#endif
   }
 
   // Wrappers for kernel_*** functions above with and without support of
@@ -955,82 +968,56 @@ class __SYCL_EXPORT handler {
   // Wrappers for kernel_single_task(...)
 
   template <typename KernelName, typename KernelType>
-  std::enable_if_t<
-      detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_single_task_wrapper(KernelType KernelFunc) {
-#else
-  kernel_single_task_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_single_task<KernelName>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename KernelType>
-  std::enable_if_t<
-      !detail::isKernelLambdaCallableWithKernelHandler<KernelType>(), void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_single_task_wrapper(KernelType KernelFunc) {
 #else
   kernel_single_task_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_single_task<KernelName>(KernelFunc);
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType>()) {
+      kernel_handler KH;
+      kernel_single_task<KernelName>(KernelFunc, KH);
+    } else {
+      kernel_single_task<KernelName>(KernelFunc);
+    }
   }
 
   // Wrappers for kernel_parallel_for(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_parallel_for_wrapper(KernelType KernelFunc) {
-#else
-  kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_wrapper(KernelType KernelFunc) {
 #else
   kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_parallel_for<KernelName, ElementType>(KernelFunc);
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType, ElementType>()) {
+      kernel_handler KH;
+      kernel_parallel_for<KernelName, ElementType>(KernelFunc, KH);
+    } else {
+      kernel_parallel_for<KernelName, ElementType>(KernelFunc);
+    }
   }
 
   // Wrappers for kernel_parallel_for_work_group(...)
 
   template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
-#ifdef __SYCL_NONCONST_FUNCTOR__
-  kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
-#else
-  kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
-#endif
-    kernel_handler KH;
-    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
-  }
-
-  template <typename KernelName, typename ElementType, typename KernelType>
-  std::enable_if_t<!detail::isKernelLambdaCallableWithKernelHandler<
-                       KernelType, ElementType>(),
-                   void>
+  void
 #ifdef __SYCL_NONCONST_FUNCTOR__
   kernel_parallel_for_work_group_wrapper(KernelType KernelFunc) {
 #else
   kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
 #endif
-    kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType, ElementType>()) {
+      kernel_handler KH;
+      kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc, KH);
+    } else {
+      kernel_parallel_for_work_group<KernelName, ElementType>(KernelFunc);
+    }
   }
-#endif
 
   std::shared_ptr<detail::kernel_bundle_impl>
   getOrInsertHandlerKernelBundle(bool Insert) const;
@@ -1192,9 +1179,8 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     kernel_single_task_wrapper<NameT>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant.
     MNDRDesc.set(range<1>{1});
@@ -1299,11 +1285,10 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkItems;
     (void)WorkItemOffset;
     kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
@@ -1332,10 +1317,9 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)ExecutionRange;
     kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(ExecutionRange);
     MNDRDesc.set(std::move(ExecutionRange));
     StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
@@ -1565,10 +1549,9 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkGroups;
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
     StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
@@ -1598,11 +1581,10 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)NumWorkGroups;
     (void)WorkGroupSize;
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     nd_range<Dims> ExecRange =
         nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
     detail::checkValueRange<Dims>(ExecRange);
@@ -1693,10 +1675,9 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     kernel_single_task<NameT>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant
     MNDRDesc.set(range<1>{1});
@@ -1734,11 +1715,10 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkItems;
     kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NumWorkItems);
     MNDRDesc.set(std::move(NumWorkItems));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
@@ -1769,12 +1749,11 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkItems;
     (void)WorkItemOffset;
     kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
@@ -1806,11 +1785,10 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NDRange;
     kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NDRange);
     MNDRDesc.set(std::move(NDRange));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
@@ -1846,11 +1824,10 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkGroups;
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(NumWorkGroups);
     MNDRDesc.setNumWorkGroups(NumWorkGroups);
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
@@ -1884,12 +1861,11 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, group<Dims>>;
-#ifdef __SYCL_DEVICE_ONLY__
     (void)Kernel;
     (void)NumWorkGroups;
     (void)WorkGroupSize;
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType>(KernelFunc);
-#else
+#ifndef __SYCL_DEVICE_ONLY__
     nd_range<Dims> ExecRange =
         nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
     detail::checkValueRange<Dims>(ExecRange);
@@ -2300,32 +2276,25 @@ class __SYCL_EXPORT handler {
 
   friend class ::MockHandler;
 
-  template <
-      typename TransformedArgType, int Dims, typename KernelType,
-      typename std::enable_if_t<detail::isKernelLambdaCallableWithKernelHandler<
-          KernelType, TransformedArgType>()> * = nullptr>
-  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
-                                   range<Dims> NumWorkItems) {
-    return [=](TransformedArgType Arg, kernel_handler KH) {
-      if (Arg[0] >= NumWorkItems[0])
-        return;
-      Arg.set_allowed_range(NumWorkItems);
-      KernelFunc(Arg, KH);
-    };
-  }
-
-  template <typename TransformedArgType, int Dims, typename KernelType,
-            typename std::enable_if_t<
-                !detail::isKernelLambdaCallableWithKernelHandler<
-                    KernelType, TransformedArgType>()> * = nullptr>
+  template <typename TransformedArgType, int Dims, typename KernelType>
   auto getRangeRoundedKernelLambda(KernelType KernelFunc,
                                    range<Dims> NumWorkItems) {
-    return [=](TransformedArgType Arg) {
-      if (Arg[0] >= NumWorkItems[0])
-        return;
-      Arg.set_allowed_range(NumWorkItems);
-      KernelFunc(Arg);
-    };
+    if constexpr (detail::isKernelLambdaCallableWithKernelHandler<
+                      KernelType, TransformedArgType>()) {
+      return [=](TransformedArgType Arg, kernel_handler KH) {
+        if (Arg[0] >= NumWorkItems[0])
+          return;
+        Arg.set_allowed_range(NumWorkItems);
+        KernelFunc(Arg, KH);
+      };
+    } else {
+      return [=](TransformedArgType Arg) {
+        if (Arg[0] >= NumWorkItems[0])
+          return;
+        Arg.set_allowed_range(NumWorkItems);
+        KernelFunc(Arg);
+      };
+    }
   }
 };
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/kernel_handler.hpp b/sycl/include/CL/sycl/kernel_handler.hpp
index 9e618976ca2b4..9a1fbaf15efb6 100644
--- a/sycl/include/CL/sycl/kernel_handler.hpp
+++ b/sycl/include/CL/sycl/kernel_handler.hpp
@@ -49,7 +49,7 @@ class kernel_handler {
       typename T = typename std::remove_reference_t<decltype(S)>::value_type,
       std::enable_if_t<std::is_fundamental_v<T>> * = nullptr>
   T getSpecializationConstantOnDevice() {
-    const char *SymbolicID = __builtin_unique_stable_name(
+    const char *SymbolicID = __builtin_sycl_unique_stable_name(
         detail::specialization_id_name_generator<S>);
     return __sycl_getScalar2020SpecConstantValue<T>(
         SymbolicID, &S, MSpecializationConstantsBuffer);
@@ -59,7 +59,7 @@ class kernel_handler {
       typename T = typename std::remove_reference_t<decltype(S)>::value_type,
       std::enable_if_t<std::is_compound_v<T>> * = nullptr>
   T getSpecializationConstantOnDevice() {
-    const char *SymbolicID = __builtin_unique_stable_name(
+    const char *SymbolicID = __builtin_sycl_unique_stable_name(
         detail::specialization_id_name_generator<S>);
     return __sycl_getComposite2020SpecConstantValue<T>(
         SymbolicID, &S, MSpecializationConstantsBuffer);
diff --git a/sycl/test/check_device_code/buffer_location_codegen.cpp b/sycl/test/check_device_code/buffer_location_codegen.cpp
index 2edff0b8fcefd..cc558f0f0eb10 100644
--- a/sycl/test/check_device_code/buffer_location_codegen.cpp
+++ b/sycl/test/check_device_code/buffer_location_codegen.cpp
@@ -1,6 +1,6 @@
 // RUN: %clangxx -fsycl -c -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
 
-// CHECK: define {{.*}}spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE15kernel_function"{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_E15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
 // CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1}
 
 #include <CL/sycl.hpp>
diff --git a/sycl/test/warnings/warnings.cpp b/sycl/test/warnings/warnings.cpp
index 97ca4877f5a00..929e4132de68f 100644
--- a/sycl/test/warnings/warnings.cpp
+++ b/sycl/test/warnings/warnings.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl --no-system-header-prefix=CL/sycl -fsyntax-only -Wall -Wextra -Wno-ignored-attributes -Wno-deprecated-declarations -Wpessimizing-move -Wunused-variable -Wmismatched-tags -Wunneeded-internal-declaration -Werror -Wno-unknown-cuda-version %s -o %t.out
+// RUN: %clangxx -fsycl --no-system-header-prefix=CL/sycl -fsyntax-only -Wall -Wextra -Wno-ignored-attributes -Wno-deprecated-declarations -Wpessimizing-move -Wunused-variable -Wmismatched-tags -Wunneeded-internal-declaration -Werror -Wno-unknown-cuda-version -Wno-unused-parameter %s -o %t.out
 
 #include <CL/sycl.hpp>