1 files changed, 102 insertions, 0 deletions
diff --git a/kernels/compiler_subgroup_image_block_write.cl b/kernels/compiler_subgroup_image_block_write.cl
index 7d97c59a..1632f929 100644
--- a/kernels/compiler_subgroup_image_block_write.cl
+++ b/kernels/compiler_subgroup_image_block_write.cl
@@ -53,3 +53,105 @@ __kernel void compiler_subgroup_image_block_write_us8(image2d_t dst, global usho
   intel_sub_group_block_write_us8(dst,coord, src[id]);
 }
 #endif
+#ifdef MEDIA_BLOCK_IO
+__kernel void compiler_subgroup_media_block_write_ui1(image2d_t dst, global uint *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(uint) + sizeof(uint) * get_group_id(0) * get_local_size(0),yid);
+  intel_sub_group_media_block_write_ui(coord, 16, 1, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_ui2(image2d_t dst, global uint2 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(uint) + sizeof(uint) * get_group_id(0) * get_local_size(0),yid*2);
+  intel_sub_group_media_block_write_ui2(coord, 16, 2, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_ui4(image2d_t dst, global uint4 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(uint) + sizeof(uint) * get_group_id(0) * get_local_size(0),yid*4);
+  intel_sub_group_media_block_write_ui4(coord, 16, 4, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_ui8(image2d_t dst, global uint8 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(uint) + sizeof(uint) * get_group_id(0) * get_local_size(0),yid*8);
+  intel_sub_group_media_block_write_ui8(coord, 16, 8, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_us1(image2d_t dst, global ushort *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(ushort) + sizeof(ushort) * get_group_id(0) * get_local_size(0),yid);
+  intel_sub_group_media_block_write_us(coord, 16, 1, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_us2(image2d_t dst, global ushort2 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(ushort) + sizeof(ushort) * get_group_id(0) * get_local_size(0),yid*2);
+  intel_sub_group_media_block_write_us2(coord, 16, 2, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_us4(image2d_t dst, global ushort4 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(ushort) + sizeof(ushort) * get_group_id(0) * get_local_size(0),yid*4);
+  intel_sub_group_media_block_write_us4(coord, 16, 4, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_us8(image2d_t dst, global ushort8 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(ushort) + sizeof(ushort) * get_group_id(0) * get_local_size(0),yid*8);
+  intel_sub_group_media_block_write_us8(coord, 16, 8, src[id], dst);
+}
+__kernel void  __attribute__((intel_reqd_sub_group_size(8)))
+compiler_subgroup_media_block_write_us16(image2d_t dst, global ushort16 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(ushort) + sizeof(ushort) * get_group_id(0) * get_local_size(0),yid*16);
+  intel_sub_group_media_block_write_us16(coord, 8, 16, src[id], dst);
+}
+
+__kernel void compiler_subgroup_media_block_write_uc1(image2d_t dst, global uchar *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(char) + sizeof(char) * get_group_id(0) * get_local_size(0),yid);
+  intel_sub_group_media_block_write_uc(coord, 16, 1, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_uc2(image2d_t dst, global uchar2 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(char) + sizeof(char) * get_group_id(0) * get_local_size(0),yid*2);
+  intel_sub_group_media_block_write_uc2(coord, 16, 2, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_uc4(image2d_t dst, global uchar4 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(char) + sizeof(char) * get_group_id(0) * get_local_size(0),yid*4);
+  intel_sub_group_media_block_write_uc4(coord, 16, 4, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_uc8(image2d_t dst, global uchar8 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(char) + sizeof(char) * get_group_id(0) * get_local_size(0),yid*8);
+  intel_sub_group_media_block_write_uc8(coord, 16, 8, src[id], dst);
+}
+__kernel void compiler_subgroup_media_block_write_uc16(image2d_t dst, global uchar16 *src)
+{
+  int id = get_global_id(0);
+  int yid = get_global_id(1);
+  int2 coord = (int2)(get_sub_group_size()*get_sub_group_id()*sizeof(char) + sizeof(char) * get_group_id(0) * get_local_size(0),yid*16);
+  intel_sub_group_media_block_write_uc16(coord, 16, 16, src[id], dst);
+}
+#endif