Allow implicit conversions between incompatible pointer types in overload resolution in C.

[george.burgess.iv]

Summary

In C, we allow pointer conversions like:

void foo(int *a);
long b;
void runFoo() { foo(&b); }

...But not when overloading:

int bar(char *a) __attribute__((overloadable));
int bar(struct{}) __attribute__((overloadable));
unsigned char d;
void runBar() { bar(&d); } // error: no matching function for call to 'bar'

This patch makes it so bar(&d) isn't an error.

A few notes:

• We'll still emit a warning about the conversion, assuming the user hasn't opted out of said warnings.
• Every such conversion is ranked as the worst possible thing. So, if there's any ambiguity at all (e.g. if bar had a void(signed char*) overload, as well), we'll still emit an error.
• We consider conversions that drop qualifiers to be equally as bad as converting between incompatible pointer types. I'm happy to make conversions that only drop qualifiers preferred over incompatible pointer types, if that would be better.

[rsmith]

lib/Sema/SemaOverload.cpp

1813–1815

Yuck, this violates our invariants: First should only ever be some kind of lvalue conversion (it should in this case be an ICK_Array_To_Pointer or ICK_Function_To_Pointer if the argument was an array or function, respectively, and otherwise ICK_Identity). Third should only ever be ICK_Identity or ICK_Qualification. It seems fine to always set Third to ICK_Identity and model the qualification change as part of the second "standard" conversion; this should still compare worse than any other sequence by the rank test.

Now I've noticed this, I'd like to see it fixed, but since this is pre-existing I don't mind whether we fix this as part of this change or separately.

5931

You should check Second here rather than First.

[george.burgess.iv]

lib/Sema/SemaOverload.cpp

1813–1815

I'll fix this in a follow-up commit; thanks for pointing it out!

[jvesely]

This change breaks OpenCL(libclc). CLC expects signed/unsigned overloads to be non-ambiguous.
For example:
atomic_max(global int *p, int val);
atomic_max(global unsigned int *p, unsigned int val);

need to work and be unambiguous (they are implemented using different instructions different instructions).

[george.burgess.iv]

Thanks for the heads-up!

I'm assuming that below is an example of code that this patch broke?

void foo(int *a, int) __attribute__((overloadable));
void foo(unsigned int *a, unsigned int) __attribute__((overloadable));

void callFoo() {
  unsigned int i;
  foo(&i, 0u); // fine.
  foo(&i, 0); // now-ambiguous overload.
}

[jvesely]

In D24113#535255, @george.burgess.iv wrote:

Thanks for the heads-up!

I'm assuming that below is an example of code that this patch broke?

void foo(int *a, int) __attribute__((overloadable));
void foo(unsigned int *a, unsigned int) __attribute__((overloadable));

void callFoo() {
  unsigned int i;
  foo(&i, 0u); // fine.
  foo(&i, 0); // now-ambiguous overload.
}

yes, you're correct. (I'm pretty sure I tried this before posting, I must have done something wrong...) I'll send libclc patch. The specs don't really talk about overloads so I think the behaviour is OK. Although I think that users will expect atomic_add(unsigned int *, 5) to work.without error.

thanks

[george.burgess.iv]

Although I think that users will expect atomic_add(unsigned int *, 5) to work.without error

Totally agree; I was just making sure that I understood the problem. :)

I think, with a small tweak in semantics, we can make everything work without needing to patch existing code. Let me play around a little.

[yaxunl]

This seems to cause regression for the following test:

__kernel void test_atomic_fn(volatile __global uint *destMemory, __global uint *oldValues)
{
    int  tid = get_global_id(0);

oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );
   atom_add( &destMemory[0], tid + 3 );

}

/tmp/AMD_17652_49/t_17652_51.cl:6:19: error: call to 'atom_add' is ambiguous
        oldValues[tid] = atom_add( &destMemory[0], tid + 3 );
                         ^~~~~~~~
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14336:12: note: candidate function
int __ovld atom_add(volatile __global int *p, int val);
           ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14337:21: note: candidate function
unsigned int __ovld atom_add(volatile __global unsigned int *p, unsigned int val);
                    ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14345:13: note: candidate function
long __ovld atom_add(volatile __global long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14346:22: note: candidate function
unsigned long __ovld atom_add(volatile __global unsigned long *p, unsigned long val);
                     ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14347:13: note: candidate function
long __ovld atom_add(volatile __local long *p, long val);
            ^
/srv/opencl_rocm/drivers/opencl/library/amdgcn/headers/build/lnx64a/B_dbg/<stdin>:14348:22: note: candidate function
unsigned long __ovld atom_add(volatile __local unsigned long *p, unsigned long val);
                     ^
/tmp/AMD_17652_49/t_17652_51.cl:7:2: error: call to 'atom_add' is ambiguous
        atom_add( &destMemory[0], tid + 3 );
        ^~~~~~~~

Is this intended?

Thanks.

include/clang/Basic/AttrDocs.td

 
 * A conversion from a pointer of type ``T*`` to a pointer of type ``U*`` is
   considered a pointer conversion (with conversion rank) if ``T`` and ``U`` are
   compatible types.
 
 * A conversion from type ``T`` to a value of type ``U`` is permitted if ``T``
   and ``U`` are compatible types.  This conversion is given "conversion" rank.
 
+* A conversion from a pointer of type ``T*`` to a pointer of type ``U*``, where
+  ``T`` and ``U`` are incompatible, is allowed, but is ranked below all other
+  types of conversions. Please note: ``U`` lacking qualifiers that are present
+  on ``T`` is sufficient for ``T`` and ``U`` to be incompatible.
+
 The declaration of ``overloadable`` functions is restricted to function
 declarations and definitions.  Most importantly, if any function with a given
 name is given the ``overloadable`` attribute, then all function declarations
 and definitions with that name (and in that scope) must have the
 ``overloadable`` attribute.  This rule even applies to redeclarations of
 functions whose original declaration had the ``overloadable`` attribute, e.g.,
 
 .. code-block:: c

include/clang/Sema/Overload.h

enum ImplicitConversionKind {
     ICK_Vector_Conversion,     ///< Vector conversions
     ICK_Vector_Splat,          ///< A vector splat from an arithmetic type
     ICK_Complex_Real,          ///< Complex-real conversions (C99 6.3.1.7)
     ICK_Block_Pointer_Conversion,    ///< Block Pointer conversions 
     ICK_TransparentUnionConversion, ///< Transparent Union Conversions
     ICK_Writeback_Conversion,  ///< Objective-C ARC writeback conversion
     ICK_Zero_Event_Conversion, ///< Zero constant to event (OpenCL1.2 6.12.10)
     ICK_C_Only_Conversion,     ///< Conversions allowed in C, but not C++
+    ICK_Incompatible_Pointer_Conversion, ///< C-only conversion between pointers
+                                         ///  with incompatible types
     ICK_Num_Conversion_Kinds,  ///< The number of conversion kinds
   };
 
   /// ImplicitConversionRank - The rank of an implicit conversion
   /// kind. The enumerator values match with Table 9 of (C++
   /// 13.3.3.1.1) and are listed such that better conversion ranks
   /// have smaller values.
   enum ImplicitConversionRank {
     ICR_Exact_Match = 0,         ///< Exact Match
     ICR_Promotion,               ///< Promotion
     ICR_Conversion,              ///< Conversion
     ICR_Complex_Real_Conversion, ///< Complex <-> Real conversion
     ICR_Writeback_Conversion,    ///< ObjC ARC writeback conversion
-    ICR_C_Conversion             ///< Conversion only allowed in the C standard.
+    ICR_C_Conversion,            ///< Conversion only allowed in the C standard.
                                  ///  (e.g. void* to char*)
+    ICR_C_Conversion_Extension   ///< Conversion not allowed by the C standard,
+                                 ///  but that we accept as an extension anyway.
   };
 
   ImplicitConversionRank GetConversionRank(ImplicitConversionKind Kind);
 
   /// NarrowingKind - The kind of narrowing conversion being performed by a
   /// standard conversion sequence according to C++11 [dcl.init.list]p7.
   enum NarrowingKind {
     /// Not a narrowing conversion.

lib/Sema/SemaExprCXX.cpp

case ICK_Zero_Event_Conversion:
     break;
 
   case ICK_Lvalue_To_Rvalue:
   case ICK_Array_To_Pointer:
   case ICK_Function_To_Pointer:
   case ICK_Qualification:
   case ICK_Num_Conversion_Kinds:
   case ICK_C_Only_Conversion:
+  case ICK_Incompatible_Pointer_Conversion:
     llvm_unreachable("Improper second standard conversion");
   }
 
   switch (SCS.Third) {
   case ICK_Identity:
     // Nothing to do.
     break;
 

lib/Sema/SemaOverload.cpp

static const ImplicitConversionRank
     ICR_Conversion,
     ICR_Complex_Real_Conversion,
     ICR_Conversion,
     ICR_Conversion,
     ICR_Writeback_Conversion,
     ICR_Exact_Match, // NOTE(gbiv): This may not be completely right --
                      // it was omitted by the patch that added
                      // ICK_Zero_Event_Conversion
-    ICR_C_Conversion
+    ICR_C_Conversion,
+    ICR_C_Conversion_Extension
   };
   return Rank[(int)Kind];
 }
 
 /// GetImplicitConversionName - Return the name of this kind of
 /// implicit conversion.
 static const char* GetImplicitConversionName(ImplicitConversionKind Kind) {
   static const char* const Name[(int)ICK_Num_Conversion_Kinds] = {
static const char* const Name[(int)ICK_Num_Conversion_Kinds] = {
     "Derived-to-base conversion",
     "Vector conversion",
     "Vector splat",
     "Complex-real conversion",
     "Block Pointer conversion",
     "Transparent Union Conversion",
     "Writeback conversion",
     "OpenCL Zero Event Conversion",
-    "C specific type conversion"
+    "C specific type conversion",
+    "Incompatible pointer conversion"
   };
   return Name[Kind];
 }
 
 /// StandardConversionSequence - Set the standard conversion
 /// sequence to the identity conversion.
 void StandardConversionSequence::setAsIdentityConversion() {
   First = ICK_Identity;
if (CanonFrom == CanonTo)
     return true;
 
   // If we have not converted the argument type to the parameter type,
   // this is a bad conversion sequence, unless we're resolving an overload in C.
   if (S.getLangOpts().CPlusPlus || !InOverloadResolution)
     return false;
 
   ExprResult ER = ExprResult{From};
-  auto Conv = S.CheckSingleAssignmentConstraints(ToType, ER,
-                                                 /*Diagnose=*/false,
-                                                 /*DiagnoseCFAudited=*/false,
-                                                 /*ConvertRHS=*/false);
-  if (Conv != Sema::Compatible)
+  Sema::AssignConvertType Conv =
+      S.CheckSingleAssignmentConstraints(ToType, ER,
+  ImplicitConversionKind ImplicitConv;
+  switch (Conv) {
+  case Sema::Compatible:
+    ImplicitConv = ICK_C_Only_Conversion;
+    break;
+  // For our purposes, discarding qualifiers is just as bad as using an
+  // incompatible pointer. Note that an IncompatiblePointer conversion can drop
+  // qualifiers, as well.
+  case Sema::CompatiblePointerDiscardsQualifiers:
+  case Sema::IncompatiblePointer:
+  case Sema::IncompatiblePointerSign:
+    ImplicitConv = ICK_Incompatible_Pointer_Conversion;
+    break;
+  default:
     return false;
+  }
 
   SCS.setAllToTypes(ToType);
   // We need to set all three because we want this conversion to rank terribly,
   // and we don't know what conversions it may overlap with.
-  SCS.First = ICK_C_Only_Conversion;
-  SCS.Second = ICK_C_Only_Conversion;
-  SCS.Third = ICK_C_Only_Conversion;
+  SCS.First = ImplicitConv;
+  SCS.Second = ImplicitConv;
+  SCS.Third = ImplicitConv;
   return true;
 }
-  
 static bool
 IsTransparentUnionStandardConversion(Sema &S, Expr* From, 
                                      QualType &ToType,
                                      bool InOverloadResolution,
                                      StandardConversionSequence &SCS,
                                      bool CStyle) {
     
   const RecordType *UT = ToType->getAsUnionType();
static bool CheckConvertedConstantConversions(Sema &S,
   case ICK_Vector_Conversion:
   case ICK_Vector_Splat:
   case ICK_Complex_Real:
   case ICK_Block_Pointer_Conversion:
   case ICK_TransparentUnionConversion:
   case ICK_Writeback_Conversion:
   case ICK_Zero_Event_Conversion:
   case ICK_C_Only_Conversion:
+  case ICK_Incompatible_Pointer_Conversion:
     return false;
 
   case ICK_Lvalue_To_Rvalue:
   case ICK_Array_To_Pointer:
   case ICK_Function_To_Pointer:
     llvm_unreachable("found a first conversion kind in Second");
 
   case ICK_Qualification:
for (unsigned i = 0; i < NumNamedArgs; i++) {
 
       ImplicitConversionSequence ConversionState
         = TryCopyInitialization(*this, argExpr, param->getType(),
                                 /*SuppressUserConversions*/false,
                                 /*InOverloadResolution=*/true,
                                 /*AllowObjCWritebackConversion=*/
                                 getLangOpts().ObjCAutoRefCount,
                                 /*AllowExplicit*/false);
-        if (ConversionState.isBad()) {
-          Match = false;
-          break;
-        }
+      // This function looks for a reasonably-exact match, so we consider
+      // incompatible pointer conversions to be a failure here.
+      if (ConversionState.isBad() ||
+          (ConversionState.isStandard() &&
+           ConversionState.Standard.First ==
+               ICK_Incompatible_Pointer_Conversion)) {
     }
     // Promote additional arguments to variadic methods.
     if (Match && Method->isVariadic()) {
       for (unsigned i = NumNamedArgs, e = Args.size(); i < e; ++i) {
         if (Args[i]->isTypeDependent()) {
           Match = false;
           break;
         }

test/CodeGen/builtins-systemz-zvector-error.c

void test_core(void) {
   vec_scatter_element(vd, vul, ptrd, -1);    // expected-error {{no matching function}}
                                              // expected-note@vecintrin.h:* 6 {{candidate function not viable}}
                                              // expected-note@vecintrin.h:* 1 {{must be a constant integer from 0 to 1}}
   vec_scatter_element(vd, vul, ptrd, 2);     // expected-error {{no matching function}}
                                              // expected-note@vecintrin.h:* 6 {{candidate function not viable}}
                                              // expected-note@vecintrin.h:* 1 {{must be a constant integer from 0 to 1}}
 
   vsc = vec_load_bndry(cptrsc, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vsc = vec_load_bndry(cptrsc, 200);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vsc = vec_load_bndry(cptrsc, 32);    // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vsc = vec_load_bndry(cptrsc, 8192);  // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vuc = vec_load_bndry(cptruc, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vss = vec_load_bndry(cptrss, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vus = vec_load_bndry(cptrus, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vsi = vec_load_bndry(cptrsi, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vui = vec_load_bndry(cptrui, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vsl = vec_load_bndry(cptrsl, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
   vul = vec_load_bndry(cptrul, idx);   // expected-error {{no matching function}}
-                                       // expected-note@vecintrin.h:* 8 {{candidate function not viable}}
-                                       // expected-note@vecintrin.h:* 1 {{must be a constant power of 2 from 64 to 4096}}
+                                       // expected-note@vecintrin.h:* 9 {{must be a constant power of 2 from 64 to 4096}}
 
   vuc = vec_genmask(idx);  // expected-error {{no matching function}}
                            // expected-note@vecintrin.h:* {{must be a constant integer}}
 
   vuc = vec_genmasks_8(0, idx);    // expected-error {{no matching function}}
                                    // expected-note@vecintrin.h:* {{must be a constant integer}}
   vuc = vec_genmasks_8(idx, 0);    // expected-error {{no matching function}}
                                    // expected-note@vecintrin.h:* {{must be a constant integer}}

test/CodeGen/overloadable.c

void foo() {
 
   // CHECK: store void (i8*)* @_Z13addrof_singlePc
   void (*p3)(char *) = &addrof_single;
   // CHECK: @_Z13addrof_singlePc
   void (*p4)(int *) = &addrof_single;
   // CHECK: @_Z13addrof_singlePc
   void *vp3 = &addrof_single;
 }
+
+
+void ovl_bar(char *) __attribute__((overloadable));
+void ovl_bar(int) __attribute__((overloadable));
+
+// CHECK-LABEL: define void @bar
+void bar() {
+  char charbuf[1];
+  unsigned char ucharbuf[1];
+
+  // CHECK: call void @_Z7ovl_barPc
+  ovl_bar(charbuf);
+  // CHECK: call void @_Z7ovl_barPc
+  ovl_bar(ucharbuf);
+}

test/Sema/overloadable.c

 int *accept_funcptr(int (*)()) __attribute__((overloadable)); //         \
   // expected-note{{candidate function}}
 float *accept_funcptr(int (*)(int, double)) __attribute__((overloadable)); //  \
   // expected-note{{candidate function}}
 
 void test_funcptr(int (*f1)(int, double),
                   int (*f2)(int, float)) {
   float *fp = accept_funcptr(f1);
-  accept_funcptr(f2); // expected-error{{no matching function for call to 'accept_funcptr'}}
+  accept_funcptr(f2); // expected-error{{call to 'accept_funcptr' is ambiguous}}
 }
 
 struct X { int x; float y; };
 struct Y { int x; float y; };
 int* accept_struct(struct X x) __attribute__((__overloadable__));
 float* accept_struct(struct Y y) __attribute__((overloadable));
 
 void test_struct(struct X x, struct Y y) {
void fn_type_conversions() {
   void disabled(char *c) __attribute__((overloadable, enable_if(1, "The function name lies.")));
   // To be clear, these should all point to the last overload of 'disabled'
   void (*dptr1)(char *c) = &disabled;
   void (*dptr2)(void *c) = &disabled; // expected-warning{{incompatible pointer types initializing 'void (*)(void *)' with an expression of type '<overloaded function type>'}} expected-note@115{{candidate function made ineligible by enable_if}} expected-note@116{{candidate function made ineligible by enable_if}} expected-note@117{{candidate function has type mismatch at 1st parameter (expected 'void *' but has 'char *')}}
   void (*dptr3)(int *c) = &disabled; // expected-warning{{incompatible pointer types initializing 'void (*)(int *)' with an expression of type '<overloaded function type>'}} expected-note@115{{candidate function made ineligible by enable_if}} expected-note@116{{candidate function made ineligible by enable_if}} expected-note@117{{candidate function has type mismatch at 1st parameter (expected 'int *' but has 'char *')}}
 
   void *specific_disabled = &disabled;
 }
+
+void incompatible_pointer_type_conversions() {
+  char charbuf[1];
+  unsigned char ucharbuf[1];
+  int intbuf[1];
+
+  void foo(char *c) __attribute__((overloadable));
+  void foo(short *c) __attribute__((overloadable));
+  foo(charbuf);
+  foo(ucharbuf); // expected-error{{call to 'foo' is ambiguous}} expected-note@131{{candidate function}} expected-note@132{{candidate function}}
+  foo(intbuf); // expected-error{{call to 'foo' is ambiguous}} expected-note@131{{candidate function}} expected-note@132{{candidate function}}
+
+  void bar(unsigned char *c) __attribute__((overloadable));
+  void bar(signed char *c) __attribute__((overloadable));
+  bar(charbuf); // expected-error{{call to 'bar' is ambiguous}} expected-note@137{{candidate function}} expected-note@138{{candidate function}}
+  bar(ucharbuf);
+  bar(intbuf); // expected-error{{call to 'bar' is ambiguous}} expected-note@137{{candidate function}} expected-note@138{{candidate function}}
+}
+
+void dropping_qualifiers_is_incompatible() {
+  const char ccharbuf[1];
+  volatile char vcharbuf[1];
+
+  void foo(char *c) __attribute__((overloadable));
+  void foo(const volatile unsigned char *c) __attribute__((overloadable));
+
+  foo(ccharbuf); // expected-error{{call to 'foo' is ambiguous}} expected-note@148{{candidate function}} expected-note@149{{candidate function}}
+  foo(vcharbuf); // expected-error{{call to 'foo' is ambiguous}} expected-note@148{{candidate function}} expected-note@149{{candidate function}}
+}

test/Sema/pass-object-size.c

void FunctionPtrs() {
   TakeFn(NotOverloaded); //expected-error{{cannot take address of function 'NotOverloaded' because parameter 1 has pass_object_size attribute}}
   TakeFn(&NotOverloaded); //expected-error{{cannot take address of function 'NotOverloaded' because parameter 1 has pass_object_size attribute}}
 
   TakeFnOvl(NotOverloaded); //expected-error{{cannot take address of function 'NotOverloaded' because parameter 1 has pass_object_size attribute}}
   TakeFnOvl(&NotOverloaded); //expected-error{{cannot take address of function 'NotOverloaded' because parameter 1 has pass_object_size attribute}}
 
   int P;
   (&NotOverloaded)(&P); //expected-error{{cannot take address of function 'NotOverloaded' because parameter 1 has pass_object_size attribute}}
-  (&IsOverloaded)(&P); //expected-error{{no matching function}} expected-note@35{{candidate address cannot be taken because parameter 1 has pass_object_size attribute}} expected-note@36{{candidate function not viable: no known conversion from 'int *' to 'char *' for 1st argument}}
+  (&IsOverloaded)(&P); //expected-warning{{incompatible pointer types passing 'int *' to parameter of type 'char *'}} expected-note@36{{passing argument to parameter 'p' here}}
 }

test/SemaOpenCL/event_t_overload.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 
-void __attribute__((overloadable)) foo(event_t, __local char *); // expected-note {{candidate function not viable: no known conversion from '__global int *' to '__local char *' for 2nd argument}}
-void __attribute__((overloadable)) foo(event_t, __local float *); // expected-note {{candidate function not viable: no known conversion from '__global int *' to '__local float *' for 2nd argument}}
+void __attribute__((overloadable)) foo(event_t, __local char *); // expected-note {{candidate function}}
+void __attribute__((overloadable)) foo(event_t, __local float *); // expected-note {{candidate function}}
 
 void kernel ker(__local char *src1, __local float *src2, __global int *src3) {
   event_t evt;
   foo(evt, src1);
   foo(0, src2);
-  foo(evt, src3); // expected-error {{no matching function for call to 'foo'}}
+  foo(evt, src3); // expected-error {{call to 'foo' is ambiguous}}
 }