[Beignet] [BUG] Built-in function get_global_size error

[Sun, Yi]

Alright, thank you Homer for your suggestion :p


Thanks
  --Sun, Yi

On Sun, 2013-06-09 at 01:46 +0000, Xing, Homer wrote:
> Hi Yi,
> 
> Your test case may support both OpenCL 1.0 and OpenCL 1.2.
> 
> You can do this by following.
> 
>    #ifdef (__CL_VERSION_1_2)
>     // Test here, obey OpenCL 1.2 way
>    #else
>     // Test here, obey OpenCL 1.1, OpenCL 1.0 way.
>    #endif
> 
> Best,
> Homer Hsing
> 
> -----Original Message-----
> From: beignet-bounces+homer.xing=intel.com at lists.freedesktop.org [mailto:beignet-bounces+homer.xing=intel.com at lists.freedesktop.org] On Behalf Of Sun, Yi
> Sent: Sunday, June 09, 2013 9:09 AM
> To: Zhigang Gong
> Cc: beignet at lists.freedesktop.org; Jin, Gordon
> Subject: Re: [Beignet] [BUG] Built-in function get_global_size error
> 
> Hi Zhigang,
> 
> So surprising here I'm!  I have found the issue that we're using the different version of OpenCL. Mine is 1.0 but I assume you're using v1.2.
> Because I compared the description of this function, in v1.2 it should return 1 while out-of range, but in v1.0 get_global_size return 0 contrarily. Attachment is the screen-shot for the function description in v1.0.
> 
> So should we use OpenCL version 1.2 specification but not v1.0?
> 
> Thanks
>   --Sun, Yi
> 
> On Sat, 2013-06-08 at 18:26 +0800, Zhigang Gong wrote:
> > Yi,
> > 
> > Thanks for the test case. One comment here, according to the Open CL spec:
> > 
> > size_t get_global_size (uint dimindx)
> > 
> > Returns the number of global work-items specified for dimension 
> > identified by dimindx. This value is given by the global_work_size 
> > argument to clEnqueueNDRangeKernel.
> > Valid values of dimindx are 0 to get_work_dim() – 1.
> > For other values of dimindx, get_global_size() returns 1.
> > For clEnqueueTask, this always returns 1.
> > 
> > For the out-of-range dim argument, it should return 1 rather than 0.
> > You may need to modify your case to comply with the spec.
> > 
> > And I took a look at the GBE side, we do have bugs here. And I wrote a 
> > patch embedded here to fix it, after you modify your test case. You 
> > can try it with my patch as below:
> > 
> > From 83239b8fd74b99a4efae222916fd10cd68bf7bce Mon Sep 17 00:00:00 2001
> > From: Zhigang Gong <zhigang.gong at linux.intel.com>
> > Date: Sat, 8 Jun 2013 18:10:11 +0800
> > Subject: [PATCH] GBE: Fix some builtin functions' return value.
> > 
> > Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
> > ---
> >  backend/src/ocl_stdlib.h | 29 ++++++++++++++++-------------
> >  1 file changed, 16 insertions(+), 13 deletions(-)
> > 
> > diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h index 
> > 013f4cc..77499d7 100644
> > --- a/backend/src/ocl_stdlib.h
> > +++ b/backend/src/ocl_stdlib.h
> > @@ -374,19 +374,22 @@ DECL_INTERNAL_WORK_ITEM_FN(get_global_offset)
> >  DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)
> >  #undef DECL_INTERNAL_WORK_ITEM_FN
> >  
> > -#define DECL_PUBLIC_WORK_ITEM_FN(NAME) \ -INLINE unsigned 
> > 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(); \
> > -  else return 0; \
> > -}
> > -DECL_PUBLIC_WORK_ITEM_FN(get_group_id)
> > -DECL_PUBLIC_WORK_ITEM_FN(get_local_id)
> > -DECL_PUBLIC_WORK_ITEM_FN(get_local_size)
> > -DECL_PUBLIC_WORK_ITEM_FN(get_global_size)
> > -DECL_PUBLIC_WORK_ITEM_FN(get_global_offset)
> > -DECL_PUBLIC_WORK_ITEM_FN(get_num_groups)
> > +#define DECL_PUBLIC_WORK_ITEM_FN(NAME, OTHER_RET)    \
> > +INLINE unsigned NAME(unsigned int dim) {             \
> > +  if (dim == 0) return __gen_ocl_##NAME##0();        \
> > +  else if (dim > 0 && dim < get_work_dim()) {        \
> > +    if (dim == 1) return __gen_ocl_##NAME##1();      \
> > +    else if (dim == 2) return __gen_ocl_##NAME##2(); \
> > +  }                                                  \
> > +  return OTHER_RET;                                  \
> > +}
> > +
> > +DECL_PUBLIC_WORK_ITEM_FN(get_group_id, 0) 
> > +DECL_PUBLIC_WORK_ITEM_FN(get_local_id, 0) 
> > +DECL_PUBLIC_WORK_ITEM_FN(get_local_size, 1) 
> > +DECL_PUBLIC_WORK_ITEM_FN(get_global_size, 1) 
> > +DECL_PUBLIC_WORK_ITEM_FN(get_global_offset, 0) 
> > +DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)
> >  #undef DECL_PUBLIC_WORK_ITEM_FN
> >  
> >  INLINE uint get_global_id(uint dim) {
>