libocl: Add three work-item built-in function

Add get global/local linear id by calculate with global/local
id, size and offset. The get_queue_local_size() and get_loal_size()
should be different when the global work group size is not uniform,
but now they are the same. We will refine these functions when we
support non-uniform work-group size.

Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Junyan He <junyan.he@linux.intel.com>

2 files changed, 33 insertions, 0 deletions

diff --git a/backend/src/libocl/include/ocl_workitem.h b/backend/src/libocl/include/ocl_workitem.h
index 84bb1fbc..c3b0bdb4 100644
--- a/backend/src/libocl/include/ocl_workitem.h
+++ b/backend/src/libocl/include/ocl_workitem.h
@@ -24,9 +24,12 @@ OVERLOADABLE uint get_work_dim(void);
 OVERLOADABLE uint get_global_size(uint dimindx);
 OVERLOADABLE uint get_global_id(uint dimindx);
 OVERLOADABLE uint get_local_size(uint dimindx);
+OVERLOADABLE uint get_enqueued_local_size(uint dimindx);
 OVERLOADABLE uint get_local_id(uint dimindx);
 OVERLOADABLE uint get_num_groups(uint dimindx);
 OVERLOADABLE uint get_group_id(uint dimindx);
 OVERLOADABLE uint get_global_offset(uint dimindx);
+OVERLOADABLE uint get_global_linear_id(void);
+OVERLOADABLE uint get_local_linear_id(void);
 
 #endif  /* __OCL_WORKITEM_H__ */
diff --git a/backend/src/libocl/src/ocl_workitem.cl b/backend/src/libocl/src/ocl_workitem.cl
index 6ddc4067..235f12b1 100644
--- a/backend/src/libocl/src/ocl_workitem.cl
+++ b/backend/src/libocl/src/ocl_workitem.cl
@@ -55,3 +55,33 @@ DECL_PUBLIC_WORK_ITEM_FN(get_num_groups, 1)
 OVERLOADABLE uint get_global_id(uint dim) {
   return get_local_id(dim) + get_local_size(dim) * get_group_id(dim) + get_global_offset(dim);
 }
+
+OVERLOADABLE uint get_enqueued_local_size (uint dimindx)
+{
+  //TODO: should be different with get_local_size when support
+  //non-uniform work-group size
+  return get_local_size(dimindx);
+}
+
+OVERLOADABLE uint get_global_linear_id(void)
+{
+  uint dim = __gen_ocl_get_work_dim();
+  if (dim == 1) return get_global_id(0) - get_global_offset(0);
+  else if (dim == 2) return (get_global_id(1) - get_global_offset(1)) * get_global_size(0) +
+                             get_global_id(0) -get_global_offset(0);
+  else if (dim == 3) return ((get_global_id(2) - get_global_offset(2)) *
+                              get_global_size(1) * get_global_size(0)) +
+                            ((get_global_id(1) - get_global_offset(1)) * get_global_size (0)) +
+                             (get_global_id(0) - get_global_offset(0));
+  else return 0;
+}
+
+OVERLOADABLE uint get_local_linear_id(void)
+{
+  uint dim = __gen_ocl_get_work_dim();
+  if (dim == 1) return get_local_id(0);
+  else if (dim == 2) return get_local_id(1) * get_local_size (0) + get_local_id(0);
+  else if (dim == 3) return (get_local_id(2) * get_local_size(1) * get_local_size(0)) +
+                            (get_local_id(1) * get_local_size(0)) + get_local_id(0);
+  else return 0;
+}