[PATCH] D58236: Make address space conversions a bit stricter.

[Bevin Hansson via Phabricator via cfe-commits]

ebevhan marked 5 inline comments as done.
ebevhan added inline comments.


================
Comment at: include/clang/Basic/DiagnosticSemaKinds.td:6996
+  "|%diff{casting $ to type $|casting between types}0,1}2"
+  " changes address space of nested pointer">;
 def err_typecheck_incompatible_ownership : Error<
----------------
Anastasia wrote:
> 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"
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.


================
Comment at: lib/Sema/SemaCast.cpp:2294
+
+  // FIXME: C++ might want to emit different errors here.
   if (Self.getLangOpts().OpenCL) {
----------------
Anastasia wrote:
> 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. :)
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.


================
Comment at: lib/Sema/SemaExpr.cpp:14199
+    // qualifiers.
+    // XXX: Canonical types?
+    const Type *SrcPtr = SrcType->getPointeeType().getTypePtr();
----------------
Anastasia wrote:
> 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.
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.


================
Comment at: test/CodeGenOpenCL/numbered-address-space.cl:17
-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);
-}
----------------
Anastasia wrote:
> Does this not compile any more?
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.


================
Comment at: test/SemaOpenCL/address-spaces.cl:89
+  __local int * __global * __private * lll;
+  lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global **' from '__global int **'}}
+}
----------------
Anastasia wrote:
> ebevhan wrote:
> > This doesn't seem entirely correct still, but I'm not sure what to do about it.
> 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.
> 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.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D58236/new/

https://reviews.llvm.org/D58236