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

[Zhigang Gong]

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) {
-- 
1.7.11.7




On Sat, Jun 08, 2013 at 09:42:31AM +0000, Sun, Yi wrote:
> Now the results for calling the function get_global_size are following.
> But the lines with sign (*) are wrong, which should be 0, according to
> the Spec chapter 6.11.
> The test case is ready as attachment.
> 
> get_global_size(-1) = 0 (dimension:1)
> get_global_size(0) = 1 (dimension:1) 
> *get_global_size(1) = 1 (dimension:1) 
> *get_global_size(2) = 1 (dimension:1) 
> 
> get_global_size(-1) = 0 (dimension:2)
> get_global_size(0) = 1 (dimension:2) 
> get_global_size(1) = 1 (dimension:2) 
> *get_global_size(2) = 1 (dimension:2) 
> get_global_size(3) = 0 (dimension:2) 
> 
> get_global_size(-1) = 0 (dimension:3)
> get_global_size(0) = 1 (dimension:3) 
> get_global_size(1) = 1 (dimension:3) 
> get_global_size(2) = 1 (dimension:3) 
> get_global_size(3) = 0 (dimension:3) 
> get_global_size(4) = 0 (dimension:3) 
> 
> 
> Thanks
>   --Sun, Yi
>