[Beignet] [PATCH 10/22 V2] Add thw workitem module into the libocl

[junyan.he at inbox.com]

From: Junyan He <junyan.he at linux.intel.com>

Signed-off-by: Junyan He <junyan.he at linux.intel.com>
---
 backend/src/libocl/include/ocl_workitem.h |   15 +++++++++++
 backend/src/libocl/src/ocl_workitem.cl    |   40 +++++++++++++++++++++++++++++
 2 files changed, 55 insertions(+)
 create mode 100644 backend/src/libocl/include/ocl_workitem.h
 create mode 100644 backend/src/libocl/src/ocl_workitem.cl

diff --git a/backend/src/libocl/include/ocl_workitem.h b/backend/src/libocl/include/ocl_workitem.h
new file mode 100644
index 0000000..c9c605e
--- /dev/null
+++ b/backend/src/libocl/include/ocl_workitem.h
@@ -0,0 +1,15 @@
+#ifndef __OCL_WORKITEM_H__
+#define __OCL_WORKITEM_H__
+
+#include "ocl_types.h"
+
+uint get_work_dim(void);
+uint get_global_size(uint dimindx);
+uint get_global_id(uint dimindx);
+uint get_local_size(uint dimindx);
+uint get_local_id(uint dimindx);
+uint get_num_groups(uint dimindx);
+uint get_group_id(uint dimindx);
+uint get_global_offset(uint dimindx);
+
+#endif  /* __OCL_WORKITEM_H__ */
diff --git a/backend/src/libocl/src/ocl_workitem.cl b/backend/src/libocl/src/ocl_workitem.cl
new file mode 100644
index 0000000..af0e895
--- /dev/null
+++ b/backend/src/libocl/src/ocl_workitem.cl
@@ -0,0 +1,40 @@
+#include "ocl_workitem.h"
+
+PURE CONST uint __gen_ocl_get_work_dim(void);
+uint get_work_dim(void)
+{
+  return __gen_ocl_get_work_dim();
+}
+
+
+#define DECL_INTERNAL_WORK_ITEM_FN(NAME) \
+PURE CONST unsigned int __gen_ocl_##NAME##0(void); \
+PURE CONST unsigned int __gen_ocl_##NAME##1(void); \
+PURE CONST unsigned int __gen_ocl_##NAME##2(void);
+DECL_INTERNAL_WORK_ITEM_FN(get_group_id)
+DECL_INTERNAL_WORK_ITEM_FN(get_local_id)
+DECL_INTERNAL_WORK_ITEM_FN(get_local_size)
+DECL_INTERNAL_WORK_ITEM_FN(get_global_size)
+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, OTHER_RET)    \
+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 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
+
+uint get_global_id(uint dim) {
+  return get_local_id(dim) + get_local_size(dim) * get_group_id(dim) + get_global_offset(dim);
+}
-- 
1.7.9.5