[cfe-dev] Got compile error 'comparison between non-overlapping address spaces' on 3.6 branch

[Anastasia Stulova]

> So, when Clang do semantic check, 0 is interpreted as integer instead of null pointer?

 

Yes, Clang will cast 0 implicitly to the null pointer of the type it is being compare to, so the address space will be taken into account correctly in this case.

 

If you use (void *)0, Clang will apply default address space to your null pointer (as it’s not provided explicitly). Therefore, pointers will appear to point to different address spaces.

 

Regards,

Anastasia 

 

From: Anastasia Stulova [mailto:anastasia.stulova at arm.com] 
Sent: Friday, February 13, 2015 6:23 PM
To: 'Sameer Sahasrabuddhe'; Song, Ruiling
Cc: cfe-dev at cs.uiuc.edu
Subject: RE: Got compile error 'comparison between non-overlapping address spaces' on 3.6 branch

 

Hi Ruiling,

 

I agree with Sameer, in general in CL2.0 your example will have a comparison between global and generic address spaces, which is allowed.

 

However, for all CL versions Clang still uses private as default address space of pointers. Some time ago, I’ve sent a patch changing this to comply with CL2.0 rules. My patch wasn’t complete though, because default address space rules are a bit more complex, i.e. program scope pointers should have global address space as default and all other pointers generic. I will see if I can find time to rework this patch.

 

In  the meantime, would this be sufficient for you:

#define NULL 0

This will result in the comparison between a pointer and an integer literal, which is allowed in OpenCL. Also this is the only solution for CL earlier than v2.0.

 

Regards,

Anastasia

 

From: Sameer Sahasrabuddhe [mailto:sameer.sahasrabuddhe at amd.com] 
Sent: 13 February 2015 06:17
To: Song, Ruiling; Anastasia Stulova
Cc: cfe-dev at cs.uiuc.edu
Subject: Re: Got compile error 'comparison between non-overlapping address spaces' on 3.6 branch

 


Hello Ruiling,

As far as I can see, the C spec defines the literal "0" as the null pointer, so the following macro seems better:

#define NULL 0

But ignoring the NULL, I think the behaviour might change between OpenCL 1.2 and OpenCL 2.0. In CL12, the "void *" is likely to be a pointer to the private address space. In that case, the comparison is undefined. In CL20, it is likely to be a pointer to the generic address space, and it should be allowed by the frontend. (Because generic address space includes global address space).

I have to use words like "should" and "likely to" because this needs confirmation from language lawyers ... the question actually pushes the boundaries of how much I understand the details of OpenCL! :)

Sameer.

On 2/13/2015 11:25 AM, Song, Ruiling wrote:

Hi Sameer, Anastasia,
 
Do you have any comment on this? Thanks!
 
Ruiling

-----Original Message-----
From: Song, Ruiling
Sent: Wednesday, February 11, 2015 10:38 AM
To: 'cfe-dev at cs.uiuc.edu'
Subject: Got compile error 'comparison between non-overlapping address
spaces' on 3.6 branch
 
I am a developer working on Beignet (OpenCL driver for Intel Graphics card).
I am trying Clang 3.6 branch, and encounter below problem. It works well in
older clang version.
We define NULL as "#define NULL (void *)0", so user could directly use it in
OpenCL language like __kernel void test(__global unsigned int *dst)
   if (dst != NULL) {do something;}
}
But I got below error when compiling using clang 3.6.
"error: comparison between  ('__global unsigned int *' and 'void *') which
are pointers to non-overlapping address spaces"
 
Is this the expected behavior or a bug?
 
Thanks!
Ruiling

 
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20150213/db7d7a6c/attachment.html>