Make address space conversions a bit stricter.

[ebevhan]

Summary

The semantics for converting nested pointers between address
spaces are not very well defined. Some conversions which do not
really carry any meaning only produce warnings, and in some cases
warnings hide invalid conversions, such as 'global int*' to
'local float*'!

This patch changes the logic in checkPointerTypesForAssignment
and checkAddressSpaceCast to fail properly on implicit conversions
that should definitely not be permitted. We also dig deeper into the
pointer types and warn on explicit conversions where the address
space in a nested pointer changes, regardless of whether the address
space is compatible with the corresponding pointer nesting level
on the destination type.

See https://bugs.llvm.org/show_bug.cgi?id=39674

[ebevhan]

test/SemaOpenCL/address-spaces.cl

89

This doesn't seem entirely correct still, but I'm not sure what to do about it.

[rjmccall]

lib/Sema/SemaCast.cpp

2308

This should if (Nested ? DestPPointee.getAddressSpace() != SrcPPointee.getAddressSpace() : !DestPPtr->isAddressSpaceOverlapping(*SrcPtr)), I think.

lib/Sema/SemaExpr.cpp

7706

The TODO here is fixed as much as anything else is.

7781

Please extract variables for cast<PointerType>(lhptee)->getPointeeType() (and for rhptee).

14208

Yeah, I think that would be more straightforward.

[ebevhan]

lib/Sema/SemaCast.cpp

2295

I'd also like to petition to remove the guard on OpenCL here. Maybe an RFC for formalizing the support for address space conversion semantics is in order?

lib/Sema/SemaExpr.cpp

14208

I ended up not making it that different from the original one. Perhaps it's not different enough?

[Anastasia]

include/clang/Basic/DiagnosticSemaKinds.td

6996

I am wondering if we could just unify with the diagnostic above?

We could add another select at the end:

" changes address space of %select{nested|}3 pointer"

lib/Sema/SemaCast.cpp

2294

I think my patch is divorcing C++ from here https://reviews.llvm.org/D58346.

Although, I might need to update it to add the same checks. I can do it once your patch is finalized. :)

2295

Yes, I'd support that! It would be good to generalize the rules as much as we can rather than adding extra checks for languages (the semantic is very similar).

As a matter of fact I tried to do the same in https://reviews.llvm.org/D58346 in TryAddressSpaceCast and some C++ tests fail. So I had to add OpenCL check for now. :( Perhaps, they could just be changed but needs agreement with the community first.

lib/Sema/SemaExpr.cpp

14199

May be, since if we are using a typedef that is a pointer type isPointerType() might not return true? I would just extend a test case with typedef to a pointer type as one of the nested levels.

test/CodeGenOpenCL/numbered-address-space.cl

17

Does this not compile any more?

test/SemaOpenCL/address-spaces.cl

89

Is it because Sema::IncompatiblePointer has priority? We might want to change that. I think it was ok before because qualifier's mismatch was only a warning but now with the address spaces we are giving an error. I wonder if adding a separate enum item for address spaces (something like Sema::IncompatibleNestedPointerAddressSpace) would simplify things.

[ebevhan]

include/clang/Basic/DiagnosticSemaKinds.td

6996

That is doable, but all of the 'typecheck' errors/warnings are made to be regular. If I add another parameter, there needs to be a special case in DiagnoseAssignmentResult for that error in particular.

lib/Sema/SemaCast.cpp

2294

Yes, I noticed that you weren't calling checkAddressSpaceCast any longer in one of the code paths.

I suspect that your patch might already be doing the same thing... though not explicitly. TryAddressSpaceCast will only check compatibility on the top level pointee, and then it goes on to check type similarity, and if the types are not similar, it bails. Types with different nested pointer address spaces are not (should not be?) considered compatible, so we should get the same outcome. We won't get the 'nested' error, though.

lib/Sema/SemaExpr.cpp

14199

It sort of depends on what checkPointerTypesForAssignment considers invalid.

I think maybe I should throw this logic away and instead simply add a new enum member to AssignmentResult.

test/CodeGenOpenCL/numbered-address-space.cl

17

No, these tests were a bit shaky. I'm not even sure what they're supposed to be testing. It's trying to pass an arbitrary AS int * pointer to a function that takes __local float *. That AS conversion is illegal (implicitly), but the illegal conversion was 'shadowed' by the 'incompatible pointer' warning, so we didn't get an error. This is one of the things this patch fixes.

Since it's a codegen test, it should be producing something, but I'm not really sure what is interesting to produce here, so I just removed it.

test/SemaOpenCL/address-spaces.cl

89

Is it because Sema::IncompatiblePointer has priority?

Sort of. The problem is that the AS pointee qualifiers match up until the 'end' of the RHS pointer chain (LHS: private->global->local, RHS: private->global), so we never get an 'incompatible address space' to begin with. We only get that if 1) the bottommost type is equal after unwrapping pointers (as far as both sides go), or 2) any of the 'shared' AS qualifiers (as far as both sides go) were different.

The idea is that stopping when either side is no longer a pointer will produce 'incompatible pointers' when you have different pointer depths, but it doesn't consider anything below the 'shallowest' side of the pointer chain, so we miss out on any AS mismatches further down.

(Not that there's anything to mismatch, really. There is no matching pointer on the other side, so what is really the error?)

What should the criteria be for when the pointer types 'run out'? I could have it keep digging through the other pointer until it hits a different AS? This would mean that this:

int **** a;
int ** b = a;

could give a different warning than it does today, though (incompatible nested qualifiers instead of incompatible pointers, which doesn't make sense...) . We would have to skip the lhptee == rhptee check if we 'kept going' despite one side not being a pointer type. So I don't know if that's the right approach in general.

Or should we be searching 'backwards' instead, starting from the innermost pointee? I don't know.

It really feels like the whole checkPointerTypesForAssignment routine and everything surrounding it is a bit messy. It relies on an implicit result from another function (typesAreCompatible) and then tries to deduce why that function thought the types weren't compatible. Then another function later on (DiagnoseAssignmentResult) tries to deduce why THIS function thought something was wrong.

I wonder if adding a separate enum item for address spaces (something like Sema::IncompatibleNestedPointerAddressSpace) would simplify things.

This would simplify the logic on the error emission side, since we don't need to duplicate the logic for determining what went wrong, but doesn't help with diagnosing the actual problem. Probably a good idea to add it anyway, I just wanted to avoid adding a new enum member since that means updating a lot of code elsewhere.

[ebevhan]

test/SemaOpenCL/address-spaces.cl

89

We only get that if 1) the bottommost type is equal after unwrapping pointers (as far as both sides go), or 2) any of the 'shared' AS qualifiers (as far as both sides go) were different.

Sorry, should only be 2) here. Was focused on the whole 'incompatible nested qualifiers' result.

[Anastasia]

include/clang/Basic/DiagnosticSemaKinds.td

6996

Oh I see... might not worth it?

test/CodeGenOpenCL/numbered-address-space.cl

17

Ok, I will just loop in @arsenm to confirm. OpenCL doesn't regulate arbitrary address spaces. And C doesn't regulate OpenCL ones. So interplay between those two has undefined behavior in my opinion. However, OpenCL code can make use of arbitrary address spaces since it's a valid Clang extension... But I am not sure what happens with this undefined behaviors.

For this specific case I would rather expect an error... but not sure it's worth testing this anyway.

May be Matt can provide us some more insights!

test/SemaOpenCL/address-spaces.cl

89

What should the criteria be for when the pointer types 'run out'? I could have it keep digging through the other pointer until it hits a different AS?

Hmm, good point! C99 spec seems to be helpless. C++ seems to imply that it checks pointers left to right as far as I interpret that from [conv.qual]. Not sure what we should do... Would it make sense to align with C++ or otherwise whatever is simpler? At least there is a diagnostic generated. So perhaps after all it's good enough for now!

I wonder if adding a separate enum item for address spaces (something like Sema::IncompatibleNestedPointerAddressSpace) would simplify things.

This would simplify the logic on the error emission side, since we don't need to duplicate the logic for determining what went wrong, but doesn't help with diagnosing the actual problem. Probably a good idea to add it anyway, I just wanted to avoid adding a new enum member since that means updating a lot of code elsewhere.

Ok, common helper function could be another solution to avoid duplication but it seems the logic is not entirely identical.

[ebevhan]

include/clang/Basic/DiagnosticSemaKinds.td

6996

I think keeping the generality makes it a bit simpler. Technically many of the errors here could be folded into one or two instead, but that hasn't been done, so...

test/SemaOpenCL/address-spaces.cl

89

Would it make sense to align with C++ or otherwise whatever is simpler?

I think C++ is a bit stricter in general, it doesn't permit this, but C is more lenient. There is a diagnostic, so I left a FIXME but it should probably be revisited.

Ok, common helper function could be another solution to avoid duplication but it seems the logic is not entirely identical.

I added a new enum member and removed the logic on the error emission side.

[Anastasia]

lib/Sema/SemaExpr.cpp

7795

I am wondering since the behavior is so specialized for the address spaces whether we should name this something like IncompatibleNestedPointerAddressSpaceMismatch.

test/SemaOpenCL/address-spaces.cl

87

Are we sure it has to be an error? May be we can change this wording to something like - unknown behavior that needs clarifying? Some similar text to your comment in the code.

[ebevhan]

test/SemaOpenCL/address-spaces.cl

87

Well... If __private int ** to __generic int ** is an error but __private int *** to __generic int ** isn't, that seems a bit odd to me...

[Anastasia]

LGTM! Thanks a lot for fixing this old bug! Btw, do you plan to look at generalizing this to C++ as well?

I don't feel we need anything for mixing OpenCL with GNU style address spaces at this point. We can always extend this in the future (in case it turns out to be useful to someone!).

test/SemaOpenCL/address-spaces.cl

87

Ok, I am still confused about this case because it's converting number of pointers. It's not exactly the same to me as converting address spaces of pointers. Since __generic is default addr space in OpenCL, in your example following 2 inner address spaces seem to match:
__private int * __generic* __generic* -> __generic int * __generic*

But in any case keeping FIXME is definitely the right thing to do here until we come up with some sort of rules.

[efriedma]

Generally, with an explicit cast, C allows any pointer cast with a reasonable interpretation, even if the underlying operation is suspicious. For example, you can cast an "long*" to a "int*" (as in "(int*)(long*)p") without any complaint, even though dereferencing the result is likely undefined behavior. (C doesn't allow explicitly writing a reinterpret_cast, but that's basically the operation in question.)

Along those lines, in general, the normal C rules should allow casting foo* to bar* for any object types foo and bar, even if foo and bar are pointers with address spaces, like __local int * and __global int *. I don't see anything in the OpenCL standard that would contradict this. It looks like this patch changes that?

I agree we shouldn't allow implicit conversions, though...

[rjmccall]

test/SemaOpenCL/address-spaces.cl

87

For reasons that have always been obscure to me, Clang is very permissive about pointer-type changes in C, far more than the C standard actually requires. It's hard to change that in general, but I think it's fine to harden specific new cases like nested address-space conversions.

[ebevhan]

In D58236#1412267, @Anastasia wrote:

LGTM! Thanks a lot for fixing this old bug! Btw, do you plan to look at generalizing this to C++ as well?

That does sound like a good idea and I will probably look into it when I have more time. I'd also like to look into making an RFC for the improved address space specification support. Again, when I have time :)

In D58236#1412324, @efriedma wrote:

Generally, with an explicit cast, C allows any pointer cast with a reasonable interpretation, even if the underlying operation is suspicious. For example, you can cast an "long*" to a "int*" (as in "(int*)(long*)p") without any complaint, even though dereferencing the result is likely undefined behavior. (C doesn't allow explicitly writing a reinterpret_cast, but that's basically the operation in question.)

That is true... The current rules are very permissive about pointer casts.

Along those lines, in general, the normal C rules should allow casting foo* to bar* for any object types foo and bar, even if foo and bar are pointers with address spaces, like __local int * and __global int *. I don't see anything in the OpenCL standard that would contradict this. It looks like this patch changes that?

I vaguely recall that when I was looking into fixing this in our downstream a while back I found that OpenCL did not have any explicit provisions for nested address spaces, only direct ones.

So should we drop the extra error on nested space mismatches, since the rules should be as permissive as possible unless the spec prohibits it? Technically the patch could be changed to only fix the incorrect warning on direct address space mismatches instead.

[Anastasia]

Along those lines, in general, the normal C rules should allow casting foo* to bar* for any object types foo and bar, even if foo and bar are pointers with address spaces, like __local int * and __global int *. I don't see anything in the OpenCL standard that would contradict this. It looks like this patch changes that?

s6.5 of OpenCL spec is very explicit about this case: "A pointer to address space A can only be assigned to a pointer to the same address space A or a pointer to the generic address space. Casting a pointer to address space A to a pointer to address space B is illegal if A and B are named address spaces and A is not the same as B". Although there is a typo with missing first occurrence of B.

The behavior with generic is explained in s6.5.5.

Btw, this patch doesn't change anything in the address space conversions of pointers in OpenCL in general but just changes the behavior of nested pointers for which spec doesn't say anything. I think trying to reject code that is doing something dangerous is a good thing!

[efriedma]

I think trying to reject code that is doing something dangerous is a good thing!

Refusing to compile code which is suspicious, but not forbidden by the specification, will likely cause compatibility issues; there are legitimate reasons to use casts which look weird.

I'm not against adding some sort of address space suspicious cast warning to catch cases where we think the user meant to do something else. But that's a separate issue, and it needs a proper cost-benefit analysis, including an analysis of the false-positive rate on existing code.

[Anastasia]

In D58236#1414069, @efriedma wrote:

I think trying to reject code that is doing something dangerous is a good thing!

Refusing to compile code which is suspicious, but not forbidden by the specification, will likely cause compatibility issues; there are legitimate reasons to use casts which look weird.

The spec dioesn't allow these conversions either, it just simply doesn't cover this corner case at all. I don't think we are changing anything in terms of compatibility. If you have any examples of such casts that can be legitimate I would like to understand them better. What I have seen so far were the examples where addrspacecast was lost in IR for the memory segments translation and therefore wrong memory areas were accessed.

I'm not against adding some sort of address space suspicious cast warning to catch cases where we think the user meant to do something else.

I simply don't see how these conversions can be useful and some are definitely indirectly forbidden (there is no precise wording however). There are other ways to perform such conversions differently (by being more explicit) where correct IR can be then generated with addrspacecast. I don't think we are loosing anything in terms of functionality.

But that's a separate issue, and it needs a proper cost-benefit analysis, including an analysis of the false-positive rate on existing code.

Do you have any suggestions how to do this in practice with such rare corner case?

[efriedma]

In D58236#1416765, @Anastasia wrote:

In D58236#1414069, @efriedma wrote:

I think trying to reject code that is doing something dangerous is a good thing!

Refusing to compile code which is suspicious, but not forbidden by the specification, will likely cause compatibility issues; there are legitimate reasons to use casts which look weird.

The spec dioesn't allow these conversions either, it just simply doesn't cover this corner case at all. I don't think we are changing anything in terms of compatibility. If you have any examples of such casts that can be legitimate I would like to understand them better. What I have seen so far were the examples where addrspacecast was lost in IR for the memory segments translation and therefore wrong memory areas were accessed.

The spec just says that the casts follow C rules... and C says you can cast a pointer to an object type to a pointer to another object type (subject to alignment restrictions). By default, a pointer to a pointer isn't special.

In practice, unusual casts tend to show up in code building a datastructure using union-like constructs. In plain C, for example, sometimes you have a pointer to a float, and sometimes you have a pointer to an int, determined dynamically. I expect similar cases show up where a pointer points to memory which contains either a pointer in the global address-space, or a pointer in the local address-space, determined dynamically. In some cases, it might be clearer to use void* in more places, but that's mostly style issue.

I'm not against adding some sort of address space suspicious cast warning to catch cases where we think the user meant to do something else.

I simply don't see how these conversions can be useful and some are definitely indirectly forbidden (there is no precise wording however). There are other ways to perform such conversions differently (by being more explicit) where correct IR can be then generated with addrspacecast. I don't think we are loosing anything in terms of functionality.

In some cases, the correct code may not involve an addrspacecast at all; the pointer could just have the "wrong" pointee type, and the code expects to explicitly fix it. In that case, how do you fix it? Write (foo*)(void*)p instead of (foo*)p?

But that's a separate issue, and it needs a proper cost-benefit analysis, including an analysis of the false-positive rate on existing code.

Do you have any suggestions how to do this in practice with such rare corner case?

If the warning never triggers on any code you have access to, that would still be a useful datapoint.

[Anastasia]

In D58236#1417429, @efriedma wrote:

In D58236#1416765, @Anastasia wrote:

In D58236#1414069, @efriedma wrote:

I think trying to reject code that is doing something dangerous is a good thing!

Refusing to compile code which is suspicious, but not forbidden by the specification, will likely cause compatibility issues; there are legitimate reasons to use casts which look weird.

The spec dioesn't allow these conversions either, it just simply doesn't cover this corner case at all. I don't think we are changing anything in terms of compatibility. If you have any examples of such casts that can be legitimate I would like to understand them better. What I have seen so far were the examples where addrspacecast was lost in IR for the memory segments translation and therefore wrong memory areas were accessed.

The spec just says that the casts follow C rules... and C says you can cast a pointer to an object type to a pointer to another object type (subject to alignment restrictions). By default, a pointer to a pointer isn't special.

In practice, unusual casts tend to show up in code building a datastructure using union-like constructs. In plain C, for example, sometimes you have a pointer to a float, and sometimes you have a pointer to an int, determined dynamically. I expect similar cases show up where a pointer points to memory which contains either a pointer in the global address-space, or a pointer in the local address-space, determined dynamically. In some cases, it might be clearer to use void* in more places, but that's mostly style issue.

Ok, do you have any example for pointers with different types by some chance? It would be very helpful because I can try to see what would be the behavior with different address spaces...

One fundamental difference between C and OpenCL C is that we wanted to give more clear semantic to address spaces. It is not possible to casts between pointers of arbitrary address spaces. We wanted to make such conversions very explicit using generic address space. If the code is written using unions (or some other way) that requires such "meaningless" conversions it can still be cast but the cast has to be written using pointer indirection with generic address spaces. That was done deliberately to prevent accidental erroneous patterns to be compiled.

But that's a separate issue, and it needs a proper cost-benefit analysis, including an analysis of the false-positive rate on existing code.

Do you have any suggestions how to do this in practice with such rare corner case?

If the warning never triggers on any code you have access to, that would still be a useful datapoint.

To my knowledge these corner cases hasn't occurred yet in any code pattern known to me up. I created this upstream bug myself while sorting out some other address space related aspects. However, it's not impossible that it might happen in the future.

I still feel this might be the right direction for OpenCL, do you think this might not be right for C? If that's the case potentially we could add OpenCL specific checks for now and then try to clarify what should be the right strategy for C... it feels in C we just converts between addr spaces freely. The only danger of this conversion with nested pointers is that it produces bitcast instead of addrspacecast, which might result in accessing the wrong memory. Therefore, we wanted to restrict this in the first place.

[ebevhan]

Any more input on this?

I could redo the patch to simply fix the bug and not make the conversions stricter, if that's preferable.

[Anastasia]

In D58236#1437690, @ebevhan wrote:

Any more input on this?

I could redo the patch to simply fix the bug and not make the conversions stricter, if that's preferable.

I was playing a bit with some examples of enum with pointer field of various address spaces and I couldn't find any case where successfully converting nested pointers was useful. Instead it could easily result in execution of incorrect code. So I am sticking to the opinion to be more strict.

Similarly, C++ is more strict with similar conversions i.e. nested pointers between Derived and Base. I still think since we are implementing new rules it is good to be more helpful rather than inheriting older logic of permitting everything like in C.

Also of course if we find later any issue we could fix them and modify the behavior.

[efriedma]

If nobody else agrees with my position on this, I'm not going to continue arguing on the explicit cast behavior. But please add a testcase showing that at least (global int**)(void*)(local int**)p works without an error.

[rjmccall]

I think C probably requires us to allow this under an explicit cast, but we can at least warn about it. It does not require us to allow this as an implicit conversion, and that should be an error.

[ebevhan]

In D58236#1443082, @rjmccall wrote:

I think C probably requires us to allow this under an explicit cast, but we can at least warn about it. It does not require us to allow this as an implicit conversion, and that should be an error.

Are you referring to casts to and from void*? I think the standard says those have to be allowed, and I don't see anything about whether or not the conversion has to be explicit.

[rjmccall]

In D58236#1443556, @ebevhan wrote:

In D58236#1443082, @rjmccall wrote:

I think C probably requires us to allow this under an explicit cast, but we can at least warn about it. It does not require us to allow this as an implicit conversion, and that should be an error.

Are you referring to casts to and from void*? I think the standard says those have to be allowed, and I don't see anything about whether or not the conversion has to be explicit.

No, I mean casts between pointer types that change the pointee address space.

[ebevhan]

Well, it doesn't seem to me like there is consensus on prohibiting nested address space conversion like this.

I can simply redo the patch to only include the bugfix on implicit conversions and drop the nesting level checks.

[Anastasia]

In D58236#1466263, @ebevhan wrote:

Well, it doesn't seem to me like there is consensus on prohibiting nested address space conversion like this.

I can simply redo the patch to only include the bugfix on implicit conversions and drop the nesting level checks.

I thought the conclusion is:

• Explicit conversions allowed for nested pointers with a warning.
• Implicit conversions are disallowed for nested pointers with an error.

Do you not see it this way?

[ebevhan]

In D58236#1467151, @Anastasia wrote:

In D58236#1466263, @ebevhan wrote:

Well, it doesn't seem to me like there is consensus on prohibiting nested address space conversion like this.

I can simply redo the patch to only include the bugfix on implicit conversions and drop the nesting level checks.

I thought the conclusion is:

• Explicit conversions allowed for nested pointers with a warning.
• Implicit conversions are disallowed for nested pointers with an error.

Do you not see it this way?

Alright, that seems reasonable to me. Some of the patch goes away, then. I'll update it.

[rjmccall]

Alright, this seems reasonable to me.

[ebevhan]

This was accepted a while ago, but never landed. I don't have commit access; could someone commit it?

[Anastasia]

In D58236#1494722, @ebevhan wrote:

This was accepted a while ago, but never landed. I don't have commit access; could someone commit it?

Sure!

include/clang/Basic/DiagnosticSemaKinds.td

def err_typecheck_incompatible_address_space : Error<
   "returning from function with different return type}0,1"
   "|%diff{converting $ to type $|converting between types}0,1"
   "|%diff{initializing $ with an expression of type $|"
   "initializing with expression of different type}0,1"
   "|%diff{sending $ to parameter of type $|"
   "sending to parameter of different type}0,1"
   "|%diff{casting $ to type $|casting between types}0,1}2"
   " changes address space of pointer">;
+def err_typecheck_incompatible_nested_address_space : Error<
+  "%select{%diff{assigning $ to $|assigning to different types}1,0"
+  "|%diff{passing $ to parameter of type $|"
+  "passing to parameter of different type}0,1"
+  "|%diff{returning $ from a function with result type $|"
+  "returning from function with different return type}0,1"
+  "|%diff{converting $ to type $|converting between types}0,1"
+  "|%diff{initializing $ with an expression of type $|"
+  "initializing with expression of different type}0,1"
+  "|%diff{sending $ to parameter of type $|"
+  "sending to parameter of different type}0,1"
+  "|%diff{casting $ to type $|casting between types}0,1}2"
+  " changes address space of nested pointer">;
 def err_typecheck_incompatible_ownership : Error<
   "%select{%diff{assigning $ to $|assigning to different types}1,0"
   "|%diff{passing $ to parameter of type $|"
   "passing to parameter of different type}0,1"
   "|%diff{returning $ from a function with result type $|"
   "returning from function with different return type}0,1"
   "|%diff{converting $ to type $|converting between types}0,1"
   "|%diff{initializing $ with an expression of type $|"

lib/Sema/SemaCast.cpp

static TryCastResult TryReinterpretCast(Sema &Self, ExprResult &SrcExpr,
   // object pointers.
   return SuccessResult;
 }
 
 void CastOperation::checkAddressSpaceCast(QualType SrcType, QualType DestType) {
   // In OpenCL only conversions between pointers to objects in overlapping
   // addr spaces are allowed. v2.0 s6.5.5 - Generic addr space overlaps
   // with any named one, except for constant.
+
+  // Converting the top level pointee addrspace is permitted for compatible
+  // addrspaces (such as 'generic int *' to 'local int *' or vice versa), but
+  // if any of the nested pointee addrspaces differ, we consider it illegal
+  // regardless of addrspace compatibility. This makes
+  //   local int ** p;
+  //   return (generic int **) p;
+  // illegal even though local -> generic is permitted.
+
+  // FIXME: C++ might want to emit different errors here.
   if (Self.getLangOpts().OpenCL) {
-    auto SrcPtrType = SrcType->getAs<PointerType>();
-    if (!SrcPtrType)
-      return;
-    auto DestPtrType = DestType->getAs<PointerType>();
-    if (!DestPtrType)
-      return;
-    if (!DestPtrType->isAddressSpaceOverlapping(*SrcPtrType)) {
-      Self.Diag(OpRange.getBegin(),
-                diag::err_typecheck_incompatible_address_space)
+    const Type *DestPtr, *SrcPtr;
+    bool Nested = false;
+    unsigned DiagID = diag::err_typecheck_incompatible_address_space;
+    DestPtr = Self.getASTContext().getCanonicalType(DestType.getTypePtr()),
+    SrcPtr  = Self.getASTContext().getCanonicalType(SrcType.getTypePtr());
+
+    while (isa<PointerType>(DestPtr) && isa<PointerType>(SrcPtr)) {
+      const PointerType *DestPPtr = cast<PointerType>(DestPtr);
+      const PointerType *SrcPPtr = cast<PointerType>(SrcPtr);
+      QualType DestPPointee = DestPPtr->getPointeeType();
+      QualType SrcPPointee = SrcPPtr->getPointeeType();
+      if (Nested ? DestPPointee.getAddressSpace() !=
+                   SrcPPointee.getAddressSpace()
+                 : !DestPPtr->isAddressSpaceOverlapping(*SrcPPtr)) {
+        Self.Diag(OpRange.getBegin(), DiagID)
             << SrcType << DestType << Sema::AA_Casting
             << SrcExpr.get()->getSourceRange();
         SrcExpr = ExprError();
+        return;
+      }
+
+      DestPtr = DestPPtr->getPointeeType().getTypePtr();
+      SrcPtr = SrcPPtr->getPointeeType().getTypePtr();
+      Nested = true;
+      DiagID = diag::err_typecheck_incompatible_nested_address_space;
     }
   }
 }
 
 void CastOperation::CheckCXXCStyleCast(bool FunctionalStyle,
                                        bool ListInitialization) {
   assert(Self.getLangOpts().CPlusPlus);
 

lib/Sema/SemaExpr.cpp

checkPointerTypesForAssignment(Sema &S, QualType LHSType, QualType RHSType) {
   if (lhq.getObjCLifetime() != rhq.getObjCLifetime() &&
       lhq.compatiblyIncludesObjCLifetime(rhq)) {
     // Ignore lifetime for further calculation.
     lhq.removeObjCLifetime();
     rhq.removeObjCLifetime();
   }
 
   if (!lhq.compatiblyIncludes(rhq)) {
-    // Treat address-space mismatches as fatal.  TODO: address subspaces
+    // Treat address-space mismatches as fatal.
     if (!lhq.isAddressSpaceSupersetOf(rhq))
-      ConvTy = Sema::IncompatiblePointerDiscardsQualifiers;
+      return Sema::IncompatiblePointerDiscardsQualifiers;
 
     // It's okay to add or remove GC or lifetime qualifiers when converting to
     // and from void*.
     else if (lhq.withoutObjCGCAttr().withoutObjCLifetime()
                         .compatiblyIncludes(
                                 rhq.withoutObjCGCAttr().withoutObjCLifetime())
              && (lhptee->isVoidType() || rhptee->isVoidType()))
       ; // keep old
if (!S.Context.typesAreCompatible(ltrans, rtrans)) {
     }
 
     // If we are a multi-level pointer, it's possible that our issue is simply
     // one of qualification - e.g. char ** -> const char ** is not allowed. If
     // the eventual target type is the same and the pointers have the same
     // level of indirection, this must be the issue.
     if (isa<PointerType>(lhptee) && isa<PointerType>(rhptee)) {
       do {
-        lhptee = cast<PointerType>(lhptee)->getPointeeType().getTypePtr();
-        rhptee = cast<PointerType>(rhptee)->getPointeeType().getTypePtr();
+        std::tie(lhptee, lhq) =
+          cast<PointerType>(lhptee)->getPointeeType().split().asPair();
+        std::tie(rhptee, rhq) =
+          cast<PointerType>(rhptee)->getPointeeType().split().asPair();
+
+        // Inconsistent address spaces at this point is invalid, even if the
+        // address spaces would be compatible.
+        if (lhq.getAddressSpace() != rhq.getAddressSpace())
+          return Sema::IncompatibleNestedPointerQualifiers;
+
       } while (isa<PointerType>(lhptee) && isa<PointerType>(rhptee));
 
       if (lhptee == rhptee)
         return Sema::IncompatibleNestedPointerQualifiers;
     }
 
     // General pointer incompatibility takes priority over qualifiers.
     return Sema::IncompatiblePointer;
   }
   if (!S.getLangOpts().CPlusPlus &&
       S.IsFunctionConversion(ltrans, rtrans, ltrans))
     return Sema::IncompatiblePointer;
   return ConvTy;
case CompatiblePointerDiscardsQualifiers:
     // expression, rather than a type), which should be done as part
     // of a larger effort to fix checkPointerTypesForAssignment for
     // C++ semantics.
     if (getLangOpts().CPlusPlus &&
         IsStringLiteralToNonConstPointerConversion(SrcExpr, DstType))
       return false;
     DiagKind = diag::ext_typecheck_convert_discards_qualifiers;
     break;
-  case IncompatibleNestedPointerQualifiers:
+  case IncompatibleNestedPointerQualifiers: {
+    // Perform array-to-pointer decay if necessary.
+    if (SrcType->isArrayType()) SrcType = Context.getArrayDecayedType(SrcType);
+
+    // Address space qualifier mismatch in nested pointers is always an error.
+    // Strip the pointers until we either hit a mismatch or hit the end of the
+    // nesting.
+    // Use pointee type here since this should only care about nested
+    // qualifiers.
+    // XXX: Canonical types?
+    const Type *SrcPtr = SrcType->getPointeeType().getTypePtr();
+    const Type *DstPtr = DstType->getPointeeType().getTypePtr();
+
     DiagKind = diag::ext_nested_pointer_qualifier_mismatch;
+    while (SrcPtr->isPointerType() && DstPtr->isPointerType()) {
+      if (SrcPtr->getPointeeType().getAddressSpace() !=
+          DstPtr->getPointeeType().getAddressSpace()) {
+        DiagKind = diag::err_typecheck_incompatible_nested_address_space;
+        isInvalid = true;
         break;
+      }
+      SrcPtr = SrcPtr->getPointeeType().getTypePtr();
+      DstPtr = DstPtr->getPointeeType().getTypePtr();
+    }
+    break;
+  }
   case IntToBlockPointer:
     DiagKind = diag::err_int_to_block_pointer;
     break;
   case IncompatibleBlockPointer:
     DiagKind = diag::err_typecheck_convert_incompatible_block_pointer;
     break;
   case IncompatibleObjCQualifiedId: {
     if (SrcType->isObjCQualifiedIdType()) {

test/CodeGenOpenCL/numbered-address-space.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -O0 -o - %s | FileCheck %s
 
 // Make sure using numbered address spaces doesn't trigger crashes when a
 // builtin has an address space parameter.
 
 // CHECK-LABEL: @test_numbered_as_to_generic(
 // CHECK: addrspacecast i32 addrspace(42)* %0 to i32*
 void test_numbered_as_to_generic(__attribute__((address_space(42))) int *arbitary_numbered_ptr) {
   generic int* generic_ptr = arbitary_numbered_ptr;
   *generic_ptr = 4;
 }
 
-// CHECK-LABEL: @test_numbered_as_to_builtin(
-// CHECK: addrspacecast i32 addrspace(42)* %0 to float addrspace(3)*
-void test_numbered_as_to_builtin(__attribute__((address_space(42))) int *arbitary_numbered_ptr, float src) {
-  volatile float result = __builtin_amdgcn_ds_fmaxf(arbitary_numbered_ptr, src, 0, 0, false);
-}
-
 // CHECK-LABEL: @test_generic_as_to_builtin_parameter_explicit_cast(
 // CHECK: addrspacecast i32 addrspace(3)* %0 to i32*
 void test_generic_as_to_builtin_parameter_explicit_cast(__local int *local_ptr, float src) {
   generic int* generic_ptr = local_ptr;
   volatile float result = __builtin_amdgcn_ds_fmaxf((__local float*) generic_ptr, src, 0, 0, false);
 }
 
 // CHECK-LABEL: @test_generic_as_to_builtin_parameter_implicit_cast(
-// CHECK: addrspacecast i32* %2 to float addrspace(3)*
+// CHECK: bitcast i32 addrspace(3)* %0 to float addrspace(3)*
 void test_generic_as_to_builtin_parameter_implicit_cast(__local int *local_ptr, float src) {
-  generic int* generic_ptr = local_ptr;
-
-  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false);
+  volatile float result = __builtin_amdgcn_ds_fmaxf(local_ptr, src, 0, 0, false);
 }
-

test/SemaOpenCL/address-spaces.cl

 }
 
 void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __local int *l2, __private int *p, __private int *p2) {
   g = (__global int *)g2;
   l = (__local int *)l2;
   p = (__private int *)p2;
 }
 
+#if !__OPENCL_CPP_VERSION__
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+  g = gg;    // expected-error {{assigning '__global int **' to '__global int *' changes address space of pointer}}
+  g = l;     // expected-error {{assigning '__local int *' to '__global int *' changes address space of pointer}}
+  g = ll;    // expected-error {{assigning '__local int **' to '__global int *' changes address space of pointer}}
+  g = gg_f;  // expected-error {{assigning '__global float **' to '__global int *' changes address space of pointer}}
+  g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}}
+
+  gg = g;    // expected-error {{assigning '__global int *' to '__global int **' changes address space of pointer}}
+  gg = l;    // expected-error {{assigning '__local int *' to '__global int **' changes address space of pointer}}
+  gg = ll;   // expected-error {{assigning '__local int **' to '__global int **' changes address space of nested pointer}}
+  gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int **' from '__global float **'}}
+  gg = (__global int * __private *)gg_f;
+
+  l = g;     // expected-error {{assigning '__global int *' to '__local int *' changes address space of pointer}}
+  l = gg;    // expected-error {{assigning '__global int **' to '__local int *' changes address space of pointer}}
+  l = ll;    // expected-error {{assigning '__local int **' to '__local int *' changes address space of pointer}}
+  l = gg_f;  // expected-error {{assigning '__global float **' to '__local int *' changes address space of pointer}}
+  l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}}
+
+  ll = g;    // expected-error {{assigning '__global int *' to '__local int **' changes address space of pointer}}
+  ll = gg;   // expected-error {{assigning '__global int **' to '__local int **' changes address space of nested pointer}}
+  ll = l;    // expected-error {{assigning '__local int *' to '__local int **' changes address space of pointer}}
+  ll = gg_f; // expected-error {{assigning '__global float **' to '__local int **' changes address space of nested pointer}}
+  ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of nested pointer}}
+
+  gg_f = g;  // expected-error {{assigning '__global int *' to '__global float **' changes address space of pointer}}
+  gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float **' from '__global int **'}}
+  gg_f = l;  // expected-error {{assigning '__local int *' to '__global float **' changes address space of pointer}}
+  gg_f = ll; // expected-error {{assigning '__local int **' to '__global float **' changes address space of nested pointer}}
+  gg_f = (__global float * __private *)gg;
+
+  // FIXME: This doesn't seem right.
+  __local int * __global * __private * lll;
+  lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+}
+#else
+void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
+  g = gg;    // expected-error {{assigning to '__global int *' from incompatible type '__global int **'}}
+  g = l;     // expected-error {{assigning to '__global int *' from incompatible type '__local int *'}}
+  g = ll;    // expected-error {{assigning to '__global int *' from incompatible type '__local int **'}}
+  g = gg_f;  // expected-error {{assigning to '__global int *' from incompatible type '__global float **'}}
+  g = (__global int *)gg_f; // expected-error {{casting '__global float **' to type '__global int *' changes address space of pointer}}
+
+  gg = g;    // expected-error {{assigning to '__global int **' from incompatible type '__global int *'; take the address with &}}
+  gg = l;    // expected-error {{assigning to '__global int **' from incompatible type '__local int *'}}
+  gg = ll;   // expected-error {{assigning to '__global int **' from incompatible type '__local int **'}}
+  gg = gg_f; // expected-error {{assigning to '__global int **' from incompatible type '__global float **'}}
+  gg = (__global int * __private *)gg_f;
+
+  l = g;     // expected-error {{assigning to '__local int *' from incompatible type '__global int *'}}
+  l = gg;    // expected-error {{assigning to '__local int *' from incompatible type '__global int **'}}
+  l = ll;    // expected-error {{assigning to '__local int *' from incompatible type '__local int **'}}
+  l = gg_f;  // expected-error {{assigning to '__local int *' from incompatible type '__global float **'}}
+  l = (__local int *)gg_f; // expected-error {{casting '__global float **' to type '__local int *' changes address space of pointer}}
+
+  ll = g;    // expected-error {{assigning to '__local int **' from incompatible type '__global int *'}}
+  ll = gg;   // expected-error {{assigning to '__local int **' from incompatible type '__global int **'}}
+  ll = l;    // expected-error {{assigning to '__local int **' from incompatible type '__local int *'; take the address with &}}
+  ll = gg_f; // expected-error {{assigning to '__local int **' from incompatible type '__global float **'}}
+  ll = (__local int * __private *)gg_f; // expected-error {{casting '__global float **' to type '__local int **' changes address space of nested pointer}}
+
+  gg_f = g;  // expected-error {{assigning to '__global float **' from incompatible type '__global int *'}}
+  gg_f = gg; // expected-error {{assigning to '__global float **' from incompatible type '__global int **'}}
+  gg_f = l;  // expected-error {{assigning to '__global float **' from incompatible type '__local int *'}}
+  gg_f = ll; // expected-error {{assigning to '__global float **' from incompatible type '__local int **'}}
+  gg_f = (__global float * __private *)gg;
+}
+#endif
+
 __private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
 __global int func_return_global(void);      //expected-error {{return value cannot be qualified with address space}}
 __local int func_return_local(void);        //expected-error {{return value cannot be qualified with address space}}
 __constant int func_return_constant(void);  //expected-error {{return value cannot be qualified with address space}}
 #if __OPENCL_C_VERSION__ >= 200
 __generic int func_return_generic(void);    //expected-error {{return value cannot be qualified with address space}}
 #endif
 

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}}
 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 {{call to 'foo' is ambiguous}}
+  foo(evt, src3); // expected-error {{no matching function for call to 'foo'}}
 }

test/SemaOpenCL/numbered-address-space.cl

 
 void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
   generic int* generic_ptr = as3_ptr; // FIXME: This should error
   volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float*) generic_ptr, src, 0, 0, false); // expected-error {{passing '__attribute__((address_space(3))) float *' to parameter of type '__local float *' changes address space of pointer}}
 }
 
 void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
   generic int* generic_ptr = as3_ptr;
-  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__local float *'}}
+  volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *' to parameter of type '__local float *' changes address space of pointer}}
 }
 

test/SemaOpenCL/queue_t_overload.cl

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