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

[Song, Ruiling]

Hi Anastasia,

Thanks for your detailed explanation. Using zero is OK enough to solve the issue.
So, when Clang do semantic check, 0 is interpreted as integer instead of null pointer?
I previously thought Clang treat 0 as null pointer here directly, which it is finally transformed into.
BTW, in C language, NULL may be defined as 0 or (void*)0 in different implementation.

Thanks!
Ruiling

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<mailto: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<mailto: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/6e894b1e/attachment.html>