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

Mon Feb 18 09:47:24 PST 2019

Anastasia added a subscriber: arsenm.
Anastasia 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<
----------------
ebevhan wrote:
> 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.
Oh I see... might not worth it?

================
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);
-}
----------------
ebevhan wrote:
> 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.
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!

================
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 **'}}
+}
----------------
ebevhan wrote:
> ebevhan wrote:
> > 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.
> > 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.
> 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.

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

https://reviews.llvm.org/D58236