[Beignet] [PATCH] libocl: Refine return type of workitem built-in functions

Yang, Rong R rong.r.yang at intel.com
Fri Apr 8 06:46:03 UTC 2016


LGTM, pushed, merge your patch "[Beignet] [PATCH] Utest: Add utest load spir for spir64" to one commit to avoid utest assert.

> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Xiuli Pan
> Sent: Tuesday, March 29, 2016 16:12
> To: beignet at lists.freedesktop.org
> Cc: Pan, Xiuli <xiuli.pan at intel.com>
> Subject: [Beignet] [PATCH] libocl: Refine return type of workitem built-in
> functions
> 
> From: Pan Xiuli <xiuli.pan at intel.com>
> 
> Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
> ---
>  backend/src/libocl/include/ocl_workitem.h | 20 ++++++++++----------
>  backend/src/libocl/src/ocl_workitem.cl    |  8 ++++----
>  2 files changed, 14 insertions(+), 14 deletions(-)
> 
> diff --git a/backend/src/libocl/include/ocl_workitem.h
> b/backend/src/libocl/include/ocl_workitem.h
> index c3b0bdb..1a96aa8 100644
> --- a/backend/src/libocl/include/ocl_workitem.h
> +++ b/backend/src/libocl/include/ocl_workitem.h
> @@ -21,15 +21,15 @@
>  #include "ocl_types.h"
> 
>  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);
> +OVERLOADABLE size_t get_global_size(uint dimindx); OVERLOADABLE
> size_t
> +get_global_id(uint dimindx); OVERLOADABLE size_t get_local_size(uint
> +dimindx); OVERLOADABLE size_t get_enqueued_local_size(uint dimindx);
> +OVERLOADABLE size_t get_local_id(uint dimindx); OVERLOADABLE size_t
> +get_num_groups(uint dimindx); OVERLOADABLE size_t get_group_id(uint
> +dimindx); OVERLOADABLE size_t get_global_offset(uint dimindx);
> +OVERLOADABLE size_t get_global_linear_id(void); OVERLOADABLE size_t
> +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 dc8fa6d..eb6210d 100644
> --- a/backend/src/libocl/src/ocl_workitem.cl
> +++ b/backend/src/libocl/src/ocl_workitem.cl
> @@ -38,7 +38,7 @@ DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)
>  #undef DECL_INTERNAL_WORK_ITEM_FN
> 
>  #define DECL_PUBLIC_WORK_ITEM_FN(NAME, OTHER_RET)    \
> -OVERLOADABLE unsigned NAME(unsigned int dim) {             \
> +OVERLOADABLE size_t NAME(unsigned int dim) {             \
>    if (dim == 0) return __gen_ocl_##NAME##0();        \
>    else if (dim == 1) return __gen_ocl_##NAME##1();   \
>    else if (dim == 2) return __gen_ocl_##NAME##2();   \
> @@ -54,11 +54,11 @@ DECL_PUBLIC_WORK_ITEM_FN(get_global_offset, 0)
> DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)  #undef
> DECL_PUBLIC_WORK_ITEM_FN
> 
> -OVERLOADABLE uint get_global_id(uint dim) {
> +OVERLOADABLE size_t get_global_id(uint dim) {
>    return get_local_id(dim) + get_enqueued_local_size(dim) *
> get_group_id(dim) + get_global_offset(dim);  }
> 
> -OVERLOADABLE uint get_global_linear_id(void)
> +OVERLOADABLE size_t get_global_linear_id(void)
>  {
>    uint dim = __gen_ocl_get_work_dim();
>    if (dim == 1) return get_global_id(0) - get_global_offset(0); @@ -71,7 +71,7
> @@ OVERLOADABLE uint get_global_linear_id(void)
>    else return 0;
>  }
> 
> -OVERLOADABLE uint get_local_linear_id(void)
> +OVERLOADABLE size_t get_local_linear_id(void)
>  {
>    uint dim = __gen_ocl_get_work_dim();
>    if (dim == 1) return get_local_id(0);
> --
> 2.5.0
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list