[Beignet] [PATCH OCL2.0 1/2] libocl: Add three work-item built-in function

[Pan, Xiuli]

Ping for push.

-----Original Message-----
From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of He Junyan
Sent: Thursday, December 3, 2015 12:36 PM
To: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH OCL2.0 1/2] libocl: Add three work-item built-in function

I think that it can also be merged to master.

On Thu, Dec 03, 2015 at 12:34:26PM +0800, He Junyan wrote:
> Date: Thu, 3 Dec 2015 12:34:26 +0800
> From: He Junyan <junyan.he at inbox.com>
> To: beignet at lists.freedesktop.org
> Subject: Re: [Beignet] [PATCH OCL2.0 1/2] libocl: Add three work-item  
> built-in function
> 
> LGTM, thanks.
> 
> On Wed, Dec 02, 2015 at 04:57:38PM +0800, Pan Xiuli wrote:
> > Date: Wed,  2 Dec 2015 16:57:38 +0800
> > From: Pan Xiuli <xiuli.pan at intel.com>
> > To: beignet at lists.freedesktop.org
> > Cc: xiuli.pan at intel.com
> > Subject: [Beignet] [PATCH OCL2.0 1/2] libocl: Add three work-item 
> > built-in  function
> > X-Mailer: git-send-email 2.1.4
> > 
> > Add get global/local linear id by calculate with global/local id, 
> > size and offset. The get_queue_local_size() and get_loal_size() 
> > should be different when the global work group size is not uniform, 
> > but now they are the same. We will refine these functions when we 
> > support non-uniform work-group size.
> > 
> > Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
> > ---
> >  backend/src/libocl/include/ocl_workitem.h |  3 +++
> >  backend/src/libocl/src/ocl_workitem.cl    | 30 ++++++++++++++++++++++++++++++
> >  2 files changed, 33 insertions(+)
> > 
> > diff --git a/backend/src/libocl/include/ocl_workitem.h 
> > b/backend/src/libocl/include/ocl_workitem.h
> > index 84bb1fb..c3b0bdb 100644
> > --- a/backend/src/libocl/include/ocl_workitem.h
> > +++ b/backend/src/libocl/include/ocl_workitem.h
> > @@ -24,9 +24,12 @@ OVERLOADABLE uint get_work_dim(void);  
> > OVERLOADABLE uint get_global_size(uint dimindx);  OVERLOADABLE uint 
> > get_global_id(uint dimindx);  OVERLOADABLE uint get_local_size(uint 
> > dimindx);
> > +OVERLOADABLE uint get_enqueued_local_size(uint dimindx);
> >  OVERLOADABLE uint get_local_id(uint dimindx);  OVERLOADABLE uint 
> > get_num_groups(uint dimindx);  OVERLOADABLE uint get_group_id(uint 
> > dimindx);  OVERLOADABLE uint get_global_offset(uint dimindx);
> > +OVERLOADABLE uint get_global_linear_id(void); OVERLOADABLE uint 
> > +get_local_linear_id(void);
> >  
> >  #endif  /* __OCL_WORKITEM_H__ */
> > diff --git a/backend/src/libocl/src/ocl_workitem.cl 
> > b/backend/src/libocl/src/ocl_workitem.cl
> > index 6ddc406..235f12b 100644
> > --- a/backend/src/libocl/src/ocl_workitem.cl
> > +++ b/backend/src/libocl/src/ocl_workitem.cl
> > @@ -55,3 +55,33 @@ DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)  
> > OVERLOADABLE uint get_global_id(uint dim) {
> >    return get_local_id(dim) + get_local_size(dim) * 
> > get_group_id(dim) + get_global_offset(dim);  }
> > +
> > +OVERLOADABLE uint get_enqueued_local_size (uint dimindx) {
> > +  //TODO: should be different with get_local_size when support
> > +  //non-uniform work-group size
> > +  return get_local_size(dimindx);
> > +}
> > +
> > +OVERLOADABLE uint get_global_linear_id(void) {
> > +  uint dim = __gen_ocl_get_work_dim();
> > +  if (dim == 1) return get_global_id(0) - get_global_offset(0);
> > +  else if (dim == 2) return (get_global_id(1) - get_global_offset(1)) * get_global_size(0) +
> > +                             get_global_id(0) 
> > +-get_global_offset(0);
> > +  else if (dim == 3) return ((get_global_id(2) - get_global_offset(2)) *
> > +                              get_global_size(1) * get_global_size(0)) +
> > +                            ((get_global_id(1) - get_global_offset(1)) * get_global_size (0)) +
> > +                             (get_global_id(0) - 
> > +get_global_offset(0));
> > +  else return 0;
> > +}
> > +
> > +OVERLOADABLE uint get_local_linear_id(void) {
> > +  uint dim = __gen_ocl_get_work_dim();
> > +  if (dim == 1) return get_local_id(0);
> > +  else if (dim == 2) return get_local_id(1) * get_local_size (0) + 
> > +get_local_id(0);
> > +  else if (dim == 3) return (get_local_id(2) * get_local_size(1) * get_local_size(0)) +
> > +                            (get_local_id(1) * get_local_size(0)) + 
> > +get_local_id(0);
> > +  else return 0;
> > +}
> > --
> > 2.1.4
> > 
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
> 
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet