diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index ac14b6dd4248a..7709e96608e34 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1689,6 +1689,16 @@ def SYCLAddIRAttributesGlobalVariable : InheritableAttr {
   let Documentation = [SYCLAddIRAttributesGlobalVariableDocs];
 }
 
+def SYCLAddIRAnnotationsMember : InheritableAttr {
+  let Spellings = [CXX11<"__sycl_detail__", "add_ir_annotations_member">];
+  let Args = [VariadicExprArgument<"Args">];
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
+  let Subjects = SubjectList<[Field], ErrorDiag>;
+  let AcceptsExprPack = 1;
+  let AdditionalMembers = SYCLAddIRAttrCommonMembers.MemberCode;
+  let Documentation = [SYCLAddIRAnnotationsMemberDocs];
+}
+
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index e43edf8b9a893..4ea8136799ad0 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3611,6 +3611,83 @@ where the last N*2 are as described above.
   }];
 }
 
+def SYCLAddIRAnnotationsMemberDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "add_ir_annotations_member";
+  let Content = [{
+This attribute can be applied to a non-static field. Access to a field with this
+attribute will generate a call to ``llvm.ptr.annotation`` with the arguments
+paired.
+The attribute must contain N*2 arguments. Each of the first N of these arguments
+must be either string literals or ``constexpr const char *``. The following N
+must be integer, floating point, character, boolean, ``const char *``, or an
+enumeration as either a literal or ``constexpr`` value. The first N arguments
+and the second N arguments are zipped into pairs when creating the call to
+``llvm.ptr.annotation``, i.e. in the constant global variable pointed to by the
+generated annotation intrinsic call, the N+1'th argument of the attribute will
+occur after the first argument of the attribute, the N+2'th argument of the
+attribute will occur after the second argument of the attribute, etc.
+The generated call to ``llvm.ptr.annotation`` will have the following arguments
+in this order:
+  * First argument is a pointer to the field.
+  * A pointer to a string literal in a constant global variable. This will
+    always be "sycl-properties".
+  * A pointer to a string literal in a constant global variable with the name of
+    the source file.
+  * The line number of the field declaration in the source file.
+  * A pointer to a constant global variable containing pointers to string
+    literals in constant global variables. These pointers to string literals
+    occur in pairs. If the second value of a pair was a ``nullptr`` or an empty
+    string then the pointer will be a null-pointer.
+A pair will not be in the call to ``llvm.ptr.annotation`` if the first value of
+the pair is an empty string.
+
+.. code-block:: c++
+
+  struct Foo {
+    int *ptr
+  #ifdef __SYCL_DEVICE_ONLY__
+        [[__sycl_detail__::add_ir_annotations_member(
+            "Attr1", "Attr2", "Attribute value", 3.14)]]
+  #endif
+        ;
+  };
+  // Accessing the 'ptr' field of 'Foo' will result in a call to
+  // 'llvm.ptr.annotation' with the second argument pointing to a constant
+  // global variable containing "sycl-properties" and the fifth argument
+  // pointing to a constant global variable with pointers to string literals
+  // "Attr1", "Attribute value", "Attr2", "3.14" in that order.
+
+Optionally, the first argument of the attribute can be an initializer list
+containing only string literals. This initializer list acts as a filter,
+allowing only pairs with the first value in the initializer list to be
+generated. If this intializer list is present, the attribute must have N*2+1
+arguments, where the last N*2 are as described above.
+
+.. code-block:: c++
+
+  struct Foo {
+    int *ptr
+  #ifdef __SYCL_DEVICE_ONLY__
+        [[__sycl_detail__::add_ir_annotations_member(
+            {"Attr2"}, "Attr1", "Attr2", "Attribute value", 3.14)]]
+  #endif
+        ;
+  };
+  // Accessing the 'ptr' field of 'Foo' will result in a call to
+  // 'llvm.ptr.annotation' with the second argument pointing to a constant
+  // global variable containing "sycl-properties" and the fifth argument
+  // pointing to a constant global variable with pointers to string literals
+  // "Attr2", "3.14" in that order.
+
+.. Note:: This attribute will only generate a call to ``llvm.ptr.annotation`` in
+  SYCL device code.
+
+.. Note:: This attribute is intended as an internal implementation detail and is
+  not intended to be used by external users.
+  }];
+}
+
 def SYCLDeviceIndirectlyCallableDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "intel::device_indirectly_callable";
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 3c2e0a863d5ad..7f448e05bfe85 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10637,6 +10637,11 @@ class Sema final {
   void AddSYCLAddIRAttributesGlobalVariableAttr(Decl *D,
                                                 const AttributeCommonInfo &CI,
                                                 MutableArrayRef<Expr *> Args);
+  SYCLAddIRAnnotationsMemberAttr *
+  MergeSYCLAddIRAnnotationsMemberAttr(Decl *D,
+                                      const SYCLAddIRAnnotationsMemberAttr &A);
+  void AddSYCLAddIRAnnotationsMemberAttr(Decl *D, const AttributeCommonInfo &CI,
+                                         MutableArrayRef<Expr *> Args);
   void AddReqdWorkGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
                                 Expr *XDim, Expr *YDim, Expr *ZDim);
   ReqdWorkGroupSizeAttr *
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 045a5a12138da..a0833d15f5787 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -4518,6 +4518,9 @@ LValue CodeGenFunction::EmitLValueForField(LValue base,
 
   // Emit attribute annotation for a field.
   if (getLangOpts().SYCLIsDevice) {
+    if (field->hasAttr<SYCLAddIRAnnotationsMemberAttr>())
+      addr = EmitFieldSYCLAnnotations(field, addr);
+
     SmallString<256> AnnotStr;
     CGM.generateIntelFPGAAnnotation(field, AnnotStr);
     if (!AnnotStr.empty())
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 900c485173bce..e2d837237d77d 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -2721,6 +2721,40 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D,
   return Address(V, Addr.getElementType(), Addr.getAlignment());
 }
 
+llvm::Value *CodeGenFunction::EmitSYCLAnnotationCall(
+    llvm::Function *AnnotationFn, llvm::Value *AnnotatedVal,
+    SourceLocation Location, const SYCLAddIRAnnotationsMemberAttr *Attr) {
+  SmallVector<llvm::Value *, 5> Args = {
+      AnnotatedVal,
+      Builder.CreateBitCast(CGM.EmitAnnotationString("sycl-properties"),
+                            Int8PtrTy),
+      Builder.CreateBitCast(CGM.EmitAnnotationUnit(Location), Int8PtrTy),
+      CGM.EmitAnnotationLineNo(Location), CGM.EmitSYCLAnnotationArgs(Attr)};
+  return Builder.CreateCall(AnnotationFn, Args);
+}
+
+Address CodeGenFunction::EmitFieldSYCLAnnotations(const FieldDecl *D,
+                                                  Address Addr) {
+  const auto *SYCLAnnotAttr = D->getAttr<SYCLAddIRAnnotationsMemberAttr>();
+  assert(SYCLAnnotAttr && "no add_ir_annotations_member attribute");
+  llvm::Value *V = Addr.getPointer();
+  llvm::Type *VTy = V->getType();
+  auto *PTy = dyn_cast<llvm::PointerType>(VTy);
+  unsigned AS = PTy ? PTy->getAddressSpace() : 0;
+  llvm::Type *IntrType = VTy;
+  if (!VTy->getPointerElementType()->isIntegerTy())
+    IntrType = llvm::PointerType::getWithSamePointeeType(CGM.Int8PtrTy, AS);
+  llvm::Function *F =
+      CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation, IntrType);
+
+  if (VTy != IntrType)
+    V = Builder.CreateBitCast(V, IntrType);
+  V = EmitSYCLAnnotationCall(F, V, D->getLocation(), SYCLAnnotAttr);
+  if (VTy != IntrType)
+    V = Builder.CreateBitCast(V, VTy);
+  return Address(V, Addr.getElementType(), Addr.getAlignment());
+}
+
 Address CodeGenFunction::EmitIntelFPGAFieldAnnotations(const FieldDecl *D,
                                                        Address Addr,
                                                        StringRef AnnotStr) {
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 005dcd3c1b483..d971f4d9fc3bc 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4504,6 +4504,16 @@ class CodeGenFunction : public CodeGenTypeCache {
   /// annotation result.
   Address EmitFieldAnnotations(const FieldDecl *D, Address V);
 
+  /// Emit a "sycl-properties" annotation call (intrinsic).
+  llvm::Value *
+  EmitSYCLAnnotationCall(llvm::Function *AnnotationFn,
+                         llvm::Value *AnnotatedVal, SourceLocation Location,
+                         const SYCLAddIRAnnotationsMemberAttr *Attr);
+
+  /// Emit sycl field annotations for given field & value. Returns the
+  /// annotation result.
+  Address EmitFieldSYCLAnnotations(const FieldDecl *D, Address V);
+
   /// Emit Intel FPGA field annotations for the given field and value. Returns
   /// the annotation result.
   Address EmitIntelFPGAFieldAnnotations(const FieldDecl *D, Address V,
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index de33c095f4b2e..51a309f403fd7 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2873,6 +2873,55 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
     Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
 }
 
+llvm::Constant *CodeGenModule::EmitSYCLAnnotationArgs(
+    const SYCLAddIRAnnotationsMemberAttr *Attr) {
+  llvm::SmallVector<std::pair<std::string, std::string>, 4>
+      AnnotationNameValPairs =
+          Attr->getFilteredAttributeNameValuePairs(getContext());
+  if (AnnotationNameValPairs.empty())
+    return llvm::ConstantPointerNull::get(GlobalsInt8PtrTy);
+
+  // For each name-value pair of the SYCL annotation attribute, create an
+  // annotation string for it. This will be the annotation arguments. If the
+  // value is the empty string, use a null-pointer instead.
+  llvm::SmallVector<llvm::Constant *, 4> LLVMArgs;
+  llvm::FoldingSetNodeID ID;
+  LLVMArgs.reserve(AnnotationNameValPairs.size() * 2);
+  for (const std::pair<std::string, std::string> &NVP :
+       AnnotationNameValPairs) {
+    llvm::Constant *NameStrC = EmitAnnotationString(NVP.first);
+    llvm::Constant *ValueStrC =
+        NVP.second == "" ? llvm::ConstantPointerNull::get(GlobalsInt8PtrTy)
+                         : EmitAnnotationString(NVP.second);
+    LLVMArgs.push_back(NameStrC);
+    LLVMArgs.push_back(ValueStrC);
+    ID.Add(NameStrC);
+    ID.Add(ValueStrC);
+  }
+
+  // If another SYCL annotation had the same arguments we can reuse the
+  // annotation value it created.
+  llvm::Constant *&LookupRef = SYCLAnnotationArgs[ID.ComputeHash()];
+  if (LookupRef)
+    return LookupRef;
+
+  // Create an anonymous struct global variable pointing to the annotation
+  // arguments in the order they were added above. This is the final constant
+  // used as the annotation value.
+  auto *Struct = llvm::ConstantStruct::getAnon(LLVMArgs);
+  auto *GV = new llvm::GlobalVariable(getModule(), Struct->getType(), true,
+                                      llvm::GlobalValue::PrivateLinkage, Struct,
+                                      ".args");
+  GV->setSection(AnnotationSection);
+  GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
+  auto *Bitcasted = llvm::ConstantExpr::getBitCast(GV, GlobalsInt8PtrTy);
+
+  // Set the look-up reference to the final annotation value for future
+  // annotations to reuse.
+  LookupRef = Bitcasted;
+  return Bitcasted;
+}
+
 void CodeGenModule::AddGlobalSYCLIRAttributes(llvm::GlobalVariable *GV,
                                               const RecordDecl *RD) {
   const auto *A = RD->getAttr<SYCLAddIRAttributesGlobalVariableAttr>();
diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h
index 57a651fb77618..22f2df4c7be06 100644
--- a/clang/lib/CodeGen/CodeGenModule.h
+++ b/clang/lib/CodeGen/CodeGenModule.h
@@ -410,6 +410,10 @@ class CodeGenModule : public CodeGenTypeCache {
   /// Used for uniquing of annotation arguments.
   llvm::DenseMap<unsigned, llvm::Constant *> AnnotationArgs;
 
+  /// Used for uniquing of SYCL annotation arguments. SYCL annotations are
+  /// handled differently than regular annotations so they cannot share map.
+  llvm::DenseMap<unsigned, llvm::Constant *> SYCLAnnotationArgs;
+
   llvm::StringMap<llvm::GlobalVariable *> CFConstantStringMap;
 
   llvm::DenseMap<llvm::Constant *, llvm::GlobalVariable *> ConstantStringMap;
@@ -1308,6 +1312,10 @@ class CodeGenModule : public CodeGenTypeCache {
   /// annotations are emitted during finalization of the LLVM code.
   void AddGlobalAnnotations(const ValueDecl *D, llvm::GlobalValue *GV);
 
+  /// Emit additional args of the annotation.
+  llvm::Constant *
+  EmitSYCLAnnotationArgs(const SYCLAddIRAnnotationsMemberAttr *Attr);
+
   /// Add attributes from add_ir_attributes_global_variable on TND to GV.
   void AddGlobalSYCLIRAttributes(llvm::GlobalVariable *GV,
                                  const RecordDecl *RD);
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index f765ad84ca7c0..5cf00574411fb 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2805,6 +2805,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
   else if (const auto *A =
                dyn_cast<SYCLAddIRAttributesGlobalVariableAttr>(Attr))
     NewAttr = S.MergeSYCLAddIRAttributesGlobalVariableAttr(D, *A);
+  else if (const auto *A = dyn_cast<SYCLAddIRAnnotationsMemberAttr>(Attr))
+    NewAttr = S.MergeSYCLAddIRAnnotationsMemberAttr(D, *A);
   else if (const auto *A = dyn_cast<ReqdWorkGroupSizeAttr>(Attr))
     NewAttr = S.MergeReqdWorkGroupSizeAttr(D, *A);
   else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 93706678dc8ca..77ec34f5ced28 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -7903,6 +7903,38 @@ static void handleSYCLAddIRAttributesGlobalVariableAttr(Sema &S, Decl *D,
   S.AddSYCLAddIRAttributesGlobalVariableAttr(D, A, Args);
 }
 
+SYCLAddIRAnnotationsMemberAttr *Sema::MergeSYCLAddIRAnnotationsMemberAttr(
+    Decl *D, const SYCLAddIRAnnotationsMemberAttr &A) {
+  if (const auto *ExistingAttr = D->getAttr<SYCLAddIRAnnotationsMemberAttr>()) {
+    checkSYCLAddIRAttributesMergeability(A, *ExistingAttr, *this);
+    return nullptr;
+  }
+  return A.clone(Context);
+}
+
+void Sema::AddSYCLAddIRAnnotationsMemberAttr(Decl *D,
+                                             const AttributeCommonInfo &CI,
+                                             MutableArrayRef<Expr *> Args) {
+  auto *Attr = SYCLAddIRAnnotationsMemberAttr::Create(Context, Args.data(),
+                                                      Args.size(), CI);
+  if (evaluateAddIRAttributesArgs(Attr->args_begin(), Attr->args_size(), *this,
+                                  CI))
+    return;
+  D->addAttr(Attr);
+}
+
+static void handleSYCLAddIRAnnotationsMemberAttr(Sema &S, Decl *D,
+                                                 const ParsedAttr &A) {
+  llvm::SmallVector<Expr *, 4> Args;
+  Args.reserve(A.getNumArgs());
+  for (unsigned I = 0; I < A.getNumArgs(); I++) {
+    assert(A.getArgAsExpr(I));
+    Args.push_back(A.getArgAsExpr(I));
+  }
+
+  S.AddSYCLAddIRAnnotationsMemberAttr(D, A, Args);
+}
+
 namespace {
 struct IntrinToName {
   uint32_t Id;
@@ -11381,6 +11413,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLAddIRAttributesGlobalVariable:
     handleSYCLAddIRAttributesGlobalVariableAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLAddIRAnnotationsMember:
+    handleSYCLAddIRAnnotationsMemberAttr(S, D, AL);
+    break;
 
   // Swift attributes.
   case ParsedAttr::AT_SwiftAsyncName:
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index edd5da5030608..8c1b5b03987e9 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -828,6 +828,18 @@ static void instantiateSYCLAddIRAttributesGlobalVariableAttr(
   S.AddSYCLAddIRAttributesGlobalVariableAttr(New, *A, Args);
 }
 
+static void instantiateSYCLAddIRAnnotationsMemberAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const SYCLAddIRAnnotationsMemberAttr *A, Decl *New) {
+  EnterExpressionEvaluationContext ConstantEvaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  SmallVector<Expr *, 4> Args;
+  if (S.SubstExprs(ArrayRef<Expr *>(A->args().begin(), A->args().end()),
+                   /*IsCall=*/false, TemplateArgs, Args))
+    return;
+  S.AddSYCLAddIRAnnotationsMemberAttr(New, *A, Args);
+}
+
 static void instantiateWorkGroupSizeHintAttr(
     Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
     const WorkGroupSizeHintAttr *A, Decl *New) {
@@ -1155,6 +1167,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
           *this, TemplateArgs, SYCLAddIRAttributesGlobalVariable, New);
       continue;
     }
+    if (const auto *SYCLAddIRAnnotationsMember =
+            dyn_cast<SYCLAddIRAnnotationsMemberAttr>(TmplAttr)) {
+      instantiateSYCLAddIRAnnotationsMemberAttr(
+          *this, TemplateArgs, SYCLAddIRAnnotationsMember, New);
+      continue;
+    }
     if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(TmplAttr)) {
       instantiateWorkGroupSizeHintAttr(*this, TemplateArgs, A, New);
       continue;
diff --git a/clang/test/AST/ast-attr-add-ir-annotations-pack.cpp b/clang/test/AST/ast-attr-add-ir-annotations-pack.cpp
new file mode 100644
index 0000000000000..05ba9fc34bb14
--- /dev/null
+++ b/clang/test/AST/ast-attr-add-ir-annotations-pack.cpp
@@ -0,0 +1,848 @@
+// RUN: not %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// Tests the AST produced from instantiating templates using the
+// __sycl_detail__::add_ir_annotations_member attribute with pack expansion
+// arguments.
+
+constexpr const char AttrName1[] = "Attr1";
+constexpr const char AttrName2[] = "Attr2";
+constexpr const char AttrName3[] = "Attr3";
+constexpr const char AttrVal1[] = "Val1";
+constexpr const char AttrVal2[] = "Val2";
+constexpr const char AttrVal3[] = "Val3";
+
+template <int... Is> struct ClassWithAnnotFieldTemplate1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", Is...)]];
+};
+template <int... Is> struct ClassWithAnnotFieldTemplate2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", "Attr2", "Attr3", Is...)]];
+};
+template <const char *...Names> struct ClassWithAnnotFieldTemplate3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Names..., 1, 2, 3)]];
+};
+template <const char *...Names> struct ClassWithAnnotFieldTemplate4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, Names..., 1, 2, 3)]];
+};
+template <const char *...Strs> struct ClassWithAnnotFieldTemplate5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Strs...)]];
+};
+template <const char *...Strs> struct ClassWithAnnotFieldTemplate6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, Strs...)]];
+};
+
+void InstantiateClassWithAnnotFieldTemplates() {
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate1
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate1 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate1
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr2"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'Is' 'int'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate1 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument integral 1
+  // CHECK-NEXT:       TemplateArgument integral 2
+  // CHECK-NEXT:       TemplateArgument integral 3
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate1
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr2"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate1 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate1 'void (const ClassWithAnnotFieldTemplate1<1, 2, 3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate1<1, 2, 3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate1 'void (ClassWithAnnotFieldTemplate1<1, 2, 3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate1<1, 2, 3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate1 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument integral 1
+  // CHECK-NEXT:       TemplateArgument integral 2
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate1
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate1 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate1 'void (const ClassWithAnnotFieldTemplate1<1, 2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate1<1, 2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate1 'void (ClassWithAnnotFieldTemplate1<1, 2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate1<1, 2> &&'
+  ClassWithAnnotFieldTemplate1<1, 2, 3> InstantiatedCWAFS1;
+  ClassWithAnnotFieldTemplate1<1, 2> InstantiatedCWAFS2;
+
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate2
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate2 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate2
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr2"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'Is' 'int'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate2 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument integral 1
+  // CHECK-NEXT:       TemplateArgument integral 2
+  // CHECK-NEXT:       TemplateArgument integral 3
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate2
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr2"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char[6]' lvalue
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'int'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 ... Is
+  // CHECK-NEXT:             IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate2 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate2 'void (const ClassWithAnnotFieldTemplate2<1, 2, 3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate2<1, 2, 3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate2 'void (ClassWithAnnotFieldTemplate2<1, 2, 3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate2<1, 2, 3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate2 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument integral 1
+  // CHECK-NEXT:       TemplateArgument integral 2
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate2
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate2 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate2 'void (const ClassWithAnnotFieldTemplate2<1, 2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate2<1, 2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate2 'void (ClassWithAnnotFieldTemplate2<1, 2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate2<1, 2> &&'
+  ClassWithAnnotFieldTemplate2<1, 2, 3> InstantiatedCWAFS3;
+  ClassWithAnnotFieldTemplate2<1, 2> InstantiatedCWAFS4;
+
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate3
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate3 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate3
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'const char *' NonTypeTemplateParm {{.*}} 'Names' 'const char *'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate3 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate3
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate3 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate3 'void (const ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate3 'void (ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate3 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate3
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate3 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate3 'void (const ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate3 'void (ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> &&'
+  ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> InstantiatedCWAFS5;
+  ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> InstantiatedCWAFS6;
+
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate4
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate4 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate4
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'const char *' NonTypeTemplateParm {{.*}} 'Names' 'const char *'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate4 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate4
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Names
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 1
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 1
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 2
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 2
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'int'
+  // CHECK-NEXT:           value: Int 3
+  // CHECK-NEXT:           IntegerLiteral {{.*}} 'int' 3
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate4 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate4 'void (const ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate4 'void (ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate4 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate4
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate4 'void () noexcept' inline default trivial
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate4 'void (const ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate4 'void (ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> &&'
+  ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> InstantiatedCWAFS7;
+  ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> InstantiatedCWAFS8;
+
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate5 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'const char *' NonTypeTemplateParm {{.*}} 'Strs' 'const char *'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate5 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate5 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (const ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate5 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate5 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate5 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal3' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal3' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate5 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate5 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate5
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate5 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate5 'void (ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &&'
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> InstantiatedCWAFS9;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> InstantiatedCWAFS10;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> InstantiatedCWAFS11;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> InstantiatedCWAFS12;
+
+  // CHECK:      ClassTemplateDecl {{.*}} ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:   NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:   CXXRecordDecl {{.*}} struct ClassWithAnnotFieldTemplate6 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:     FieldDecl {{.*}} ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         PackExpansionExpr {{.*}} '<dependent type>'
+  // CHECK-NEXT:           DeclRefExpr {{.*}} 'const char *' NonTypeTemplateParm {{.*}} 'Strs' 'const char *'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate6 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate6 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (const ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate6 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} col:26 referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate6 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate6 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal3' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:       SYCLAddIRAnnotationsMemberAttr
+  // CHECK-NEXT:         InitListExpr {{.*}} 'void'
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr1"
+  // CHECK-NEXT:           StringLiteral {{.*}} 'const char[6]' lvalue "Attr3"
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[6]' lvalue Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:         ConstantExpr {{.*}} 'const char *'
+  // CHECK-NEXT:           value: LValue
+  // CHECK-NEXT:           SubstNonTypeTemplateParmExpr {{.*}} 'const char *'
+  // CHECK-NEXT:             NonTypeTemplateParmDecl {{.*}} referenced 'const char *' depth 0 index 0 ... Strs
+  // CHECK-NEXT:             ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+  // CHECK-NEXT:               DeclRefExpr {{.*}} 'const char[5]' lvalue Var {{.*}} 'AttrVal3' 'const char[5]'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate6 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> &&'
+  // CHECK-NEXT:   ClassTemplateSpecializationDecl {{.*}} struct ClassWithAnnotFieldTemplate6 definition
+  // CHECK-NEXT:     DefinitionData
+  // CHECK-NEXT:       DefaultConstructor
+  // CHECK-NEXT:       CopyConstructor
+  // CHECK-NEXT:       MoveConstructor
+  // CHECK-NEXT:       CopyAssignment
+  // CHECK-NEXT:       MoveAssignment
+  // CHECK-NEXT:       Destructor
+  // CHECK-NEXT:     TemplateArgument pack
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName1' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName2' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrName3' 'const char[6]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal1' 'const char[5]'
+  // CHECK-NEXT:       TemplateArgument decl
+  // CHECK-NEXT:         Var {{.*}} 'AttrVal2' 'const char[5]'
+  // CHECK-NEXT:     CXXRecordDecl {{.*}} implicit struct ClassWithAnnotFieldTemplate6
+  // CHECK-NEXT:     FieldDecl {{.*}} referenced ptr 'int *'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit used ClassWithAnnotFieldTemplate6 'void () noexcept'
+  // CHECK-NEXT:       CompoundStmt
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'const ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &'
+  // CHECK-NEXT:     CXXConstructorDecl {{.*}} implicit constexpr ClassWithAnnotFieldTemplate6 'void (ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &&)'
+  // CHECK-NEXT:       ParmVarDecl {{.*}} 'ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> &&'
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> InstantiatedCWAFS13;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> InstantiatedCWAFS14;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> InstantiatedCWAFS15;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> InstantiatedCWAFS16;
+
+  (void)*InstantiatedCWAFS1.ptr;
+  (void)*InstantiatedCWAFS2.ptr;
+  (void)*InstantiatedCWAFS3.ptr;
+  (void)*InstantiatedCWAFS4.ptr;
+  (void)*InstantiatedCWAFS5.ptr;
+  (void)*InstantiatedCWAFS6.ptr;
+  (void)*InstantiatedCWAFS7.ptr;
+  (void)*InstantiatedCWAFS8.ptr;
+  (void)*InstantiatedCWAFS9.ptr;
+  (void)*InstantiatedCWAFS10.ptr;
+  (void)*InstantiatedCWAFS11.ptr;
+  (void)*InstantiatedCWAFS12.ptr;
+  (void)*InstantiatedCWAFS13.ptr;
+  (void)*InstantiatedCWAFS14.ptr;
+  (void)*InstantiatedCWAFS15.ptr;
+  (void)*InstantiatedCWAFS16.ptr;
+}
diff --git a/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp b/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp
new file mode 100644
index 0000000000000..e254934db8c0d
--- /dev/null
+++ b/clang/test/CodeGenSYCL/add_ir_annotations_member.cpp
@@ -0,0 +1,138 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -DTEST_SCALAR -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+
+// Tests the generation of IR annotation calls from
+// __sycl_detail__::add_ir_annotations_member attributes.
+
+#include "mock_properties.hpp"
+#include "sycl.hpp"
+
+#ifdef TEST_SCALAR
+#define TEST_T char
+#else
+#define TEST_T int *
+#endif
+
+template <typename... Properties> class g {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(Properties::name..., Properties::value...)]]
+#endif
+      ;
+
+  g() : x() {}
+  g(TEST_T _x) : x(_x) {}
+};
+
+class h {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17",
+          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]]
+#endif
+      ;
+
+  h() : x() {}
+  h(TEST_T _x) : x(_x) {}
+};
+
+template <typename... Properties> class gh {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          Properties::name..., "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17",
+          Properties::value..., "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2)]]
+#endif
+      ;
+
+  gh() : x() {}
+  gh(TEST_T _x) : x(_x) {}
+};
+
+template <typename... Properties> class hg {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          "Prop11", "Prop12", "Prop13", "Prop14", "Prop15", "Prop16", "Prop17", Properties::name...,
+          "Another property string", 2, false, TestEnum::Enum1, nullptr, nullptr, ScopedTestEnum::ScopedEnum2, Properties::value...)]]
+#endif
+      ;
+
+  hg() : x() {}
+  hg(TEST_T _x) : x(_x) {}
+};
+
+int main() {
+  sycl::queue q;
+  g<prop1, prop2, prop3, prop4, prop5, prop6, prop7> a;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel1>(
+        [=]() {
+          (void)a.x;
+        });
+  });
+  h b;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel2>(
+        [=]() {
+          (void)b.x;
+        });
+  });
+  gh<prop1, prop2, prop3, prop4, prop5, prop6, prop7> c;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel3>(
+        [=]() {
+          (void)c.x;
+        });
+  });
+  hg<prop1, prop2, prop3, prop4, prop5, prop6, prop7> d;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel4>(
+        [=]() {
+          (void)d.x;
+        });
+  });
+}
+
+// CHECK-DAG: @[[AnnotName:.*]] = private unnamed_addr constant [16 x i8] c"sycl-properties\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop1\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop2Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop2\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop3Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop3\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop4Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop4\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop5Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop5\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop6Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop6\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop7Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop7\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop11Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop11\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop12Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop12\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop13Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop13\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop14Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop14\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop15Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop15\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop16Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop16\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop17Name:.*]] = private unnamed_addr constant [7 x i8] c"Prop17\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Value:.*]] = private unnamed_addr constant [16 x i8] c"Property string\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop2_7_14Value:.*]] = private unnamed_addr constant [2 x i8] c"1\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop3Value:.*]] = private unnamed_addr constant [5 x i8] c"true\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop4_12_17Value:.*]] = private unnamed_addr constant [2 x i8] c"2\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop11Value:.*]] = private unnamed_addr constant [24 x i8] c"Another property string\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop13Value:.*]] = private unnamed_addr constant [6 x i8] c"false\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[GArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2_7_14Value]], [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]], [6 x i8]* @[[Prop4Name]], [2 x i8]* @[[Prop4_12_17Value]], [6 x i8]* @[[Prop5Name]], i8* null, [6 x i8]* @[[Prop6Name]], i8* null, [6 x i8]* @[[Prop7Name]], [2 x i8]* @[[Prop2_7_14Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[HArgs:.*]] = private unnamed_addr constant { [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]* } { [7 x i8]* @[[Prop11Name]], [24 x i8]* @[[Prop11Value]], [7 x i8]* @[[Prop12Name]], [2 x i8]* @[[Prop4_12_17Value]], [7 x i8]* @[[Prop13Name]], [6 x i8]* @[[Prop13Value]], [7 x i8]* @[[Prop14Name]], [2 x i8]* @[[Prop2_7_14Value]], [7 x i8]* @[[Prop15Name]], i8* null, [7 x i8]* @[[Prop16Name]], i8* null, [7 x i8]* @[[Prop17Name]], [2 x i8]* @[[Prop4_12_17Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[GHArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]*, [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2_7_14Value]], [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]], [6 x i8]* @[[Prop4Name]], [2 x i8]* @[[Prop4_12_17Value]], [6 x i8]* @[[Prop5Name]], i8* null, [6 x i8]* @[[Prop6Name]], i8* null, [6 x i8]* @[[Prop7Name]], [2 x i8]* @[[Prop2_7_14Value]], [7 x i8]* @[[Prop11Name]], [24 x i8]* @[[Prop11Value]], [7 x i8]* @[[Prop12Name]], [2 x i8]* @[[Prop4_12_17Value]], [7 x i8]* @[[Prop13Name]], [6 x i8]* @[[Prop13Value]], [7 x i8]* @[[Prop14Name]], [2 x i8]* @[[Prop2_7_14Value]], [7 x i8]* @[[Prop15Name]], i8* null, [7 x i8]* @[[Prop16Name]], i8* null, [7 x i8]* @[[Prop17Name]], [2 x i8]* @[[Prop4_12_17Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[HGArgs:.*]] = private unnamed_addr constant { [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]*, [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* } { [7 x i8]* @[[Prop11Name]], [24 x i8]* @[[Prop11Value]], [7 x i8]* @[[Prop12Name]], [2 x i8]* @[[Prop4_12_17Value]], [7 x i8]* @[[Prop13Name]], [6 x i8]* @[[Prop13Value]], [7 x i8]* @[[Prop14Name]], [2 x i8]* @[[Prop2_7_14Value]], [7 x i8]* @[[Prop15Name]], i8* null, [7 x i8]* @[[Prop16Name]], i8* null, [7 x i8]* @[[Prop17Name]], [2 x i8]* @[[Prop4_12_17Value]], [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2_7_14Value]], [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]], [6 x i8]* @[[Prop4Name]], [2 x i8]* @[[Prop4_12_17Value]], [6 x i8]* @[[Prop5Name]], i8* null, [6 x i8]* @[[Prop6Name]], i8* null, [6 x i8]* @[[Prop7Name]], [2 x i8]* @[[Prop2_7_14Value]] }, section "llvm.metadata"
+
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* }* @[[GArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]* }* @[[HArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]*, [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]* }* @[[GHArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [7 x i8]*, [24 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, [6 x i8]*, [7 x i8]*, [2 x i8]*, [7 x i8]*, i8*, [7 x i8]*, i8*, [7 x i8]*, [2 x i8]*, [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, i8*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* }* @[[HGArgs]] to i8*))
diff --git a/clang/test/CodeGenSYCL/add_ir_annotations_member_filter.cpp b/clang/test/CodeGenSYCL/add_ir_annotations_member_filter.cpp
new file mode 100644
index 0000000000000..d9c2000761a29
--- /dev/null
+++ b/clang/test/CodeGenSYCL/add_ir_annotations_member_filter.cpp
@@ -0,0 +1,63 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -DTEST_SCALAR -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+
+// Tests the optional filter parameter of
+// __sycl_detail__::add_ir_annotations_member attributes.
+
+#include "mock_properties.hpp"
+#include "sycl.hpp"
+
+#ifdef TEST_SCALAR
+#define TEST_T char
+#else
+#define TEST_T int *
+#endif
+
+template <typename... Properties> class g {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          {"Prop1", "Prop7", "Prop5"},
+          Properties::name..., Properties::value...)]]
+#endif
+      ;
+
+  g() : x() {}
+  g(TEST_T _x) : x(_x) {}
+};
+
+int main() {
+  sycl::queue q;
+  g<prop1, prop2, prop3, prop4, prop5, prop6, prop7> a;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel1>(
+        [=]() {
+          (void)a.x;
+        });
+  });
+}
+
+// CHECK-DAG: @[[AnnotName:.*]] = private unnamed_addr constant [16 x i8] c"sycl-properties\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop1\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop5Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop5\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop7Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop7\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Value:.*]] = private unnamed_addr constant [16 x i8] c"Property string\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop2_7Value:.*]] = private unnamed_addr constant [2 x i8] c"1\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[GArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop5Name]], i8* null, [6 x i8]* @[[Prop7Name]], [2 x i8]* @[[Prop2_7Value]] }, section "llvm.metadata"
+
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, i8*, [6 x i8]*, [2 x i8]* }* @[[GArgs]] to i8*))
+
+// CHECK-NOT: @[[Prop2Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop2\00", section "llvm.metadata"
+// CHECK-NOT: @[[Prop3Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop3\00", section "llvm.metadata"
+// CHECK-NOT: @[[Prop4Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop4\00", section "llvm.metadata"
+// CHECK-NOT: @[[Prop6Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop6\00", section "llvm.metadata"
+// CHECK-NOT: @[{{.*}} = private unnamed_addr constant [5 x i8] c"true\00", section "llvm.metadata"
+// CHECK-NOT: @{{.*}} = private unnamed_addr constant [2 x i8] c"2\00", section "llvm.metadata"
diff --git a/clang/test/CodeGenSYCL/add_ir_annotations_member_reuse.cpp b/clang/test/CodeGenSYCL/add_ir_annotations_member_reuse.cpp
new file mode 100644
index 0000000000000..ca5bcdd4dffc0
--- /dev/null
+++ b/clang/test/CodeGenSYCL/add_ir_annotations_member_reuse.cpp
@@ -0,0 +1,112 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -internal-isystem %S/Inputs -disable-llvm-passes \
+// RUN:    -triple spir64-unknown-unknown -fsycl-is-device -DTEST_SCALAR -S \
+// RUN:    -emit-llvm %s -o - | FileCheck %s
+
+// Tests the reuse of generated annotation value global variables for
+// __sycl_detail__::add_ir_annotations_member attributes.
+
+#include "mock_properties.hpp"
+#include "sycl.hpp"
+
+#ifdef TEST_SCALAR
+#define TEST_T char
+#else
+#define TEST_T int *
+#endif
+
+template <typename... Properties> class g {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          Properties::name..., Properties::value...)]]
+#endif
+      ;
+
+  g() : x() {}
+  g(TEST_T _x) : x(_x) {}
+};
+
+template <typename... Properties> class h {
+public:
+  TEST_T x
+#ifdef __SYCL_DEVICE_ONLY__
+      [[__sycl_detail__::add_ir_annotations_member(
+          {"Prop1", "Prop2", "Prop3"},
+          Properties::name..., Properties::value...)]]
+#endif
+      ;
+
+  h() : x() {}
+  h(TEST_T _x) : x(_x) {}
+};
+
+int main() {
+  sycl::queue q;
+  g<prop1, prop2, prop3> a;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel1>(
+        [=]() {
+          (void)a.x;
+        });
+  });
+  g<prop1, prop2, prop3> b;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel2>(
+        [=]() {
+          (void)b.x;
+        });
+  });
+  h<prop1, prop2, prop4, prop3> c;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel3>(
+        [=]() {
+          (void)c.x;
+        });
+  });
+  g<prop1, prop2> d;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel4>(
+        [=]() {
+          (void)d.x;
+        });
+  });
+  g<prop1, prop2, prop3, prop5> e;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel5>(
+        [=]() {
+          (void)e.x;
+        });
+  });
+  g<prop3, prop2, prop1> f;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel6>(
+        [=]() {
+          (void)f.x;
+        });
+  });
+}
+
+// CHECK-DAG: @[[AnnotName:.*]] = private unnamed_addr constant [16 x i8] c"sycl-properties\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop1\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop2Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop2\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop3Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop3\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop5Name:.*]] = private unnamed_addr constant [6 x i8] c"Prop5\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[Prop1Value:.*]] = private unnamed_addr constant [16 x i8] c"Property string\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop2Value:.*]] = private unnamed_addr constant [2 x i8] c"1\00", section "llvm.metadata"
+// CHECK-DAG: @[[Prop3Value:.*]] = private unnamed_addr constant [5 x i8] c"true\00", section "llvm.metadata"
+
+// CHECK-DAG: @[[ReusedArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2Value]], [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[DArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2Value]] }, section "llvm.metadata"
+// CHECK-DAG: @[[EArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, i8* } { [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2Value]], [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]], [6 x i8]* @[[Prop5Name]], i8* null }, section "llvm.metadata"
+// CHECK-DAG: @[[FArgs:.*]] = private unnamed_addr constant { [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [16 x i8]* } { [6 x i8]* @[[Prop3Name]], [5 x i8]* @[[Prop3Value]], [6 x i8]* @[[Prop2Name]], [2 x i8]* @[[Prop2Value]], [6 x i8]* @[[Prop1Name]], [16 x i8]* @[[Prop1Value]] }, section "llvm.metadata"
+
+// CHECK-COUNT-3: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]* }* @[[ReusedArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]* }* @[[DArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [16 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [5 x i8]*, [6 x i8]*, i8* }* @[[EArgs]] to i8*))
+// CHECK-DAG: %{{.*}} = call i8 addrspace(4)* @llvm.ptr.annotation.p4i8(i8 {{.*}}, i8* getelementptr inbounds ([16 x i8], [16 x i8]* @[[AnnotName]], i32 0, i32 0), i8* getelementptr inbounds {{.*}}, i32 {{.*}}, i8* bitcast ({ [6 x i8]*, [5 x i8]*, [6 x i8]*, [2 x i8]*, [6 x i8]*, [16 x i8]* }* @[[FArgs]] to i8*))
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index dd2516beefb60..16057b939d904 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -156,6 +156,7 @@
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLAddIRAnnotationsMember (SubjectMatchRule_field)
 // CHECK-NEXT: SYCLAddIRAttributesFunction (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLAddIRAttributesGlobalVariable (SubjectMatchRule_record)
 // CHECK-NEXT: SYCLAddIRAttributesKernelParameter (SubjectMatchRule_variable_is_parameter)
diff --git a/clang/test/SemaSYCL/attr-add-ir-annotations-packs.cpp b/clang/test/SemaSYCL/attr-add-ir-annotations-packs.cpp
new file mode 100644
index 0000000000000..33e80417bb390
--- /dev/null
+++ b/clang/test/SemaSYCL/attr-add-ir-annotations-packs.cpp
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
+
+// Tests that __sycl_detail__::add_ir_annotations_member allows pack expansions
+// in its arguments.
+
+constexpr const char AttrName1[] = "Attr1";
+constexpr const char AttrName2[] = "Attr2";
+constexpr const char AttrName3[] = "Attr3";
+constexpr const char AttrVal1[] = "Val1";
+constexpr const char AttrVal2[] = "Val2";
+constexpr const char AttrVal3[] = "Val3";
+
+template <int... Is> struct ClassWithAnnotFieldTemplate1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", Is...)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+template <int... Is> struct ClassWithAnnotFieldTemplate2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", "Attr2", "Attr3", Is...)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+template <const char *...Names> struct ClassWithAnnotFieldTemplate3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Names..., 1, 2, 3)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+template <const char *...Names> struct ClassWithAnnotFieldTemplate4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, Names..., 1, 2, 3)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+template <const char *...Strs> struct ClassWithAnnotFieldTemplate5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Strs...)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+template <const char *...Strs> struct ClassWithAnnotFieldTemplate6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, Strs...)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+
+void InstantiateClassWithAnnotFieldTemplates() {
+  ClassWithAnnotFieldTemplate1<1, 2, 3> InstantiatedCWAFS1;
+  ClassWithAnnotFieldTemplate1<1, 2> InstantiatedCWAFS2; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate1<1, 2>' requested here}}
+
+  ClassWithAnnotFieldTemplate2<1, 2, 3> InstantiatedCWAFS3;
+  ClassWithAnnotFieldTemplate2<1, 2> InstantiatedCWAFS4; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate2<1, 2>' requested here}}
+
+  ClassWithAnnotFieldTemplate3<AttrName1, AttrName2, AttrName3> InstantiatedCWAFS5;
+  ClassWithAnnotFieldTemplate3<AttrName1, AttrName2> InstantiatedCWAFS6; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate3<AttrName1, AttrName2>' requested here}}
+
+  ClassWithAnnotFieldTemplate4<AttrName1, AttrName2, AttrName3> InstantiatedCWAFS7;
+  ClassWithAnnotFieldTemplate4<AttrName1, AttrName2> InstantiatedCWAFS8; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate4<AttrName1, AttrName2>' requested here}}
+
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrVal1> InstantiatedCWAFS9;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrVal1, AttrVal2> InstantiatedCWAFS10;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> InstantiatedCWAFS11;
+  ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> InstantiatedCWAFS12; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate5<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2>' requested here}}
+
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrVal1> InstantiatedCWAFS13;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrVal1, AttrVal2> InstantiatedCWAFS14;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2, AttrVal3> InstantiatedCWAFS15;
+  ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2> InstantiatedCWAFS16; // expected-note {{in instantiation of template class 'ClassWithAnnotFieldTemplate6<AttrName1, AttrName2, AttrName3, AttrVal1, AttrVal2>' requested here}}
+
+  (void)*InstantiatedCWAFS1.ptr;
+  (void)*InstantiatedCWAFS2.ptr;
+  (void)*InstantiatedCWAFS3.ptr;
+  (void)*InstantiatedCWAFS4.ptr;
+  (void)*InstantiatedCWAFS5.ptr;
+  (void)*InstantiatedCWAFS6.ptr;
+  (void)*InstantiatedCWAFS7.ptr;
+  (void)*InstantiatedCWAFS8.ptr;
+  (void)*InstantiatedCWAFS9.ptr;
+  (void)*InstantiatedCWAFS10.ptr;
+  (void)*InstantiatedCWAFS11.ptr;
+  (void)*InstantiatedCWAFS12.ptr;
+  (void)*InstantiatedCWAFS13.ptr;
+  (void)*InstantiatedCWAFS14.ptr;
+  (void)*InstantiatedCWAFS15.ptr;
+  (void)*InstantiatedCWAFS16.ptr;
+}
diff --git a/clang/test/SemaSYCL/attr-add-ir-annotations.cpp b/clang/test/SemaSYCL/attr-add-ir-annotations.cpp
new file mode 100644
index 0000000000000..f14b362b19bb8
--- /dev/null
+++ b/clang/test/SemaSYCL/attr-add-ir-annotations.cpp
@@ -0,0 +1,579 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+
+// Tests valid and invalid arguments in
+// __sycl_detail__::add_ir_annotations_member attributes.
+
+enum TestEnum {
+  EnumVal1,
+  EnumVal2
+};
+
+enum class ScopedTestEnum : short {
+  ScopedEnumVal1,
+  ScopedEnumVal2
+};
+
+constexpr decltype(nullptr) CENullptr = nullptr;
+constexpr const char CEStr[] = "Text";
+constexpr int CEInt = 1;
+constexpr float CEFloat = 3.14;
+constexpr bool CETrue = true;
+constexpr bool CEFalse = false;
+constexpr TestEnum CEEnum = TestEnum::EnumVal1;
+constexpr char CEChar = 'F';
+constexpr ScopedTestEnum CESEnum = ScopedTestEnum::ScopedEnumVal2;
+
+constexpr const char CEAttrName1[] = "CEAttr1";
+constexpr const char CEAttrName2[] = "CEAttr2";
+constexpr const char CEAttrName3[] = "CEAttr3";
+constexpr const char CEAttrName4[] = "CEAttr4";
+constexpr const char CEAttrName5[] = "CEAttr5";
+constexpr const char CEAttrName6[] = "CEAttr6";
+constexpr const char CEAttrName7[] = "CEAttr7";
+constexpr const char CEAttrName8[] = "CEAttr8";
+constexpr const char CEAttrName9[] = "CEAttr9";
+
+struct ClassWithAnnotFieldLiteral1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", nullptr)]];
+};
+struct ClassWithAnnotFieldLiteral2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Text")]];
+};
+struct ClassWithAnnotFieldLiteral3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", 1)]];
+};
+struct ClassWithAnnotFieldLiteral4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", 3.14)]];
+};
+struct ClassWithAnnotFieldLiteral5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", true)]];
+};
+struct ClassWithAnnotFieldLiteral6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", false)]];
+};
+struct ClassWithAnnotFieldLiteral7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", TestEnum::EnumVal1)]];
+};
+struct ClassWithAnnotFieldLiteral8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", 'F')]];
+};
+struct ClassWithAnnotFieldLiteral9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", "Attr4", "Attr5", "Attr6", "Attr7", "Attr8", "Attr9", nullptr, "Text", 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldLiteral10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", nullptr)]];
+};
+struct ClassWithAnnotFieldLiteral11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", "Text")]];
+};
+struct ClassWithAnnotFieldLiteral12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", 1)]];
+};
+struct ClassWithAnnotFieldLiteral13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", 3.14)]];
+};
+struct ClassWithAnnotFieldLiteral14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", true)]];
+};
+struct ClassWithAnnotFieldLiteral15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", false)]];
+};
+struct ClassWithAnnotFieldLiteral16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", TestEnum::EnumVal1)]];
+};
+struct ClassWithAnnotFieldLiteral17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", 'F')]];
+};
+struct ClassWithAnnotFieldLiteral18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", "Attr2", "Attr3", "Attr4", "Attr5", "Attr6", "Attr7", "Attr8", "Attr9", nullptr, "Text", 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldLiteral19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldLiteral20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", ScopedTestEnum::ScopedEnumVal2)]];
+};
+
+struct ClassWithAnnotFieldCEVal1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CENullptr)]];
+};
+struct ClassWithAnnotFieldCEVal2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEStr)]];
+};
+struct ClassWithAnnotFieldCEVal3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEInt)]];
+};
+struct ClassWithAnnotFieldCEVal4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEFloat)]];
+};
+struct ClassWithAnnotFieldCEVal5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CETrue)]];
+};
+struct ClassWithAnnotFieldCEVal6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEFalse)]];
+};
+struct ClassWithAnnotFieldCEVal7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEEnum)]];
+};
+struct ClassWithAnnotFieldCEVal8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEChar)]];
+};
+struct ClassWithAnnotFieldCEVal9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", "Attr4", "Attr5", "Attr6", "Attr7", "Attr8", "Attr9", CENullptr, CEStr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]];
+};
+struct ClassWithAnnotFieldCEVal10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CENullptr)]];
+};
+struct ClassWithAnnotFieldCEVal11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEStr)]];
+};
+struct ClassWithAnnotFieldCEVal12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEInt)]];
+};
+struct ClassWithAnnotFieldCEVal13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEFloat)]];
+};
+struct ClassWithAnnotFieldCEVal14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CETrue)]];
+};
+struct ClassWithAnnotFieldCEVal15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEFalse)]];
+};
+struct ClassWithAnnotFieldCEVal16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEEnum)]];
+};
+struct ClassWithAnnotFieldCEVal17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CEChar)]];
+};
+struct ClassWithAnnotFieldCEVal18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", "Attr2", "Attr3", "Attr4", "Attr5", "Attr6", "Attr7", "Attr8", "Attr9", CENullptr, CEStr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]];
+};
+struct ClassWithAnnotFieldCEVal19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CESEnum)]];
+};
+struct ClassWithAnnotFieldCEVal20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, "Attr1", CESEnum)]];
+};
+
+struct ClassWithAnnotFieldCEName1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, nullptr)]];
+};
+struct ClassWithAnnotFieldCEName2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, "Text")]];
+};
+struct ClassWithAnnotFieldCEName3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, 1)]];
+};
+struct ClassWithAnnotFieldCEName4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, 3.14)]];
+};
+struct ClassWithAnnotFieldCEName5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, true)]];
+};
+struct ClassWithAnnotFieldCEName6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, false)]];
+};
+struct ClassWithAnnotFieldCEName7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, TestEnum::EnumVal1)]];
+};
+struct ClassWithAnnotFieldCEName8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, 'F')]];
+};
+struct ClassWithAnnotFieldCEName9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9, nullptr, "Text", 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldCEName10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, nullptr)]];
+};
+struct ClassWithAnnotFieldCEName11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, "Text")]];
+};
+struct ClassWithAnnotFieldCEName12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, 1)]];
+};
+struct ClassWithAnnotFieldCEName13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, 3.14)]];
+};
+struct ClassWithAnnotFieldCEName14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, true)]];
+};
+struct ClassWithAnnotFieldCEName15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, false)]];
+};
+struct ClassWithAnnotFieldCEName16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, TestEnum::EnumVal1)]];
+};
+struct ClassWithAnnotFieldCEName17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, 'F')]];
+};
+struct ClassWithAnnotFieldCEName18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9, nullptr, "Text", 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldCEName19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, ScopedTestEnum::ScopedEnumVal2)]];
+};
+struct ClassWithAnnotFieldCEName20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, ScopedTestEnum::ScopedEnumVal2)]];
+};
+
+struct ClassWithAnnotFieldCE1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CENullptr)]];
+};
+struct ClassWithAnnotFieldCE2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEStr)]];
+};
+struct ClassWithAnnotFieldCE3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEInt)]];
+};
+struct ClassWithAnnotFieldCE4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEFloat)]];
+};
+struct ClassWithAnnotFieldCE5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CETrue)]];
+};
+struct ClassWithAnnotFieldCE6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEFalse)]];
+};
+struct ClassWithAnnotFieldCE7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEEnum)]];
+};
+struct ClassWithAnnotFieldCE8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEChar)]];
+};
+struct ClassWithAnnotFieldCE9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9, CENullptr, CEStr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]];
+};
+struct ClassWithAnnotFieldCE10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CENullptr)]];
+};
+struct ClassWithAnnotFieldCE11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEStr)]];
+};
+struct ClassWithAnnotFieldCE12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEInt)]];
+};
+struct ClassWithAnnotFieldCE13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEFloat)]];
+};
+struct ClassWithAnnotFieldCE14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CETrue)]];
+};
+struct ClassWithAnnotFieldCE15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEFalse)]];
+};
+struct ClassWithAnnotFieldCE16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEEnum)]];
+};
+struct ClassWithAnnotFieldCE17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEChar)]];
+};
+struct ClassWithAnnotFieldCE18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9, CENullptr, CEStr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]];
+};
+struct ClassWithAnnotFieldCE19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CESEnum)]];
+};
+struct ClassWithAnnotFieldCE20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr1", "Attr3"}, CEAttrName1, CESEnum)]];
+};
+
+template <decltype(nullptr) Null> struct ClassWithAnnotFieldTemplate1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", Null)]];
+};
+template <const char *Str> struct ClassWithAnnotFieldTemplate2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", Str)]];
+};
+template <int I> struct ClassWithAnnotFieldTemplate3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", I)]];
+};
+template <bool B> struct ClassWithAnnotFieldTemplate4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", B)]];
+};
+template <TestEnum E> struct ClassWithAnnotFieldTemplate5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", E)]];
+};
+template <char C> struct ClassWithAnnotFieldTemplate6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", C)]];
+};
+template <decltype(nullptr) Null, const char *Str, int I, bool B, TestEnum E, char C, ScopedTestEnum SE> struct ClassWithAnnotFieldTemplate7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", "Attr4", "Attr5", "Attr6", "Attr7", Null, Str, I, B, E, C, SE)]];
+};
+template <decltype(nullptr) Null> struct ClassWithAnnotFieldTemplate8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, Null)]];
+};
+template <const char *Str> struct ClassWithAnnotFieldTemplate9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, Str)]];
+};
+template <int I> struct ClassWithAnnotFieldTemplate10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, I)]];
+};
+template <bool B> struct ClassWithAnnotFieldTemplate11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, B)]];
+};
+template <TestEnum E> struct ClassWithAnnotFieldTemplate12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, E)]];
+};
+template <char C> struct ClassWithAnnotFieldTemplate13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, C)]];
+};
+template <decltype(nullptr) Null, const char *Str, int I, bool B, TestEnum E, char C, ScopedTestEnum SE> struct ClassWithAnnotFieldTemplate14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, Null, Str, I, B, E, C, SE)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, nullptr)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, "Text")]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, 1)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, 3.14)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, true)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, false)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate21 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, TestEnum::EnumVal1)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate22 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, 'F')]];
+};
+template <const char *Name1, const char *Name2, const char *Name3, const char *Name4, const char *Name5, const char *Name6, const char *Name7, const char *Name8, const char *Name9> struct ClassWithAnnotFieldTemplate23 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name1, Name2, Name3, Name4, Name5, Name6, Name7, Name8, Name9, nullptr, "Text", 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate24 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CENullptr)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate25 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEStr)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate26 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEInt)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate27 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEFloat)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate28 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CETrue)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate29 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEFalse)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate30 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEEnum)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate31 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CEChar)]];
+};
+template <const char *Name1, const char *Name2, const char *Name3, const char *Name4, const char *Name5, const char *Name6, const char *Name7, const char *Name8, const char *Name9> struct ClassWithAnnotFieldTemplate32 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name1, Name2, Name3, Name4, Name5, Name6, Name7, Name8, Name9, CENullptr, CEStr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]];
+};
+template <ScopedTestEnum SE> struct ClassWithAnnotFieldTemplate33 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", SE)]];
+};
+template <ScopedTestEnum SE> struct ClassWithAnnotFieldTemplate34 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEAttrName1, SE)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate35 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, ScopedTestEnum::ScopedEnumVal2)]];
+};
+template <const char *Name> struct ClassWithAnnotFieldTemplate36 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(Name, CESEnum)]];
+};
+void InstantiateClassWithAnnotFieldTemplates() {
+  ClassWithAnnotFieldTemplate1<nullptr> InstantiatedCWAFS1;
+  ClassWithAnnotFieldTemplate1<CENullptr> InstantiatedCWAFS2;
+  ClassWithAnnotFieldTemplate2<CEStr> InstantiatedCWAFS3;
+  ClassWithAnnotFieldTemplate3<1> InstantiatedCWAFS4;
+  ClassWithAnnotFieldTemplate3<CEInt> InstantiatedCWAFS5;
+  ClassWithAnnotFieldTemplate4<true> InstantiatedCWAFS6;
+  ClassWithAnnotFieldTemplate4<CETrue> InstantiatedCWAFS7;
+  ClassWithAnnotFieldTemplate4<false> InstantiatedCWAFS8;
+  ClassWithAnnotFieldTemplate4<CEFalse> InstantiatedCWAFS9;
+  ClassWithAnnotFieldTemplate5<TestEnum::EnumVal1> InstantiatedCWAFS10;
+  ClassWithAnnotFieldTemplate5<CEEnum> InstantiatedCWAFS11;
+  ClassWithAnnotFieldTemplate6<'F'> InstantiatedCWAFS12;
+  ClassWithAnnotFieldTemplate6<CEChar> InstantiatedCWAFS13;
+  ClassWithAnnotFieldTemplate7<nullptr, CEStr, 1, true, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2> InstantiatedCWAFS14;
+  ClassWithAnnotFieldTemplate7<CENullptr, CEStr, CEInt, CETrue, CEEnum, CEChar, CESEnum> InstantiatedCWAFS15;
+  ClassWithAnnotFieldTemplate8<nullptr> InstantiatedCWAFS16;
+  ClassWithAnnotFieldTemplate8<CENullptr> InstantiatedCWAFS17;
+  ClassWithAnnotFieldTemplate9<CEStr> InstantiatedCWAFS18;
+  ClassWithAnnotFieldTemplate10<1> InstantiatedCWAFS19;
+  ClassWithAnnotFieldTemplate10<CEInt> InstantiatedCWAFS20;
+  ClassWithAnnotFieldTemplate11<true> InstantiatedCWAFS21;
+  ClassWithAnnotFieldTemplate11<CETrue> InstantiatedCWAFS22;
+  ClassWithAnnotFieldTemplate11<false> InstantiatedCWAFS23;
+  ClassWithAnnotFieldTemplate11<CEFalse> InstantiatedCWAFS24;
+  ClassWithAnnotFieldTemplate12<TestEnum::EnumVal1> InstantiatedCWAFS25;
+  ClassWithAnnotFieldTemplate12<CEEnum> InstantiatedCWAFS26;
+  ClassWithAnnotFieldTemplate13<'F'> InstantiatedCWAFS27;
+  ClassWithAnnotFieldTemplate13<CEChar> InstantiatedCWAFS28;
+  ClassWithAnnotFieldTemplate14<nullptr, CEStr, 1, true, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2> InstantiatedCWAFS29;
+  ClassWithAnnotFieldTemplate14<CENullptr, CEStr, CEInt, CETrue, CEEnum, CEChar, CESEnum> InstantiatedCWAFS30;
+  ClassWithAnnotFieldTemplate15<CEAttrName1> InstantiatedCWAFS31;
+  ClassWithAnnotFieldTemplate16<CEAttrName1> InstantiatedCWAFS32;
+  ClassWithAnnotFieldTemplate17<CEAttrName1> InstantiatedCWAFS33;
+  ClassWithAnnotFieldTemplate18<CEAttrName1> InstantiatedCWAFS34;
+  ClassWithAnnotFieldTemplate19<CEAttrName1> InstantiatedCWAFS35;
+  ClassWithAnnotFieldTemplate20<CEAttrName1> InstantiatedCWAFS36;
+  ClassWithAnnotFieldTemplate21<CEAttrName1> InstantiatedCWAFS37;
+  ClassWithAnnotFieldTemplate22<CEAttrName1> InstantiatedCWAFS38;
+  ClassWithAnnotFieldTemplate23<CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9> InstantiatedCWAFS39;
+  ClassWithAnnotFieldTemplate24<CEAttrName1> InstantiatedCWAFS40;
+  ClassWithAnnotFieldTemplate25<CEAttrName1> InstantiatedCWAFS41;
+  ClassWithAnnotFieldTemplate26<CEAttrName1> InstantiatedCWAFS42;
+  ClassWithAnnotFieldTemplate27<CEAttrName1> InstantiatedCWAFS43;
+  ClassWithAnnotFieldTemplate28<CEAttrName1> InstantiatedCWAFS44;
+  ClassWithAnnotFieldTemplate29<CEAttrName1> InstantiatedCWAFS45;
+  ClassWithAnnotFieldTemplate30<CEAttrName1> InstantiatedCWAFS46;
+  ClassWithAnnotFieldTemplate31<CEAttrName1> InstantiatedCWAFS47;
+  ClassWithAnnotFieldTemplate32<CEAttrName1, CEAttrName2, CEAttrName3, CEAttrName4, CEAttrName5, CEAttrName6, CEAttrName7, CEAttrName8, CEAttrName9> InstantiatedCWAFS48;
+  ClassWithAnnotFieldTemplate33<ScopedTestEnum::ScopedEnumVal2> InstantiatedCWAFS49;
+  ClassWithAnnotFieldTemplate34<ScopedTestEnum::ScopedEnumVal2> InstantiatedCWAFS50;
+  ClassWithAnnotFieldTemplate35<CEAttrName1> InstantiatedCWAFS51;
+  ClassWithAnnotFieldTemplate36<CEAttrName1> InstantiatedCWAFS52;
+
+  (void)*InstantiatedCWAFS1.ptr;
+  (void)*InstantiatedCWAFS2.ptr;
+  (void)*InstantiatedCWAFS3.ptr;
+  (void)*InstantiatedCWAFS4.ptr;
+  (void)*InstantiatedCWAFS5.ptr;
+  (void)*InstantiatedCWAFS6.ptr;
+  (void)*InstantiatedCWAFS7.ptr;
+  (void)*InstantiatedCWAFS8.ptr;
+  (void)*InstantiatedCWAFS9.ptr;
+  (void)*InstantiatedCWAFS10.ptr;
+  (void)*InstantiatedCWAFS11.ptr;
+  (void)*InstantiatedCWAFS12.ptr;
+  (void)*InstantiatedCWAFS13.ptr;
+  (void)*InstantiatedCWAFS14.ptr;
+  (void)*InstantiatedCWAFS15.ptr;
+  (void)*InstantiatedCWAFS16.ptr;
+  (void)*InstantiatedCWAFS17.ptr;
+  (void)*InstantiatedCWAFS18.ptr;
+  (void)*InstantiatedCWAFS19.ptr;
+  (void)*InstantiatedCWAFS20.ptr;
+  (void)*InstantiatedCWAFS21.ptr;
+  (void)*InstantiatedCWAFS22.ptr;
+  (void)*InstantiatedCWAFS23.ptr;
+  (void)*InstantiatedCWAFS24.ptr;
+  (void)*InstantiatedCWAFS25.ptr;
+  (void)*InstantiatedCWAFS26.ptr;
+  (void)*InstantiatedCWAFS27.ptr;
+  (void)*InstantiatedCWAFS28.ptr;
+  (void)*InstantiatedCWAFS29.ptr;
+  (void)*InstantiatedCWAFS30.ptr;
+  (void)*InstantiatedCWAFS31.ptr;
+  (void)*InstantiatedCWAFS32.ptr;
+  (void)*InstantiatedCWAFS33.ptr;
+  (void)*InstantiatedCWAFS34.ptr;
+  (void)*InstantiatedCWAFS35.ptr;
+  (void)*InstantiatedCWAFS36.ptr;
+}
+
+struct InvalidClassWithAnnotField1 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1")]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField2 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", nullptr, "Attr2")]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField3 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", nullptr)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField4 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr5", "Attr3"}, "Attr1")]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField5 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr5", "Attr3"}, "Attr1", nullptr, "Attr2")]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField6 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr5", "Attr3"}, "Attr1", "Attr2", nullptr)]]; // expected-error {{attribute 'add_ir_annotations_member' must specify a value for each specified name in the argument list}}
+};
+struct InvalidClassWithAnnotField7 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", {"Attr5", "Attr3"}, nullptr)]]; // expected-error {{only the first argument of attribute 'add_ir_annotations_member' can be an initializer list}}
+};
+struct InvalidClassWithAnnotField8 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", nullptr, {"Attr5", "Attr3"})]]; // expected-error {{only the first argument of attribute 'add_ir_annotations_member' can be an initializer list}}
+};
+struct InvalidClassWithAnnotField9 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", {"Attr5", "Attr3"}, nullptr, "Text")]]; // expected-error {{only the first argument of attribute 'add_ir_annotations_member' can be an initializer list}}
+};
+struct InvalidClassWithAnnotField10 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({1}, "Attr1", nullptr)]]; // expected-error {{initializer list in the first argument of 'add_ir_annotations_member' must contain only string literals}}
+};
+struct InvalidClassWithAnnotField11 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({true, "Attr3"}, "Attr1", nullptr)]]; // expected-error {{initializer list in the first argument of 'add_ir_annotations_member' must contain only string literals}}
+};
+struct InvalidClassWithAnnotField12 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member({"Attr3", 'c'}, "Attr1", nullptr)]]; // expected-error {{initializer list in the first argument of 'add_ir_annotations_member' must contain only string literals}}
+};
+struct InvalidClassWithAnnotField13 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(nullptr, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField14 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(1, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField15 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(3.14, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField16 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(true, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField17 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(false, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField18 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(TestEnum::EnumVal1, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField19 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member('F', "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField20 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(nullptr, 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2, nullptr, 1, 3.14, true, false, TestEnum::EnumVal1, 'F', ScopedTestEnum::ScopedEnumVal2)]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField21 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", 3.14, "Attr3", 1, 3.14, true)]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField22 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CENullptr, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField23 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEInt, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField24 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEFloat, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField25 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CETrue, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField26 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEFalse, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField27 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEEnum, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField28 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CEChar, "Attr1")]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField29 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member(CENullptr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum, CENullptr, CEInt, CEFloat, CETrue, CEFalse, CEEnum, CEChar, CESEnum)]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField30 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", CEFloat, "Attr3", CEInt, CEFloat, CETrue)]]; // expected-error {{each name argument in 'add_ir_annotations_member' must be a 'const char *' usable in a constant expression}}
+};
+struct InvalidClassWithAnnotField31 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", &CEInt)]]; // expected-error {{each value argument in 'add_ir_annotations_member' must be an integer, a floating point, a character, a boolean, 'const char *', or an enumerator usable as a constant expression}}
+};
+struct InvalidClassWithAnnotField32 {
+  int *ptr [[__sycl_detail__::add_ir_annotations_member("Attr1", "Attr2", "Attr3", 1, &CEInt, CEInt)]]; // expected-error {{each value argument in 'add_ir_annotations_member' must be an integer, a floating point, a character, a boolean, 'const char *', or an enumerator usable as a constant expression}}
+};
+
+struct [[__sycl_detail__::add_ir_annotations_member("Attr1", 1)]] InvalidAnnotationsMemberSubjectStruct;     // expected-error {{'add_ir_annotations_member' attribute only applies to non-static data members}}
+[[__sycl_detail__::add_ir_annotations_member("Attr1", 1)]] void InvalidAnnotationsMemberSubjectFunction() {} // expected-error {{'add_ir_annotations_member' attribute only applies to non-static data members}}
+[[__sycl_detail__::add_ir_annotations_member("Attr1", 1)]] int InvalidAnnotationsMemberSubjectVar;           // expected-error {{'add_ir_annotations_member' attribute only applies to non-static data members}}