[Sema] Improve diagnostic about addr spaces for overload candidates

[Anastasia]

Summary

As discussed in https://reviews.llvm.org/D69938#inline-629726 this commit improves the diagnostic of addr spaces. The approach is currently reusing diagnostic streaming of Qualifiers.

There are a number of issues however:

1. Address spaces don't always have representation i.e. in C++ it is just a number or Default address space doesn't correspond to anything at all. For the former case with the current approach we will get '__attribute__((address_space(N)))' printed into the diagnostic. For the latter one it will print unqualified.

2. In OpenCL we threat __private and Default in the same way and therefore it doesn't get printed. With this patch therefore unqualified will appear even if __private was specified in the original source. There is a bug open to fix that however: https://bugs.llvm.org/show_bug.cgi?id=43295. Perhaps it should be fixed as soon as this gets committed.

[rjmccall]

Hmm. How about:

// For the `this` argument
candidate function not viable: 'this' object is in '__private' address space, but method expects object in '__global' address space

// For pointer arguments
candidate function not viable: cannot pass pointer to '__private' address space as a pointer to '__global' address space in 1st argument

// For reference arguments
candidate function not viable: cannot bind reference in '__global' address space to object in '__private' address space in 1st argument

This would require you to render a string describing the address space yourself. I would suggest "generic address space" for the default AS outside of OpenCL, "'foo' address space" for a language-specific AS (including implicit __private in OpenCL), or "address space N" for a numbered address space.

What do you think?

[Anastasia]

• Allow sending address spaces into diagnostics
• Change wording of diagnostics for address spaces in overloading

[rjmccall]

clang/include/clang/Basic/DiagnosticSemaKinds.td

3996

I think you need to move "address space" into the diagnostic rendering for ak_addrspace so that you can say "generic address space" instead of "address space generic".

clang/lib/AST/TypePrinter.cpp

1777

Can you just make these two cases in the switch? opencl_private already seems to be there.

[Anastasia]

• Moved "address space" printing into diagnostic engine
• Moved LangAS::Default into switch/case statement.

[rjmccall]

Thanks, that looks much better.

[zatrazz]

Hi,

I am investigating a recurring regression on aarch64-linux bots and bisecting the commits on the build [1] that introduced the regression it points to this one. I don't understand exactly what is triggering the issue, but it only happen on the stage2 build with stage1 clang built with different compilers (gcc 5.4, gcc-7.4.0, and gcc-8.3.0). The issue is on clang itself while trying to execute the Analysis/uninit-sometimes.cpp test:

******************** TEST 'Clang :: Analysis/uninit-sometimes.cpp' FAILED ********************
Script:
--
: 'RUN: at line 1';   /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/clang -cc1 -internal-isystem /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/lib/clang/10.0.0/include -nostdsysteminc -std=gnu++11 -Wsometimes-uninitialized -verify /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp
: 'RUN: at line 2';   /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/clang -cc1 -internal-isystem /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/lib/clang/10.0.0/include -nostdsysteminc -std=gnu++11 -Wsometimes-uninitialized -fdiagnostics-parseable-fixits /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp 2>&1 | /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/FileCheck /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp
--
Exit Code: 134

Command Output (stderr):
--
clang: /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/lib/Basic/Diagnostic.cpp:591: void HandleSelectModifier(const clang::Diagnostic &, unsigned int, const char *, unsigned int, SmallVectorImpl<char> &): Assertion `NextVal != ArgumentEnd && "Value for integer select modifier was" " larger than the number of options in the diagnostic string!"' failed.
Stack dump:
0.	Program arguments: /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/clang -cc1 -internal-isystem /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/lib/clang/10.0.0/include -nostdsysteminc -std=gnu++11 -Wsometimes-uninitialized -verify /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp 
1.	/home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp:161:1: current parser token 'int'
2.	/home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp:146:32: parsing function body 'test_for_range_true'
#0 0x000000000182be04 PrintStackTraceSignalHandler(void*) (/home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/clang+0x182be04)
/home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/tools/clang/test/Analysis/Output/uninit-sometimes.cpp.script: line 2: 63297 Aborted                 (core dumped) /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/bin/clang -cc1 -internal-isystem /home/buildslave/buildslave/clang-cmake-aarch64-full/stage2/lib/clang/10.0.0/include -nostdsysteminc -std=gnu++11 -Wsometimes-uninitialized -verify /home/buildslave/buildslave/clang-cmake-aarch64-full/llvm/clang/test/Analysis/uninit-sometimes.cpp

The bots run a ubuntu 16.04 container and I can't reproduce on a different one, I am still trying to come up with a better reproduce case. Also, moving the both DiagnosticBuilder::AddString and DiagnosticBuilder::AddTaggedVal outside clang/include/clang/Basic/Diagnostic.h to clang/lib/Basic/Diagnostic.cpp "fixes" the issue (indicating the something is wrong at code generation). This issue has been shown for some time on aarch64-linux bot [2]. Any idea if this is an environmental issue or a real issue?

[1] http://lab.llvm.org:8011/builders/clang-cmake-aarch64-full/builds/8562
[2] http://lab.llvm.org:8011/builders/clang-cmake-aarch64-full

clang/include/clang/AST/Type.h

public:
   friend Qualifiers operator-(Qualifiers L, Qualifiers R) {
     L -= R;
     return L;
   }
 
   std::string getAsString() const;
   std::string getAsString(const PrintingPolicy &Policy) const;
 
+  static std::string getAddrSpaceAsString(LangAS AS);
+
   bool isEmptyWhenPrinted(const PrintingPolicy &Policy) const;
   void print(raw_ostream &OS, const PrintingPolicy &Policy,
              bool appendSpaceIfNonEmpty = false) const;
 
   void Profile(llvm::FoldingSetNodeID &ID) const {
     ID.AddInteger(Mask);
   }
 

 inline const Type *Type::getPointeeOrArrayElementType() const {
   const Type *type = this;
   if (type->isAnyPointerType())
     return type->getPointeeType().getTypePtr();
   else if (type->isArrayType())
     return type->getBaseElementTypeUnsafe();
   return type;
 }
+/// Insertion operator for diagnostics. This allows sending address spaces into
+/// a diagnostic with <<.
+inline const DiagnosticBuilder &operator<<(const DiagnosticBuilder &DB,
+                                           LangAS AS) {
+  DB.AddTaggedVal(static_cast<std::underlying_type_t<LangAS>>(AS),
+                  DiagnosticsEngine::ArgumentKind::ak_addrspace);
+  return DB;
+}
+
+/// Insertion operator for partial diagnostics. This allows sending adress
+/// spaces into a diagnostic with <<.
+inline const PartialDiagnostic &operator<<(const PartialDiagnostic &PD,
+                                           LangAS AS) {
+  PD.AddTaggedVal(static_cast<std::underlying_type_t<LangAS>>(AS),
+                  DiagnosticsEngine::ArgumentKind::ak_addrspace);
+  return PD;
+}
 
 /// Insertion operator for diagnostics. This allows sending Qualifiers into a
 /// diagnostic with <<.
 inline const DiagnosticBuilder &operator<<(const DiagnosticBuilder &DB,
                                            Qualifiers Q) {
   DB.AddTaggedVal(Q.getAsOpaqueValue(),
                   DiagnosticsEngine::ArgumentKind::ak_qual);
   return DB;

clang/include/clang/Basic/Diagnostic.h

enum ArgumentKind {
     ak_uint,
 
     /// enum TokenKind : unsigned
     ak_tokenkind,
 
     /// IdentifierInfo
     ak_identifierinfo,
 
+    /// address space
+    ak_addrspace,
+
     /// Qualifiers
     ak_qual,
 
     /// QualType
     ak_qualtype,
 
     /// DeclarationName
     ak_declarationname,

clang/include/clang/Basic/DiagnosticSemaKinds.td

def note_ovl_candidate_bad_arc_conv : Note<
     "%diff{of type $ to $|type to parameter type}3,4 for "
     "%select{%ordinal6 argument|object argument}5 under ARC">;
 def note_ovl_candidate_bad_lvalue : Note<
     "candidate %sub{select_ovl_candidate_kind}0,1,2 not viable: "
     "expects an l-value for "
     "%select{%ordinal4 argument|object argument}3">;
 def note_ovl_candidate_bad_addrspace : Note<
     "candidate %sub{select_ovl_candidate_kind}0,1,2 not viable: "
-    "address space mismatch in %select{%ordinal6|'this'}5 argument (%3), "
-    "parameter type must be %4">;
+    "cannot %select{pass pointer to|bind reference in}5 %3 "
+    "%select{as a pointer to|to object in}5 %4 in %ordinal6 "
+    "argument">;
+def note_ovl_candidate_bad_addrspace_this : Note<
+    "candidate %sub{select_ovl_candidate_kind}0,1,2 not viable: "
+    "'this' object is in %3, but method expects object in %4">;
 def note_ovl_candidate_bad_gc : Note<
     "candidate %sub{select_ovl_candidate_kind}0,1,2 not viable: "
     "%select{%ordinal7|'this'}6 argument (%3) has %select{no|__weak|__strong}4 "
     "ownership, but parameter has %select{no|__weak|__strong}5 ownership">;
 def note_ovl_candidate_bad_ownership : Note<
     "candidate %sub{select_ovl_candidate_kind}0,1,2 not viable: "
     "%select{%ordinal7|'this'}6 argument (%3) has "
     "%select{no|__unsafe_unretained|__strong|__weak|__autoreleasing}4 ownership,"
     " but parameter has %select{no|__unsafe_unretained|__strong|__weak|"

clang/lib/AST/ASTDiagnostic.cpp

void clang::FormatASTNodeDiagnosticArgument(
   ASTContext &Context = *static_cast<ASTContext*>(Cookie);
 
   size_t OldEnd = Output.size();
   llvm::raw_svector_ostream OS(Output);
   bool NeedQuotes = true;
 
   switch (Kind) {
     default: llvm_unreachable("unknown ArgumentKind");
+    case DiagnosticsEngine::ak_addrspace: {
+      assert(Modifier.empty() && Argument.empty() &&
+             "Invalid modifier for Qualfiers argument");
+
+      auto S = Qualifiers::getAddrSpaceAsString(static_cast<LangAS>(Val));
+      if (S.empty()) {
+        OS << (Context.getLangOpts().OpenCL ? "default" : "generic");
+        OS << " address space";
+      } else {
+        OS << "address space";
+        OS << " '" << S << "'";
+      }
+      NeedQuotes = false;
+      break;
+    }
     case DiagnosticsEngine::ak_qual: {
       assert(Modifier.empty() && Argument.empty() &&
              "Invalid modifier for Qualfiers argument");
 
       Qualifiers Q(Qualifiers::fromOpaqueValue(Val));
       auto S = Q.getAsString();
       if (S.empty()) {
         OS << "unqualified";
         NeedQuotes = false;
       } else {
-        OS << Q.getAsString();
+        OS << S;
       }
       break;
     }
     case DiagnosticsEngine::ak_qualtype_pair: {
       TemplateDiffTypes &TDT = *reinterpret_cast<TemplateDiffTypes*>(Val);
       QualType FromType =
           QualType::getFromOpaquePtr(reinterpret_cast<void*>(TDT.FromType));
       QualType ToType =

clang/lib/AST/TypePrinter.cpp

bool Qualifiers::isEmptyWhenPrinted(const PrintingPolicy &Policy) const {
 
   if (Qualifiers::ObjCLifetime lifetime = getObjCLifetime())
     if (!(lifetime == Qualifiers::OCL_Strong && Policy.SuppressStrongLifetime))
       return false;
 
   return true;
 }
 
+std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
+  switch (AS) {
+  case LangAS::Default:
+    return "";
+  case LangAS::opencl_global:
+    return "__global";
+  case LangAS::opencl_local:
+    return "__local";
+  case LangAS::opencl_private:
+    return "";
+  case LangAS::opencl_constant:
+    return "__constant";
+  case LangAS::opencl_generic:
+    return "__generic";
+  case LangAS::cuda_device:
+    return "__device__";
+  case LangAS::cuda_constant:
+    return "__constant__";
+  case LangAS::cuda_shared:
+    return "__shared__";
+  default:
+    return std::to_string(toTargetAddressSpace(AS));
+  }
+}
+
 // Appends qualifiers to the given string, separated by spaces.  Will
 // prefix a space if the string is non-empty.  Will not append a final
 // space.
 void Qualifiers::print(raw_ostream &OS, const PrintingPolicy& Policy,
                        bool appendSpaceIfNonEmpty) const {
   bool addSpace = false;
 
   unsigned quals = getCVRQualifiers();
   if (quals) {
     AppendTypeQualList(OS, quals, Policy.Restrict);
     addSpace = true;
   }
   if (hasUnaligned()) {
     if (addSpace)
       OS << ' ';
     OS << "__unaligned";
     addSpace = true;
   }
-  LangAS addrspace = getAddressSpace();
-  if (addrspace != LangAS::Default) {
-    if (addrspace != LangAS::opencl_private) {
-      if (addSpace)
-        OS << ' ';
-      addSpace = true;
-      switch (addrspace) {
-      case LangAS::opencl_global:
-        OS << "__global";
-        break;
-      case LangAS::opencl_local:
-        OS << "__local";
-        break;
-      case LangAS::opencl_private:
-        break;
-      case LangAS::opencl_constant:
-        OS << "__constant";
-        break;
-      case LangAS::opencl_generic:
-        OS << "__generic";
-        break;
-      case LangAS::cuda_device:
-        OS << "__device__";
-        break;
-      case LangAS::cuda_constant:
-        OS << "__constant__";
-        break;
-      case LangAS::cuda_shared:
-        OS << "__shared__";
-        break;
-      default:
-        OS << "__attribute__((address_space(";
-        OS << toTargetAddressSpace(addrspace);
-        OS << ")))";
-      }
-    }
+  auto ASStr = getAddrSpaceAsString(getAddressSpace());
+  if (!ASStr.empty()) {
+    // Wrap target address space into an attribute syntax
+    if (isTargetAddressSpace(getAddressSpace()))
+      OS << "__attribute__((address_space(" << ASStr << ")))";
+    else
+      OS << ASStr;
   }
+
   if (Qualifiers::GC gc = getObjCGCAttr()) {
     if (addSpace)
       OS << ' ';
     addSpace = true;
     if (gc == Qualifiers::Weak)
       OS << "__weak";
     else
       OS << "__strong";

clang/lib/Basic/Diagnostic.cpp

case DiagnosticsEngine::ak_identifierinfo: {
         const char *S = "(null)";
         OutStr.append(S, S + strlen(S));
         continue;
       }
 
       llvm::raw_svector_ostream(OutStr) << '\'' << II->getName() << '\'';
       break;
     }
+    case DiagnosticsEngine::ak_addrspace:
     case DiagnosticsEngine::ak_qual:
     case DiagnosticsEngine::ak_qualtype:
     case DiagnosticsEngine::ak_declarationname:
     case DiagnosticsEngine::ak_nameddecl:
     case DiagnosticsEngine::ak_nestednamespec:
     case DiagnosticsEngine::ak_declcontext:
     case DiagnosticsEngine::ak_attr:
       getDiags()->ConvertArgToString(Kind, getRawArg(ArgNo),

clang/lib/Sema/SemaOverload.cpp

static void DiagnoseBadConversion(Sema &S, OverloadCandidate *Cand,
   }
 
   if (CToTy.getUnqualifiedType() == CFromTy.getUnqualifiedType() &&
       !CToTy.isAtLeastAsQualifiedAs(CFromTy)) {
     Qualifiers FromQs = CFromTy.getQualifiers();
     Qualifiers ToQs = CToTy.getQualifiers();
 
     if (FromQs.getAddressSpace() != ToQs.getAddressSpace()) {
+      if (isObjectArgument)
+        S.Diag(Fn->getLocation(), diag::note_ovl_candidate_bad_addrspace_this)
+            << (unsigned)FnKindPair.first << (unsigned)FnKindPair.second
+            << FnDesc << (FromExpr ? FromExpr->getSourceRange() : SourceRange())
+            << FromQs.getAddressSpace() << ToQs.getAddressSpace();
+      else
         S.Diag(Fn->getLocation(), diag::note_ovl_candidate_bad_addrspace)
-          << (unsigned)FnKindPair.first << (unsigned)FnKindPair.second << FnDesc
-          << (FromExpr ? FromExpr->getSourceRange() : SourceRange()) << FromTy
-          << ToTy << (unsigned)isObjectArgument << I + 1;
+            << (unsigned)FnKindPair.first << (unsigned)FnKindPair.second
+            << FnDesc << (FromExpr ? FromExpr->getSourceRange() : SourceRange())
+            << FromQs.getAddressSpace() << ToQs.getAddressSpace()
+            << ToTy->isReferenceType() << I + 1;
       MaybeEmitInheritedConstructorNote(S, Cand->FoundDecl);
       return;
     }
 
     if (FromQs.getObjCLifetime() != ToQs.getObjCLifetime()) {
       S.Diag(Fn->getLocation(), diag::note_ovl_candidate_bad_ownership)
           << (unsigned)FnKindPair.first << (unsigned)FnKindPair.second << FnDesc
           << (FromExpr ? FromExpr->getSourceRange() : SourceRange()) << FromTy

clang/test/SemaCXX/address-space-references.cpp

 // RUN: %clang_cc1 -fsyntax-only -verify %s
 
 typedef int __attribute__((address_space(1))) int_1;
 typedef int __attribute__((address_space(2))) int_2;
 
-void f0(int_1 &); // expected-note{{candidate function not viable: address space mismatch in 1st argument ('int'), parameter type must be 'int_1 &' (aka '__attribute__((address_space(1))) int &')}} \
-// expected-note{{candidate function not viable: address space mismatch in 1st argument ('int_2' (aka '__attribute__((address_space(2))) int')), parameter type must be 'int_1 &' (aka '__attribute__((address_space(1))) int &')}}
-void f0(const int_1 &); // expected-note{{candidate function not viable: address space mismatch in 1st argument ('int'), parameter type must be 'const int_1 &' (aka 'const __attribute__((address_space(1))) int &')}} \
-// expected-note{{candidate function not viable: address space mismatch in 1st argument ('int_2' (aka '__attribute__((address_space(2))) int')), parameter type must be 'const int_1 &' (aka 'const __attribute__((address_space(1))) int &')}}
+void f0(int_1 &);       // expected-note{{candidate function not viable: cannot bind reference in generic address space to object in address space '1' in 1st argument}} \
+// expected-note{{candidate function not viable: cannot bind reference in address space '2' to object in address space '1' in 1st argument}}
+void f0(const int_1 &); // expected-note{{candidate function not viable: cannot bind reference in generic address space to object in address space '1' in 1st argument}} \
+// expected-note{{candidate function not viable: cannot bind reference in address space '2' to object in address space '1' in 1st argument}}
 
 void test_f0() {
   int i;
   static int_1 i1;
   static int_2 i2;
 
   f0(i); // expected-error{{no matching function for call to 'f0'}}
   f0(i1);
   f0(i2); // expected-error{{no matching function for call to 'f0'}}
 }

clang/test/SemaOpenCL/address-spaces-conversions-cl2.0.cl

 #define AS_INCOMP __global
 #endif
 
 void f_glob(__global int *arg_glob) {}
 #ifndef GLOBAL
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-3{{passing argument to parameter 'arg_glob' here}}
 #else
-// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{generic|constant}} int *'), parameter type must be '__global int *'}}
+// expected-note-re@-5{{candidate function not viable: cannot pass pointer to address space '__{{generic|constant}}' as a pointer to address space '__global' in 1st argument}}
 #endif
 #endif
 
 void f_loc(__local int *arg_loc) {}
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-2{{passing argument to parameter 'arg_loc' here}}
 #else
-// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be '__local int *'}}
+// expected-note-re@-4{{candidate function not viable: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to address space '__local' in 1st argument}}
 #endif
 
 void f_const(__constant int *arg_const) {}
 #ifndef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-3{{passing argument to parameter 'arg_const' here}}
 #else
-// expected-note-re@-5{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic}} int *'), parameter type must be '__constant int *'}}
+// expected-note-re@-5{{candidate function not viable: cannot pass pointer to address space '__{{global|generic}}' as a pointer to address space '__constant' in 1st argument}}
 #endif
 #endif
 
 void f_priv(__private int *arg_priv) {}
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-2{{passing argument to parameter 'arg_priv' here}}
 #else
-// expected-note-re@-4{{candidate function not viable: address space mismatch in 1st argument ('__{{global|generic|constant}} int *'), parameter type must be 'int *'}}
+// expected-note-re@-4{{candidate function not viable: cannot pass pointer to address space '__{{global|generic|constant}}' as a pointer to default address space in 1st argument}}
 #endif
 
 void f_gen(__generic int *arg_gen) {}
 #ifdef CONSTANT
 #if !__OPENCL_CPP_VERSION__
 // expected-note@-3{{passing argument to parameter 'arg_gen' here}}
 #else
-// expected-note@-5{{candidate function not viable: address space mismatch in 1st argument ('__constant int *'), parameter type must be '__generic int *'}}
+// expected-note@-5{{candidate function not viable: cannot pass pointer to address space '__constant' as a pointer to address space '__generic' in 1st argument}}
 #endif
 #endif
 
 void test_conversion(__global int *arg_glob, __local int *arg_loc,
                      __constant int *arg_const, __private int *arg_priv,
                      __generic int *arg_gen) {
 
   AS int *var_init1 = arg_glob;

clang/test/SemaOpenCLCXX/address-space-lambda.cl

 //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
 
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic'
 auto glambda = [](auto a) { return a; };
 
 __kernel void test() {
   int i;
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   auto  llambda = [&]() {i++;};
   llambda();
   glambda(1);
   // Test lambda with default parameters
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   [&] {i++;} ();
-  __constant auto err = [&]() {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}})'}}
+  __constant auto err = [&]() {}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}}
   err();                          //expected-error-re{{no matching function for call to object of type '__constant (lambda at {{.*}})'}}
   // FIXME: There is very limited addr space functionality
   // we can test when taking lambda type from the object.
   // The limitation is due to addr spaces being added to all
   // objects in OpenCL. Once we add metaprogramming utility
   // for removing address spaces from a type we can enhance
   // testing here.
   (*(__constant decltype(llambda) *)nullptr)(); //expected-error{{multiple address spaces specified for type}}
   (*(decltype(llambda) *)nullptr)();
 }
 
 __kernel void test_qual() {
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const'
   auto priv1 = []() __private {};
   priv1();
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   auto priv2 = []() __generic {};
   priv2();
-  auto priv3 = []() __global {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const __global (lambda at {{.*}})'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   priv3(); //expected-error{{no matching function for call to object of type}}
 
-  __constant auto const1 = []() __private{}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const (lambda at {{.*}}'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in default address space}} //expected-note{{conversion candidate of type 'void (*)()'}}
   const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
-  __constant auto const2 = []() __generic{}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}}'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  __constant auto const2 = []() __generic{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __constant'
   __constant auto const3 = []() __constant{};
   const3();
 
-  [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const __global (lambda at {{.*}})'}}
-  [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const (lambda at {{.*}})'}}
+  [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}}
+  [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in default address space}}
 
   [&] __private {} (); //expected-error{{lambda requires '()' before attribute specifier}} expected-error{{expected body of lambda expression}}
 
   [&] () mutable __private {} ();
   [&] () __private mutable {} (); //expected-error{{expected body of lambda expression}}
 }
 

clang/test/SemaOpenCLCXX/address-space-of-this-class-scope.cl

 //RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify
 
 struct C {
   auto fGlob() __global -> decltype(this);
   auto fGen() -> decltype(this);
   auto fErr() __global __local -> decltype(this); //expected-error{{multiple address spaces specified for type}}
 };
 
 void bar(__local C*);
-// expected-note@-1{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka '__global C *')), parameter type must be '__local C *'}}
-// expected-note@-2{{candidate function not viable: address space mismatch in 1st argument ('decltype(this)' (aka 'C *')), parameter type must be '__local C *'}}
+// expected-note@-1{{candidate function not viable: cannot pass pointer to address space '__global' as a pointer to address space '__local' in 1st argument}}
+// expected-note@-2{{candidate function not viable: cannot pass pointer to default address space as a pointer to address space '__local' in 1st argument}}
 
 __global C Glob;
 void foo(){
 bar(Glob.fGlob()); // expected-error{{no matching function for call to 'bar'}}
 // FIXME: AS of 'this' below should be correctly deduced to generic
 bar(Glob.fGen()); // expected-error{{no matching function for call to 'bar'}}
 }