From 5a288032ab2361030e7fa24bcbc3c5e8f5b07d3a Mon Sep 17 00:00:00 2001 From: Luo Xionghu Date: Wed, 14 Jun 2017 00:54:14 +0800 Subject: add utest compiler_intra_prediction for extenstion cl_intel_device_side_avc_motion_estimation. fix build warnings. Signed-off-by: Chuanbo Weng Signed-off-by: Xionghu Luo Reviewed-by: Yang Rong --- kernels/compiler_intra_prediction.cl | 91 +++++++++++++++++++++++++++ utests/CMakeLists.txt | 3 +- utests/compiler_intra_prediction.cpp | 118 +++++++++++++++++++++++++++++++++++ 3 files changed, 211 insertions(+), 1 deletion(-) create mode 100644 kernels/compiler_intra_prediction.cl create mode 100644 utests/compiler_intra_prediction.cpp 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 + +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); -- cgit v1.2.1