[cfe-commits] [llvm-commits] Patch: Add get_global_id builtin/intrinsic

[Tom Stellard]

On Fri, Feb 03, 2012 at 10:07:03PM +0000, Rotem, Nadav wrote:
> Tom, 
> 
> Our OpenCL implementation of get_global_id is not target specific and we don't resolve it in the backend. I think that get_global_id should be implemented as a simple library call.  However, this is something that needs to be discussed with Tanya Lattner, Peter Collingbourne, Anton Lokhmotov, etc. 
> 
> Nadav

Nadav,

Sorry, I guess I should be more clear.  When I say target specific I'm
talking about GPU targets.  The get_global_id() implementation
on the GPUs we've written a backend for (Evergreen, Northern Islands)
requires reading values from special registers that are preloaded by the
hardware.  I'm guessing other GPUs do something similar, so I do think
it is something that would need to be resolved in the backend.

-Tom

> 
> -----Original Message-----
> From: Tom Stellard [mailto:thomas.stellard at amd.com] 
> Sent: Friday, February 03, 2012 23:56
> To: Rotem, Nadav
> Cc: llvm-commits at cs.uiuc.edu; cfe-commits at cs.uiuc.edu
> Subject: Re: [llvm-commits] Patch: Add get_global_id builtin/intrinsic
> 
> Hi Nadav,
> 
> On Fri, Feb 03, 2012 at 09:26:14PM +0000, Rotem, Nadav wrote:
> > Tom,
> > 
> > I can see a number of problems with this patch.  First, get_global_id does not return llvm_i32_ty on 64bit systems. It returns size_t (which is i64 on x86_64).
> 
> Ok, so I guess I should use llvm_anyint_ty instead?
> 
> 
> >Second, I am not sure that this is the correct approach for implementing this. What's wrong with having get_global_id as a standard library call ? Why do we need a special intrinsic for it ?
> 
> I think it should be an intrinsic because the implementation of it is always target specific, and it seems like it would be better to have the target specific implementations live in the backends rather than a common library.
> 
> Thanks,
> Tom
> > 
> > Thanks,
> > Nadav
> > 
> 
> 
> > -----Original Message-----
> > From: llvm-commits-bounces at cs.uiuc.edu [mailto:llvm-commits-bounces at cs.uiuc.edu] On Behalf Of Tom Stellard
> > Sent: Friday, February 03, 2012 23:19
> > To: llvm-commits at cs.uiuc.edu; cfe-commits at cs.uiuc.edu
> > Subject: [llvm-commits] Patch: Add get_global_id builtin/intrinsic
> > 
> > Hi,
> > 
> > I've attached two patches, one for llvm and one for clang that add support for the OpenCL C builtin function get_global_id().  I would like to eventually add support for more OpenCL builtins to clang/llvm, but this initial patch is just to make sure I'm doing it the right way.
> > 
> > Please review.
> > 
> > Thanks,
> > Tom Stellard
> > 
> > ---------------------------------------------------------------------
> > Intel Israel (74) Limited
> > 
> > This e-mail and any attachments may contain confidential material for
> > the sole use of the intended recipient(s). Any review or distribution
> > by others is strictly prohibited. If you are not the intended
> > recipient, please contact the sender and delete all copies.
> > 
> > 
> 
> ---------------------------------------------------------------------
> Intel Israel (74) Limited
> 
> This e-mail and any attachments may contain confidential material for
> the sole use of the intended recipient(s). Any review or distribution
> by others is strictly prohibited. If you are not the intended
> recipient, please contact the sender and delete all copies.
> 
>