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

[Zhigang Gong]

Yi,

I thought our current target is OpenCL 1.1, so I just checked 1.1's specification.
Maybe we need to synch with each other here.
Nanhai, we may need your confirmation here. Which version should we focus on for
next release?

On Sun, Jun 09, 2013 at 01:09:25AM +0000, Sun, Yi wrote:
> 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) {
> 


> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet