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

Sun, Yi yi.sun at intel.com
Mon Jun 17 06:09:01 PDT 2013


Hi Zhigang,

Thank you for your comments. :p
I updated the patched again, as attached, please review.

Thanks
  --Sun, Yi

On Mon, 2013-06-17 at 18:04 +0800, Zhigang Gong wrote:
> Hi Yi,
> 
> I just took a look at the patch and found two minor problems:
> 
> 1. The indentation is not fully comply with Beignet. We always use 2 space indent, and
> never use tab. I found you mix-use 2 space and 4 space indent. And also found some tabs there.
> 
> 2. Another problem is the version check method. I made the following modification on top
> of you patch. If we compile this test case with OpenCL 1.1 header file, it will fail at
> #if CL_VERSION_1_2 | CL_VERSION_1_1.
> As CL_VERSION_1_2 is not defined at all.
> 
> -        #if CL_VERSION_1_2 | CL_VERSION_1_1
> +        #if defined(CL_VERSION_1_2) | defined(CL_VERSION_1_1)
>              OCL_ASSERT( global_size == 1);
> -        #elif CL_VERSION_1_0
> +        #elif defined(CL_VERSION_1_0)
>              OCL_ASSERT( global_size == 0);
> -               #else
> -                       OCL_ASSERT( global_size == 1);
> +        #else
> +            OCL_ASSERT( global_size == 1);
>          #endif
> 
> Any comment?
> 
> On Mon, Jun 17, 2013 at 05:40:40PM +0800, Zhigang Gong wrote:
> > Thanks for the patch. Will push it latter.
> > 
> > On Mon, Jun 17, 2013 at 08:52:21AM +0000, Sun, Yi wrote:
> > > Updated the test case to handle different version.
> > > 
> > > Thanks
> > >   --Sun, Yi
> > > 
> > > On Thu, 2013-06-13 at 07:21 +0000, Sun, Yi wrote:
> > > > 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
> > > > 
> > > > _______________________________________________
> > > > Beignet mailing list
> > > > Beignet at lists.freedesktop.org
> > > > http://lists.freedesktop.org/mailman/listinfo/beignet
> > > 
> > 
> > > From 57c4e1ed628a2ab7ebc1021a4406346f61f3f086 Mon Sep 17 00:00:00 2001
> > > From: Yi Sun <yi.sun at intel.com>
> > > Date: Thu, 13 Jun 2013 15:12:44 +0800
> > > Subject: [PATCH] Utest: Add a test case for validating built-in function
> > >  get_global_size()
> > > 
> > > v1:
> > > According to the OpenCL v1.1 & v1.2 chapter 6.11, the behavior of function get_global_size should be as following:
> > > 
> > > get_global_size(-1) = 1 (dimension:1)
> > > get_global_size(0) = 3 (dimension:1)
> > > get_global_size(1) = 1 (dimension:1)
> > > get_global_size(2) = 1 (dimension:1)
> > > 
> > > get_global_size(-1) = 1 (dimension:2)
> > > get_global_size(0) = 3 (dimension:2)
> > > get_global_size(1) = 4 (dimension:2)
> > > get_global_size(2) = 1 (dimension:2)
> > > get_global_size(3) = 1 (dimension:2)
> > > 
> > > get_global_size(-1) = 1 (dimension:3)
> > > get_global_size(0) = 3 (dimension:3)
> > > get_global_size(1) = 4 (dimension:3)
> > > get_global_size(2) = 5 (dimension:3)
> > > get_global_size(3) = 1 (dimension:3)
> > > get_global_size(4) = 1 (dimension:3)
> > > 
> > > if defined:
> > >   globals[0] = 3;
> > >   globals[1] = 4;
> > >   globals[2] = 5;
> > > 
> > > v2:
> > >     Handle different version.
> > > 
> > >     Add #if and #elif to make the test case be suitable to different version.
> > > 
> > > Signed-off-by: Yi Sun <yi.sun at intel.com>
> > > ---
> > >  kernels/builtin_global_size.cl |    3 +
> > >  utests/CMakeLists.txt          |    1 +
> > >  utests/builtin_global_size.cpp |  108 ++++++++++++++++++++++++++++++++++++++++
> > >  3 files changed, 112 insertions(+), 0 deletions(-)
> > >  create mode 100644 kernels/builtin_global_size.cl
> > >  create mode 100644 utests/builtin_global_size.cpp
> > > 
> > > diff --git a/kernels/builtin_global_size.cl b/kernels/builtin_global_size.cl
> > > new file mode 100644
> > > index 0000000..e6ddb2f
> > > --- /dev/null
> > > +++ b/kernels/builtin_global_size.cl
> > > @@ -0,0 +1,3 @@
> > > +kernel void builtin_global_size( __global int *ret, __global int *i_dim ) {
> > > +  *ret = get_global_size( *i_dim);
> > > +}
> > > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> > > index e5c03ee..c8df444 100644
> > > --- a/utests/CMakeLists.txt
> > > +++ b/utests/CMakeLists.txt
> > > @@ -82,6 +82,7 @@ set (utests_sources
> > >    compiler_vector_load_store.cpp
> > >    compiler_cl_finish.cpp
> > >    buildin_work_dim.cpp
> > > +  builtin_global_size.cpp
> > >    runtime_createcontext.cpp
> > >    runtime_null_kernel_arg.cpp
> > >    utest_assert.cpp
> > > diff --git a/utests/builtin_global_size.cpp b/utests/builtin_global_size.cpp
> > > new file mode 100644
> > > index 0000000..8518c8d
> > > --- /dev/null
> > > +++ b/utests/builtin_global_size.cpp
> > > @@ -0,0 +1,108 @@
> > > +/*
> > > +According to the OpenCL v1.1 & v1.2 chapter 6.11, the behavior of function get_global_size should be as following:
> > > +
> > > +  globals[0] = 3;
> > > +  globals[1] = 4;
> > > +  globals[2] = 5;
> > > +
> > > +#ifdef CL_VERSION_1_2 | CL_VERSION_1_1:
> > > +get_global_size(-1) = 1 (dimension:1)
> > > +get_global_size(0) = 3 (dimension:1)
> > > +get_global_size(1) = 1 (dimension:1)
> > > +get_global_size(2) = 1 (dimension:1)
> > > +
> > > +get_global_size(-1) = 1 (dimension:2)
> > > +get_global_size(0) = 3 (dimension:2)
> > > +get_global_size(1) = 4 (dimension:2)
> > > +get_global_size(2) = 1 (dimension:2)
> > > +get_global_size(3) = 1 (dimension:2)
> > > +
> > > +get_global_size(-1) = 1 (dimension:3)
> > > +get_global_size(0) = 3 (dimension:3)
> > > +get_global_size(1) = 4 (dimension:3)
> > > +get_global_size(2) = 5 (dimension:3)
> > > +get_global_size(3) = 1 (dimension:3)
> > > +get_global_size(4) = 1 (dimension:3)
> > > +
> > > +#ifdef CL_VERSION_1_0:
> > > +get_global_size(-1) = 0 (dimension:1)
> > > +get_global_size(0) = 3 (dimension:1)
> > > +get_global_size(1) = 0 (dimension:1)
> > > +get_global_size(2) = 0 (dimension:1)
> > > +
> > > +get_global_size(-1) = 0 (dimension:2)
> > > +get_global_size(0) = 3 (dimension:2)
> > > +get_global_size(1) = 4 (dimension:2)
> > > +get_global_size(2) = 0 (dimension:2)
> > > +get_global_size(3) = 1 (dimension:2)
> > > +
> > > +get_global_size(-1) = 0 (dimension:3)
> > > +get_global_size(0) = 3 (dimension:3)
> > > +get_global_size(1) = 4 (dimension:3)
> > > +get_global_size(2) = 5 (dimension:3)
> > > +get_global_size(3) = 0 (dimension:3)
> > > +get_global_size(4) = 0 (dimension:3)
> > > +
> > > +*/
> > > +#include "utest_helper.hpp"
> > > +static void builtin_global_size(void)
> > > +{
> > > +
> > > +  // Setup kernel and buffers
> > > +  int dim, dim_arg_global, global_size, err;
> > > +  OCL_CREATE_KERNEL("builtin_global_size");
> > > +
> > > +  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int), NULL);
> > > +  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, sizeof(int), NULL);
> > > +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> > > +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> > > +
> > > +  globals[0] = 3;
> > > +  globals[1] = 4;
> > > +  globals[2] = 5;
> > > +  locals[0] = 1;
> > > +  locals[1] = 1;
> > > +  locals[2] = 1;
> > > +
> > > +  for( dim=1; dim <= 3; dim++ )
> > > +  {
> > > +
> > > +    for( dim_arg_global = -1; dim_arg_global <= dim + 1; dim_arg_global++ )
> > > +    {
> > > +
> > > +        err = clEnqueueWriteBuffer( queue, buf[1], CL_TRUE, 0, sizeof(int), &dim_arg_global, 0, NULL, NULL);
> > > +        if (err != CL_SUCCESS)
> > > +        {
> > > +            printf("Error: Failed to write to source array!\n");
> > > +            exit(1);
> > > +        }
> > > +
> > > +        // Run the kernel
> > > +        OCL_NDRANGE( dim );
> > > +
> > > +        err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int), &global_size, 0, NULL, NULL);
> > > +        if (err != CL_SUCCESS)
> > > +        {
> > > +           printf("Error: Failed to read output array! %d\n", err);
> > > +           exit(1);
> > > +        }
> > > +
> > > +        //printf("get_global_size(%d) = %d (dimension:%d)\n", dim_arg_global, global_size, dim);
> > > +
> > > +        if ( dim_arg_global >= 0 && dim_arg_global < dim)
> > > +            OCL_ASSERT( global_size == dim_arg_global + 3);
> > > +        else
> > > +        {
> > > +        #if CL_VERSION_1_2 | CL_VERSION_1_1
> > > +            OCL_ASSERT( global_size == 1);
> > > +        #elif CL_VERSION_1_0
> > > +            OCL_ASSERT( global_size == 0);
> > > +		#else
> > > +			OCL_ASSERT( global_size == 1);
> > > +        #endif
> > > +        }
> > > +    }
> > > +  }
> > > +}
> > > +
> > > +MAKE_UTEST_FROM_FUNCTION(builtin_global_size);
> > > -- 
> > > 1.7.6.4
> > > 
> > 
> > _______________________________________________
> > 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: 5676 bytes
Desc: 0001-Utest-Add-a-test-case-for-validating-built-in-functi.patch
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130617/96d6a82f/attachment-0001.bin>


More information about the Beignet mailing list