Utest: Add tset case for block read/write image

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_image_block_read.cl b/kernels/compiler_subgroup_image_block_read.cl
new file mode 100644
index 00000000..d5df6dba
--- /dev/null
+++ b/kernels/compiler_subgroup_image_block_read.cl
@@ -0,0 +1,31 @@
+__kernel void compiler_subgroup_image_block_read1(image2d_t src, global uint *dst)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  uint tmp = intel_sub_group_block_read(src,coord);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_image_block_read2(image2d_t src, global uint2 *dst)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  uint2 tmp = intel_sub_group_block_read2(src,coord);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_image_block_read4(image2d_t src, global uint4 *dst)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  uint4 tmp = intel_sub_group_block_read4(src,coord);
+  dst[id] = tmp;
+}
+
+__kernel void compiler_subgroup_image_block_read8(image2d_t src, global uint8 *dst)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  uint8 tmp = intel_sub_group_block_read8(src,coord);
+  dst[id] = tmp;
+}
diff --git a/kernels/compiler_subgroup_image_block_write.cl b/kernels/compiler_subgroup_image_block_write.cl
new file mode 100644
index 00000000..d9b3717f
--- /dev/null
+++ b/kernels/compiler_subgroup_image_block_write.cl
@@ -0,0 +1,27 @@
+__kernel void compiler_subgroup_image_block_write1(image2d_t dst, global uint *src)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  intel_sub_group_block_write(dst,coord, src[id]);
+}
+
+__kernel void compiler_subgroup_image_block_write2(image2d_t dst, global uint2 *src)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  intel_sub_group_block_write2(dst,coord, src[id]);
+}
+
+__kernel void compiler_subgroup_image_block_write4(image2d_t dst, global uint4 *src)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  intel_sub_group_block_write4(dst,coord, src[id]);
+}
+
+__kernel void compiler_subgroup_image_block_write8(image2d_t dst, global uint8 *src)
+{
+  int id = get_global_id(0);
+  int2 coord = (int2)(get_simd_size()*get_sub_group_id()*sizeof(uint),0);
+  intel_sub_group_block_write8(dst,coord, src[id]);
+}