Ping for push.
-----Original Message-----
From: Beignet [mailto:beignet-boun...@lists.freedesktop.org] On Behalf Of He
Junyan
Sent: Thursday, December 3, 2015 12:36 PM
To: beignet@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...@inbox.com>
> To: beignet@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....@intel.com>
> > To: beignet@lists.freedesktop.org
> > Cc: xiuli....@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....@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@lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
>
>
> _______________________________________________
> Beignet mailing list
> Beignet@lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
