summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--kernels/compiler_intra_prediction.cl91
-rw-r--r--utests/CMakeLists.txt3
-rw-r--r--utests/compiler_intra_prediction.cpp118
3 files changed, 211 insertions, 1 deletions
diff --git a/kernels/compiler_intra_prediction.cl b/kernels/compiler_intra_prediction.cl
new file mode 100644
index 00000000..28e81e52
--- /dev/null
+++ b/kernels/compiler_intra_prediction.cl
@@ -0,0 +1,91 @@
+
+__kernel __attribute__((intel_reqd_sub_group_size(16)))
+void compiler_intra_prediction(
+ __read_only image2d_t srcImg,
+ __global uchar *luma_mode,
+ __global ushort *luma_distortion,
+ __global uchar *luma_shape,
+ __global uint* dwo_buffer,
+ __global uint* pld_buffer){
+
+ int gr_id0 = get_group_id(0);
+ int gr_id1 = get_group_id(1);
+
+ ushort2 src_coord;
+ /*src_coord.x = gr_id0 * 16;
+ src_coord.y = gr_id1 * 16;*/
+ src_coord.x = 2 * 16;
+ src_coord.y = 1 * 16;
+
+ intel_sub_group_avc_sic_payload_t payload = intel_sub_group_avc_sic_initialize(src_coord);
+
+ uchar sad_adjustment = CLK_AVC_ME_SAD_ADJUST_MODE_NONE_INTEL;
+ uchar intra_partition_mask = CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_16x16_INTEL;
+//XXX: Different from official value?
+#undef CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_RIGHT_MASK_ENABLE_INTEL
+#undef CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_LEFT_MASK_ENABLE_INTEL
+#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_RIGHT_MASK_ENABLE_INTEL 0x4
+#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_LEFT_MASK_ENABLE_INTEL 0x8
+ uint nb_avail = CLK_AVC_ME_INTRA_NEIGHBOR_LEFT_MASK_ENABLE_INTEL |
+ CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_MASK_ENABLE_INTEL |
+ CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_LEFT_MASK_ENABLE_INTEL |
+ CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_RIGHT_MASK_ENABLE_INTEL;
+
+ uint sgl_id = get_sub_group_local_id();
+ int2 nb_coord;
+ float4 color;
+
+ nb_coord.x = src_coord.x - 1;
+ nb_coord.y = src_coord.y + sgl_id;
+ color = read_imagef(srcImg, nb_coord);
+ uchar left_edge = color.s0 * 255;
+
+ nb_coord.x = src_coord.x - 1;
+ nb_coord.y = src_coord.y - 1;
+ color = read_imagef(srcImg, nb_coord);
+ uchar upper_left_corner = color.s0 * 255;
+
+ nb_coord.x = src_coord.x + sgl_id;
+ nb_coord.y = src_coord.y - 1;
+ color = read_imagef(srcImg, nb_coord);
+ uchar upper_edge = color.s0 * 255;
+
+ uchar upper_right_edge = 0;
+ if(sgl_id < 8){
+ nb_coord.x = src_coord.x + 16 + sgl_id;
+ nb_coord.y = src_coord.y - 1;
+ color = read_imagef(srcImg, nb_coord);
+ upper_right_edge = color.s0 * 255;
+ }
+ payload = intel_sub_group_avc_sic_configure_ipe(
+ intra_partition_mask, nb_avail, left_edge, upper_left_corner, upper_edge,
+ upper_right_edge, sad_adjustment, payload);
+
+ uchar shape_cost_16_16 = (1 << 4) | 5;
+ uchar shape_cost_8_8 = (1 << 4) | 4;
+ uchar shape_cost_4_4 = (1 << 4) | 3;
+ uint intra_shape_cost = (shape_cost_4_4 << 24) | (shape_cost_8_8 << 16) | (shape_cost_16_16 << 8) | (0x0);
+ payload = intel_sub_group_avc_sic_set_intra_luma_shape_penalty(intra_shape_cost, payload);
+
+ sampler_t vs = 0;
+ intel_sub_group_avc_sic_result_t result =
+ intel_sub_group_avc_sic_evaluate_ipe(srcImg, vs, payload);
+
+ uchar shape = intel_sub_group_avc_sic_get_ipe_luma_shape(result);
+ ushort dist = intel_sub_group_avc_sic_get_best_ipe_luma_distortion(result);
+ ulong modes = intel_sub_group_avc_sic_get_packed_ipe_luma_modes(result);
+
+ int lid_x = get_local_id(0);
+ int mb_idx = gr_id0 + gr_id1 * get_num_groups(0);
+ if (lid_x == 0) {
+ luma_shape[mb_idx] = shape;
+ luma_distortion[mb_idx] = dist;
+ uchar mode = modes & 0xF;
+ luma_mode[mb_idx] = mode;
+ }
+
+ dwo_buffer[mb_idx*16*4 + lid_x + 16*0] = result.s0;
+ dwo_buffer[mb_idx*16*4 + lid_x + 16*1] = result.s1;
+ dwo_buffer[mb_idx*16*4 + lid_x + 16*2] = result.s2;
+ dwo_buffer[mb_idx*16*4 + lid_x + 16*3] = result.s3;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 655a314b..ef34e475 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -307,7 +307,8 @@ set (utests_sources
compiler_device_enqueue.cpp
compiler_sqrt_div.cpp
compiler_remove_negative_add.cpp
- compiler_fdiv2rcp.cpp)
+ compiler_fdiv2rcp.cpp
+ compiler_intra_prediction.cpp)
if (LLVM_VERSION_NODOT VERSION_GREATER 34)
SET(utests_sources
diff --git a/utests/compiler_intra_prediction.cpp b/utests/compiler_intra_prediction.cpp
new file mode 100644
index 00000000..5cb5814b
--- /dev/null
+++ b/utests/compiler_intra_prediction.cpp
@@ -0,0 +1,118 @@
+#include "utest_helper.hpp"
+#include <string.h>
+
+void compiler_intra_prediction(void)
+{
+ if (!cl_check_device_side_avc_motion_estimation()) {
+ return;
+ }
+ if (!cl_check_reqd_subgroup())
+ return;
+
+ OCL_CREATE_KERNEL("compiler_intra_prediction");
+
+ const size_t w = 80;
+ const size_t h = 48;
+ const size_t mv_w = (w + 15) / 16;
+ const size_t mv_h = (h + 15) / 16;
+
+ cl_image_format format;
+ cl_image_desc desc;
+
+ memset(&desc, 0x0, sizeof(cl_image_desc));
+ memset(&format, 0x0, sizeof(cl_image_format));
+
+ uint8_t *image_data1 = (uint8_t *)malloc(w * h); // src
+ for (size_t j = 0; j < h; j++) {
+ for (size_t i = 0; i < w; i++) {
+ if (i >= 32 && i <= 47 && j >= 16 && j <= 31)
+ image_data1[w * j + i] = 2;
+ else
+ image_data1[w * j + i] = 1;
+ }
+ }
+
+ format.image_channel_order = CL_R;
+ format.image_channel_data_type = CL_UNORM_INT8;
+ desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+ desc.image_width = w;
+ desc.image_height = h;
+ desc.image_row_pitch = 0;
+ OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1); // src
+
+ OCL_CREATE_BUFFER(buf[1], 0, mv_w * mv_h * sizeof(uint8_t), NULL);
+ OCL_CREATE_BUFFER(buf[2], 0, mv_w * mv_h * sizeof(uint16_t), NULL);
+ OCL_CREATE_BUFFER(buf[3], 0, mv_w * mv_h * sizeof(uint8_t), NULL);
+ OCL_CREATE_BUFFER(buf[4], 0, mv_w * mv_h * sizeof(uint32_t) * 16 * 8, NULL);
+ OCL_CREATE_BUFFER(buf[5], 0, mv_w * mv_h * sizeof(uint32_t) * 8 * 8, NULL);
+
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+ OCL_SET_ARG(3, sizeof(cl_mem), &buf[3]);
+ OCL_SET_ARG(4, sizeof(cl_mem), &buf[4]);
+ OCL_SET_ARG(5, sizeof(cl_mem), &buf[5]);
+
+ globals[0] = w;
+ globals[1] = h / 16;
+ locals[0] = 16;
+ locals[1] = 1;
+ OCL_NDRANGE(2);
+
+ OCL_MAP_BUFFER(1);
+ OCL_MAP_BUFFER(2);
+ OCL_MAP_BUFFER(3);
+ OCL_MAP_BUFFER(4);
+ OCL_MAP_BUFFER(5);
+ uint8_t *modes = (uint8_t *)buf_data[1];
+ uint16_t *residual = (uint16_t *)buf_data[2];
+ uint8_t *shape = (uint8_t *)buf_data[3];
+#define VME_DEBUG 0
+#if VME_DEBUG
+ uint32_t *dwo = (uint32_t *)buf_data[4];
+ uint32_t *pld = (uint32_t *)buf_data[5];
+ std::cout << std::endl;
+ for (uint32_t j = 0; j <= mv_h - 1; ++j) {
+ for (uint32_t i = 0; i <= mv_w - 1; ++i) {
+ uint32_t mv_num = j * mv_w + i;
+ std::cout << "******* mv num = " << mv_num << ": " << std::endl;
+ std::cout << "payload register result: " << std::endl;
+ for (uint32_t row_num = 0; row_num < 8; row_num++) {
+ for (int32_t idx = 7; idx >= 0; idx--)
+ printf("%.8x ", pld[mv_num * 64 + row_num * 8 + idx]);
+ printf("\n");
+ }
+ std::cout << std::endl;
+ std::cout << "writeback register result: " << std::endl;
+ for (uint32_t row_num = 0; row_num < 4; row_num++) {
+ for (int32_t wi = 7; wi >= 0; wi--)
+ printf("%.8x ", dwo[mv_num * 16 * 4 + row_num * 16 + wi]);
+ printf("\n");
+ for (int32_t wi = 15; wi >= 8; wi--)
+ printf("%.8x ", dwo[mv_num * 16 * 4 + row_num * 16 + wi]);
+ printf("\n");
+ }
+ std::cout << std::endl;
+ printf("modes: %u\n", modes[mv_num]);
+ std::cout << std::endl;
+ std::cout << "residual: " << residual[mv_num] << std::endl;
+ std::cout << std::endl;
+ printf("shape: %u\n", shape[mv_num]);
+ std::cout << std::endl;
+ }
+ }
+#endif
+ OCL_ASSERT(modes[7] == 2);
+ OCL_ASSERT(residual[7] == 266);
+ OCL_ASSERT(shape[7] == 0);
+
+ OCL_UNMAP_BUFFER(1);
+ OCL_UNMAP_BUFFER(2);
+ OCL_UNMAP_BUFFER(3);
+ OCL_UNMAP_BUFFER(4);
+ OCL_UNMAP_BUFFER(5);
+
+ free(image_data1);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_intra_prediction);
generated by cgit v1.2.3 (git 2.39.1) at 2024-11-21 13:38:24 +0000