Ok, but is actually based on existed get_global/local_id/size/offset. It should be correct if the depend on functions are correst. And I see there is missing a get_global_offset utest, maybe I can also write am utest for it, too.
-----Original Message----- From: Beignet [mailto:[email protected]] On Behalf Of He Junyan Sent: Wednesday, December 2, 2015 2:23 PM To: [email protected] Subject: Re: [Beignet] [PATCH OCL2.0] libocl: Add three work-item built-in function This patch is OK. I appreciate that you can also add a utest case for it. On Wed, Dec 02, 2015 at 01:37:41PM +0800, Pan Xiuli wrote: > Date: Wed, 2 Dec 2015 13:37:41 +0800 > From: Pan Xiuli <[email protected]> > To: [email protected] > Cc: [email protected] > Subject: [Beignet] [PATCH OCL2.0] 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 <[email protected]> > --- > 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 > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
