[Beignet] [BUG] Built-in function get_global_size error
Sun, Yi
yi.sun at intel.com
Thu Jun 13 00:21:06 PDT 2013
Hi Zhigang,
Your patch works well, verified.
And Updated the test case to comply the spec v1.1 and v1.2
Thanks
--Sun, Yi
On Sun, 2013-06-09 at 11:22 +0800, Zhigang Gong wrote:
> 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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-Utest-Add-a-test-case-for-validating-built-in-functi.patch
Type: text/x-patch
Size: 4671 bytes
Desc: 0001-Utest-Add-a-test-case-for-validating-built-in-functi.patch
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130613/091ec71f/attachment-0001.bin>
More information about the Beignet
mailing list