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



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to