Utest: Add tset case for block read/write buffer

V2: Rename test case to buffer block read/write test

Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

2 files changed, 58 insertions, 0 deletions

diff --git a/kernels/compiler_subgroup_buffer_block_read.cl b/kernels/compiler_subgroup_buffer_block_read.cl
new file mode 100644
index 00000000..9edaa2ed
--- /dev/null
+++ b/kernels/compiler_subgroup_buffer_block_read.cl
@@ -0,0 +1,31 @@
+__kernel void compiler_subgroup_buffer_block_read1(global uint *src, global uint *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = src + get_sub_group_id() * get_max_sub_group_size();
+  uint tmp = intel_sub_group_block_read(p);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_buffer_block_read2(global uint *src, global uint2 *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*2;
+  uint2 tmp = intel_sub_group_block_read2(p);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_buffer_block_read4(global uint *src, global uint4 *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*4;
+  uint4 tmp = intel_sub_group_block_read4(p);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_buffer_block_read8(global uint *src, global uint8 *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = src + get_sub_group_id() * get_max_sub_group_size()*8;
+  uint8 tmp = intel_sub_group_block_read8(p);
+  dst[id] = tmp;
+}
diff --git a/kernels/compiler_subgroup_buffer_block_write.cl b/kernels/compiler_subgroup_buffer_block_write.cl
new file mode 100644
index 00000000..f7358550
--- /dev/null
+++ b/kernels/compiler_subgroup_buffer_block_write.cl
@@ -0,0 +1,27 @@
+__kernel void compiler_subgroup_buffer_block_write1(global uint *src, global uint *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = dst + get_sub_group_id() * get_max_sub_group_size();
+  intel_sub_group_block_write(p,src[id]);
+}
+
+__kernel void compiler_subgroup_buffer_block_write2(global uint2 *src, global uint *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = dst + get_sub_group_id() * get_max_sub_group_size()*2;
+  intel_sub_group_block_write2(p,src[id]);
+}
+
+__kernel void compiler_subgroup_buffer_block_write4(global uint4 *src, global uint *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = dst + get_sub_group_id() * get_max_sub_group_size()*4;
+  intel_sub_group_block_write4(p,src[id]);
+}
+
+__kernel void compiler_subgroup_buffer_block_write8(global uint8 *src, global uint *dst)
+{
+  int id = get_global_id(0);
+  global uint * p = dst + get_sub_group_id() * get_max_sub_group_size()*8;
+  intel_sub_group_block_write8(p,src[id]);
+}