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

[He Junyan]

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