From 6e1e3a0764c4336e35b9001d4fb9b30e40072136 Mon Sep 17 00:00:00 2001
From: Junyan He <junyan.he@linux.intel.com>
Date: Mon, 1 Sep 2014 10:06:09 +0800
Subject: [PATCH] Add the async module into the libocl

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@intel.com>
---
 backend/src/libocl/include/ocl_async.h | 49 ++++++++++++++++++++++++
 backend/src/libocl/src/ocl_async.cl    | 69 ++++++++++++++++++++++++++++++++++
 2 files changed, 118 insertions(+)
 create mode 100644 backend/src/libocl/include/ocl_async.h
 create mode 100644 backend/src/libocl/src/ocl_async.cl

diff --git a/backend/src/libocl/include/ocl_async.h b/backend/src/libocl/include/ocl_async.h
new file mode 100644
index 0000000..eccd381
--- /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/src/ocl_async.cl b/backend/src/libocl/src/ocl_async.cl
new file mode 100644
index 0000000..57d6859
--- /dev/null
+++ b/backend/src/libocl/src/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
-- 
2.7.4