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:[email protected]] > Sent: Friday, February 03, 2012 23:56 > To: Rotem, Nadav > Cc: [email protected]; [email protected] > 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: [email protected] > > [mailto:[email protected]] On Behalf Of Tom Stellard > > Sent: Friday, February 03, 2012 23:19 > > To: [email protected]; [email protected] > > 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. > > _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
