Add the kernels used by clEnqueueBufferFill API

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

5 files changed, 47 insertions, 0 deletions

diff --git a/src/kernels/cl_internal_fill_buf_align128.cl b/src/kernels/cl_internal_fill_buf_align128.cl
new file mode 100644
index 00000000..552820cc
--- /dev/null
+++ b/src/kernels/cl_internal_fill_buf_align128.cl
@@ -0,0 +1,9 @@
+kernel void __cl_fill_region_align128 ( global float16* dst, float16 pattern0,
+                                        unsigned int offset, unsigned int size, float16 pattern1)
+{
+    int i = get_global_id(0);
+    if (i < size) {
+        dst[i*2+offset] = pattern0;
+        dst[i*2+offset+1] = pattern1;
+    }
+}
diff --git a/src/kernels/cl_internal_fill_buf_align2.cl b/src/kernels/cl_internal_fill_buf_align2.cl
new file mode 100644
index 00000000..0b9a4cfd
--- /dev/null
+++ b/src/kernels/cl_internal_fill_buf_align2.cl
@@ -0,0 +1,8 @@
+kernel void __cl_fill_region_align2 ( global char2 * dst, char2 pattern,
+			             unsigned int offset, unsigned int size)
+{
+    int i = get_global_id(0);
+    if (i < size) {
+        dst[i+offset] = pattern;
+    }
+}
diff --git a/src/kernels/cl_internal_fill_buf_align4.cl b/src/kernels/cl_internal_fill_buf_align4.cl
new file mode 100644
index 00000000..aefd92fc
--- /dev/null
+++ b/src/kernels/cl_internal_fill_buf_align4.cl
@@ -0,0 +1,8 @@
+kernel void __cl_fill_region_align4 ( global float* dst, float pattern,
+			             unsigned int offset, unsigned int size)
+{
+    int i = get_global_id(0);
+    if (i < size) {
+        dst[i+offset] = pattern;
+    }
+}
diff --git a/src/kernels/cl_internal_fill_buf_align8.cl b/src/kernels/cl_internal_fill_buf_align8.cl
new file mode 100644
index 00000000..edaff772
--- /dev/null
+++ b/src/kernels/cl_internal_fill_buf_align8.cl
@@ -0,0 +1,14 @@
+#define COMPILER_ABS_FUNC_N(N) \
+    kernel void __cl_fill_region_align8_##N ( global float##N* dst, float##N pattern, \
+                                              unsigned int offset, unsigned int size) { \
+         int i = get_global_id(0); \
+         if (i < size) { \
+             dst[i+offset] = pattern; \
+         }  \
+    }
+
+
+COMPILER_ABS_FUNC_N(2)
+COMPILER_ABS_FUNC_N(4)
+COMPILER_ABS_FUNC_N(8)
+COMPILER_ABS_FUNC_N(16)
diff --git a/src/kernels/cl_internal_fill_buf_unalign.cl b/src/kernels/cl_internal_fill_buf_unalign.cl
new file mode 100644
index 00000000..90762b00
--- /dev/null
+++ b/src/kernels/cl_internal_fill_buf_unalign.cl
@@ -0,0 +1,8 @@
+kernel void __cl_fill_region_unalign ( global char * dst, char pattern,
+			               unsigned int offset, unsigned int size)
+{
+    int i = get_global_id(0);
+    if (i < size) {
+        dst[i+offset] = pattern;
+    }
+}