[Beignet] [BUG] Built-in function get_global_size error
Zhigang Gong
zhigang.gong at linux.intel.com
Sat Jun 8 20:22:31 PDT 2013
Homer,
I just greped __CL_VERSION_XXX, and failed to get anything, so ....
Now, I just took a look at the cl.h and found the following marcos' definitions:
/* OpenCL Version */
#define CL_VERSION_1_0 1
#define CL_VERSION_1_1 1
#define CL_VERSION_1_2 1
So the macro name is not __CL_VERSION_XXX, and it defines all the macros for
the three versions.
And as we already use OpenCL 1.2's header file in our project, if we use thes
macros to determine the version, we always go to the OpenCL 1.2 path.
If we don't fallback the OpenCL header file (for opencl 1.2) to version 1.1,
we may have to focus on OpenCL version 1.2 from now on.
On the other hand, if we want to choose 1.1 now, we need to fallback the header
files, and we already implemented/used some new APIs introduced in 1.2. This will
lead to some effort to fix them.
Nanhai, what's your opinion?
On Sun, Jun 09, 2013 at 03:04:00AM +0000, Xing, Homer wrote:
> Hi,
>
> CL_VERSION_1_X macros are also defined in CL/cl.h. They can take effect in user C program, not just inside OpenCL kernels. :)
>
> -----Original Message-----
> From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
> Sent: Sunday, June 09, 2013 10:59 AM
> To: Xing, Homer
> Cc: Sun, Yi; beignet at lists.freedesktop.org; Jin, Gordon
> Subject: Re: [Beignet] [BUG] Built-in function get_global_size error
>
> Homer,
>
> I'm afraid this may not work. As __CL_VERSION_XXX is a predefined OpenCL macros, it should only take effect in a OpenCL kernel from my understanding. If I am wrong, please point it out.
> Thanks.
>
> On Sun, Jun 09, 2013 at 01:46:43AM +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) {
> >
> > _______________________________________________
> > Beignet mailing list
> > Beignet at lists.freedesktop.org
> > http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list