OpenCL parsing / type support

classic Classic list List threaded Threaded
15 messages Options
Reply | Threaded
Open this post in threaded view
|

OpenCL parsing / type support

brdavs
Hi,

I know many vendors use clang/llvm for their implementations of OpenCL.

I am wondering if OpenCL parser and types are supported in standard OpenCL distributions or they are strictly proprietary modules.

Thanks,

Marko


     
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Chris Lattner

On Mar 8, 2010, at 11:38 PM, brdavs wrote:

> Hi,
>
> I know many vendors use clang/llvm for their implementations of OpenCL.
>
> I am wondering if OpenCL parser and types are supported in standard OpenCL distributions or they are strictly proprietary modules.

Hi Marko,

I'm not an expert, but clang mainline does a bunch of OpenCL features, such as the vector extensions:
http://clang.llvm.org/docs/LanguageExtensions.html#vectors

AFAIK, common opencl implementations have various hacks on top of mainline clang, so I doubt clang mainline will be a compliant implementation out of the box.

-Chris
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

brdavs

> AFAIK, common opencl implementations have various hacks on
> top of mainline clang, so I doubt clang mainline will be a
> compliant implementation out of the box.

That's how it appears. For parsing only, what's missing is:

1. address space qualifiers (eg., __global float*, __constant int,
__local char*)
2. function qualifiers (eg, __kernel void foo(...) ),
some builtin types (float4, float8, float16, image2d_t, etc.)
3. image access qualifiers (eg, __read_only, __write_only,
__read_write).

It also seems that clang-cc had some OpenCL support (at least it was listed in the command line parser as an acceptable language), but don't know what was supported.

Marko



     
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Chris Lattner

On Mar 9, 2010, at 11:42 PM, brdavs wrote:

>
>> AFAIK, common opencl implementations have various hacks on
>> top of mainline clang, so I doubt clang mainline will be a
>> compliant implementation out of the box.
>
> That's how it appears. For parsing only, what's missing is:
>
> 1. address space qualifiers (eg., __global float*, __constant int,
> __local char*)
> 2. function qualifiers (eg, __kernel void foo(...) ),
> some builtin types (float4, float8, float16, image2d_t, etc.)

These are not magic parser features, in the opencl implementations I'm aware of, these are macros that come from an implicit #include.  These are #defines for various attributes, __global typically turns into __attribute__((address_space.  __kernel turns into attribute(annotate)

> 3. image access qualifiers (eg, __read_only, __write_only,
> __read_write).

I don't know what these do, not familiar enough with OpenCL.

-Chris


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Arnaud Allard de Grandmaison
>> 1. address space qualifiers (eg., __global float*, __constant int,
>> __local char*)
>> 2. function qualifiers (eg, __kernel void foo(...) ),
>> some builtin types (float4, float8, float16, image2d_t, etc.)
>These are not magic parser features, in the opencl implementations I'm
> aware of, these are macros that come from an implicit #include.  These are > #defines for various attributes, __global typically turns into
> __attribute__((address_space.  __kernel turns into attribute(annotate)

Although not magic parser feature per se, in the case of the address_space attributes this requires modifying clang to accept them nicely in the function prototypes. I have a patch for this for llvm/clang-2.6, and I plan to provided a patch for clang-head soon.

--
Arnaud de Grandmaison

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Arnaud Allard de Grandmaison
To be more specific, this is about intrinsic prototypes with address spaces, not function prototypes.
--
Arnaud de Grandmaison
________________________________________
From: [hidden email] [[hidden email]] On Behalf Of Arnaud Allard de Grandmaison [[hidden email]]
Sent: Wednesday, March 10, 2010 7:34 PM
To: [hidden email]
Subject: Re: [cfe-dev] OpenCL parsing / type support

>> 1. address space qualifiers (eg., __global float*, __constant int,
>> __local char*)
>> 2. function qualifiers (eg, __kernel void foo(...) ),
>> some builtin types (float4, float8, float16, image2d_t, etc.)
>These are not magic parser features, in the opencl implementations I'm
> aware of, these are macros that come from an implicit #include.  These are > #defines for various attributes, __global typically turns into
> __attribute__((address_space.  __kernel turns into attribute(annotate)

Although not magic parser feature per se, in the case of the address_space attributes this requires modifying clang to accept them nicely in the function prototypes. I have a patch for this for llvm/clang-2.6, and I plan to provided a patch for clang-head soon.

--
Arnaud de Grandmaison

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Anton Lokhmotov
In reply to this post by brdavs
I am also looking into OpenCL support in Clang, and am learning quite a few
of hacks Chris and Arnaud are referring to.  It would be good, however, to
implement OpenCL features properly, so that setting LangOpts.OpenCL=1 would
just work.  I suggest we make a conscious community-wide effort here so that
various implementations do not reinvent the wheel. (So Arnaud it would be
great if you could contribute to Clang head your patch for address space and
kernel attributes!)

Right now I'm staring into TokenKinds.def.  One issue there is the use of
KEYALL for

a) non-C99 (and hence automatically non-OpenCL) keywords (e.g. __func__)
b) GNU extensions (e.g. _Decimal32, __builtin_choose_expr)
c) MS extensions (e.g. __forceinline and __declspec).  

This means that a Clang-based OpenCL front-end may not accept these as valid
identifiers, although the standard does not list them as keywords!  It's
easy to make existing keywords valid for OpenCL but it's not so easy to
exclude them (e.g. valid in almost all C variants but invalid in OpenCL).

I appreciate many people are busy with preparing for release 2.7 but this is
something we should address sooner rather than later.

Anton L.


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Sebastian Redl

On Thu, 11 Mar 2010 11:17:32 -0000, "Anton Lokhmotov"
<[hidden email]> wrote:
> Right now I'm staring into TokenKinds.def.  One issue there is the use
of
> KEYALL for
>
> a) non-C99 (and hence automatically non-OpenCL) keywords (e.g. __func__)
> b) GNU extensions (e.g. _Decimal32, __builtin_choose_expr)
> c) MS extensions (e.g. __forceinline and __declspec).  
>
> This means that a Clang-based OpenCL front-end may not accept these as
> valid
> identifiers, although the standard does not list them as keywords!

This shouldn't matter. Identifiers containing double underscores or
starting with an underscore and an uppercase letter are not valid
identifiers for user code under any circumstances.

Sebastian

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Arnaud Allard de Grandmaison
In reply to this post by Anton Lokhmotov
I just have a patch for the address space not for kernels.

The goal is to be able to qualify the pointer parameters of intrinsics function with an address space. This is done by adding a number as a postfix to the '*' modifier in clang/Basic/Builtins.def. I have not invented anything there, as I have seen this in a opencl/llvm presentation.

This way, you can now have a memcpy-like function, transferring blocks from address space 0 to 1, with the following declaration in your clang/include/clang/Basic/BuiltinsZZZZ.def:

BUILTIN(__builtin_memxfer, "v*1v*1vC*z", "n")

Clang/head is currently broken (building it fails with 'llvm::DIFile' has not been declared), so I have not been able to test it, but the patch applies and the patched files compile just fine.
--
Arnaud de Grandmaison

-----Original Message-----
From: [hidden email] [mailto:[hidden email]] On Behalf Of Anton Lokhmotov
Sent: Thursday, March 11, 2010 12:18 PM
To: [hidden email]
Subject: Re: [cfe-dev] OpenCL parsing / type support

I am also looking into OpenCL support in Clang, and am learning quite a few
of hacks Chris and Arnaud are referring to.  It would be good, however, to
implement OpenCL features properly, so that setting LangOpts.OpenCL=1 would
just work.  I suggest we make a conscious community-wide effort here so that
various implementations do not reinvent the wheel. (So Arnaud it would be
great if you could contribute to Clang head your patch for address space and
kernel attributes!)

Right now I'm staring into TokenKinds.def.  One issue there is the use of
KEYALL for

a) non-C99 (and hence automatically non-OpenCL) keywords (e.g. __func__)
b) GNU extensions (e.g. _Decimal32, __builtin_choose_expr)
c) MS extensions (e.g. __forceinline and __declspec).

This means that a Clang-based OpenCL front-end may not accept these as valid
identifiers, although the standard does not list them as keywords!  It's
easy to make existing keywords valid for OpenCL but it's not so easy to
exclude them (e.g. valid in almost all C variants but invalid in OpenCL).

I appreciate many people are busy with preparing for release 2.7 but this is
something we should address sooner rather than later.

Anton L.


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev

intrinsics+addrSpace.patch (1K) Download Attachment
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Chris Lattner
In reply to this post by brdavs

On Mar 11, 2010, at 3:17 AM, Anton Lokhmotov wrote:

> I am also looking into OpenCL support in Clang, and am learning quite a few
> of hacks Chris and Arnaud are referring to.  It would be good, however, to
> implement OpenCL features properly, so that setting LangOpts.OpenCL=1 would
> just work.  I suggest we make a conscious community-wide effort here so that
> various implementations do not reinvent the wheel. (So Arnaud it would be
> great if you could contribute to Clang head your patch for address space and
> kernel attributes!)

I'm fine with improving infrastructure to support this, but there is no reason to add first class language support for things like __kernel.  It just makes the compiler more complex for no reason.  If you're interested in driving this, a better thing to focus on would be to provide an implementation of the *implicit include* which provides the #defines etc.

-Chris

>
> Right now I'm staring into TokenKinds.def.  One issue there is the use of
> KEYALL for
>
> a) non-C99 (and hence automatically non-OpenCL) keywords (e.g. __func__)
> b) GNU extensions (e.g. _Decimal32, __builtin_choose_expr)
> c) MS extensions (e.g. __forceinline and __declspec).  
>
> This means that a Clang-based OpenCL front-end may not accept these as valid
> identifiers, although the standard does not list them as keywords!  It's
> easy to make existing keywords valid for OpenCL but it's not so easy to
> exclude them (e.g. valid in almost all C variants but invalid in OpenCL).
>
> I appreciate many people are busy with preparing for release 2.7 but this is
> something we should address sooner rather than later.
>
> Anton L.
>
>
> _______________________________________________
> cfe-dev mailing list
> [hidden email]
> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

John McCall
In reply to this post by Arnaud Allard de Grandmaison
On Mar 11, 2010, at 7:04 AM, Arnaud Allard de Grandmaison wrote:
> This way, you can now have a memcpy-like function, transferring blocks from address space 0 to 1, with the following declaration in your clang/include/clang/Basic/BuiltinsZZZZ.def:
>
> BUILTIN(__builtin_memxfer, "v*1v*1vC*z", "n")

Seems reasonable to me;  I'll go ahead and commit this when I get a chance.

> Clang/head is currently broken (building it fails with 'llvm::DIFile' has not been declared), so I have not been able to test it, but the patch applies and the patched files compile just fine.

DIFile was added to LLVM on Monday in r97994.  In general Clang/head is tightly coupled to LLVM/head, and you should assume that missing declarations (particularly in the llvm:: namespace) are due to such changes.

John.
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Anton Lokhmotov
In reply to this post by Sebastian Redl
> From: Chris Lattner
> Sent: 11 March 2010 18:05
>
> I'm fine with improving infrastructure to support this, but there is no
> reason to add first class language support for things like __kernel.
> It just makes the compiler more complex for no reason.  If you're
> interested in driving this, a better thing to focus on would be to
> provide an implementation of the *implicit include* which provides the
> #defines etc.
Yes, I'm interested in driving this.  I didn't mean to suggest that an implicit header is necessarily a bad idea.  I agree that the things that can be supported this way, such as vector types and address (and hopefully function) qualifiers, should be supported this way.  Where do you think such a header should live (e.g. which subdirectory of Clang)?

It appears, however, that certain things cannot be supported only with headers.  For example, the OpenCL standard restricts usage of the half type (16-bit floating point) to declaring a pointer to a buffer that contains half values; such pointers cannot be dereferenced.  In other words, half behaves similar to void.  However, sizeof(half) must be 2.  (Obviously, if someone knows an ingenious quirk to solve this conundrum, I'd be interested to hear.)

The standard also specifies many restrictions which the compiler should check as semantic analyses, possibly in a few separate passes.  Also, there are some alignment restrictions which perhaps call for a special OpenCL triple (since the target should be OpenCL-compliant)?

I'm new to Clang, so I'd appreciate your thoughts on how to implement OpenCL support properly.


> From: Sebastian Redl
> Sent: 11 March 2010 14:34

> > One issue there is the use of KEYALL for
> > a) non-C99 (and hence automatically non-OpenCL) keywords (e.g. __func__)
> > b) GNU extensions (e.g. _Decimal32, __builtin_choose_expr)
> > c) MS extensions (e.g. __forceinline and __declspec).
> >
> > This means that a Clang-based OpenCL front-end may not accept these
> > as valid identifiers, although the standard does not list them as keywords!
>
> This shouldn't matter. Identifiers containing double underscores or
> starting with an underscore and an uppercase letter are not valid
> identifiers for user code under any circumstances.

Perhaps I've exaggerated the problem a bit but still extensions' keywords should not be listed as KEYALL ;).

Cheers,
Anton.



_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Chris Lattner
In reply to this post by Sebastian Redl
On Mar 12, 2010, at 2:59 AM, Anton Lokhmotov wrote:

>> From: Chris Lattner
>> Sent: 11 March 2010 18:05
>>
>> I'm fine with improving infrastructure to support this, but there is no
>> reason to add first class language support for things like __kernel.
>> It just makes the compiler more complex for no reason.  If you're
>> interested in driving this, a better thing to focus on would be to
>> provide an implementation of the *implicit include* which provides the
>> #defines etc.
> Yes, I'm interested in driving this.  I didn't mean to suggest that an implicit header is necessarily a bad idea.  I agree that the things that can be supported this way, such as vector types and address (and hopefully function) qualifiers, should be supported this way.  Where do you think such a header should live (e.g. which subdirectory of Clang)?

clang/lib/Headers would make sense to me.

> It appears, however, that certain things cannot be supported only with headers.  For example, the OpenCL standard restricts usage of the half type (16-bit floating point) to declaring a pointer to a buffer that contains half values; such pointers cannot be dereferenced.  In other words, half behaves similar to void.  However, sizeof(half) must be 2.  (Obviously, if someone knows an ingenious quirk to solve this conundrum, I'd be interested to hear.)

You're right, 'half float' is something that clang will need to support directly.  I think that Anton has some patches coming to add at least backend support for half float.
>
> The standard also specifies many restrictions which the compiler should check as semantic analyses, possibly in a few separate passes.  Also, there are some alignment restrictions which perhaps call for a special OpenCL triple (since the target should be OpenCL-compliant)?

Yep, sounds right.  These checks should be based on the existing LangOptions.OpenCL bit, and that bit should be set when using an OpenCL target triple.  I'm not sure how much of that is wired up.

-Chris
_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Arnaud Allard de Grandmaison
In reply to this post by John McCall
Thanks John. I have seen you had time to improve the patch as well.

This is also what I assume, as I update llvm and clang at the same time, but for historical reasons, I am working with a git clone for llvm, and an svn repository for clang. Yes, this is _ugly_. And this can create such discrepancies from time to time. Just dreaming of the day when it will be possible to directly clone llvm and clang form an official llvm.org git repository... This would definitely help me to maintain my company branch and help as well the patch flow.

--
Arnaud de Grandmaison

-----Original Message-----
From: John McCall [mailto:[hidden email]]
Sent: Friday, March 12, 2010 1:20 AM
To: Arnaud Allard de Grandmaison
Cc: [hidden email]
Subject: Re: [cfe-dev] OpenCL parsing / type support

On Mar 11, 2010, at 7:04 AM, Arnaud Allard de Grandmaison wrote:
> This way, you can now have a memcpy-like function, transferring blocks from address space 0 to 1, with the following declaration in your clang/include/clang/Basic/BuiltinsZZZZ.def:
>
> BUILTIN(__builtin_memxfer, "v*1v*1vC*z", "n")

Seems reasonable to me;  I'll go ahead and commit this when I get a chance.

> Clang/head is currently broken (building it fails with 'llvm::DIFile' has not been declared), so I have not been able to test it, but the patch applies and the patched files compile just fine.

DIFile was added to LLVM on Monday in r97994.  In general Clang/head is tightly coupled to LLVM/head, and you should assume that missing declarations (particularly in the llvm:: namespace) are due to such changes.

John.

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
Reply | Threaded
Open this post in threaded view
|

Re: OpenCL parsing / type support

Sumesh Udayakumaran
On a related topic,  clang does not seem to accept  vector compares - If a < b where a, b are of vector type. However, the assembly language does allow result of vector types for compares. So is this unsupported yet ?  Open CL also allows for comparison of vectors. 


thanks
Sumesh



On Mon, Mar 15, 2010 at 2:13 AM, Arnaud Allard de Grandmaison <[hidden email]> wrote:
Thanks John. I have seen you had time to improve the patch as well.

This is also what I assume, as I update llvm and clang at the same time, but for historical reasons, I am working with a git clone for llvm, and an svn repository for clang. Yes, this is _ugly_. And this can create such discrepancies from time to time. Just dreaming of the day when it will be possible to directly clone llvm and clang form an official llvm.org git repository... This would definitely help me to maintain my company branch and help as well the patch flow.

--
Arnaud de Grandmaison

-----Original Message-----
From: John McCall [mailto:[hidden email]]
Sent: Friday, March 12, 2010 1:20 AM
To: Arnaud Allard de Grandmaison
Cc: [hidden email]
Subject: Re: [cfe-dev] OpenCL parsing / type support

On Mar 11, 2010, at 7:04 AM, Arnaud Allard de Grandmaison wrote:
> This way, you can now have a memcpy-like function, transferring blocks from address space 0 to 1, with the following declaration in your clang/include/clang/Basic/BuiltinsZZZZ.def:
>
> BUILTIN(__builtin_memxfer, "v*1v*1vC*z", "n")

Seems reasonable to me;  I'll go ahead and commit this when I get a chance.

> Clang/head is currently broken (building it fails with 'llvm::DIFile' has not been declared), so I have not been able to test it, but the patch applies and the patched files compile just fine.

DIFile was added to LLVM on Monday in r97994.  In general Clang/head is tightly coupled to LLVM/head, and you should assume that missing declarations (particularly in the llvm:: namespace) are due to such changes.

John.

_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev


_______________________________________________
cfe-dev mailing list
[hidden email]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev