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

[Pan, Xiuli]

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:beignet-bounces at lists.freedesktop.org] On Behalf Of He Junyan
Sent: Wednesday, December 2, 2015 2:23 PM
To: beignet at lists.freedesktop.org
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 <xiuli.pan at intel.com>
> To: beignet at lists.freedesktop.org
> Cc: xiuli.pan at intel.com
> 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 <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