[Beignet] [PATCH 02/18] Add the workitems functions into lib ocl

[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/Makefile.in            | 14 ++++++++---
 backend/src/libocl/include/ocl_workitem.h | 15 ++++++++++++
 backend/src/libocl/lib/ocl_workitem.cl    | 40 +++++++++++++++++++++++++++++++
 3 files changed, 66 insertions(+), 3 deletions(-)
 create mode 100644 backend/src/libocl/include/ocl_workitem.h
 create mode 100644 backend/src/libocl/lib/ocl_workitem.cl

diff --git a/backend/src/libocl/Makefile.in b/backend/src/libocl/Makefile.in
index 628d6c0..cc74fcb 100644
--- a/backend/src/libocl/Makefile.in
+++ b/backend/src/libocl/Makefile.in
@@ -7,7 +7,7 @@ GENERATED_FILES=
 GENERATED_HEADERS=ocl_defines.h
 GENERATED_CL_SRCS=$(addprefix lib/, $(GENERATED_FILES))
 GENERATED_CL_HEADERS=$(addprefix include/, $(GENERATED_HEADERS))
-CL_FILE_NAMES= $(GENERATED_FILES)
+CL_FILE_NAMES=ocl_workitem.cl $(GENERATED_FILES)
 LL_FILE_NAMES=
 CL_SRCS=$(addprefix lib/, $(CL_FILE_NAMES))
 LL_SRCS=$(addprefix lib/, $(LL_FILE_NAMES))
@@ -20,7 +20,7 @@ CLANG_OCL_FLAGS=-I./include -fno-builtin -Dcl_khr_fp64 -ffp-contract=off
 all:$(GENERATED_CL_HEADERS) $(GENERATED_CL_SRCS) lib/beignet.bc
 
 lib/beignet.bc:$(GENERATED_CL_HEADERS) $(GENERATED_CL_SRCS) $(CL_BITCODES) $(LL_BITCODES)
-	touch $@
+	llvm-link -o $@ $(CL_BITCODES) $(LL_BITCODES)
 
 include/ocl_defines.h:include/ocl_defines.inh
 	@echo "Generate the header:  $@"
@@ -28,12 +28,20 @@ include/ocl_defines.h:include/ocl_defines.inh
 	@cat $< > $@
 	@cat ../ocl_common_defines.h >> $@
 
+%.bc:%.cl $(GENERATED_CL_HEADERS)
+	clang -cc1 $(CLANG_OCL_FLAGS) -emit-llvm-bc -triple spir -o $@ -x cl $<
+
 install:all
 	mkdir -p $(HEADER_INSTALL_PREFIX)
 	mkdir -p $(BITCODE_INSTALL_PREFIX)
 	cp lib/beignet.bc $(BITCODE_INSTALL_PREFIX)
 	cp include/*.h $(HEADER_INSTALL_PREFIX)
 
+%.d:%.cl $(GENERATED_CL_SRCS) $(GENERATED_CL_HEADERS)
+	@clang -MM -MF $@ $(CLANG_OCL_FLAGS) $<
+	@sed -i '1s/\(^.*\)\.o/lib\/\1\.bc/' $@
+
+-include $(CL_DEPENDS)
 
 clean:
-	rm -fr lib/beignet.bc
+	rm -fr $(CL_BITCODES) $(LL_BITCODES) $(CL_DEPENDS) $(GENERATED_CL_HEADERS) $(GENERATED_CL_SRCS) lib/beignet.bc
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/lib/ocl_workitem.cl b/backend/src/libocl/lib/ocl_workitem.cl
new file mode 100644
index 0000000..af0e895
--- /dev/null
+++ b/backend/src/libocl/lib/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.8.3.2