[Beignet] [PATCH 06/18] Add the sync and async functions into 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/Makefile.in         |  2 +-
 backend/src/libocl/include/ocl_async.h | 49 ++++++++++++++++++++++++
 backend/src/libocl/include/ocl_sync.h  | 18 +++++++++
 backend/src/libocl/lib/ocl_async.cl    | 69 ++++++++++++++++++++++++++++++++++
 backend/src/libocl/lib/ocl_sync.cl     | 14 +++++++
 5 files changed, 151 insertions(+), 1 deletion(-)
 create mode 100644 backend/src/libocl/include/ocl_async.h
 create mode 100644 backend/src/libocl/include/ocl_sync.h
 create mode 100644 backend/src/libocl/lib/ocl_async.cl
 create mode 100644 backend/src/libocl/lib/ocl_sync.cl

diff --git a/backend/src/libocl/Makefile.in b/backend/src/libocl/Makefile.in
index 06adc08..5e020ab 100644
--- a/backend/src/libocl/Makefile.in
+++ b/backend/src/libocl/Makefile.in
@@ -7,7 +7,7 @@ GENERATED_FILES=ocl_as.cl ocl_convert.cl
 GENERATED_HEADERS=ocl_defines.h ocl_as.h ocl_convert.h
 GENERATED_CL_SRCS=$(addprefix lib/, $(GENERATED_FILES))
 GENERATED_CL_HEADERS=$(addprefix include/, $(GENERATED_HEADERS))
-CL_FILE_NAMES=ocl_workitem.cl ocl_atom.cl $(GENERATED_FILES)
+CL_FILE_NAMES=ocl_workitem.cl ocl_atom.cl ocl_async.cl ocl_sync.cl $(GENERATED_FILES)
 LL_FILE_NAMES=
 CL_SRCS=$(addprefix lib/, $(CL_FILE_NAMES))
 LL_SRCS=$(addprefix lib/, $(LL_FILE_NAMES))
diff --git a/backend/src/libocl/include/ocl_async.h b/backend/src/libocl/include/ocl_async.h
new file mode 100644
index 0000000..833bc21
--- /dev/null
+++ b/backend/src/libocl/include/ocl_async.h
@@ -0,0 +1,49 @@
+#ifndef __OCL_ASYNC_H__
+#define __OCL_ASYNC_H__ 
+
+#include "ocl_types.h"
+
+#define DEFN(TYPE) \
+OVERLOADABLE event_t async_work_group_copy (local TYPE *dst,  const global TYPE *src, \
+							 size_t num, event_t event); \
+OVERLOADABLE event_t async_work_group_copy (global TYPE *dst,  const local TYPE *src, \
+							  size_t num, event_t event); \
+OVERLOADABLE event_t async_work_group_strided_copy (local TYPE *dst,  const global TYPE *src, \
+								 size_t num, size_t src_stride, event_t event); \
+OVERLOADABLE event_t async_work_group_strided_copy (global TYPE *dst,  const local TYPE *src, \
+								  size_t num, size_t dst_stride, event_t event); \
+
+#define DEF(TYPE) \
+  DEFN(TYPE); DEFN(TYPE##2); DEFN(TYPE##3); DEFN(TYPE##4); DEFN(TYPE##8); DEFN(TYPE##16);
+DEF(char)
+DEF(uchar)
+DEF(short)
+DEF(ushort)
+DEF(int)
+DEF(uint)
+DEF(long)
+DEF(ulong)
+DEF(float)
+DEF(double)
+#undef DEFN
+#undef DEF
+
+void wait_group_events (int num_events, event_t *event_list);
+
+#define DEFN(TYPE) \
+OVERLOADABLE void prefetch(const global TYPE *p, size_t num);
+#define DEF(TYPE) \
+DEFN(TYPE); DEFN(TYPE##2); DEFN(TYPE##3); DEFN(TYPE##4); DEFN(TYPE##8); DEFN(TYPE##16)
+DEF(char);
+DEF(uchar);
+DEF(short);
+DEF(ushort);
+DEF(int);
+DEF(uint);
+DEF(long);
+DEF(ulong);
+DEF(float);
+#undef DEFN
+#undef DEF
+
+#endif
diff --git a/backend/src/libocl/include/ocl_sync.h b/backend/src/libocl/include/ocl_sync.h
new file mode 100644
index 0000000..f983824
--- /dev/null
+++ b/backend/src/libocl/include/ocl_sync.h
@@ -0,0 +1,18 @@
+#ifndef __OCL_SYNC_H__
+#define __OCL_SYNC_H__
+
+#include "ocl_types.h"
+
+/////////////////////////////////////////////////////////////////////////////
+// Synchronization functions
+/////////////////////////////////////////////////////////////////////////////
+#define CLK_LOCAL_MEM_FENCE  (1 << 0)
+#define CLK_GLOBAL_MEM_FENCE (1 << 1)
+
+typedef uint cl_mem_fence_flags;
+void barrier(cl_mem_fence_flags flags);
+void mem_fence(cl_mem_fence_flags flags);
+void read_mem_fence(cl_mem_fence_flags flags);
+void write_mem_fence(cl_mem_fence_flags flags);
+
+#endif  /* __OCL_SYNC_H__ */
diff --git a/backend/src/libocl/lib/ocl_async.cl b/backend/src/libocl/lib/ocl_async.cl
new file mode 100644
index 0000000..57d6859
--- /dev/null
+++ b/backend/src/libocl/lib/ocl_async.cl
@@ -0,0 +1,69 @@
+#include "ocl_async.h"
+#include "ocl_sync.h"
+#include "ocl_workitem.h"
+
+#define BODY(SRC_STRIDE, DST_STRIDE) \
+  uint size = get_local_size(2) * get_local_size(1) * get_local_size(0); \
+  uint count = num / size;  \
+  uint offset = get_local_id(2) * get_local_size(1) + get_local_id(1);  \
+  offset = offset * get_local_size(0) + get_local_id(0); \
+  for(uint i=0; i<count; i+=1) { \
+    *(dst + offset * DST_STRIDE) = *(src + offset * SRC_STRIDE); \
+    offset += size;                                 \
+  } \
+  if(offset < num) \
+    *(dst + offset * DST_STRIDE) = *(src + offset * SRC_STRIDE); \
+  return 0;
+
+#define DEFN(TYPE) \
+OVERLOADABLE event_t async_work_group_copy (local TYPE *dst,  const global TYPE *src, \
+							 size_t num, event_t event) { \
+  BODY(1, 1); \
+} \
+OVERLOADABLE event_t async_work_group_copy (global TYPE *dst,  const local TYPE *src, \
+							  size_t num, event_t event) { \
+  BODY(1, 1); \
+} \
+OVERLOADABLE event_t async_work_group_strided_copy (local TYPE *dst,  const global TYPE *src, \
+								 size_t num, size_t src_stride, event_t event) { \
+  BODY(src_stride, 1); \
+} \
+OVERLOADABLE event_t async_work_group_strided_copy (global TYPE *dst,  const local TYPE *src, \
+								  size_t num, size_t dst_stride, event_t event) { \
+  BODY(1, dst_stride); \
+}
+#define DEF(TYPE) \
+  DEFN(TYPE); DEFN(TYPE##2); DEFN(TYPE##3); DEFN(TYPE##4); DEFN(TYPE##8); DEFN(TYPE##16);
+DEF(char)
+DEF(uchar)
+DEF(short)
+DEF(ushort)
+DEF(int)
+DEF(uint)
+DEF(long)
+DEF(ulong)
+DEF(float)
+DEF(double)
+#undef BODY
+#undef DEFN
+#undef DEF
+
+void wait_group_events (int num_events, event_t *event_list) {
+  barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
+}
+
+#define DEFN(TYPE) \
+OVERLOADABLE void prefetch(const global TYPE *p, size_t num) { }
+#define DEF(TYPE) \
+DEFN(TYPE); DEFN(TYPE##2); DEFN(TYPE##3); DEFN(TYPE##4); DEFN(TYPE##8); DEFN(TYPE##16)
+DEF(char);
+DEF(uchar);
+DEF(short);
+DEF(ushort);
+DEF(int);
+DEF(uint);
+DEF(long);
+DEF(ulong);
+DEF(float);
+#undef DEFN
+#undef DEF
diff --git a/backend/src/libocl/lib/ocl_sync.cl b/backend/src/libocl/lib/ocl_sync.cl
new file mode 100644
index 0000000..3489450
--- /dev/null
+++ b/backend/src/libocl/lib/ocl_sync.cl
@@ -0,0 +1,14 @@
+#include "ocl_sync.h"
+
+void __gen_ocl_barrier_local(void);
+void __gen_ocl_barrier_global(void);
+void __gen_ocl_barrier_local_and_global(void);
+
+void mem_fence(cl_mem_fence_flags flags) {
+}
+
+void read_mem_fence(cl_mem_fence_flags flags) {
+}
+
+void write_mem_fence(cl_mem_fence_flags flags) {
+}
-- 
1.8.3.2