Re: Ideas for Google Summer of Code

Re: Ideas for Google Summer of Code

[Ross Ridge]

Paolo Bonzini writes:
>Regarding the NVIDIA GPU backend, I think NVIDIA is not yet distributing
>details about the instruction set unlike ATI, is it?  In this case, I
>think ATI would be a better match.

I think a GPU backend would be well beyond the scope of a Summer of
Code project.  GPUs don't have normal stacks and addressing support
is limitted. 

>Another possibility is to analyze OpenCL C and try to integrate its
>features in GCC as much as possible.  This would include
>
>1) masking/swizzling support for GCC's "generic vector" language extension;

A project that started and ended here would give GCC, in particular
GCC's Cell SPU port, the only major required functionality in the OpenCL
language, outside the runtime, that GCC is missing.

>2) half-precision floats;

Do you mean just conversion only support, like Sandra Loosemore's
proposed ARM patch, or full arithmetic support like any other scalar or
vector type?

					Ross Ridge

Re: Ideas for Google Summer of Code

[Joseph S. Myers]

On Mon, 30 Mar 2009, Ross Ridge wrote:

> >2) half-precision floats;
> 
> Do you mean just conversion only support, like Sandra Loosemore's
> proposed ARM patch, or full arithmetic support like any other scalar or
> vector type?

Arithmetic via converting to float, doing arithmetic on that and 
converting back to half-precision gives you the same semantics as IEEE 
arithmetic directly on half-precision, except for the special case of 
converting double to half-precision where you can get double rounding if 
you go via float.  (This is because the precision of float is strictly 
more than twice that of half-precision.  Of course the above only applies 
to single arithmetic operations; more complicated combinations would need 
truncation to half-precision at each step rather than doing everything in 
float.)

So if you have hardware float operations but not hardware half-precision 
you probably want to use those hardware float operations for arithmetic 
rather than having any separate half-precision arithmetic implementation.  
There are at least three possibilities for semantics: the ARM semantics 
where promotion to float happens at the language level; full IEEE 
arithmetic with every operation having a result in the range and precision 
of half-precision; or C99 excess precision where the semantic type is 
half-precision but the values are expressed to the range and precision of 
float.

-- 
Joseph S. Myers
joseph@codesourcery.com

Re: Ideas for Google Summer of Code

[Joe Buck]

On Mon, Mar 30, 2009 at 01:36:14PM -0700, Ross Ridge wrote:
> Joe Buck writes:
> >I'm having trouble finding that document, I don't see a link to it
> >on that page.  Maybe I'm missing something obvious?
> 
> Sticking "nvidia ptx" into Google turned up this document:
> 
>         http://www.nvidia.com/object/io_1195170102263.html
> 
> It's an intermediate language, so isn't tied to any particular NVIDIA GPU.
> I beleive there's something similar for AMD/ATI GPUs.
> 
> btw. The computational power of Intel's integrated GPUs is pretty dismal,
> so I don't think GCC port targetting them would be very useful.

True for what they are shipping now, but Larrabee should change that.

Re: Ideas for Google Summer of Code

[phil++]

Paolo,

Thanks for the feedback, I am not very experienced in compilers so it
is hard to judge how long a task will take...

By sharing I meant sharing of code between NVIDIA and GCC.  It
probably won't happen I guess.

Here is my proposal for an OpenCL runtime with a target runtime as
well.  If you think it is too ambitious or not
ambitious enough, I will change it.
=================================================================================

Project Title:
Make the OpenCL Platform Layer API and Runtime API for the Cell
Processor and CPUs.

Project Synopsis:
The aim of this project is to create an implementation that supports
the Platform Layer API
and Runtime API of OpenCL that can target the Cell Processor and CPUs.
The Platform
Layer API is:
-A hardware abstraction layer over diverse computational resources
-An interface to query, select and initialize compute devices
-An interface to create compute devices and work-queues
The Runtime API is:
-Able to execute compute kernels
-Able to manage scheduling, compute and memory resources (I am
confused as to the wording
of this, does it mean: manage the scheduling of compute and memory resources?)
(Source http://www.khronos.org/developers/library/overview/opencl_overview.pdf,
page 13).

This project will use the existing gcc and ppu-gcc/spu-gcc compilers for offline
compilation of binary programs.

Project Details:
(Part 1)
In this project I will make a C library and runtime that supports some
of the functions listed
here: http://www.khronos.org/registry/cl/api/1.0/cl.h
Specifically I will add support for:
clGetPlatformInfo - Get info about OpenCL
clGetDeviceIDs - Get what devices are supported on system
clGetDeviceInfo - Get info about a specific device
clCreateContext - Create an OpenCL context
clReleaseContext - Release an OpenCL context
clCreateCommandQueue - Create a command-queue on a specific device
clReleaseCommandQueue - Release a command-queue
clCreateBuffer - Create a buffer object
clEnqueueReadBuffer - Enqueue a read
clEnqueueWriteBuffer - Enqueue a write
clCreateProgramWithBinary - Create a program object from a pre-compiled binary.
clReleaseProgram - Release a program object
clCreateKernel - Create a kernel object
clReleaseKernel - Release a kernel object
clSetKernelArg - Set the kernel arguments
clEnqueueNDRangeKernel - Enqueue a command to execute a kernel on a device
clEnqueueTask - Enqueue a single work item
clWaitForEvents - Wait for events to complete
clReleaseEvent - Release an event

This will allow for rudimentary launches of CPU and Cell kernels in a
common interface.  Any functions
that are required for the above to work will also be added.  The
OpenCL compiler will not be implemented.

(Part 2)
Also, a runtime library for the target (CPU or Cell) must be created
that includes the following intrinsics:
Information Functions: (section 6.11.1 of
http://www.khronos.org/registry/cl/specs/opencl-1.0.33.pdf)
uint get_work_dim ()
size_t get_global_size (uint dimindx)
size_t get_global_id (uint dimindx)
size_t get_local_size (uint dimindx)
size_t get_local_id (uint dimindx)
size_t get_num_groups (uint dimindx)
size_t get_group_id (uint dimindx)
Synchronization Functions: (sections 6.11.9 - 6.11.10 of
http://www.khronos.org/registry/cl/specs/opencl-1.0.33.pdf)
void barrier (cl_mem_fence_flags flags)
void mem_fence (cl_mem_fence_flags flags)
void read_mem_fence (cl_mem_fence_flags flags)
void write_mem_fence (cl_mem_fence_flags flags)
Async Copies to/from Memory: (section 6.11.11 of
http://www.khronos.org/registry/cl/specs/opencl-1.0.33.pdf)
event_t async_work_group_copy (__local gentype *dst, const __global
gentype *src,size_t num_elements, event_t event)
event_t async_work_group_copy (__global gentype *dst,const __local
gentype *src,size_t num_elements, event_t event)
void wait_group_events (int num_events, event_t *event_list)

Developer Details:
I am a Ph.D. student in Computer Science at Syracuse University in
Syracuse, NY, USA.  I have been
programming prefessionally in C since I was about 15 years old (My
first job was a Win32 device driver
for a furnace controller ISA add-in card.  This device driver has run
for 10 years, 351 days a year, 24 hours
a day.  There have been no bugs or updates ever to the driver).  I
have lately been fixing bugs in open
source software and I have been on the mailing list of gcc for about 6
months.  In short, programming
is my life, I love it.

Experience with the Cell Processor:
I have been working with the Cell Processor on the playstation 3 for
about a year now.  I am familiar
with using ppu-gcc and spu-gcc, creation of shared memory between the
local stores and main memory,
explicit mfc_put, mfc_get calls, mailboxes, etc.

Success Criteria:
1. There is a working OpenCL runtime for the CPU and Cell Processor
that supports the limited functionality above (Part 1)
2. There is a working runtime library for the target that supports the
limitied functionality above (Part 2)

Road Map:
April 20th - Start getting up to speed and get to know mentors.  Prepare design.
May 23rd - Begin coding.
June 30th - Target runtime library is done.
July 30th - OpenCL runtime is done.
August 10th - Support Documentation is done.

Benefit to GCC:
This project will bring GCC a step closer to supporting OpenCL for the
CPU and Cell
targets.  The target runtime library will be an intermediate
deliverable that can be
used until the OpenCL C compiler is done.
=========================================================================================

Thanks for all of the support and interest so far everyone!

Sincerely,
Phil Pratt-Szeliga

Re: Ideas for Google Summer of Code

[Ross Ridge]

Joe Buck writes:
>I'm having trouble finding that document, I don't see a link to it
>on that page.  Maybe I'm missing something obvious?

Sticking "nvidia ptx" into Google turned up this document:

	http://www.nvidia.com/object/io_1195170102263.html

It's an intermediate language, so isn't tied to any particular NVIDIA GPU.
I beleive there's something similar for AMD/ATI GPUs.

btw. The computational power of Intel's integrated GPUs is pretty dismal,
so I don't think GCC port targetting them would be very useful.

					Ross Ridge

Re: Ideas for Google Summer of Code

[Philip Pratt-Szeliga]

Joe Buck,

You have to choose an operating system and then a table appears below
it to download the CUDA toolkit.

Phil



On Mon, Mar 30, 2009 at 1:32 PM, Joe Buck <Joe.Buck@synopsys.com> wrote:
> On Mon, Mar 30, 2009 at 10:10:24AM -0700, phil++ wrote:
>> Hi All,
>>
>> Thanks for the feedback.  As far as the NVIDIA ISA details, I have the
>> pdf that describes it.  In the CUDA toolkit, available on NVIDIA's
>> website there is a pdf named ptx_isa_1.3.pdf under the doc directory.
>> (Get it from here http://www.nvidia.com/object/cuda_get.html).
>
> I'm having trouble finding that document, I don't see a link to it
> on that page.  Maybe I'm missing something obvious?
>

Re: Ideas for Google Summer of Code

[Paolo Bonzini]

> So we can do Intel, ATI and NVIDIA GPU backends.  NVIDIA already has
> an implementation of OpenCL working.
> http://www.nvidia.com/object/cuda_opencl.html.  Would there be any
> sharing involved with them??

If you mean between backends, 1) do not underestimate the time needed to
write a new GCC backend; 2) probably nothing can be shared (see for
example the zero-sharing between PPU and SPU).  I think that your 2) and
3) projects are way more viable.

> I am working on my proposal now and I will post it to this list before
> final submission (I've got to hurry, they are due April 3rd).  I will
> mainly focus on this pdf:
> http://www.khronos.org/developers/library/overview/opencl_overview.pdf,

Note that this *is* different from the backends you mentioned above, and
as I said I think it is more viable.

Make sure that what you propose is implementable without an OpenCL C
compiler (I think it is), or discuss to what extent the library will be
functional.  For example, do you need another runtime library
implementing the intrinsics used by kernels?

Thanks,

Paolo

Re: Ideas for Google Summer of Code

[Joe Buck]

On Mon, Mar 30, 2009 at 10:10:24AM -0700, phil++ wrote:
> Hi All,
> 
> Thanks for the feedback.  As far as the NVIDIA ISA details, I have the
> pdf that describes it.  In the CUDA toolkit, available on NVIDIA's
> website there is a pdf named ptx_isa_1.3.pdf under the doc directory.
> (Get it from here http://www.nvidia.com/object/cuda_get.html).

I'm having trouble finding that document, I don't see a link to it
on that page.  Maybe I'm missing something obvious?

Re: Ideas for Google Summer of Code

[phil++]

Hi All,

Thanks for the feedback.  As far as the NVIDIA ISA details, I have the
pdf that describes it.  In the CUDA toolkit, available on NVIDIA's
website there is a pdf named ptx_isa_1.3.pdf under the doc directory.
(Get it from here http://www.nvidia.com/object/cuda_get.html).

So we can do Intel, ATI and NVIDIA GPU backends.  NVIDIA already has
an implementation of OpenCL working.
http://www.nvidia.com/object/cuda_opencl.html.  Would there be any
sharing involved with them??

I am working on my proposal now and I will post it to this list before
final submission (I've got to hurry, they are due April 3rd).  I will
mainly focus on this pdf:
http://www.khronos.org/developers/library/overview/opencl_overview.pdf,
page 13

Phil Pratt-Szeliga


On Mon, Mar 30, 2009 at 12:39 PM, Joe Buck <Joe.Buck@synopsys.com> wrote:
> On Mon, Mar 30, 2009 at 01:09:56AM -0700, Paolo Bonzini wrote:
>> > I am a PhD student who has been working with CUDA for the GPU and also
>> > gcc for Cell BE for about a year now. (By work I mean developing
>> > applications).  I am looking to bring GCC closer to being able to
>> > support OpenCL as a Google Summer of Code.
>>
>> This is very interesting and I'm willing to help with mentoring.
>> However I think your projects are more ahead than what is actually in
>> GCC right now!
>>
>> Regarding the NVIDIA GPU backend, I think NVIDIA is not yet distributing
>> details about the instruction set unlike ATI, is it?  In this case, I
>> think ATI would be a better match.
>
> Intel GPUs as well.  I think that at this stage, there's only sufficient
> documentation available to do a GPU back end for Intel and ATI, not nVidia
> (though I don't know if any progress has been made via reverse
> engineering, perhaps by the nouveau project).  I think anyone who pursues
> this angle would be well-advised to look at more than one architecture.
>

Re: Ideas for Google Summer of Code

[Joe Buck]

On Mon, Mar 30, 2009 at 01:09:56AM -0700, Paolo Bonzini wrote:
> > I am a PhD student who has been working with CUDA for the GPU and also
> > gcc for Cell BE for about a year now. (By work I mean developing
> > applications).  I am looking to bring GCC closer to being able to
> > support OpenCL as a Google Summer of Code.
> 
> This is very interesting and I'm willing to help with mentoring.
> However I think your projects are more ahead than what is actually in
> GCC right now!
> 
> Regarding the NVIDIA GPU backend, I think NVIDIA is not yet distributing
> details about the instruction set unlike ATI, is it?  In this case, I
> think ATI would be a better match.

Intel GPUs as well.  I think that at this stage, there's only sufficient
documentation available to do a GPU back end for Intel and ATI, not nVidia
(though I don't know if any progress has been made via reverse
engineering, perhaps by the nouveau project).  I think anyone who pursues
this angle would be well-advised to look at more than one architecture.

Re: Ideas for Google Summer of Code

[Paolo Bonzini]

Joseph S. Myers wrote:
> On Mon, 30 Mar 2009, Paolo Bonzini wrote:
> 
>> 2) half-precision floats;
> 
> I expect Sandra will be merging half-precision support for ARM to trunk 
> fairly soon; see her announcement 
> <http://gcc.gnu.org/ml/gcc-patches/2009-01/msg01062.html>.  This includes 
> generic support for the IEEE half-precision format (plus an ARM-specific 
> variant) and lots of testcases that are likely to be useful for any other 
> targets wanting half-precision as well, as well as ARM-specific pieces 
> including the Cortex-A9 half-precision instructions.

Great news!  I suppose this should be an incentive for the OpenCL C
project then. :-)

Paolo

Re: Ideas for Google Summer of Code

[Joseph S. Myers]

On Mon, 30 Mar 2009, Paolo Bonzini wrote:

> 2) half-precision floats;

I expect Sandra will be merging half-precision support for ARM to trunk 
fairly soon; see her announcement 
<http://gcc.gnu.org/ml/gcc-patches/2009-01/msg01062.html>.  This includes 
generic support for the IEEE half-precision format (plus an ARM-specific 
variant) and lots of testcases that are likely to be useful for any other 
targets wanting half-precision as well, as well as ARM-specific pieces 
including the Cortex-A9 half-precision instructions.

-- 
Joseph S. Myers
joseph@codesourcery.com

Re: Ideas for Google Summer of Code

[Paolo Bonzini]

> I am a PhD student who has been working with CUDA for the GPU and also
> gcc for Cell BE for about a year now. (By work I mean developing
> applications).  I am looking to bring GCC closer to being able to
> support OpenCL as a Google Summer of Code.

This is very interesting and I'm willing to help with mentoring.
However I think your projects are more ahead than what is actually in
GCC right now!

Regarding the NVIDIA GPU backend, I think NVIDIA is not yet distributing
details about the instruction set unlike ATI, is it?  In this case, I
think ATI would be a better match.

Another possibility is to analyze OpenCL C and try to integrate its
features in GCC as much as possible.  This would include

1) masking/swizzling support for GCC's "generic vector" language extension;

2) half-precision floats;

3) having all the rounding conversion operators of OpenCL C as builtins.
 Subproject: the i386 has a lot of fast implementations for
floating-point rounding operators (things like x+2^52-2^52), move them
to target-independent code.

4) ???


I'm not extremely familiar with OpenCL, but if I recall correctly you're
right that having a runtime for it does not strictly *need* an OpenCL C
compiler.  However, if you want to pursue the OpenCL runtime project,
you should make sure to specify in the proposal how the runtime and
compiler relate (or why they are independent).

In any case, I would be happy to mentor you. :-)

Paolo

Re: Ideas for Google Summer of Code

[Daniel Kraft]

Hi,

phil++ wrote:
> I am a PhD student who has been working with CUDA for the GPU and also
> gcc for Cell BE for about a year now. (By work I mean developing
> applications).  I am looking to bring GCC closer to being able to
> support OpenCL as a Google Summer of Code.  Here are some of my ideas:
> 
> 1. Make an NVIDIA GPU backend
> 2. Make the OpenCL runtime for the Cell Processor
> 3. Make the OpenCL runtime for a homogeneous processor

I'm no mentor and thus don't take my comment too seriously, but I've 
started to work with CUDA recently (and may be doing so for my 
Bachelor's thesis) and would *love* to see some gcc support in this 
direction!

Regarding your points, I tend to favour point 1.  Thinking about it, 
when adding support for Co-Arrays to the Fortran front-end, maybe we 
could allow NVIDIA cards as backend for Co-Array'ed programs?  In 
general, it could be interesting to allow, for instance, writing kernels 
in Fortran, so existing code could be reused (or for the ones that like 
Fortran :D).  Maybe some basic support for that could be achieved, too? 
  (That is, gcc compiled kernels in some way.)

Yours,
Daniel

Ideas for Google Summer of Code

[phil++]

Hello All,

I am a PhD student who has been working with CUDA for the GPU and also
gcc for Cell BE for about a year now. (By work I mean developing
applications).  I am looking to bring GCC closer to being able to
support OpenCL as a Google Summer of Code.  Here are some of my ideas:

1. Make an NVIDIA GPU backend
2. Make the OpenCL runtime for the Cell Processor
3. Make the OpenCL runtime for a homogeneous processor

What idea is most valuable to the gcc community?  Are there any other
more valuable ideas?

About me:  I have been reading the gcc mailing list for about 6 months
now, I have submitted my first patch to a simple package that comes
with Ubuntu (glife).  I have been programming in C for about 10 years
and C++ for less.

Phil Pratt-Szeliga