Name String | |
cl_intel_device_side_avc_motion_estimation | |
Contributors | |
Biju George, Intel | |
Raghukrishnan Embar, Intel | |
Ryan Lei, Intel | |
Ben Ashbaugh, Intel | |
Bartosz Sochacki, Intel | |
Krishna Madaparambil, Intel | |
Mateusz Tabaka, Intel | |
Contact | |
Biju George, Intel (biju.george 'at' intel.com) | |
Version | |
Version 5, November 9, 2018 | |
Number | |
OpenCL Extension #50 | |
Status | |
First Draft | |
Dependencies | |
OpenCL 1.2 is required. | |
The OpenCL Intel vendor extension cl_intel_subgroups is required. The | |
VME built-in functions are an extension of the subgroup functions | |
defined in cl_intel_subgroups. | |
The built-in functions which perform intra estimation depend on the | |
OpenCL Intel vendor extension cl_intel_media_block_io in order to | |
read-in the neighboring macroblock edge pixels. | |
The built-in functions which perform chroma based intra estimation | |
operations depend on the OpenCL Intel vendor extension | |
cl_intel_planar_yuv in order to create NV12 source images. | |
This extension is written against revision 29 of the OpenCL 2.0 API | |
specification, against revision 33 of the OpenCL 2.0 OpenCL C | |
specification, and against revision 32 of the OpenCL 2.0 extension | |
specification. | |
Overview | |
Video motion estimation (VME) is defined as of set motion estimation | |
operations that are used to determine the motion vectors, intra | |
estimation angles and macroblock partitioning combination that best | |
describe the transformation to the source macroblock, from blocks in | |
one or more previous reference pictures (inter-prediction), or from | |
other blocks in the same source picture (intra-prediction). It does | |
this by searching for spatial and temporal patterns on the current and | |
various forward and backward reference pictures. | |
The goal of this extension is to provide programmers with a fine- | |
grained interface to the AVC VME media sampler in Intel graphics | |
processors. It describes the specification of low-level built-in | |
functions, callable from OpenCL kernels, that facilitate the | |
programming of the VME media sampler to evaluate specific AVC motion | |
estimation operations. If only a coarser-level interface at the level | |
of built-in kernels suffices, then the Intel vendor extensions | |
cl_intel_motion_estimation and cl_advanced_motion_estimation may be | |
considered. | |
Built-in functions are defined for all the major operations of the VME | |
media sampler. The major operations of the AVC VME media sampler in | |
Intel Graphics Processors can be described as follows: | |
1. Integer motion estimation (IME) | |
Perform motion estimation on a given source macroblock in | |
a source image over a single or dual reference window in a | |
reference image, at full-pixel resolution, to determine the | |
best integer motion vectors and their associated distortions, | |
and the best macroblock shape partitioning combination. | |
2. Motion estimation refinement (REF) | |
Perform refinement operations on the results of IME. The two sub- | |
operations are: | |
Fractional motion estimation (FME) | |
Perform sub-pixel refinement on the results of an IME operation. | |
half-pixel (HPEL) or quarter-pixel (QPEL) refinements are performed | |
to determine the best sub-pixel motion vectors and their associated | |
distortions. | |
Bidirectional motion estimation (BME) | |
Perform bidirectional refinement on the results of an IME | |
operation using two reference images to check if the bidirectional | |
mode using two references yields lesser distortions. An FME can | |
optionally be performed implicitly as part of a bidirectional | |
refinement. | |
3. Skip and Intra check (SIC) | |
Performs the following two sub-operations: | |
Skip check (SKC) | |
Compute the pixel distortion of a user-specified shape and motion | |
vector combination. The VME media sampler fetches necessary pixels, | |
performs fractional and bidirectional filtering (as necessary), and | |
then computes the distortion between the derived reference and | |
source. The skip decision can optionally be enhanced to include a | |
4x4 forward transform, the results of which are compared against a | |
user specified threshold to emulate the effects of the forward | |
quantization zeroing effect. | |
Intra prediction estimation (IPE) | |
Perform intra prediction on a given source macroblock to determine | |
the best intra prediction modes and the best shape partitioning | |
combination. | |
New API Enums | |
Accepted as arguments to clGetDeviceInfo | |
CL_DEVICE_AVC_ME_VERSION_INTEL 0x410B | |
CL_DEVICE_AVC_ME_SUPPORTS_TEXTURE_SAMPLER_USE_INTEL 0x410C | |
CL_DEVICE_AVC_ME_SUPPORTS_PREEMPTION_INTEL 0x410D | |
Additional valid constant values returned by clGetDeviceInfo: | |
CL_AVC_ME_VERSION_1_INTEL 0x1 | |
Additions to Chapter 4 of the OpenCL Specification: | |
Modify the description of function clGetDeviceInfo(..) | |
Table 4.3 is extended to include the following enumeration constants. | |
+--------------------------+-----------+----------------------------+ | |
|cl_device_info |Return Type|Description | | |
+--------------------------+-----------+----------------------------+ | |
|CL_DEVICE_AVC_ME_VERSION_ |cl_uint |The AVC device-side motion | | |
|INTEL | |estimation API version | | |
| | |number supported by the | | |
| | |device and driver. | | |
+--------------------------+-----------+----------------------------+ | |
|CL_DEVICE_AVC_ME_SUPPORTS_|cl_bool |Is CL_TRUE if built-in | | |
|TEXTURE_SAMPLER_USE_INTEL | |functions using the texture | | |
| | |sampler may be used along | | |
| | |with AVC VME built-in | | |
| | |functions using the media | | |
| | |sampler on the device, and | | |
| | |CL_FALSE otherwise. | | |
+--------------------------+-----------+----------------------------+ | |
|CL_DEVICE_AVC_ME_SUPPORTS_|cl_bool |Is CL_TRUE if an enqueue of | | |
|PREEMPTION_INTEL | |an OpenCL kernel that uses | | |
| | |AVC VME built-in functions | | |
| | |supports preemption on the | | |
| | |device, and CL_FALSE | | |
| | |otherwise. (This may be | | |
| | |useful on platforms that | | |
| | |have strict latency | | |
| | |requirements on the GPU). | | |
+--------------------------+-----------+----------------------------+ | |
This extension requires a value of CL_AVC_ME_VERSION_1_INTEL returned | |
by clGetDeviceInfo for CL_DEVICE_AVC_ME_VERSION_INTEL. | |
Terms, Acronyms and Definitions | |
Adds a new section 6.13.14.X "VME built-in functions" to the OpenCL | |
2.0 C specification. | |
" | |
The following terms, acronyms and definitions are used in and | |
provide context for the VME built-in functions. | |
Macro-block (MB): | |
----------------- | |
An image is partitioned into macro-blocks of size 16x16 pixels. It is | |
the basic unit of processing for AVC video motion estimation | |
operations. | |
Shape: | |
------ | |
A MB may be partitioned into sub-blocks of one of the major shapes. A | |
sub-block with an 8x8 major shape may be further independently | |
partitioned into sub-blocks of one of the minor shapes. It is | |
represented by pre-defined shape enumeration values. | |
Major Shapes: | |
------------- | |
Shapes of 16x16, 16x8, 8x16, or 8x8 partitions of a MB. A 16x16 major | |
shape merely indicates that the MB was not further partitioned. | |
Minor Shapes: | |
------------- | |
Shapes of 8x8, 8x4, 4x8, or 4x4 sub-partitions of an 8x8 partition. A | |
8x8 minor shape merely indicates that the 8x8 major partition was not | |
further sub-partitioned. | |
Block: | |
------ | |
A sub-block of a MB with one of the major or minor shapes. | |
Reference Image: | |
---------------- | |
An image (typically from the previously decoded buffer in an encoder | |
pipeline) from which motion estimation predictions are made. | |
Source Image: | |
------------- | |
The current image for which motion estimation predictions are made. | |
Source Macro-block Offset: | |
-------------------------- | |
The 2D offset of the top left corner of the source MB in pixel units. | |
It is represented by a pair of unsigned 16-bit integers. | |
Reference Window Offset: | |
------------------------ | |
The 2D offset of the top left corner reference search window w.r.t to | |
the top left corner of the source MB in pixel units. It is represented | |
by a pair of signed 16-bit integers in the range [-2048, 2047]. | |
Reference identifier: | |
--------------------- | |
Reference identifiers are associated to pairs of forward(L0)/ | |
backward(L1) reference image parameters. Up to 16 pairs of reference | |
pairs of reference image parameters are permitted, with the permitted | |
values of reference identifiers ranging from 0 to 15. The reference | |
identifers are assigned in increasing order in which the reference | |
image parameter pairs are declared in the OpenCL kernel parameter | |
list. See new section 6.13.14.X "VME built-in functions" described | |
below for more details. | |
Motion Vector (MV): | |
------------------- | |
A 2D vector used for inter motion estimation that provides an offset | |
from the top left corner of a block in the source image to the top | |
left corner an identically sized block in the reference image. | |
Generally it is used to represent the best match of a block in the | |
reference image to a block in the source image. The best match is | |
determined as the block minimizing the distortion. MVs are specified | |
in QPEL resolution with the 2 LSB representing the fractional part | |
of the offset. It is represented by a pair of signed 16-bit signed | |
integers. | |
Packed Motion Vector: | |
-------------------------- | |
A motion vector represented as a packed 32-bit unsigned integer. The | |
lower 16 bits contains the X coordinate and the upper 16 bits contains | |
the Y coordinate. | |
Bidirectional Motion Vector (BMV): | |
---------------------------------- | |
A pair of MVs for the forward(L0) and backward(L1) images. Depending | |
on how the VME operation is configured only the forward or the | |
backward MV or both may be valid. | |
Packed Bidirectional Motion Vector: | |
----------------------------------- | |
A bidirectional MV represented as a packed 64-bit unsigned integer. | |
The lower 32-bits contain the forward packed MV, and the upper 32-bits | |
contain the backward packed MV. | |
Sum Of Absolute Difference (SAD): | |
--------------------------------- | |
The sum of absolute differences of every full/sub-pixel location in | |
the source block w.r.t every corresponding full/sub pixel in the | |
reference block as specified by a given MV. The sum of absolute | |
differences may be optionally Haar transform adjusted. It is | |
represented by an unsigned 16-bit integer value. | |
Haar Transform (HAAR): | |
---------------------- | |
A simple wavelet transform that is used to refine the distortion | |
measure of SAD. The per pixel difference goes through a 4x4 Haar | |
transform. Then the SAD is replaced by the sum of the absolute | |
values of the transform domain coefficients in the distortion. Haar | |
transform is used as a coarse estimation of the integer transform. | |
Motion Vector Cost Center (CC): | |
------------------------------- | |
A MV has an associated cost w.r.t a cost center coordinate. The | |
further away from the cost center, the larger will be the cost | |
associated with the MV. Cost centers are specified in QPEL resolution | |
with the 2 LSB representing the fractional part of the offset. | |
Motion Vector Cost Center Delta (CCD): | |
-------------------------------------- | |
The 2D offset of the cost center relative to the top left corner of | |
the source MB. Cost center deltas are specified in QPEL resolution | |
with the 2 LSB representing the fractional part of the offset. | |
It is represented by a pair of signed 16-bit integers. | |
Packed Motion Vector Cost Center Delta: | |
--------------------------------------- | |
A motion vector cost center delta represented as a packed 32-bit | |
unsigned integer. The lower 16 bits contains the X coordinate and the | |
upper 16 bits contains the Y coordinate. | |
Bidirectional Motion Vector Cost Center Delta: | |
---------------------------------------------- | |
A pair of cost center deltas for the forward and backward images. | |
Packed Bidirectional Motion Vector Cost Center Delta: | |
----------------------------------------------------- | |
A packed bidirectional motion vector cost center delta represented as | |
a 64-bit unsigned integer. The lower 32-bits contain the forward | |
packed CCD, and the upper 32-bits contain the backward packed CCD. | |
Motion Vector Cost: | |
------------------- | |
The MV cost is determined using a cost function described by a cost | |
table that is indexed based on power-of-two distances from the user | |
specified cost center, with a user specified precision (or unit) of | |
the distances from the cost center. | |
U4U4 Byte Format: | |
----------------- | |
Represents a value of (B<<S), where B, called base, is the 4 bit | |
LSB of the byte and S, called shift, is the 4 bit MSB of the byte. | |
Motion Vector Cost Table: | |
------------------------- | |
A table which specifies the cost penalties at 8 control points. The | |
first 7 control points represent the distances from cost center at | |
powers-of-two locations (2^0 to 2^6), and the last control point | |
represents the base penalty for distances that are out of range of the | |
cost function curve. It is represented by a packed array of 8 U4U4 | |
unsigned integer values. | |
Motion Vector Cost Precision: | |
----------------------------- | |
The precision (or unit) of the control points in the MV cost table. | |
It can be used to control the precision and range of the cost | |
function. It is represented by pre-defined cost precision enumeration | |
values. | |
Shape Cost: | |
----------- | |
The cost associated with encoding a particular partition shape using | |
inter or intra prediction. It is represented by a packed array of 10 | |
U4U4 unsigned integer values. | |
Distortion: | |
----------- | |
The distortion is the sum of SAD, MV cost, shape cost and multi- | |
reference cost for inter estimation, and the sum of SAD, mode cost, | |
shape cost and non-dc cost for intra estimation. It is a measure of | |
the cost of encoding a block and is represented by an unsigned 16-bit | |
integer value. | |
Intra Mode: | |
----------- | |
An intra-prediction angle which provides a prediction for the current | |
block from the edge pixels in its neighboring blocks. It is | |
represented by pre-defined intra mode enumeration values. | |
Intra Mode Cost: | |
---------------- | |
The cost associated with a computed intra mode for a block w.r.t a | |
predicted intra mode based on the computed intra modes for its | |
neighboring blocks. | |
Mode: | |
----- | |
The decision whether the inter-prediction or intra-prediction | |
minimizes distortion of a given MB. | |
Search unit (SU): | |
----------------- | |
The basic unit of searching. Possible reference search locations are | |
grouped in a predefined 4x4 pattern, and all locations within the same | |
group must be completely chosen or completely skipped. These | |
predefined groups are called search units. | |
Search Path (SP): | |
----------------- | |
The path taken during searching in a reference window. The steps taken | |
in a search path are in units of SUs. The search path must lie within | |
the defined search window. | |
Luma: | |
----- | |
Luma refers to either the Y-plane of a NV12 image or a regular image | |
with the image_channel_order and image_data_type restricted as CL_R | |
and CL_UNORM_INT8. | |
Chroma: | |
------ | |
Chroma refer to the UV-plane of a NV12 image. | |
Search Window (SW): | |
------------------- | |
The search area that will be covered during searching. The area of the | |
search window is limited to 2K luma pixels. | |
Search Window Configuration: | |
---------------------------- | |
The configuration of a search window which is a combination of the | |
search path and search window. | |
The pre-defined search window configurations are: | |
+--------------+-----------------------------------------------------+ | |
| EXHAUSTIVE | 48x40 SW with exhaustive single reference search (or| | |
| | 32x32 dual SW for exhaustive dual-reference search);| | |
| | an exhaustive search means that all SU within the | | |
| | search window are searched in a spiral pattern with | | |
| | the search center being the middle of the search | | |
| | window | | |
+--------------+-----------------------------------------------------+ | |
| SMALL | 28x28 SW with exhaustive search | | |
+--------------+-----------------------------------------------------+ | |
| TINY | 24x24 SW with exhaustive search | | |
+--------------+-----------------------------------------------------+ | |
| EXTRA TINY | 20x20 SW with exhaustive search | | |
+--------------+-----------------------------------------------------+ | |
| DIAMOND | 48x40 SW with diamond single reference search (or | | |
| | 32x32 dual SW for diamond dual-reference search); a | | |
| | diamond pattern search path is used for the first 16| | |
| | (or 7 per reference for dual reference search) SUs, | | |
| | and then gradient based searching is used for up to | | |
| | a maximum of 57 search unit. | | |
+--------------+-----------------------------------------------------+ | |
| LARGE DIAMOND| 48x40 SW with large diamond single reference search | | |
| | (or 32x32 dual SW for large diamond dual-reference | | |
| | search); a diamond pattern search pattern is used | | |
| | for the first 32 (or 10 per reference for dual | | |
| | reference search) SUs, and then gradient based | | |
| | searching is used for up to a maximum of 57 search | | |
| | units. | | |
+--------------+-----------------------------------------------------+ | |
Inter Estimation: | |
----------------- | |
The process of determining motion vectors and shapes that best | |
describe the transformation from 2D images from previously decoded | |
images in a video sequence to the currently processed image. | |
Intra-Prediction Estimation (IPE): | |
---------------------------------- | |
The process of determining prediction angles and shapes that best | |
describe the transformation from neighboring MBs in an image to the | |
currently processed MB in the same image. | |
Luma Mode: | |
---------- | |
The prediction angle returned by IPE for the luma component for a | |
block. It is represented by an unsigned 8-bit integer with the upper 4 | |
bits set to zero. | |
Integer Motion Estimation (IME): | |
------------------------------- | |
Inter-motion estimation in integer pixel resolution. | |
Fractional Motion Estimation (FME): | |
----------------------------------- | |
Inter-motion estimation in sub-pixel resolution. The result of integer | |
motion estimation on a reference image is used to perform fractional | |
refinement. | |
Bidirectional Motion Estimation (BME): | |
-------------------------------------- | |
The process of determining if the bi-directional prediction minimizes | |
the distortion w.r.t to unidirectional prediction. The results of IME | |
on forward(L0) and backward(L1) reference images are used to perform | |
bi-directional refinement. BME can be performed in integer or sub- | |
pixel resolution. If performed in sub-pixel resolution an implicit FME | |
operation is done before performing the BME. | |
Refinement (REF): | |
----------------- | |
A FME and/or BME refinement operation. | |
Skip/Spot Check (SKC): | |
---------------------- | |
The operation determining the distortion associated with a given | |
(uni or bidirectional) MV in a reference image(s) w.r.t a source | |
image. | |
Skip and Intra Check (SIC): | |
--------------------------- | |
The process of performing both SKC and IPE in the same operation. | |
Motion Check or Estimation (MCE): | |
--------------------------------- | |
A generic IME, REF, or SIC operation. | |
Forward Transform (FT): | |
----------------------- | |
An 8x8 or 4x4 integer transform used to transform the residual to the | |
frequency domain. | |
" | |
New OpenCL C Enums | |
Adds the following enums to the OpenCL C language. The corresponding | |
API enums have the same name with the prefix "CLK_" replaced with | |
"CL_". | |
These enums are used as arguments and return values of some of the | |
VME built-in function parameters and return values described in the | |
new section 6.13.14.X "VME built-in functions". | |
Interlaced image field polarity values: | |
--------------------------------------- | |
#define CLK_AVC_ME_INTERLACED_SCAN_TOP_FIELD_INTEL 0x0 | |
#define CLK_AVC_ME_INTERLACED_SCAN_BOTTOM_FIELD_INTEL 0x1 | |
Inter macro-block major shape values: | |
------------------------------------- | |
#define CLK_AVC_ME_MAJOR_16x16_INTEL 0x0 | |
#define CLK_AVC_ME_MAJOR_16x8_INTEL 0x1 | |
#define CLK_AVC_ME_MAJOR_8x16_INTEL 0x2 | |
#define CLK_AVC_ME_MAJOR_8x8_INTEL 0x3 | |
Inter macro-block minor shape values: | |
------------------------------------- | |
#define CLK_AVC_ME_MINOR_8x8_INTEL 0x0 | |
#define CLK_AVC_ME_MINOR_8x4_INTEL 0x1 | |
#define CLK_AVC_ME_MINOR_4x8_INTEL 0x2 | |
#define CLK_AVC_ME_MINOR_4x4_INTEL 0x3 | |
Inter macro-block major direction values: | |
----------------------------------------- | |
#define CLK_AVC_ME_MAJOR_FORWARD_INTEL 0x0 | |
#define CLK_AVC_ME_MAJOR_BACKWARD_INTEL 0x1 | |
#define CLK_AVC_ME_MAJOR_BIDIRECTIONAL_INTEL 0x2 | |
Inter (IME) partition mask values: | |
---------------------------------- | |
#define CLK_AVC_ME_PARTITION_MASK_ALL_INTEL 0x0 | |
#define CLK_AVC_ME_PARTITION_MASK_16x16_INTEL 0x7E | |
#define CLK_AVC_ME_PARTITION_MASK_16x8_INTEL 0x7D | |
#define CLK_AVC_ME_PARTITION_MASK_8x16_INTEL 0x7B | |
#define CLK_AVC_ME_PARTITION_MASK_8x8_INTEL 0x77 | |
#define CLK_AVC_ME_PARTITION_MASK_8x4_INTEL 0x6F | |
#define CLK_AVC_ME_PARTITION_MASK_4x8_INTEL 0x5F | |
#define CLK_AVC_ME_PARTITION_MASK_4x4_INTEL 0x3F | |
Slice type values: | |
------------------ | |
#define CLK_AVC_ME_SLICE_TYPE_PRED_INTEL 0x0 | |
#define CLK_AVC_ME_SLICE_TYPE_BPRED_INTEL 0x1 | |
#define CLK_AVC_ME_SLICE_TYPE_INTRA_INTEL 0x2 | |
Search window configuration: | |
---------------------------- | |
#define CLK_AVC_ME_SEARCH_WINDOW_EXHAUSTIVE_INTEL 0x0 | |
#define CLK_AVC_ME_SEARCH_WINDOW_SMALL_INTEL 0x1 | |
#define CLK_AVC_ME_SEARCH_WINDOW_TINY_INTEL 0x2 | |
#define CLK_AVC_ME_SEARCH_WINDOW_EXTRA_TINY_INTEL 0x3 | |
#define CLK_AVC_ME_SEARCH_WINDOW_DIAMOND_INTEL 0x4 | |
#define CLK_AVC_ME_SEARCH_WINDOW_LARGE_DIAMOND_INTEL 0x5 | |
#define CLK_AVC_ME_SEARCH_WINDOW_RESERVED0_INTEL 0x6 | |
#define CLK_AVC_ME_SEARCH_WINDOW_RESERVED1_INTEL 0x7 | |
SAD adjustment mode: | |
-------------------- | |
#define CLK_AVC_ME_SAD_ADJUST_MODE_NONE_INTEL 0x0 | |
#define CLK_AVC_ME_SAD_ADJUST_MODE_HAAR_INTEL 0x2 | |
Pixel resolution: | |
----------------- | |
#define CLK_AVC_ME_SUBPIXEL_MODE_INTEGER_INTEL 0x0 | |
#define CLK_AVC_ME_SUBPIXEL_MODE_HPEL_INTEL 0x1 | |
#define CLK_AVC_ME_SUBPIXEL_MODE_QPEL_INTEL 0x3 | |
Cost precision values: | |
---------------------- | |
#define CLK_AVC_ME_COST_PRECISION_QPEL_INTEL 0x0 | |
#define CLK_AVC_ME_COST_PRECISION_HPEL_INTEL 0x1 | |
#define CLK_AVC_ME_COST_PRECISION_PEL_INTEL 0x2 | |
#define CLK_AVC_ME_COST_PRECISION_DPEL_INTEL 0x3 | |
Inter bidirectional weights: | |
---------------------------- | |
#define CLK_AVC_ME_BIDIR_WEIGHT_QUARTER_INTEL 0x10 | |
#define CLK_AVC_ME_BIDIR_WEIGHT_THIRD_INTEL 0x15 | |
#define CLK_AVC_ME_BIDIR_WEIGHT_HALF_INTEL 0x20 | |
#define CLK_AVC_ME_BIDIR_WEIGHT_TWO_THIRD_INTEL 0x2B | |
#define CLK_AVC_ME_BIDIR_WEIGHT_THREE_QUARTER_INTEL 0x30 | |
Inter border reached values: | |
---------------------------- | |
#define CLK_AVC_ME_BORDER_REACHED_LEFT_INTEL 0x0 | |
#define CLK_AVC_ME_BORDER_REACHED_RIGHT_INTEL 0x2 | |
#define CLK_AVC_ME_BORDER_REACHED_TOP_INTEL 0x4 | |
#define CLK_AVC_ME_BORDER_REACHED_BOTTOM_INTEL 0x8 | |
Intra macro-block shape values: | |
------------------------------- | |
#define CLK_AVC_ME_INTRA_16x16_INTEL 0x0 | |
#define CLK_AVC_ME_INTRA_8x8_INTEL 0x1 | |
#define CLK_AVC_ME_INTRA_4x4_INTEL 0x2 | |
Inter skip block partition type: | |
-------------------------------- | |
#define CLK_AVC_ME_SKIP_BLOCK_PARTITION_16x16_INTEL 0x0 | |
#define CLK_AVC_ME_SKIP_BLOCK_PARTITION_8x8_INTEL 0x04000 | |
Inter skip motion vector mask: | |
------------------------------ | |
#define CLK_AVC_ME_SKIP_BLOCK_16x16_FORWARD_ENABLE_INTEL (0x1<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_16x16_BACKWARD_ ENABLE_INTEL (0x2<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_16x16_DUAL_ENABLE_INTEL (0x3<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_FORWARD_ENABLE_INTEL (0x55<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_BACKWARD_ENABLE_INTEL (0xAA<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_DUAL_ENABLE_INTEL (0xFF<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_0_FORWARD_ENABLE_INTEL (0x1<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_0_BACKWARD_ENABLE_INTEL (0x2<<24) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_1_FORWARD_ENABLE_INTEL (0x1<<26) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_1_BACKWARD_ENABLE_INTEL (0x2<<26) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_2_FORWARD_ENABLE_INTEL (0x1<<28) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_2_BACKWARD_ENABLE_INTEL (0x2<<28) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_3_FORWARD_ENABLE_INTEL (0x1<<30) | |
#define CLK_AVC_ME_SKIP_BLOCK_8x8_3_BACKWARD_ENABLE_INTEL (0x2<<30) | |
Block based skip type values: | |
----------------------------- | |
#define CLK_AVC_ME_BLOCK_BASED_SKIP_4x4_INTEL 0x0 | |
#define CLK_AVC_ME_BLOCK_BASED_SKIP_8x8_INTEL 0x80 | |
Luma intra partition mask values: | |
--------------------------------- | |
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_ALL_INTEL 0x0 | |
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_16x16_INTEL 0x6 | |
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_8x8_INTEL 0x5 | |
#define CLK_AVC_ME_INTRA_LUMA_PARTITION_MASK_4x4_INTEL 0x3 | |
Intra neighbor availability mask values: | |
---------------------------------------- | |
#define CLK_AVC_ME_INTRA_NEIGHBOR_LEFT_MASK_ENABLE_INTEL 0x60 | |
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_MASK_ENABLE_INTEL 0x10 | |
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_RIGHT_MASK_ENABLE_INTEL 0x8 | |
#define CLK_AVC_ME_INTRA_NEIGHBOR_UPPER_LEFT_MASK_ENABLE_INTEL 0x4 | |
Luma intra modes: | |
----------------- | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_INTEL 0x0 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_INTEL 0x1 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DC_INTEL 0x2 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DIAGONAL_DOWN_LEFT_INTEL 0x3 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_DIAGONAL_DOWN_RIGHT_INTEL 0x4 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_PLANE_INTEL 0x4 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_RIGHT_INTEL 0x5 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_DOWN_INTEL 0x6 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_VERTICAL_LEFT_INTEL 0x7 | |
#define CLK_AVC_ME_LUMA_PREDICTOR_MODE_HORIZONTAL_UP_INTEL 0x8 | |
Chroma intra modes: | |
------------------- | |
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_DC_INTEL 0x0 | |
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_HORIZONTAL_INTEL 0x1 | |
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_VERTICAL_INTEL 0x2 | |
#define CLK_AVC_ME_CHROMA_PREDICTOR_MODE_PLANE_INTEL 0x3 | |
Reference image select values: | |
------------------------------ | |
#define CLK_AVC_ME_FRAME_FORWARD_INTEL 0x1 | |
#define CLK_AVC_ME_FRAME_BACKWARD_INTEL 0x2 | |
#define CLK_AVC_ME_FRAME_DUAL_INTEL 0x3 | |
VME media sampler type default initialization literal: | |
------------------------------------------------------ | |
Pre-defined enumeration CLK_AVC_ME_INITIALIZE_INTEL. | |
IME payload type default initialization literal: | |
------------------------------------------------ | |
Pre-defined enumeration CLK_AVC_IME_PAYLOAD_INITIALIZE_INTEL. | |
REF payload type default initialization literal: | |
------------------------------------------------ | |
Pre-defined enumeration CLK_AVC_REF_PAYLOAD_INITIALIZE_INTEL. | |
SIC payload type default initialization literal: | |
------------------------------------------------ | |
Pre-defined enumeration CLK_AVC_SIC_PAYLOAD_INITIALIZE_INTEL. | |
IME result type default initialization literal: | |
----------------------------------------------- | |
Pre-defined enumeration CLK_AVC_IME_RESULT_INITIALIZE_INTEL. | |
IME result single reference streamout type default initialization literal: | |
-------------------------------------------------------------------------- | |
Pre-defined enumeration CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMOUT_INITIALIZE_INTEL. | |
IME result dual reference streamout type default initialization literal: | |
------------------------------------------------------------------------ | |
Pre-defined enumeration CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMOUT_INITIALIZE_INTEL. | |
IME result single reference streamin type default initialization literal: | |
------------------------------------------------------------------------- | |
Pre-defined enumeration CLK_AVC_IME_RESULT_SINGLE_REFERENCE_STREAMIN_INITIALIZE_INTEL. | |
IME result dual reference streamin type default initialization literal: | |
----------------------------------------------------------------------- | |
Pre-defined enumeration CLK_AVC_IME_RESULT_DUAL_REFERENCE_STREAMIN_INITIALIZE_INTEL. | |
REF result type default initialization literal: | |
----------------------------------------------- | |
Pre-defined enumeration CLK_AVC_REF_RESULT_INITIALIZE_INTEL. | |
SIC result type default initialization literal: | |
----------------------------------------------- | |
Pre-defined enumeration CLK_AVC_SIC_RESULT_INITIALIZE_INTEL. | |
New OpenCL C Types | |
Adds the following built-in types to the OpenCL C language. Appends | |
the following built-in type definitions to Table 6.3 in section 6.1.3 | |
"Other Built-in Data Types". | |
These built-in types are used to declare the types of some of the VME | |
built-in function parameters and return values in the new section | |
6.13.14.X "VME built-in functions". | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_mce_payload_t |This is a parameter and/or | | |
| |return type of a generic AVC MCE| | |
| |operation. It is used in MCE | | |
| |built-in functions and | | |
| |represents the payload for a | | |
| |basic IME/REF/SIC operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_payload_t |This is a parameter and/or | | |
| |return type of a basic AVC IME | | |
| |operation. It is used in IME | | |
| |built-in functions and | | |
| |represents the payload for an | | |
| |IME operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ref_payload_t |This is a parameter and/or | | |
| |return type of an AVC REF | | |
| |operation. It is used in REF | | |
| |built-in functions and | | |
| |represents the payload for a REF| | |
| |operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_sic_payload_t |This is a parameter and/or | | |
| |return type of an AVC SIC | | |
| |operation. It is used in SIC | | |
| |built-in functions and | | |
| |represents the payload for a SIC| | |
| |operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_mce_result_t |This is a parameter and/or | | |
| |return type of a MCE operation. | | |
| |It is used in MCE built-in | | |
| |functions and represents the | | |
| |evaluation result for a basic | | |
| |IME/REF/SIC operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_result_t |This is a parameter and/or | | |
| |return type of a basic AVC IME | | |
| |operation not using the stream- | | |
| |in/streamout functionality.It is| | |
| |used in IME built-in functions | | |
| |and represents the result for a | | |
| |basic IME evaluation operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_result_ |This is a parameter and/or | | |
| single_reference_streamout_t |return type of an AVC single | | |
| |reference IME operation using | | |
| |the streamout functionality. It | | |
| |is used in specific IME built- | | |
| |in functions and represents the | | |
| |evaluation result of an IME | | |
| |streamout operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_result_ |This is a parameter and/or | | |
| dual_reference_streamout_t |return type of an AVC dual | | |
| |reference IME operation using | | |
| |the streamout functionality. It | | |
| |is used in specific IME built-in| | |
| |functions and represents the | | |
| |evaluation result for an IME | | |
| |streamout operation. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_ |This is a parameter and/or | | |
| single_reference_streamin_t |return type of an AVC single IME| | |
| |operation using the stream-in | | |
| |reference functionality. It is | | |
| |used in specific IME built-in | | |
| |functions. It represents the | | |
| |additional results from the | | |
| |result of an IME evaluation | | |
| |using the streamout | | |
| |functionality that may be | | |
| |streamed-in in a subsequent IME | | |
| |streamin call. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ime_ |This is a parameter and/or | | |
| dual_reference_streamin_t |return type of an AVC dual | | |
| |reference IME operation using | | |
| |the streamin functionality. It | | |
| |is used in specific IME built- | | |
| |in functions. It represents the | | |
| |additional results from the | | |
| |result of an IME evaluation | | |
| |using the streamout | | |
| |functionality that may be | | |
| |streamed-in in a subsequent IME | | |
| |call. | | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_ref_result_t |This is a parameter and or/ | | |
| |return type of an AVC REF | | |
| |operation. It is used in REF | | |
| |built-in functions and | | |
| |represents the evaluation result| | |
| |from a REF evaluation operation.| | |
+-----------------------------------+--------------------------------+ | |
| intel_sub_group_avc_sic_result_t |This is a parameter and/or | | |
| |return type of an AVC SKC, IPE, | | |
| |or SIC operation. It is used in | | |
| |SIC built-in functions and | | |
| |represents the evaluation result| | |
| |from a SIC evaluation operation.| | |
+-----------------------------------+--------------------------------+ | |
Amends the following built-in type definition in Table 6.3 in section | |
6.1.3 "Other Built-in Data Types". | |
+---------+----------------------------------------------------------+ | |
|sampler_t|A texture or media sampler type. Refer to section 6.13.14 | | |
| |for a detailed description the built-in functions that use| | |
| |of this type. | | |
| | | | |
| |In general, the use of the word "sampler" by itself refers| | |
| |to a texture sampler. A media sampler will be explicitly | | |
| |refered to as "media sampler". | | |
| | | | |
| |In general, sampler_t parameters of built-in functions | | |
| |refer to texture sampler parameter. Media sampler | | |
| |parameters are exclusively used in the VME built-in | | |
| |functions (described in new section 6.13.14.X "VME | | |
| |built-in functions"). | | |
+---------+----------------------------------------------------------+ | |
Amends the restrictions defined for sampler types in (b) section 6.9 | |
"Restrictions". | |
"... | |
The texture sampler type (sampler_t) can only be used as the type of | |
a function argument or a variable declared in the program scope or | |
the outermost scope of a kernel function. The behavior of a sampler | |
variable declared in a non-outermost scope of a kernel function is | |
implementation-defined. The media sampler type cannot be used as a | |
function argument of a kernel function, but otherwise can be used as | |
the type of a function argument or a variable declared in any scope. | |
... " | |
Amends the following section 6.13.14.1 "Samplers". | |
Replace all strings "sampler(s)" with "texture sampler(s)". | |
Appends the following paragraphs in section 6.13.14.1 "Samplers". | |
"The VME evaluation phase built-in functions (described in new section | |
6.13.14.X "VME built-in functions" take a VME media sampler argument. | |
The media sampler can be declared as a constant variable in any scope, | |
and must be initialized with the 32-bit unsigned integer constant | |
CLK_AVC_ME_INITIALIZE_INTEL." | |
Note that VME media samplers declared using the constant qualifier are | |
not counted towards the maximum size of constant address space allowed | |
per device (i.e. CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE as described in | |
table 4.3) | |
Amends the following section 6.13.14.2 "Built-in Image Read | |
Functions". | |
"The following built-in function calls to read images with a texture | |
sampler are supported." | |
Amends the following section 6.13.14.3 "Built-in Image Sampler-less | |
Read Functions". | |
"The sampler-less read image functions behave exactly as the | |
corresponding read image functions described in section 6.13.14.2 | |
that take integer coordinates and a texture sampler with filter mode | |
set to CLK_FILTER_NEAREST, normalized coordinates set to | |
CLK_NORMALIZED_COORDS_FALSE and addressing mode to CLK_ADDRESS_NONE" | |
New OpenCL C built-in functions | |
Appends the following to the new section 6.13.14.X "VME built-in | |
functions". | |
" | |
The OpenCL C programming language implements the following built-in | |
functions to evaluate VME operations as a block operation by all work- | |
items in a subgroup. Furthermore they are defined only for subgroup | |
size of 16, and thus using these built-in functions in a kernel will | |
force a subgroup size of 16. The use of kernel attribute | |
intel_reqd_sub_group_size with a value other than 16 as described in | |
the Intel vendor extension cl_intel_required_subgroup_size is not | |
allowed. | |
These built-in functions are defined to operate on a 16x16 MB and must | |
be encountered by all work items in a subgroup executing the kernel, | |
otherwise the behavior is undefined (i.e. they can only be used only | |
in convergent control flow where all the work-items in the subgroups | |
are enabled). | |
______________________________________________________________________ | |
If VME built-in functions are used in an OpenCL kernel, then the | |
following additional requirements must be satisfied. | |
1. The VME source and reference images must either be 8-bit luma | |
2D images, or planar YUV images with NV12 format. If chroma based | |
intra estimation operations are being performed, then NV12 source | |
images are required, otherwise either 8-bit luma or NV12 images may | |
be used. For 8-bit luma images, the image_channel_order and the | |
image_data_type are restricted as CL_R and CL_UNORM_INT8 | |
respectively. NV12 images are described in the Intel vendor | |
extension cl_intel_planar_yuv. Images created from buffers are not | |
allowed as VME source and reference images. | |
2. OpenCL image kernel parameters that are used as VME source and | |
reference images in the VME evaluation built-in functions must | |
exclusively be used only by VME built-in functions. If the same | |
image needs to be used for other image operations within the same | |
kernel, then an additional image parameter may be used that is | |
bound to the same image object as for the VME image parameter in | |
the OpenCL host application for the kernel enqueue API call. | |
3. OpenCL image kernel parameters that are used as VME source and | |
reference images in the VME evaluation built-in functions must be | |
declared in a specific order in the kernel parameter list of a | |
kernel that calls any of the VME built-in functions as shown below; | |
otherwise it will result in a complier reported error. | |
+-----------------------------------------------+ | |
| __kernel void vme_custom( | | |
| read_only image2d_t src_image, | | |
| read_only image2d_t fwd0_ref_image, | | |
| read_only image2d_t bwd0_ref_image, | | |
| read_only image2d_t fwd1_ref_image, | | |
| read_only image2d_t bwd1_ref_image, | | |
| ... | | |
| read_only image2d_t fwd15_ref_image, | | |
| read_only image2d_t bwd15_ref_image, | | |
| ...) | | |
+-----------------------------------------------+ | |
The first forward(L0) reference image parameter must immediately | |
follow the source image parameter, and the first backward(L1) | |
reference image parameter must immediately follow the first forward | |
reference image parameter. Up to 16 consecutive pairs of forward | |
and backward reference image parameters can be provided, with the | |
reference images closer to the source image in timing order | |
appearing earlier in the parameter list. | |
______________________________________________________________________ | |
The VME built-in functions are organized into categories of functions | |
based on the major VME operations as follows: | |
1. MCE built-in functions | |
-------------------------- | |
A set of generic built-in functions which are common across IME, REF, | |
or SIC operations. Certain functions may not be applicable for some | |
specific operation as indicated in their descriptions. | |
2. IME built-in functions | |
------------------------- | |
A set of built-in functions to initialize, configure and evaluate an | |
integer motion estimation result. | |
3. REF built-in functions | |
------------------------- | |
A set of built-in functions to initialize, configure and evaluate a | |
fractional and/or bidirectional refinement operation on the results | |
of IME built-in functions. | |
4. SIC built-in functions | |
------------------------- | |
A set of built-in functions to initialize, configure and evaluate a | |
skip check or intra estimation operation result. | |
______________________________________________________________________ | |
A set of ordered phases of built-in functions need to be called to | |
evaluate each category of built-in functions. Some phases may not be | |
specified for certain function categories, but if present it must be | |
called in the correct order, otherwise the behavior is undefined. | |
1. Initialization (required phase) | |
--------------------------------- | |
This creates and initializes the payload for the VME operation and | |
may perform some basic configuration that invariably must be set. | |
2. Operation configuration (required phase, if present) | |
------------------------------------------------------ | |
This configures important sets of parameters for the VME operation | |
that were not set in the initialization phase in the payload. | |
3. Inter reference frame cost configuration (optional phase) | |
----------------------------------------------------------- | |
This configures the cost assigned to matching blocks in reference | |
frames based on timing distance from the source frame. | |
4. Inter motion vector/Intra mode cost configuration (optional phase) | |
-------------------------------------------------------------------- | |
This configures the inter motion vector cost function or the intra | |
mode cost function to be used for the VME operation in the payload. | |
5. Inter/Intra shape cost configuration (optional phase) | |
------------------------------------------------------- | |
This configures the inter/intra shape cost penalty to be used for the | |
VME operation in the payload. | |
6. Miscellaneous property configuration (optional phase) | |
------------------------------------------------------- | |
This configures miscellaneous operation sub-functions in the VME | |
operation payload. | |
7. Evaluation (required phase) | |
----------------------------- | |
This phase performs the evaluation of the VME operation using the | |
payload configured in the previous phases by issuing it to the VME | |
unit for processing. | |
8. Result processing (required phase) | |
------------------------------------ | |
The phase performs the extraction of various motion estimation | |
operation component results from the result of the evaluation phase. | |
The result components may be distributed among multiple work-items. | |
The results components of one VME operation may be used as input to | |
another VME operation at phase (1) or phase (7). | |
A set of built-in functions are defined for each of the major VME | |
operations of IME, REF, and SIC (SKC and IPE). These are categorized | |
as IME, REF, and SIC built-in functions. Functions that are common | |
across all categories are classified as MCE (motion check and | |
estimation) functions. Several restrictions for function argument | |
values are stated in the descriptions of the built-in function | |
descriptions. In general, if the restrictions are not satisfied the | |
behavior is undefined, unless otherwise stated. | |
______________________________________________________________________ | |
MCE built-in functions | |
---------------------- | |
A set of generic MCE operations which may be called for IME, REF, or | |
SIC operations with the restrictions as stated in their descriptions. | |
They can be called only during specific phases of these operations as | |
indicated in the description of the built-in functions. | |
Multi-reference cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable multi-reference image costing. They | |
allow for the configuration of the payloads to favorably bias the | |
major partitions coming from reference images that are closer to the | |
source image, than the ones coming from reference images that are | |
further away. The distance of the reference images, in the timing | |
order, from the source image is implied based on the order in which | |
the reference images are declared in the kernel parameter list. | |
+-------------------------------------+-----------------------------+ | |
|uchar |Get the default base | | |
|intel_sub_group_avc_mce_get_default_ |multi-reference cost penalty | | |
|inter_base_multi_reference_penalty( |in U4U4 format when HW | | |
| uchar slice_type, |assisted multi-reference | | |
| uchar qp ) |search is used. | | |
| | | | |
| |The value of slice_type must | | |
| |be a valid slice type | | |
| |enumeration value. | | |
| | | | |
| |The value of qp must be a | | |
| |valid quantization parameter | | |
| |value between 0 and 51. | | |
+-------------------------------------+-----------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Set the multi-reference base | | |
|intel_sub_group_avc_mce_set_inter_ |penalty when HW assisted | | |
|base_multi_reference_penalty( |multi-reference search is | | |
| uchar reference_base_penalty, |performed. | | |
| intel_sub_group_avc_mce_payload_t | | | |
| payload ) |Reference major partitions | | |
| |get associated with a penalty| | |
| |based on its distance from | | |
| |the source image. The | | |
| |reference_base_penalty is | | |
| |scaled using a scaling factor| | |
| |based on the implied distance| | |
| |of the reference image from | | |
| |the source image as shown | | |
| |below. | | |
| | | | |
| | 0 => 0x | | |
| | 1 to 2 => 1x | | |
| | 3 to 6 => 2x | | |
| | 7 to 15 => 3x | | |
| | | | |
| |The value of | | |
| |reference_base_penalty is | | |
| |specified in U4U4 format and | | |
| |the integer value must fit | | |
| |within 12 bits. | | |
+-------------------------------------+-----------------------------+ | |
Inter shape and direction cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable shape costing for inter estimation. | |
They allow for the configuration of payloads for the biasing of | |
certain shapes over others based on the configured parameters. | |
+------------------------------------+-------------------------------+ | |
|ulong intel_sub_group_avc_mce_get_ |Get the default packed shape | | |
|default_inter_shape_penalty( |cost for inter estimation in | | |
| uchar slice_type, |U4U4 format. | | |
| uchar qp ) | | | |
| |The value of slice_type must be| | |
| |a valid slice type enumeration | | |
| |value. | | |
| | | | |
| |The value of qp must be a valid| | |
| |quantization parameter value | | |
| |between 0 and 51. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Set the shape penalty for inter| | |
|intel_sub_group_avc_mce_set_inter_ |motion estimation. | | |
|shape_penalty( | | | |
| ulong packed_shape_penalty, |The value of packed_shape_cost | | |
| intel_sub_group_avc_mce_payload_t|is an unsigned long integer | | |
| payload ) |value with the following bits | | |
| |specifying the shape penalty in| | |
| |U4U4 format as follows: | | |
| | | | |
| | 7:0 => 16x8 and 8x16 | | |
| | 15:8 => 8x8 | | |
| | 23:16 => 8x4 and 4x8 | | |
| | 31:24 => 4x4 | | |
| | 39:32 => 16x16 | | |
| | 63:40 => must be zero | | |
| | | | |
| |The U4U4 decoded integer values| | |
| |for byte 0 and byte 4 must bit | | |
| |fit in 12 bits, while the U4U4 | | |
| |decoded integer values for the | | |
| |other bytes must fit within 10 | | |
| |bits. | | |
+------------------------------------+-------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_ |Get the default direction | | |
|default_inter_direction_penalty( |penalty for inter estimation in| | |
| uchar slice_type, |U4U4 format. | | |
| uchar qp ) | | | |
| |The value of slice_type must be| | |
| |a valid slice type enumeration | | |
| |value. | | |
| | | | |
| |The value of qp must be a valid| | |
| |quantization parameter value | | |
| |between 0 and 51. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Set the direction penalty for | | |
|intel_sub_group_avc_mce_set_ |backward images used in inter | | |
|inter_direction_penalty( |motion estimation. | | |
| uchar direction_cost, | | | |
| intel_sub_group_avc_mce_payload_t|The value of direction_cost | | |
| payload ) |must be in U4U4 format and its | | |
| |decoded integer value must bit | | |
| |fit in 12 bits. | | |
+------------------------------------+-------------------------------+ | |
Intra shape cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable shape costing for intra estimation. | |
They allow for the configuration of payloads for biasing of certain | |
shapes over others based on the configured parameters. Only the built- | |
in function providing the default shape penalty is specified as an MCE | |
function. The function which actually configures the payload for the | |
intra estimation operation is specified as a SIC built-in function. | |
+----------------------------------+--------------------------------+ | |
|uint intel_sub_group_avc_mce_get_ |Get the default packed luma | | |
|default_intra_luma_shape_penalty( |intra shape penalty in U4U4 | | |
| uchar slice_type, |format. | | |
| uchar qp ) | | | |
+----------------------------------+--------------------------------+ | |
Inter motion vector cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable motion vector costing for inter | |
estimation. The distortion measure is augmented to favor motion | |
vectors closer to the cost-center considered in conjunction with the | |
primary objective of minimizing the SAD between the source and | |
reference blocks. | |
+------------------------------------+-------------------------------+ | |
|uint2 intel_sub_group_avc_mce_get_ |Get the default inter motion | | |
|default_inter_motion_vector_ |vector cost table for the | | |
|cost_table( |pre-defined control points in | | |
| uchar slice_type, |U4U4 format for the input Qp | | |
| uchar qp ) |and slice type. | | |
| | | | |
| |The value of slice_type must be| | |
| |a valid slice type enumeration | | |
| |value. | | |
| | | | |
| |The value of qp must be a valid| | |
| |quantization parameter value | | |
| |between 0 and 51. | | |
+------------------------------------+-------------------------------+ | |
|uint2 intel_sub_group_avc_mce_get_ |Get the default predefined | | |
|default_high_penalty_cost_table( |packed U4U4 format high cost | | |
| void) |table for high Qp. This may be | | |
| |more appropriate for frame | | |
| |sequences with high motion. | | |
+------------------------------------+-------------------------------+ | |
|uint2 intel_sub_group_avc_mce_get_ |Get the default predefined | | |
|default_medium_penalty_cost_table( |packed U4U4 format medium cost | | |
| void) |table for medium Qp. This may | | |
| |be more appropriate for frame | | |
| |sequences with normal motion. | | |
+------------------------------------+-------------------------------+ | |
|uint2 intel_sub_group_avc_mce_get_ |Get the default predefined | | |
|default_low_penalty_cost_table( |packed U4U4 format low cost | | |
| void) |table for low Qp. This may be | | |
| |more appropriate for frame | | |
| |sequences with low motion. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Update the input payload to set| | |
|intel_sub_group_avc_mce_set_ |the cost precision along with | | |
|motion_vector_cost_function( |the cost center and cost table | | |
| ulong packed_cost_center_delta, |and return it. | | |
| uint2 packed_cost_table, | | | |
| uchar cost_precision, |The value of | | |
| intel_sub_group_avc_mce_payload_t|packed_cost_center_delta is the| | |
| payload ) |packed bidirectional cost | | |
| |center delta value relative to | | |
| |the source macroblock, which | | |
| |specifies the 4 bidirectional | | |
| |cost centers of each of the 8x8| | |
| |partitions of the reference | | |
| |image. If only unidirectional | | |
| |search is performed then the | | |
| |values of the backward | | |
| |reference cost centers must be | | |
| |zero. Work-item 'n' provides | | |
| |the value of cost center | | |
| |'n'. It is specified in QPEL | | |
| |units. For 16x16 partitions | | |
| |work-item '0' provides the cost| | |
| |center. For 8x16 partitions | | |
| |work-items '0' and '1' provide | | |
| |the cost centers. For 16x8 | | |
| |partitions work-items '0' and | | |
| |'2' provide the cost | | |
| |centers. The X and Y | | |
| |coordinates of each cost center| | |
| |delta must be in the range | | |
| |[-2048, 2047] and [-512.00 to | | |
| |511.75] respectively, otherwise| | |
| |the results are undefined. | | |
| | | | |
| |The packed cost table specifies| | |
| |the cost penalties for | | |
| |pre-defined control points in | | |
| |U4U4 format in the cost | | |
| |function curve. The first 7 | | |
| |bytes specify 7 control points | | |
| |representing consecutive | | |
| |powers-of-two delta units (2^0 | | |
| |to 2^6). Each delta unit, dx, | | |
| |is the distance of a motion | | |
| |vector, mv, from the specified | | |
| |cost center, cc | | |
| |(dx=|mv-cc|). The cost penalty | | |
| |values at in-between control | | |
| |points are linearly | | |
| |interpolated. The range of the | | |
| |cost function is defined to be | | |
| |from 2^0 to 2^6 delta | | |
| |units. The 8th byte of the | | |
| |packed cost table specifies the| | |
| |penalty base factor (over_cost)| | |
| |for dx distances that are | | |
| |out-of-range. The penalty of | | |
| |out-of-range cost dx distances | | |
| |is computed as min(over_cost + | | |
| |int(dx) - 64, 255). | | |
| | | | |
| |The value of the cost precision| | |
| |parameter must be a valid cost | | |
| |precision enumeration value, | | |
| |and specifies the precision of | | |
| |the delta units from the cost | | |
| |center, dx. This effectively | | |
| |can be used to control the | | |
| |range of the cost function as | | |
| |follows: | | |
| | | | |
| |PRECISION DELTAS PIXEL RANGE | | |
| |--------- ------ ----------- | | |
| | | | |
| |PEL pixel 0-64 | | |
| |DPEL dual 0-127 | | |
| | pixel | | |
| |HALF half 0-31 | | |
| | pixel | | |
| |QUARTER quarter 0-15 | | |
| | pixel | | |
| | | | |
| |The inter distortion for a | | |
| |block can be described by the | | |
| |following formula: | | |
| | | | |
| |Distortion = | | |
| | SAD (or HAAR) + | | |
| | MV_Cost_Penalty + | | |
| | Shape_Penalty + | | |
| | Direction_Cost + | | |
| | Multi_Reference_Penalty | | |
+------------------------------------+-------------------------------+ | |
Intra mode cost configuration phase functions | |
+++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable mode costing for intra estimation. | |
They allow for the configuration of payloads to bias the computed | |
intra modes to be closer to their configured neighbor modes. This form | |
of costing is similar to the inter motion vector costing. Only the | |
built-in functions providing the defaults mode costs are specified as | |
MCE functions. The remaining functions which actually configure the | |
payload for the intra estimation operation is specified as SIC built- | |
in functions. | |
+----------------------------------+---------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_|Get the default intra mode cost | | |
|default_intra_luma_mode_penalty( |penalty for in U4U4 format when | | |
| uchar slice_type, |the estimated mode differs from | | |
| uchar qp ) |its predicted mode from its | | |
| |neighbors. | | |
| | | | |
| |The value of slice_type must be a| | |
| |valid slice type enumeration | | |
| |value. | | |
| | | | |
| |The value of qp must be a valid | | |
| |quantization parameter value | | |
| |between 0 and 51. | | |
+----------------------------------+---------------------------------+ | |
|uint intel_sub_group_avc_mce_get_ |Get the default intra non-dc cost| | |
|default_non_dc_luma_intra_penalty(|penalty for intra luma estimation| | |
| void ) |in packed 32-bit integer format. | | |
+----------------------------------+---------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_|Get the default chroma mode base | | |
|default_intra_chroma_mode_base_ |penalty in U4U4 format. | | |
|penalty( | | | |
| void ) | | | |
+----------------------------------+---------------------------------+ | |
Miscellaneous property configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These built-in functions enable miscellaneous MCE properties settings. | |
+--------------------------------------+-----------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Update the input payload to | | |
|intel_sub_group_avc_mce_set_ |enable an AC only HAAR SAD | | |
|ac_only_haar( |mode and return it. It | | |
| intel_sub_group_avc_mce_payload_t |overrides any previous | | |
| payload ) |setting for sad adjustment. | | |
| | | | |
| |This feature is mainly | | |
| |intended for improved block | | |
| |matching in frame-rate | | |
| |conversion (FRC) kernels. | | |
+--------------------------------------+-----------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Update the input payload to | | |
|intel_sub_group_avc_mce_set_source_ |specify the field polarities | | |
|interlaced_field_polarity( |for interlaced source images | | |
| uchar src_field_polarity, |used for inter or intra | | |
| intel_sub_group_avc_mce_payload_t |operations. | | |
| payload ) | | | |
| |The value of | | |
| |src_field_polarity must be a | | |
| |valid field polarity | | |
| |enumeration value indicating | | |
| |the field polarity for the | | |
| |source image. | | |
+--------------------------------------+-----------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Update the input payload to | | |
|intel_sub_group_avc_mce_set_single_ |specify the field polarities | | |
|reference_interlaced_field_polarity( |for interlaced reference | | |
| uchar ref_field_polarity, |images used for single | | |
| intel_sub_group_avc_mce_payload_t |reference inter search or | | |
| payload ) |check operation. | | |
| | | | |
| |The value of | | |
| |ref_field_polarity must be a | | |
| |valid field polarity | | |
| |enumeration value indicating | | |
| |the field polarity for the | | |
| |reference image. | | |
+--------------------------------------+-----------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Update the input payload to | | |
|intel_sub_group_avc_mce_set_dual_ |specify the field polarities | | |
|reference_interlaced_field_polarities(|for interlaced reference | | |
| uchar fwd_ref_field_polarity, |images used for dual | | |
| uchar bwd_ref_field_polarity, |reference inter search or | | |
| intel_sub_group_avc_mce_payload_t |check operation. | | |
| payload ) | | | |
| |The value of | | |
| |fwd_ref_field_polarity must | | |
| |be a valid field polarity | | |
| |enumeration value indicating | | |
| |the field polarity for the | | |
| |forward image. | | |
| | | | |
| |The value of | | |
| |bwd_ref_field_polarity must | | |
| |be a valid field polarity | | |
| |enumeration value indicating | | |
| |the field polarity for the | | |
| |backward image. | | |
+--------------------------------------+-----------------------------+ | |
Result processing phase functions | |
+++++++++++++++++++++++++++++++++ | |
These built-in functions facilitate the extraction of components of | |
the result from VME unit. | |
+-----------------------------------+--------------------------------+ | |
|ulong intel_sub_group_avc_mce_get_ |Get the MCE packed BMVs result. | | |
|motion_vectors( | | | |
| intel_sub_group_avc_mce_result_t|Up to 16 packed BMVs are | | |
| result ) |returned, one per work-item. If | | |
| |the MCE search operation's | | |
| |payload was setup for | | |
| |unidirectional search then only | | |
| |the forward packed MV will be | | |
| |valid in each BMV, otherwise | | |
| |both packed MVs will be | | |
| |valid. The BMVs have to be | | |
| |selected by their respective | | |
| |work-items based on the result | | |
| |block major and minor shapes. | | |
| | | | |
| |If the major shape is: | | |
| | - 16x16, then one BMV is | | |
| | returned by work-item 0 | | |
| | - 16x8, or 8x16, then two | | |
| | BMVs are returned by work- | | |
| | items 0 and 8 | | |
| | - 8x8, then four sets of BMVs| | |
| | corresponding to the four | | |
| | partitions in traditional | | |
| | Z-order are returned by | | |
| | work-items in the ranges | | |
| | [0, 3], [4, 7], [8, 11], | | |
| | and [12, 15]; the minor | | |
| | shape will determine | | |
| | exactly which work-items | | |
| | in the reserved inclusive | | |
| | range for the partition | | |
| | returns the BMVs for that | | |
| | partition | | |
| | | | |
| |If the range of work-items for | | |
| |the 8x8 major partition is [n, | | |
| |n+3] and the minor shape is: | | |
| | - 8x8, then work-item 'n' | | |
| | returns the BMV for each | | |
| | minor partition | | |
| | - 8x4 or 4x8, then work-items| | |
| | 'n' and 'n+2' returns the | | |
| | BMVs for each minor | | |
| | partition | | |
| | - 4x4, then all work-items in| | |
| | [n, n+3] return the BMVs | | |
| | for each minor partition | | |
| | in traditional Z-order | | |
| | | | |
| |NOTES: | | |
| | | | |
| |(1) All sub-block BMVs get | | |
| |replicated for each partition. | | |
| |For example, for a 16x16 | | |
| |partition, all smaller sub-block| | |
| |BMVs are replicated to the same | | |
| |BMV, and for 8x8 partition, each| | |
| |8x8 must have its respective | | |
| |sub-block BMVs replicated. This | | |
| |is not important to extract the | | |
| |component BMVs itself, but is | | |
| |needed if the result of this | | |
| |function is used to initialize | | |
| |the input motion vectors of a | | |
| |REF initialization function. | | |
| | | | |
| |(2) With interlaced images, the | | |
| |MBs for the top field MBs are | | |
| |considered as logically | | |
| |overlapping with the bottom MBs.| | |
+-----------------------------------+--------------------------------+ | |
|ushort intel_sub_group_avc_mce_get_|Get the MCE inter distortions | | |
|inter_distortions( |result corresponding to the BMVs| | |
| intel_sub_group_avc_mce_result_t|returned by intel_sub_group_avc_| | |
| result ) |mce_get_motion_vectors(..). The | | |
| |MCE inter directions result | | |
| |returned by intel_sub_group_avc_| | |
| |mce_get_inter_directions(..) | | |
| |will specify if the distortion | | |
| |corresponds to the forward MV, | | |
| |backward MV, or the | | |
| |bidirectional MV in the BMV. Up | | |
| |to 16 distortions are returned, | | |
| |one per work-item. | | |
| | | | |
| |The distortions have to be | | |
| |selected by their respective | | |
| |work-items based on the result | | |
| |block major and minor shapes | | |
| |just as for the result MVs as | | |
| |described above. | | |
+-----------------------------------+--------------------------------+ | |
|ushort intel_sub_group_avc_mce_get_|Get the best inter distortion | | |
|best_inter_distortion( |for the whole MB. | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) | | | |
+-----------------------------------+--------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_ |Get the MCE inter MB major | | |
|inter_major_shape( |partition shape. | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) |The returned values are as per | | |
| |the inter-MB major shapes | | |
| |enumeration values. | | |
| | | | |
| |This can only be called as part | | |
| |of an IME or REF operation | | |
| |evaluation. | | |
+-----------------------------------+--------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_ |Get the MCE inter MB minor | | |
|inter_minor_shapes( |partition shapes. | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) |It returns a bit field with the | | |
| |minor shapes for the 4 8x8 | | |
| |sub-partitions in traditional Z | | |
| |order. Two bits are reserved for| | |
| |each of the four sub-partitions | | |
| |in row-major order. The | | |
| |returned 2-bit values are as per| | |
| |the inter-MB minor shapes | | |
| |enumeration values. | | |
| | | | |
| |This function returns valid | | |
| |results only if the major shape | | |
| |is 8x8, otherwise the results | | |
| |are undefined. | | |
| | | | |
| |This can only be called as part | | |
| |of an IME or REF operation | | |
| |evaluation. | | |
+-----------------------------------+--------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_ |Get the MCE inter MB major | | |
|inter_directions( |partition directions. | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) |It returns a bit field with the | | |
| |direction for up to 4 major | | |
| |sub-partitions in traditional Z | | |
| |order. Two bits are reserved for| | |
| |each of the four | | |
| |sub-partitions. The returned | | |
| |2-bit values are as per the | | |
| |inter-MB major shape direction | | |
| |enumeration values. | | |
| | | | |
| |If the major partition is: | | |
| | | | |
| |- 16x16, then bits in the range | | |
| | [0, 1] contains the direction | | |
| |- 16x8 or 8x16, then bits in the| | |
| | ranges [0, 1] and [2, 3] | | |
| | contains the two partitions | | |
| | directions | | |
| |- 8x8, then bits in the ranges | | |
| | [0, 1], [2, 3], [4, 5], and | | |
| | [6, 7] contains the four | | |
| | partitions directions | | |
| | | | |
| |The returned values are as per | | |
| |the inter direction enumeration | | |
| |values. | | |
| | | | |
| |This can only be called as part | | |
| |of an IME or REF operation | | |
| |evaluation. | | |
+-----------------------------------+--------------------------------+ | |
|uchar intel_sub_group_avc_mce_get_ |Get the count of motion vectors | | |
|inter_motion_vector_count( |(based on the partitioning | | |
| intel_sub_group_avc_mce_result_t|decision) returned by the search| | |
| result ) |operation. | | |
| | | | |
| |This can only be called as part | | |
| |of an IME or REF operation | | |
| |evaluation. | | |
+-----------------------------------+--------------------------------+ | |
|uint intel_sub_group_avc_mce_get_ |Get the MCE inter MB reference | | |
|inter_reference_ids( |identifiers in a packed integer | | |
| intel_sub_group_avc_mce_result_t|format, with the following bits | | |
| result ) |specifying the reference | | |
| |identifiers for the major | | |
| |partitions. | | |
| | | | |
| |3:0 => Fwd reference block 0 | | |
| |7:4 => Bwd reference block 0 | | |
| |11:8 => Fwd reference block 1 | | |
| |15:12 => Bwd reference block 1 | | |
| |19:16 => Fwd reference block 2 | | |
| |23:20 => Bwd reference block 2 | | |
| |27:24 => Fwd reference block 3 | | |
| |31:28 => Bwd reference block 3 | | |
| | | | |
| |The values of each individual | | |
| |4-bit reference identifier range| | |
| |from 0 to 15, with each value | | |
| |identifying the distance of | | |
| |ordered pair of forward/backward| | |
| |reference images as declared in | | |
| |the VME kernel parameter | | |
| |interface list. | | |
| | | | |
| |If the dual-reference evaluation| | |
| |functions are not used, then the| | |
| |values of the backward reference| | |
| |identifiers are undefined. | | |
| | | | |
| |The blocks are numbered using | | |
| |the traditional Z order. For | | |
| |larger block sizes, the | | |
| |sub-block reference identifier | | |
| |pairs are replicated. For | | |
| |example, for a 16x16 block all | | |
| |four pairs of reference | | |
| |identifiers are replicated to | | |
| |the value of the first pair for | | |
| |block 0. | | |
| | | | |
| |NOTE: Unless HW assisted | | |
| |multi-reference search was | | |
| |performed using the IME | | |
| |streamin/streamout evaluation | | |
| |functions, the individual 4-bit | | |
| |reference identifier pair values| | |
| |will all be the same (pointing | | |
| |to the same pair for | | |
| |forward/backward reference | | |
| |images). | | |
+-----------------------------------+--------------------------------+ | |
|uchar |Get the MCE inter MB reference | | |
|intel_sub_group_avc_mce_get_inter_ |field polarities for the | | |
|reference_interlaced_field_ |corresponding reference | | |
|polarities( |identifiers returned by | | |
| uint packed_reference_ids, |intel_sub_group_avc_ | | |
| uint |mce_get_inter_reference_ids(..) | | |
| packed_reference_parameter_ |in a packed integer format, with| | |
| field_polarities, |the following bits specifying | | |
| intel_sub_group_avc_mce_result_t|the reference field polarities | | |
| result ) |for the major partitions. | | |
| | | | |
| | 0 : Fwd reference block 0 | | |
| | 1 : Fwd reference block 1 | | |
| | 2 : Fwd reference block 2 | | |
| | 3 : Fwd reference block 3 | | |
| | 4 : Bwd reference block 0 | | |
| | 5 : Bwd reference block 1 | | |
| | 6 : Bwd reference block 2 | | |
| | 7 : Bwd reference block 3 | | |
| | | | |
| |If the dual-reference evaluation| | |
| |functions are not used, then the| | |
| |values of the backward reference| | |
| |field polarities are undefined. | | |
| | | | |
| |The blocks are numbered using | | |
| |the traditional Z order. For | | |
| |larger block sizes, the | | |
| |sub-block reference field | | |
| |polarities are replicated. For | | |
| |example, for a 16x16 block all | | |
| |four pairs of reference field | | |
| |polarities are replicated to the| | |
| |value of the first pair for | | |
| |block 0. | | |
| | | | |
| |The value of | | |
| |packed_reference_ids is as | | |
| |defined by the return value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_reference_ids(..). | | |
| | | | |
| |The value of | | |
| |packed_reference_parameter_ | | |
| |field_polarities specifies the | | |
| |packed bit field of field | | |
| |polarities for each of the (up | | |
| |to 16) forward/backward | | |
| |interleaved pairs of reference | | |
| |images in the same order as | | |
| |specified in the kernel | | |
| |parameter list, as used for the | | |
| |inter search operation. If less | | |
| |than 16 pairs are used then the | | |
| |corresponding bit field values | | |
| |are ignored. | | |
| | | | |
| |NOTE: An important restriction | | |
| |is that when multiple IME | | |
| |operations are performed for a | | |
| |HW multi-assisted | | |
| |multi-reference search operation| | |
| |using the streamin/streamout | | |
| |capabilities, the same reference| | |
| |image parameter cannot be used | | |
| |with different polarities in the| | |
| |sequence of IME operations used | | |
| |for a HW-assisted search | | |
| |operation. In other words, the | | |
| |field polarities for reference | | |
| |image parameters must be used | | |
| |consistently across IME | | |
| |operations used in a HW assisted| | |
| |multi-reference search | | |
| |operation. | | |
+-----------------------------------+--------------------------------+ | |
______________________________________________________________________ | |
IME built-in functions | |
---------------------- | |
A set of ordered phases of functions are required to be called to | |
evaluate an integer motion estimation result. | |
Initialization phase functions | |
++++++++++++++++++++++++++++++ | |
These built-in functions create a properly initialized payload that | |
can be used for further configured for evaluating IME operations. This | |
is a required initial phase. | |
+---------------------------------+----------------------------------+ | |
|intel_sub_group_avc_ime_payload_t|Return an initialized payload for | | |
|intel_sub_group_avc_ime_ |a VME integer search (IME) | | |
|initialize( |operation. | | |
| ushort2 src_coord, | | | |
| uchar partition_mask, |The payload is initialized for | | |
| uchar sad_adjustment ) |progressive frame operations, and | | |
| |the cost configuration values and | | |
| |the miscellaneous property values | | |
| |are all initialized to zero. The | | |
| |cost configuration and the | | |
| |miscellaneous property | | |
| |configuration phase functions must| | |
| |be used to override the initial | | |
| |configurations in the payload. | | |
| | | | |
| |The src_coord value represents the| | |
| |2D offset of the top left corner | | |
| |of the source MB in pixel units in| | |
| |the source image. Source MBs at | | |
| |the image borders are allowed to | | |
| |be partial, but the top-left | | |
| |corner must be within the image. | | |
| | | | |
| |If the source image is an | | |
| |interlaced scan image, then the | | |
| |bottom field lines are considered | | |
| |as logically overlapping with the | | |
| |top field lines (i.e. the top | | |
| |field MBs are considered as | | |
| |logically overlapping with the | | |
| |bottom MBs) for the purposes for | | |
| |specifying the src_coord value. | | |
| | | | |
| |The legal values for | | |
| |partition_mask can be composed by | | |
| |setting the appropriate bit fields| | |
| |specified by partition mask | | |
| |enumeration values using the '&' | | |
| |operator. | | |
| | | | |
| |The legal values for | | |
| |sad_adjustment is specified by its| | |
| |respective enumeration values. | | |
| | | | |
| |If the sad_adjustment is set to | | |
| |CLK_AVC_ME_SAD_ADJUST_MODE_HAAR_ | | |
| |INTEL, a simple wavelet transform,| | |
| |Haar transform, is used to refine | | |
| |the distortion measure of | | |
| |SAD. Haar transform here is used | | |
| |as a coarse estimation of the | | |
| |integer transform. | | |
+---------------------------------+----------------------------------+ | |
Configuration phase functions | |
+++++++++++++++++++++++++++++ | |
These built-in functions allow for configuration of the search window. | |
A call to either intel_sub_group_avc_ime_set_single_reference(..) or | |
intel_sub_group_avc_ime_set_dual_reference(..) is required. This is a | |
required phase immediately following the initialization phase | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Update the input payload for a | | |
|intel_sub_group_avc_ime_set_ |VME single-reference search | | |
|single_reference( |with the configuration for the | | |
| short2 ref_offset, |reference window search region,| | |
| uchar search_window_config, |and return it. | | |
| intel_sub_group_avc_ime_payload_t| | | |
| payload ) |The 2D ref_offset specifies the| | |
| |reference window offset. The X | | |
| |and Y coordinates must be in | | |
| |the range [-2048, 2047], | | |
| |otherwise the results are | | |
| |undefined. The reference window| | |
| |is allowed to be partially | | |
| |outside the image. Pixel | | |
| |replication is applied to | | |
| |generate out-of-bound reference| | |
| |pixels. It is specified in PEL | | |
| |units. Results are undefined in| | |
| |the reference region is | | |
| |completely outside the image. | | |
| | | | |
| |If the reference image is an | | |
| |interlaced scan image, then the| | |
| |top field lines are considered | | |
| |as logically overlapping with | | |
| |the bottom field lines | | |
| |(i.e. the top field MBs are | | |
| |considered as logically | | |
| |overlapping with the bottom | | |
| |MBs) for the purposes for | | |
| |specifying the ref_offset | | |
| |value. | | |
| | | | |
| |The parameter | | |
| |search_window_config must be a | | |
| |compile-time constant. | | |
| | | | |
| |The value of | | |
| |search_window_config must be | | |
| |one of the unreserved search | | |
| |window configuration | | |
| |enumeration values. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Update the input payload for a | | |
|intel_sub_group_avc_ime_ |VME dual-reference search with | | |
|set_dual_reference( |the configurations for the | | |
| short2 fwd_ref_offset, |reference window search | | |
| short2 bwd_ref_offset, |regions, and return it. | | |
| uchar search_window_config, | | | |
| intel_sub_group_avc_ime_payload_t|The 2D | | |
| payload ) |fwd_ref_offset/bwd_ref_offset | | |
| |specifies the forward/backward | | |
| |reference window offsets. The X| | |
| |and Y coordinates must be in | | |
| |the range [-2048, 2047], | | |
| |otherwise the results are | | |
| |undefined. The reference | | |
| |windows are allowed to be | | |
| |partially outside the | | |
| |image. Pixel replication is | | |
| |applied to generate | | |
| |out-of-bound reference | | |
| |pixels. It is specified in PEL | | |
| |units. | | |
| | | | |
| |If a reference image is an | | |
| |interlaced scan image, then the| | |
| |top field lines are considered | | |
| |as logically overlapping with | | |
| |the bottom field lines | | |
| |(i.e. the top field MBs are | | |
| |considered as logically | | |
| |overlapping with the bottom | | |
| |MBs) for the purposes for | | |
| |specifying the corresponding | | |
| |fwd_ref_offset and/or | | |
| |bwd_ref_offset values. | | |
| | | | |
| |The parameter | | |
| |search_window_config must be a | | |
| |compile-time constant and must | | |
| |be one of the unreserved search| | |
| |window configuration | | |
| |enumeration values. | | |
+------------------------------------+-------------------------------+ | |
|ushort2 intel_sub_group_avc_ime_ |Get the 2D size of the | | |
|ref_window_size( |reference window in pixel | | |
| uchar search_window_config, |units. | | |
| char dual_ref ) | | | |
| |The value of | | |
|// deprecated |search_window_config must be | | |
|ushort2 intel_sub_group_ime_ |one of the unreserved search | | |
|ref_window_size( |window configuration | | |
| uchar search_window_config, |enumeration values. | | |
| char dual_ref ) | | | |
| |The value of dual_ref must be | | |
| |set to zero for a single | | |
| |reference search window and one| | |
| |for a dual-reference search | | |
| |window. | | |
+------------------------------------+-------------------------------+ | |
|short2 |If the input 2D reference | | |
|intel_sub_group_avc_ime_ |window offset, ref_offset, | | |
|adjust_ref_offset( |causes the reference window to | | |
| short2 ref_offset, |be fully out-of-bound of the | | |
| ushort2 src_coord, |reference image, adjust it such| | |
| ushort2 ref_window_size, |that the reference window is | | |
| ushort2 image_size ) |within bounds of the reference | | |
| |image. | | |
| | | | |
| |The 2D ref_offset specifies the| | |
| |reference window offset. The X | | |
| |and Y coordinates must be in | | |
| |the range [-2048, 2047], | | |
| |otherwise the results are | | |
| |undefined. | | |
| | | | |
| |The src_coord value represents | | |
| |the 2D offset of the top left | | |
| |corner of the source MB in | | |
| |pixel units in the source | | |
| |image. Source MBs at the image | | |
| |borders are allowed to be | | |
| |partial. Results are undefined | | |
| |if the source MB is completely | | |
| |outside the image. | | |
| | | | |
| |The ref_window_size specifies | | |
| |the 2D size of the reference | | |
| |window in pixel units. | | |
| | | | |
| |The image_size specifies the 2D| | |
| |size of the progressive scan, | | |
| |or top or bottom fields, of the| | |
| |interlaced scan image in pixel | | |
| |units. | | |
| | | | |
| |If the reference image is an | | |
| |interlaced scan image, then the| | |
| |bottom field lines are | | |
| |considered as logically | | |
| |overlapping with the top field | | |
| |lines the purposes for | | |
| |specifying the image_size | | |
| |value. Since, the actual layout| | |
| |of the top and bottom fields in| | |
| |the reference image is in an | | |
| |interleaved fashion, the height| | |
| |of the top or bottom fields | | |
| |should be exactly half of the | | |
| |actual reference image height. | | |
| | | | |
| |A call to | | |
| |intel_sub_group_avc_ime_ | | |
| |adjust_ref_coord(..) is | | |
| |optional. It is required only | | |
| |if the reference window offsets| | |
| |inputs to | | |
| |intel_sub_group_avc_ime_set_ | | |
| |single_reference(..) or | | |
| |intel_sub_group_avc_ime_set_ | | |
| |dual_reference (..) is | | |
| |potentially out-of-bounds and | | |
| |need to be adjusted. | | |
+------------------------------------+-------------------------------+ | |
Payload type conversion functions | |
+++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
search configuration phase to convert IME payload to MCE payloads and | |
vice-versa. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Convert the IME payload to a | | |
|intel_sub_group_avc_ime_convert_ |generic MCE payload. | | |
|to_mce_payload( | | | |
| intel_sub_group_avc_ime_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Convert the generic MCE payload| | |
|intel_sub_group_avc_ |to a IME payload. | | |
|mce_convert_to_ime_payload( | | | |
| intel_sub_group_avc_mce_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
Multi-reference cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
search configuration phase to enable multi-reference image costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_ | | |
|inter_base_multi_reference_penalty( |set_inter_base_multi_reference_| | |
| uchar reference_base_penalty, |penalty(..) the payload | | |
| intel_sub_group_avc_ime_payload_t|conversions with to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
Inter shape and direction cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
search configuration phase to enable shape costing. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_shape_penalty( |inter_shape_penalty(..) | | |
| ulong packed_shape_cost, |with the payload conversions | | |
| intel_sub_group_avc_ime_payload_t |to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
| | | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_direction_penalty( |inter_direction_penalty(..) | | |
| uchar direction_cost, |with the payload conversions | | |
| intel_sub_group_avc_ime_payload_t|to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+-------------------------------------+------------------------------+ | |
Inter motion vector cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
search configuration phase to enable motion vector costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ | | |
|motion_vector_cost_function( |motion_vector_cost_function(..)| | |
| ulong packed_cost_center_delta, |with the payload conversions | | |
| uint2 packed_cost_table, |to/from MCE types. See MCE | | |
| uchar cost_precision, |version for description. | | |
| intel_sub_group_avc_ime_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
Miscellaneous property configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
search configuration phase to enable miscellaneous properties setting | |
in the payload. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ | | |
|source_interlaced_field_polarity( |source_interlaced_field_ | | |
| uchar src_field_polarity, |polarity(..) with the result | | |
| intel_sub_group_avc_ime_payload_t|conversions to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_ | | |
|single_reference_interlaced_field_ |set_single_reference_ | | |
|polarity( |interlaced_field_polarity(..) | | |
| uchar ref_field_polarity, |with the result conversions | | |
| intel_sub_group_avc_ime_payload_t|to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ | | |
|dual_reference_interlaced_field_ |dual_reference_interlaced_ | | |
|polarities( |field_polarities(..) with the | | |
| uchar fwd_ref_field_polarity, |result conversions to/from MCE | | |
| uchar bwd_ref_field_polarity, |types. See MCE version for | | |
| intel_sub_group_avc_ime_payload_t|description. | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Specify the maximum number of | | |
|intel_sub_group_avc_ime_set_ |motion vectors allowed for the | | |
|max_motion_vector_count( |current MB. The default setting| | |
| uchar max_motion_vector_count, |is 32. Any other value may | | |
| intel_sub_group_avc_ime_payload_t|alter the MB partitioning | | |
| payload ) |decision. The IME operation | | |
| |will compute the best allowed | | |
| |partitioning such that the | | |
| |number of sub-block motion | | |
| |vectors will not exceed | | |
| |max_motion_vector_count. | | |
| | | | |
| |The value of max_motion_vector_| | |
| |count(..) specifies the maximum| | |
| |number of motion vectors | | |
| |allowed for the current MB. It | | |
| |must be in the range [1, 32], | | |
| |otherwise the results are | | |
| |undefined. | | |
| | | | |
| |NOTE: This can be used to | | |
| |handle the restriction for | | |
| |certain profiles for AVC in | | |
| |that the maximum number of | | |
| |motion vectors allowed for two | | |
| |consecutive MBs can only be 16.| | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Update the input payload to | | |
|intel_sub_group_avc_ime_set_ |disable a mix of forward and | | |
|unidirectional_mix_disable( |backward MVs in the result. | | |
| intel_sub_group_avc_ime_payload_t| | | |
| payload ) |Default is to enable it. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Specifies the threshold value | | |
|intel_sub_group_avc_ime_set_early_ |of a distortion compute of a | | |
|search_termination_threshold( |16x16 partition of a MB for a | | |
| uchar threshold, |single-reference search, below | | |
| intel_sub_group_avc_ime_payload_t|which no more searching is | | |
| payload ) |performed for the MB. | | |
| | | | |
| |The value of threshold is | | |
| |specified in U4U4 format and | | |
| |the integer value must fit | | |
| |within 14 bits. | | |
| | | | |
| |The input payload must have | | |
| |been configured for a single- | | |
| |reference search with the 16x16| | |
| |partition enabled for this | | |
| |threshold to be set, or else | | |
| |the results are undefined. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |Set the (16) SAD weights for | | |
|intel_sub_group_avc_ime_set_ |each 4x4 sub-block. | | |
|weighted_sad( | | | |
| uint packed_sad_weights, |These values are used to | | |
| intel_sub_group_avc_ime_payload_t|decrease the SAD magnitude of | | |
| payload ) |each 4x4 sub-block by dividing | | |
| |the SAD of 4x4 sub-block of the| | |
| |source MB by its mapped weight.| | |
| | | | |
| |It requires a partition_mask of| | |
| |16x16 and forward search window| | |
| |configuration. | | |
| | | | |
| |The weighting pattern used is | | |
| |the traditional Z order for | | |
| |each 4x4 block. Weighted-SAD | | |
| |Control Mapping: | | |
| | | | |
| | 0 1 4 5 | | |
| | 2 3 6 7 | | |
| | 8 9 C D | | |
| | A B E F | | |
| | | | |
| |Each weight is of 2 bits | | |
| |represented in a packed format | | |
| |in packed_sad_weights for each | | |
| |of the 4x4 blocks in Z order. | | |
| | | | |
| |A prior call to | | |
| |intel_sub_group_avc_ime_ | | |
| |set_single_reference(..) set up| | |
| |for the forward reference image| | |
| |is required. | | |
| | | | |
| |This feature is mainly intended| | |
| |for improved block matching in | | |
| |image-rate conversion (FRC) | | |
| |kernels. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ime_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ime_set_ |intel_sub_group_avc_mce_set_ac_| | |
|ac_only_haar( |only_haar(..) with the payload | | |
| intel_sub_group_avc_ime_payload_t|conversions to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
Evaluation phase functions | |
++++++++++++++++++++++++++ | |
These built-in functions perform the evaluation of the IME operation | |
configured in the payload with a VME media sampler and return the | |
results. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Evaluate the basic IME | | |
|intel_sub_group_avc_ime_evaluate_ |operation with a single | | |
|with_single_reference( |reference and return its | | |
| read_only image2d_t src_image, |results. The IME payload must | | |
| read_only image2d_t ref_image, |have been configured with | | |
| sampler_t vme_media_sampler, |intel_sub_group_avc_ime_set_ | | |
| intel_sub_group_avc_ime_payload_t |single_reference(..). | | |
| payload ) | | | |
| |The parameter ref_image must | | |
| |be a valid forward image | | |
| |kernel parameter per the | | |
| |ordering conventions for the | | |
| |kernel parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Evaluate the basic IME | | |
|intel_sub_group_avc_ime_evaluate_ |operation with dual reference | | |
|with_dual_reference( |and return its results. The | | |
| read_only image2d_t src_image, |IME payload must have been | | |
| read_only image2d_t fwd_ref_image,|configured with | | |
| read_only image2d_t bwd_ref_image,|intel_sub_group_avc_ime_set_ | | |
| sampler_t vme_media_sampler, |dual_reference(..). | | |
| intel_sub_group_avc_ime_payload_t | | | |
| payload ) |The parameter fwd_ref_image[ | | |
| |bwd_ref_image] must be a valid| | |
| |forward[backward] image kernel| | |
| |parameter per the ordering | | |
| |conventions for the kernel | | |
| |parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_ |Evaluate the single reference | | |
|single_reference_streamout_t |IME operation with streamout | | |
|intel_sub_group_avc_ime_evaluate_ |and return its results. The | | |
|with_single_reference_streamout( |IME payload must have been | | |
| read_only image2d_t src_image, |configured with | | |
| read_only image2d_t ref_image, |intel_sub_group_avc_ime_set_ | | |
| sampler_t vme_media_sampler, |single_reference(..). | | |
| intel_sub_group_avc_ime_payload_t | | | |
| payload ) |The parameter ref_image must | | |
| |be a valid forward image | | |
| |kernel parameter per the | | |
| |ordering conventions for the | | |
| |kernel parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_ |Evaluate the dual reference | | |
|dual_reference_streamout_t |IME operation with streamout | | |
|intel_sub_group_avc_ime_evaluate_ |and return its results. The | | |
|with_dual_reference_streamout( |IME payload must have been | | |
| read_only image2d_t src_image, |configured with | | |
| read_only image2d_t fwd_ref_image,|intel_sub_group_avc_ime_set_ | | |
| read_only image2d_t bwd_ref_image,|dual_reference(..). | | |
| sampler_t vme_media_sampler, | | | |
| intel_sub_group_avc_ime_payload_t |The parameter fwd_ref_image[ | | |
| payload ) |bwd_ref_image] must be a valid| | |
| |forward[backward] image kernel| | |
| |parameter per the ordering | | |
| |conventions for the kernel | | |
| |parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Evaluate the single reference | | |
|intel_sub_group_avc_ime_evaluate_ |IME operation with streamin | | |
|with_single_reference_streamin( |and return its results. The | | |
| read_only image2d_t src_image, |IME payload must have been | | |
| read_only image2d_t ref_image, |configured with | | |
| sampler_t vme_media_sampler, |intel_sub_group_avc_ime_set_ | | |
| intel_sub_group_avc_ime_payload_t|single_reference(..). | | |
| payload, | | | |
| intel_sub_group_avc_ime_single_ |The parameter ref_image must | | |
| reference_streamin_t |be a valid forward image | | |
| streamin_components ) |kernel parameter per the | | |
| |ordering conventions for the | | |
| |kernel parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Evaluate the dual reference | | |
|intel_sub_group_avc_ime_evaluate_ |IME operation with streamin | | |
|with_dual_reference_streamin( |and return its results. The | | |
| read_only image2d_t src_image, |IME payload must have been | | |
| read_only image2d_t fwd_ref_image,|configured with | | |
| read_only image2d_t bwd_ref_image,|intel_sub_group_avc_ime_set_ | | |
| sampler_t vme_media_sampler, |dual_reference(..). | | |
| intel_sub_group_avc_ime_payload_t | | | |
| payload, |The parameter fwd_ref_image[ | | |
| intel_sub_group_avc_ime_dual_ |bwd_ref_image] must be a valid| | |
| reference_streamin_t |forward[backward] image kernel| | |
| streamin_components |parameter per the ordering | | |
| ) |conventions for the kernel | | |
| |parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_ |Evaluate the single reference | | |
|single_reference_streamout_t |IME operation with streamin | | |
|intel_sub_group_avc_ime_evaluate_ |and streamout and return its | | |
|with_single_reference_streaminout( |results. The IME payload must | | |
| read_only image2d_t src_image, |have been configured with | | |
| read_only image2d_t ref_image, |intel_sub_group_avc_ime_set_ | | |
| sampler_t vme_media_sampler, |single_reference(..). | | |
| intel_sub_group_avc_ime_payload_t | | | |
| payload, |The parameter ref_image must | | |
| intel_sub_group_avc_ime_single_ |be a valid forward image | | |
| reference_streamin_t |kernel parameter per the | | |
| streamin_components |ordering conventions for the | | |
| ) |kernel parameter list. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_ |Evaluate the dual reference | | |
|dual_reference_streamout_t |IME operation with streamin | | |
|intel_sub_group_avc_ime_evaluate_ |and streamout and return its | | |
|with_dual_reference_streaminout( |results. The IME payload must | | |
| read_only image2d_t src_image, |have been configured with | | |
| read_only image2d_t fwd_ref_image,|intel_sub_ | | |
| read_only image2d_t bwd_ref_image,|group_avc_ime_set_dual_ | | |
| sampler_t vme_media_sampler, |reference(..). | | |
| intel_sub_group_avc_ime_payload_t | | | |
| payload, |The parameter fwd_ref_image[ | | |
| intel_sub_group_avc_ime_dual_ |bwd_ref_image] must be a valid| | |
| reference_streamin_t |forward[backward] image kernel| | |
| streamin_components |parameter per the ordering | | |
| ) |conventions for the kernel | | |
| |parameter list. | | |
+-------------------------------------+------------------------------+ | |
Result type conversion functions | |
+++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
evaluation phase to convert IME results to MCE results and vice-versa. | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_mce_result_t |Convert the IME result into a | | |
|intel_sub_group_avc_ime_ |MCE result. | | |
|convert_to_mce_result( | | | |
| intel_sub_group_avc_ime_result_t| | | |
| result ) | | | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Convert the MCE result into an | | |
|intel_sub_group_avc_mce_ |IME result. | | |
|convert_to_ime_result( | | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) | | | |
+-----------------------------------+--------------------------------+ | |
Result processing phase functions | |
+++++++++++++++++++++++++++++++++ | |
These built-in functions are called following the evaluation phase to | |
extract the various result components from an IME evaluation result. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_ |Return the streamed out BMVs | | |
|single_reference_streamin_t |and distortions from the input| | |
|intel_sub_group_avc_ime_ |result from a single reference| | |
|get_single_reference_streamin( |streamout IME operation that | | |
| intel_sub_group_avc_ime_result_ |can be used as streamin input | | |
| single_reference_streamout_t |for a subsequent IME | | |
| result ) |operation. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_ |Return the streamed out BMVs | | |
|dual_reference_streamin_t |and distortions from the input| | |
|intel_sub_group_avc_ime_get_ |result from a dual reference | | |
|dual_reference_streamin( |streamout IME operation that | | |
| intel_sub_group_avc_ime_result_ |can be used as streamin input | | |
| dual_reference_streamout_t |for a subsequent IME | | |
| result ) |operation. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Strip out the single reference| | |
|intel_sub_group_avc_ime_strip_ |streamout BMVs and distortions| | |
|single_reference_streamout( |from the streamout results and| | |
| intel_sub_group_avc_ime_result_ |return the rest. | | |
| single_reference_streamout_t | | | |
| result ) | | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ime_result_t |Strip out the dual reference | | |
|intel_sub_group_avc_ime_strip_ |streamout MVs and distortions | | |
|dual_reference_streamout( |from the streamout results and| | |
| intel_sub_group_avc_ime_result_ |return the rest. | | |
| dual_reference_streamout_t | | | |
| result ) | | | |
+-------------------------------------+------------------------------+ | |
|uint intel_sub_group_avc_ime_get_ |Get the packed motion vectors | | |
|streamout_major_shape_motion_vectors(|for the input major shape from| | |
| intel_sub_group_avc_ime_result_ |the IME single reference | | |
| single_reference_streamout_t |streamout results. | | |
| result, | | | |
| uchar major_shape ) |The parameter major_shape must| | |
| |be a valid inter macro-block | | |
| |major shape enumeration value | | |
| |and a compile-time constant. | | |
| | | | |
| |Up to 4 packed MVs are | | |
| |returned, one per | | |
| |work-item. If the major shape | | |
| |is: | | |
| |- 16x6, then one packed MV | | |
| | is returned by work-item 0 | | |
| |- 16x8, or 8x16, then two | | |
| | packed MVs are returned by | | |
| | work-items 0 and 1 | | |
| |- 8x8, then four packed MVs | | |
| | are returned by work-items | | |
| | 0 to 3. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ime_get_ |Get the distortions for the | | |
|streamout_major_shape_distortions( |input major shape from the IME| | |
| intel_sub_group_avc_ime_result_ |single reference streamout | | |
| single_reference_streamout_t |results. | | |
| result, | | | |
| uchar major_shape ) |The parameter major_shape must| | |
| |be a valid inter macro-block | | |
| |major shape enumeration value | | |
| |and a compile-time constant. | | |
| | | | |
| |Up to 4 distortions are | | |
| |returned, one per work-item in| | |
| |the same format as for motion | | |
| |vectors as described above. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |Get the reference identifiers | | |
|streamout_major_shape_reference_ids( |for the input major shape and | | |
| intel_sub_group_avc_ime_result_ |direction from the IME dual | | |
| single_reference_streamout_t |reference streamout results. | | |
| result, | | | |
| uchar major_shape ) |The parameter major_shape must| | |
| |be a valid inter macro-block | | |
| |major shape enumeration value | | |
| |and a compile-time constant. | | |
| | | | |
| |Up to 4 reference identifiers | | |
| |are returned, one per | | |
| |work-item in the same format | | |
| |as for motion vectors as | | |
| |described above. | | |
+-------------------------------------+------------------------------+ | |
|uint intel_sub_group_avc_ime_get_ |Get the packed motion vectors | | |
|streamout_major_shape_motion_vectors(|for the input major shape and | | |
| intel_sub_group_avc_ime_result_ |direction from the IME dual | | |
| dual_reference_streamout_t result,|reference streamout results. | | |
| uchar major_shape, | | | |
| uchar direction ) |The parameter major_shape must| | |
| |be a valid inter macro-block | | |
| |major shape enumeration value | | |
| |and a compile-time constant. | | |
| | | | |
| |The parameter direction must | | |
| |be a valid unidirectional | | |
| |inter macro-block major | | |
| |direction value and a | | |
| |compile-time constant. | | |
| | | | |
| |Up to 4 packed MVs are | | |
| |returned, one per work-item in| | |
| |the same format as for motion | | |
| |vectors for single reference | | |
| |streamout as described above. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ime_get_ |Get the distortions for the | | |
|streamout_major_shape_distortions( |input major shape and | | |
| intel_sub_group_avc_ime_result_ |direction from the IME dual | | |
| dual_reference_streamout_t result,|reference streamout results. | | |
| uchar major_shape, | | | |
| uchar direction ) |The parameter major_shape must| | |
| |be a valid inter macro-block | | |
| |major shape enumeration value | | |
| |and a compile-time constant. | | |
| | | | |
| |The parameter direction must | | |
| |be a valid unidirectional | | |
| |inter macro-block major | | |
| |direction value and a | | |
| |compile-time constant. | | |
| | | | |
| |Up to 4 distortions are | | |
| |returned, one per work-item in| | |
| |the same format as for motion | | |
| |vectors as described above. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |Get the reference identifiers | | |
|streamout_major_shape_reference_ids( |for the input major shape and | | |
| intel_sub_group_avc_ime_result_ |direction from the IME dual | | |
| dual_reference_streamout_t result,|reference streamout results. | | |
| uchar major_shape, | | | |
| uchar direction ) | | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |Get the bitmask indicating | | |
|border_reached( |whether any border of | | |
| uchar image_select, |forward/backward reference | | |
| intel_sub_group_avc_ime_result_t |image is reached by one or | | |
| result ) |more MVs in the winning inter | | |
| |shape. The bitmask values are | | |
| |as per the inter border | | |
| |reached enumeration values. | | |
| | | | |
| |The search window must have | | |
| |been configured for a forward | | |
| |reference if image_select is | | |
| |set as | | |
| |CLK_AVC_ME_FRAME_FORWARD_INTEL| | |
| |and with a backward reference | | |
| |if image_select is set as | | |
| |CLK_AVC_ME_FRAME_BACKWARD_ | | |
| |INTEL. | | |
| | | | |
| |The value of image_select is | | |
| |either | | |
| |CLK_AVC_ME_FRAME_FORWARD_INTEL| | |
| |or CLK_AVC_ME_FRAME_BACKWARD_ | | |
| |INTEL and must be a | | |
| |compile-time constant. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |Get the indication that the | | |
|truncated_search_indication( |search operation was prevented| | |
| intel_sub_group_avc_ime_result_t |from providing the lowest | | |
| result ) |distortion solution due to the| | |
| |tighter constraints on the | | |
| |maximum number of MB motion | | |
| |vectors configured for the | | |
| |search operation. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |Get the indication that | | |
|unidirectional_early_search_ |unidirectional search | | |
|termination( |operation terminated early | | |
| intel_sub_group_avc_ime_result_t |because the configured | | |
| result ) |distortion threshold was met. | | |
+-------------------------------------+------------------------------+ | |
|uint intel_sub_group_avc_ime_get_ |Get the 16x16 motion vector | | |
|weighting_pattern_minimum_ |corresponding to the minimum | | |
|motion_vector( |16x16 distortion when applying| | |
| intel_sub_group_avc_ime_result_t |the traditional Z-order SAD | | |
| result ) |weighting pattern. | | |
| | | | |
| |This can only be called if a | | |
| |SAD weighting pattern was set | | |
| |prior to evaluation using | | |
| |intel_sub_group_avc_ime_set_ | | |
| |weighted_sad(..). | | |
| | | | |
| |The argument for parameter | | |
| |"pattern_id" should be a | | |
| |compile-time constant. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ime_get_ |Get the minimum 16x16 | | |
|weighting_pattern_minimum_distortion(|distortion when applying the | | |
| intel_sub_group_avc_ime_result_t |traditional Z-order SAD | | |
| result ) |weighting pattern. | | |
+-------------------------------------+------------------------------+ | |
|ulong intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|motion_vectors( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |motion_vectors(..) with the | | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_distortions( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_distortions(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|best_inter_distortion( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |best_inter_distortion(..)with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_major_shape( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_major_shape(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_minor_shapes( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_minor_shapes(..) with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_directions( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_directions(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_motion_vector_count( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_motion_vector_count(..) | | |
| result ) |with the result conversions | | |
| |to/from MCE types. See MCE | | |
| |version for description. | | |
+-------------------------------------+------------------------------+ | |
|uint intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_reference_ids( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ime_result_t |inter_reference_ids(..) with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ime_get_ |This is a wrapper for | | |
|inter_reference_interlaced_ |intel_sub_group_avc_mce_get_ | | |
|field_polarities( |inter_reference_interlaced_ | | |
| uint packed_reference_ids, |field_polarities(..) with the| | |
| uint packed_reference_parameter_ |result conversions to/from MCE| | |
| field_polarities, |types. See MCE version for | | |
| intel_sub_group_avc_ime_result_t |description. | | |
| result ) | | | |
+-------------------------------------+------------------------------+ | |
______________________________________________________________________ | |
REF built-in functions | |
---------------------- | |
A set of ordered phases of functions are required to be called to | |
evaluate a refinement operation result. | |
Initialization phase functions | |
++++++++++++++++++++++++++++++ | |
These built-in functions create a properly initialized payload that | |
can be used for further configured for evaluating REF operations. This | |
is a required initial phase. A call to either | |
intel_sub_group_avc_fme_initialize or | |
intel_sub_group_avc_bme_initialize is required. | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |Return an initialized payload | | |
|intel_sub_group_avc_fme_initialize(|for a VME fractional motion | | |
| ushort2 src_coord, |estimation operation (FME). | | |
| ulong motion_vectors, | | | |
| uchar major_shapes, |The payload is initialized for | | |
| uchar minor_shapes, |progressive frame operations, | | |
| uchar directions, |and the cost configuration | | |
| uchar pixel_resolution, |values and the miscellaneous | | |
| uchar sad_adjustment ) |property values are all | | |
| |initialized to zero. The cost | | |
| |configuration and the | | |
| |miscellaneous property | | |
| |configuration phase functions | | |
| |must be used to override the | | |
| |initial configurations in the | | |
| |payload. | | |
| | | | |
| |The src_coord value represents | | |
| |the 2D offset of the top left | | |
| |corner of the source MB in pixel| | |
| |units in the source image. | | |
| | | | |
| |If the source image is an | | |
| |interlaced scan image, then the | | |
| |bottom field lines are | | |
| |considered as logically | | |
| |overlapping with the top field | | |
| |lines (i.e. the top field MBs | | |
| |are considered as logically | | |
| |overlapping with the bottom MBs)| | |
| |for the purposes for specifying | | |
| |the src_coord value. | | |
| | | | |
| |The parameter motion_vectors | | |
| |contains the BMVs returned by an| | |
| |IME in the same format as | | |
| |returned by | | |
| |intel_sub_group_avc_mce_get_ | | |
| |motion_vectors(..). The MVs are | | |
| |in QPEL units. The X and Y | | |
| |coordinates of each MV must be | | |
| |in the range [-2048.00, | | |
| |2047.75), otherwise the results | | |
| |are undefined. | | |
| | | | |
| |(If this argument value is | | |
| |manually composed, all sub-block| | |
| |MVs must be replicated per its | | |
| |format for each partition. For | | |
| |example for 16x16 partition, all| | |
| |sub-block MVs must be replicated| | |
| |to the same MV, and for 8x8 | | |
| |partition, each 8x8 must have | | |
| |its respective sub-block MVs | | |
| |replicated.) | | |
| | | | |
| |Legal values and format for | | |
| |major_shapes are as per the | | |
| |return value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_major_shape(..). | | |
| | | | |
| |Legal values and format for | | |
| |minor_shapes are as per the | | |
| |return value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_minor_shapes(..). | | |
| | | | |
| |Legal values and format for | | |
| |directions are as per the return| | |
| |value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_directions(..). | | |
| | | | |
| |Legal values for | | |
| |pixel_resolution is either | | |
| |CLK_AVC_ME_SUBPIXEL_MODE_HPEL_ | | |
| |INTEL or | | |
| |CLK_AVC_ME_SUBPIXEL_MODE_QPEL_ | | |
| |INTEL. | | |
| | | | |
| |The legal values for | | |
| |sad_adjustment is specified by | | |
| |its respective enumeration | | |
| |values. | | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |Return an initialized payload | | |
|intel_sub_group_avc_bme_initialize(|for a VME bidirectional motion | | |
| ushort2 src_coord, |estimation (BME) operation. | | |
| ulong motion_vectors, | | | |
| uchar major_shapes, |The payload is initialized for | | |
| uchar minor_shapes, |progressive frame operations, | | |
| uchar directions, |and the cost configuration | | |
| uchar pixel_resolution, |values and the miscellaneous | | |
| uchar bidirectional_weight, |property values are all | | |
| uchar sad_adjustment ) |initialized to zero. The cost | | |
| |configuration and the | | |
| |miscellaneous property | | |
| |configuration phase functions | | |
| |must be used to override the | | |
| |initial configurations in the | | |
| |payload. | | |
| | | | |
| |If the specified | | |
| |pixel_resolution is a sub-pixel | | |
| |resolution then an implicit FME | | |
| |operation is performed with the | | |
| |BME operation. | | |
| | | | |
| |The src_coord value represents | | |
| |the 2D offset of the top left | | |
| |corner of the source MB in pixel| | |
| |units in the source image. | | |
| | | | |
| |If the source image is an | | |
| |interlaced scan image, then the | | |
| |bottom field lines are | | |
| |considered as logically | | |
| |overlapping with the top field | | |
| |lines (i.e. the top field MBs | | |
| |are considered as logically | | |
| |overlapping with the bottom MBs)| | |
| |for the purposes for specifying | | |
| |the src_coord value. | | |
| | | | |
| |The parameter motion_vectors | | |
| |contains the BMVs returned by an| | |
| |IME in the same format as | | |
| |returned by | | |
| |intel_sub_group_avc_mce_get_ | | |
| |motion_vectors(..). | | |
| | | | |
| |(If this argument value is | | |
| |manually composed, all sub-block| | |
| |MVs must be replicated per its | | |
| |format for each partition. For | | |
| |example for 16x16 partition, all| | |
| |sub-block MVs must be replicated| | |
| |to the same MV, and for 8x8 | | |
| |partition, each 8x8 must have | | |
| |its respective sub-block MVs | | |
| |replicated.) | | |
| | | | |
| |Legal values and format for | | |
| |major_shapes is as per the | | |
| |return value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_major_shape(..). | | |
| | | | |
| |Legal values and format for | | |
| |minor_shapes is as per the | | |
| |return value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_minor_shapes(..). | | |
| | | | |
| |Legal values and format for | | |
| |directions are as per the return| | |
| |value of | | |
| |intel_sub_group_avc_mce_get_ | | |
| |inter_directions(..). | | |
| | | | |
| |Legal values for | | |
| |pixel_resolution is either | | |
| |CLK_AVC_ME_SUBPIXEL_MODE_ | | |
| |INTEGER_INTEL, | | |
| |CLK_AVC_ME_SUBPIXEL_MODE_ | | |
| |HPEL_INTEL or | | |
| |CLK_AVC_ME_SUBPIXEL_MODE_ | | |
| |QPEL_INTEL. | | |
| | | | |
| |Legal values for | | |
| |bidirectional_weight and sad | | |
| |adjustment is as per their | | |
| |respective enumeration values. | | |
| | | | |
| |The legal values for | | |
| |sad_adjustment is specified by | | |
| |its respective enumeration | | |
| |values. | | |
+-----------------------------------+--------------------------------+ | |
Payload type conversion functions | |
+++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
initialization phase to convert REF payload to MCE payloads and | |
vice-versa. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Convert the REF payload to a | | |
|intel_sub_group_avc_ref_ |generic ME payload. | | |
|convert_to_mce_payload( | | | |
| intel_sub_group_avc_ref_payload_t | | | |
| payload ) | | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |Convert the generic ME payload| | |
|intel_sub_group_avc_mce_ |to an FBR payload. | | |
|convert_to_ref_payload( | | | |
| intel_sub_group_avc_mce_payload_t | | | |
| payload ) | | | |
+-------------------------------------+------------------------------+ | |
Multi-reference cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
initialization phase to enable multi-reference image costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_ | | |
|inter_base_multi_reference_penalty( |set_inter_base_multi_reference_| | |
| uchar reference_base_penalty, |penalty(..) the payload | | |
| intel_sub_group_avc_ref_payload_t|conversions with to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
Inter shape and direction cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
initialization phase to enable shape costing. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_shape_penalty( |inter_shape_penalty(..) with | | |
| ulong packed_shape_cost, |the payload conversions | | |
| intel_sub_group_avc_ref_payload_t |to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_direction_penalty( |inter_direction_penalty(..) | | |
| uchar direction_cost, |with the payload conversions | | |
| intel_sub_group_avc_ref_payload_t |to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+-------------------------------------+------------------------------+ | |
Inter motion vector cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
initialization phase to enable motion vector costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|motion_vector_cost_function( |motion_vector_cost_function | | |
| ulong packed_cost_center_delta, |with the payload conversions | | |
| uint2 packed_cost_table, |to/from MCE types. See MCE | | |
| uchar cost_precision, |version for description. | | |
| intel_sub_group_avc_ref_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
Miscellaneous property configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
initialization phase to enable miscellaneous properties setting in the | |
payload. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|source_interlaced_field_polarity( |source_interlaced_field_ | | |
| uchar src_field_polarity, |polarity(..) with the result | | |
| intel_sub_group_avc_ref_payload_t|conversions to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|single_reference_interlaced_ |single_reference_ | | |
|field_polarity( |interlaced_field_polarity(..) | | |
| uchar ref_field_polarity, |with the result conversions | | |
| intel_sub_group_avc_ref_payload_t|to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|dual_reference_interlaced_ |dual_reference_interlaced_ | | |
|field_polarities( |field_polarities(..) with the | | |
| uchar fwd_ref_field_polarity, |result conversions to/from MCE | | |
| uchar bwd_ref_field_polarity, |types. See MCE version for | | |
| intel_sub_group_avc_ref_payload_t|description. | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |Update the input payload to | | |
|intel_sub_group_avc_ref_set_ |disable a mix of bidirectional | | |
|bidirectional_mix_disable( |and unidirectional MVs in the | | |
| intel_sub_group_avc_ref_payload_t|result. | | |
| payload ) | | | |
| |Default is to enable it. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |Update the input payload to do | | |
|intel_sub_group_avc_ref_set_ |enable bilinear filter | | |
|bilinear_filter_enable( |interpolation instead of 4-tap | | |
| intel_sub_group_avc_ref_payload_t|filter interpolation. Default | | |
| payload ) |is 4-tap filter interpolation. | | |
| | | | |
| |This should not be called if | | |
| |the payload was initialized | | |
| |with integer pixel resolution. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_ref_set_ |intel_sub_group_avc_mce_set_ | | |
|ac_only_haar( |ac_only_haar(..) with the | | |
| intel_sub_group_avc_ref_payload_t|payload conversions to/from MCE| | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
Evaluation phase functions | |
++++++++++++++++++++++++++ | |
These built-in functions perform the evaluation of the REF operation | |
configured in the payload with a VME media sampler and return the | |
results. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_result_t |Evaluate the basic REF | | |
|intel_sub_group_avc_ref_ |operation with single reference| | |
|evaluate_with_single_reference( |and return its results. | | |
| read_only image2d_t src_image, | | | |
| read_only image2d_t ref_image, |The parameter ref_image must be| | |
| sampler_t vme_media_sampler, |a valid forward image kernel | | |
| intel_sub_group_avc_ref_payload_t|parameter per the ordering | | |
| payload ) |conventions for the kernel | | |
| |parameter list. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_result_t |Evaluate the basic REF | | |
|intel_sub_group_avc_ref_ |operation with dual reference | | |
|evaluate_with_dual_reference( |and return its results. | | |
| read_only image2d_t src_image, | | | |
| read_only image2d_t fwd_ref_image,|The parameter | | |
| read_only image2d_t bwd_ref_image,|fwd_ref_image[bwd_ref_image] | | |
| sampler_t vme_media_sampler, |must be a valid | | |
| intel_sub_group_avc_ref_payload_t |forward[backward] image kernel | | |
| payload ) |parameter per the ordering | | |
| |conventions for the kernel | | |
| |parameter list. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_result_t |Evaluate the basic REF | | |
|intel_sub_group_avc_ref_ |operation with multi references| | |
|evaluate_with_multi_reference( |and return its results. | | |
| read_only image2d_t src_image, | | | |
| uint packed_reference_ids, |A unique pair of reference | | |
| sampler_t vme_media_sampler, |identifiers (indicating unique | | |
| intel_sub_group_avc_ref_payload_t|forward/backward reference | | |
| payload ) |images) may be specified for | | |
| |each of the allowed major | | |
| |partitions using | | |
| |packed_dual_ref_ids. | | |
| | | | |
| |The value of | | |
| |packed_reference_ids is a | | |
| |integer with the following bits| | |
| |specifying the values for the | | |
| |pair of reference images for | | |
| |each major partition. | | |
| | | | |
| |3:0 => Fwd reference block 0 | | |
| |7:4 => Bwd reference block 0 | | |
| |11:8 => Fwd reference block 1 | | |
| |15:12 => Bwd reference block 1 | | |
| |19:16 => Fwd reference block 2 | | |
| |23:20 => Bwd reference block 2 | | |
| |27:24 => Fwd reference block 3 | | |
| |31:28 => Bwd reference block 3 | | |
| | | | |
| |A forward[backward] reference | | |
| |idenitifer value of 'n' | | |
| |indicates the forward[backward]| | |
| |image from the 'n'th pair of | | |
| |forward/backward reference | | |
| |images, with the value of 'n' | | |
| |ranging from 0 to 15. | | |
| | | | |
| |If the REF operation is | | |
| |configured with only forward | | |
| |reference images then, the | | |
| |values of the backward | | |
| |reference identifiers are not | | |
| |used. | | |
| | | | |
| |The blocks are numbered using | | |
| |the traditional Z order. For | | |
| |larger block sizes, the | | |
| |sub-block reference identifier | | |
| |pairs must be replicated. For | | |
| |example, for a 16x16 block, all| | |
| |four pair of reference | | |
| |identifiers must be replicated | | |
| |to the value of the first pair | | |
| |for block 0. | | |
| | | | |
| |The value for the | | |
| |packed_reference_ids argument | | |
| |is obtained by calling | | |
| |intel_sub_group_avc_ime_get_ | | |
| |inter_reference_ids(..) for | | |
| |the preceding IME operation's | | |
| |result. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_ref_result_t |Evaluate the basic REF | | |
|intel_sub_group_avc_ref_ |operation with multi references| | |
|evaluate_with_multi_reference( |and return its results. This | | |
| read_only image2d_t src_image, |is used for interlaced source | | |
| uint packed_reference_ids, |and reference images. | | |
| uchar packed_reference_ | | | |
| field_polarities, |A unique pair of reference | | |
| sampler_t vme_media_sampler, |identifiers (indicating unique | | |
| intel_sub_group_avc_ref_payload_t|forward/backward reference | | |
| payload ) |images) may be specified for | | |
| |each of the allowed major | | |
| |partitions using | | |
| |packed_reference_ids. | | |
| | | | |
| |The value of | | |
| |packed_reference_ids is an | | |
| |integer with the following bits| | |
| |specifying the values for the | | |
| |pair of reference images for | | |
| |each major partition. | | |
| | | | |
| |3:0 => Fwd reference block 0 | | |
| |7:4 => Bwd reference block 0 | | |
| |11:8 => Fwd reference block 1 | | |
| |15:12 => Bwd reference block 1 | | |
| |19:16 => Fwd reference block 2 | | |
| |23:20 => Bwd reference block 2 | | |
| |27:24 => Fwd reference block 3 | | |
| |31:28 => Bwd reference block 3 | | |
| | | | |
| |A forward[backward] reference | | |
| |idenitifer value of 'n' | | |
| |indicates the forward[backward]| | |
| |image from the 'n'th pair of | | |
| |forward/backward reference | | |
| |images, with the value of 'n' | | |
| |ranging from 0 to 15. | | |
| | | | |
| |If the REF operation is | | |
| |configured with only forward | | |
| |reference images then, the | | |
| |values of the backward | | |
| |reference identifiers are not | | |
| |used. | | |
| | | | |
| |The blocks are numbered using | | |
| |the traditional Z order. For | | |
| |larger block sizes, the | | |
| |sub-block reference identifier | | |
| |pairs must be replicated. For | | |
| |example, for a 16x16 block, all| | |
| |four pair of reference | | |
| |identifiers must be replicated | | |
| |to the value of the first pair | | |
| |for block 0. | | |
| | | | |
| |The value for the | | |
| |packed_reference_ids argument | | |
| |is obtained by calling | | |
| |intel_sub_group_avc_ime_get_ | | |
| |inter_reference_ids(..) for | | |
| |the preceding IME operation's | | |
| |result. | | |
| | | | |
| |Reference field polarities for | | |
| |forward and backward reference | | |
| |images are specified for each | | |
| |of the allowed major partitions| | |
| |using packed_reference_field_ | | |
| |polarities. | | |
| | | | |
| |The value of | | |
| |packed_reference_field_ | | |
| |polarities is an integer with | | |
| |the following bits specifying | | |
| |the reference field polarities | | |
| |for the major partitions. | | |
| | | | |
| |0 : Fwd reference block 0 | | |
| |1 : Fwd reference block 1 | | |
| |2 : Fwd reference block 2 | | |
| |3 : Fwd reference block 3 | | |
| |4 : Bwd reference block 0 | | |
| |5 : Bwd reference block 1 | | |
| |6 : Bwd reference block 2 | | |
| |7 : Bwd reference block 3 | | |
| | | | |
| |If the dual-reference | | |
| |evaluation functions are not | | |
| |used, then the values of the | | |
| |backward reference field | | |
| |polarities are not used. | | |
| | | | |
| |The blocks are numbered using | | |
| |the traditional Z order. For | | |
| |larger block sizes, the | | |
| |sub-block reference field | | |
| |polarities are replicated. For | | |
| |example, for a 16x16 block all | | |
| |four pairs of reference field | | |
| |polarities are replicated to | | |
| |the value of the first pair for| | |
| |block 0. | | |
+------------------------------------+-------------------------------+ | |
Result type conversion functions | |
+++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
evaluation phase to convert REF results to MCE results and vice-versa. | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_mce_result_t |Convert the REF result into a | | |
|intel_sub_group_avc_ref_ |MCE result. | | |
|convert_to_mce_result( | | | |
| intel_sub_group_avc_ref_result_t| | | |
| result ) | | | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_ref_result_t |Convert the MCE result into an | | |
|intel_sub_group_avc_mce_ |REF result. | | |
|convert_to_ref_result( | | | |
| intel_sub_group_avc_mce_result_t| | | |
| result ) | | | |
+-----------------------------------+--------------------------------+ | |
Result processing phase functions | |
+++++++++++++++++++++++++++++++++ | |
These built-in functions are called following the evaluation phase to | |
extract the various result components from an REF evaluation result. | |
+-------------------------------------+------------------------------+ | |
|ulong intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|motion_vectors( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |motion_vectors(..) with the | | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_distortions( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_distortions(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|ushort intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|best_inter_distortion( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |best_inter_distortion(..)with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_major_shape( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_major_shape(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_minor_shapes( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_minor_shapes(..) with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_directions( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_directions(..) with the| | |
| result ) |result conversions to/from MCE| | |
| |types. See MCE version for | | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_motion_vector_count( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_motion_vector_count(..) | | |
| result ) |with the result conversions | | |
| |to/from MCE types. See MCE | | |
| |version for description. | | |
+-------------------------------------+------------------------------+ | |
|uint intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_reference_ids( |intel_sub_group_avc_mce_get_ | | |
| intel_sub_group_avc_ref_result_t |inter_reference_ids(..) with | | |
| result ) |the result conversions to/from| | |
| |MCE types. See MCE version for| | |
| |description. | | |
+-------------------------------------+------------------------------+ | |
|uchar intel_sub_group_avc_ref_get_ |This is a wrapper for | | |
|inter_reference_interlaced_ |intel_sub_group_avc_mce_get_ | | |
|field_polarities( |inter_reference_interlaced_ | | |
| uint packed_reference_ids, |field_polarities(..) with the| | |
| uint packed_reference_parameter_ |result conversions to/from MCE| | |
| field_polarities, |types. See MCE version for | | |
| intel_sub_group_avc_ref_result_t |description. | | |
| result ) | | | |
+-------------------------------------+------------------------------+ | |
______________________________________________________________________ | |
SIC built-in functions | |
---------------------- | |
A set of ordered phases of functions are required to be called to | |
evaluate a skip check or intra estimation operation result. | |
Initialization phase functions | |
++++++++++++++++++++++++++++++ | |
These built-in functions create a properly initialized payload that | |
can be used for further configured for evaluating SIC operations. This | |
is a required initial phase. | |
+-----------------------------------+--------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Return an initialized payload | | |
|intel_sub_group_avc_sic_initialize(|for a VME SIC operation. | | |
| ushort2 src_coord ) | | | |
| | | | |
| |If the source image is an | | |
| |interlaced scan image, then the | | |
| |bottom field lines are | | |
| |considered as logically | | |
| |overlapping with the top field | | |
| |lines (i.e. the top field MBs | | |
| |are considered as logically | | |
| |overlapping with the bottom MBs)| | |
| |for the purposes for specifying | | |
| |the src_coord value. | | |
| | | | |
| |If the SIC operation is being | | |
| |configured for chroma based | | |
| |intra estimation, then the x and| | |
| |y coordinates of src_coord must | | |
| |be multiples of 2. | | |
+-----------------------------------+--------------------------------+ | |
Configuration phase functions | |
+++++++++++++++++++++++++++++ | |
These built-in functions allow for configuration of a skip check | |
and/or intra estimation operation. This is a required phase following | |
the initialization phase to configure the skip check or intra | |
estimation operation. A call to intel_sub_group_avc_sic_configure_ | |
skc(..) or intel_sub_group_avc_sic_configure_ipe(..) is required. Both | |
intel_sub_group_avc_sic_configure_skc(..) and intel_sub_group_avc_ | |
sic_configure_ipe(..) may be called to initialize the payload to | |
perform both skip checks and intra estimation as part for the same SIC | |
evaluation, but if called together intel_sub_group_avc_sic_configure_ | |
skc(..) must be called first in the call sequence. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Configure the SIC payload for | | |
|intel_sub_group_avc_sic_ |(uni or bi-directional) skip | | |
|configure_skc( |checks. | | |
| uint skip_block_partition_type, | | | |
| uint skip_motion_vector_mask, |The legal values for | | |
| ulong motion_vectors, |skip_block_partition_type must | | |
| uchar bidirectional_weight, |be one of the specified | | |
| uchar skip_sad_adjustment, |partition mask enumeration | | |
| intel_sub_group_avc_sic_payload_t|values. | | |
| payload ) | | | |
| |Legal values for | | |
| |skip_motion_vector_mask can be | | |
| |composed using the '|' operator| | |
| |from the enumeration values | | |
| |defined for it; both | | |
| |unidirectional and | | |
| |bidirectional skip vectors can | | |
| |be specified uniquely for each | | |
| |major partition (16x16 or 8x8) | | |
| |by an appropriate selection of | | |
| |the skip motion vector mask | | |
| |enumeration values. If the | | |
| |16x16 skip_block_partition_type| | |
| |is specified, then only the | | |
| |16x16 enumeration values may be| | |
| |used, else only the 8x8 | | |
| |enumeration values may be used.| | |
| | | | |
| |For convenience, one of the | | |
| |following two mechanisms can be| | |
| |used to obtain the value for | | |
| |the skip_motion_vector_mask | | |
| |argument. | | |
| | | | |
| |If the shape and directions are| | |
| |all compile-time constants, | | |
| |then either one of the two | | |
| |following macros may be used to| | |
| |set this argument's value for | | |
| |16x16 or 8x8 major partitions | | |
| |respectively: | | |
| |- CLK_AVC_ME_SKIP_BLOCK_16x16_ | | |
| | INTEL(DIRECTION) | | |
| |- CLK_AVC_ME_SKIP_BLOCK_8x8_ | | |
| | INTEL(DIRECTION0, DIRECTION1,| | |
| | DIRECTION2, DIRECTION3) | | |
| | | | |
| |The DIRECTION{0-3} argument | | |
| |values must be one of the inter| | |
| |macro-block major direction | | |
| |values. The directions are | | |
| |specified for each major | | |
| |partition in traditional | | |
| |Z-order. | | |
| | | | |
| |If the shape or directions are | | |
| |not compile-time constants, | | |
| |then the helper function | | |
| |intel_sub_group_avc_sic_get_ | | |
| |motion_vector_mask(..) may be | | |
| |used to set this argument's | | |
| |value. | | |
| | | | |
| |The parameter motion_vectors | | |
| |specifies the input packed | | |
| |BMVs. Either the forward or | | |
| |backward is ignored if the | | |
| |setting in | | |
| |skip_motion_vector_mask is | | |
| |backward or forward | | |
| |respectively. If the setting is| | |
| |bidirectional, then both the | | |
| |forward and backward motion | | |
| |vectors will be used. If the | | |
| |skip_block_partition_type is | | |
| |16x16, work-item 0 in the | | |
| |subgroup provides the BMV, and | | |
| |if the | | |
| |skip_block_partition_type is | | |
| |8x8, work-items 0 to 4 in the | | |
| |subgroup provide the four | | |
| |BMVs. The MVs are in QPEL | | |
| |units. The X and Y coordinates | | |
| |of each MV must be in the range| | |
| |[-2048.00, 2047.75] and | | |
| |[-512.00 to 511.75] | | |
| |respectively, otherwise the | | |
| |results are undefined. | | |
| | | | |
| |Legal values for | | |
| |bidirectional_weight and | | |
| |skip_sad_adjustment are as per | | |
| |their respective enumeration | | |
| |values. If the setting is | | |
| |unidirectional, then the | | |
| |bidirectional_weight parameter | | |
| |value is ignored and can be set| | |
| |to the value '0 '. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Return an initialized payload | | |
|intel_sub_group_avc_sic_ |for a VME luma intra prediction| | |
|configure_ipe( |estimation (IPE) operation. | | |
| uchar luma_intra_partition_mask, | | | |
| uchar |The luma_intra_partition_mask | | |
| intra_neighbour_availabilty, |can be composed from their | | |
| uchar left_edge_luma_pixels, |respective enumeration values | | |
| uchar |using the '&' operator. | | |
| upper_left_corner_luma_pixel, | | | |
| uchar upper_edge_luma_pixels, |Legal values for | | |
| uchar |intra_neighbour_availabilty and| | |
| upper_right_edge_luma_pixels, |intra_sad_adjustment are as per| | |
| uchar intra_sad_adjustment , |their respective enumeration | | |
| intel_sub_group_avc_sic_payload_t|values, based on the intra | | |
| payload ) |neighboring macroblock's to | | |
| |consider in intra mode | | |
| |estimation. | | |
| | | | |
| |The other parameters specify | | |
| |the neighbor edge pixels for | | |
| |the left, top-left corner, top | | |
| |and top right edges with each | | |
| |work-item providing each pixel | | |
| |value. These pixels values are | | |
| |used to perform the intra mode | | |
| |estimation. | | |
| | | | |
| |For the left and top edge | | |
| |pixels, successive subgroup | | |
| |work-items 0 to 15 provide the | | |
| |successive edge pixels. For the| | |
| |top-right edge, successive | | |
| |work-items 0 to 7 provide the | | |
| |successive edge pixels; the | | |
| |pixel values in work-items 8 to| | |
| |15 are ignored. The top-left | | |
| |corner pixel is a uniform pixel| | |
| |value with each work-item | | |
| |providing the same corner | | |
| |pixel. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Return an initialized payload | | |
|intel_sub_group_avc_sic_ |for a VME luma and chroma intra| | |
|configure_ipe( |prediction estimation (IPE) | | |
| uchar luma_intra_partition_mask, |operation. | | |
| uchar intra_neighbour_ | | | |
| availabilty, |The luma_intra_partition_mask | | |
| uchar left_edge_luma_pixels, |can be composed from their | | |
| uchar upper_left_corner_luma_ |respective enumeration values | | |
| pixel, |using the '&' operator. Chroma | | |
| uchar upper_edge_luma_ |intra prediction estimation is | | |
| pixels, |performed on the 8x8 CbCr | | |
| uchar upper_right_edge_luma_ |macroblock. | | |
| pixels, | | | |
| ushort left_edge_chroma_pixels, |Legal values for | | |
| ushort upper_left_corner_chroma_ |intra_neighbour_availabilty and| | |
| pixel, |intra_sad_adjustment are as per| | |
| ushort upper_edge_chroma_pixels, |their respective enumeration | | |
| uchar intra_sad_adjustment, |values, based on the intra | | |
| intel_sub_group_avc_sic_payload_t|neighboring macroblock's to | | |
| payload ) |consider in intra mode | | |
| |estimation. | | |
| | | | |
| |The other parameters specify | | |
| |the neighbor luma and chroma | | |
| |edge pixels for the left, | | |
| |top-left corner, top and | | |
| |top-right (luma only) edges | | |
| |with each work-item providing | | |
| |each pixel value. These pixels | | |
| |values are used to perform the | | |
| |intra mode estimation. | | |
| | | | |
| |For the left and top luma edge | | |
| |pixels successive, subgroup | | |
| |work-items 0 to 15 provide the | | |
| |successive edge pixels. For the| | |
| |top-right luma edge, successive| | |
| |work-items 0 to 7 provide the | | |
| |successive edge pixels; the | | |
| |pixel values in work-items 8 to| | |
| |15 are ignored. The top-left | | |
| |corner pixel is a uniform pixel| | |
| |value with each work-item | | |
| |providing the same corner | | |
| |pixel. | | |
| | | | |
| |For the left and top chroma | | |
| |CbCr pixels, successive | | |
| |subgroup work-items 0 to 7 | | |
| |provide the successive CbCr | | |
| |pixels; the pixel values in | | |
| |work-items 8 to 15 are | | |
| |ignored. The top-left corner | | |
| |pixel is a uniform CbCr pixel | | |
| |value with each work-item | | |
| |providing the same corner CbCr | | |
| |pixel. | | |
+------------------------------------+-------------------------------+ | |
|uint intel_sub_group_avc_sic_get_ |Compose the value for the input| | |
|motion_vector_mask( |argument | | |
| uint skip_block_partition_type, |skip_motion_vector_mask for the| | |
| uchar direction ) |input skip_block_partition_type| | |
| |and direction. | | |
| | | | |
| |The legal values for | | |
| |skip_block_partition_type must | | |
| |be one of the specified | | |
| |partition mask enumeration | | |
| |values (16x16 or 8x8). | | |
| | | | |
| |The direction parameter is a | | |
| |bit field with the directions | | |
| |for the 4 8x8 sub-partitions in| | |
| |traditional Z order, or for | | |
| |only the 16x16 partition. Two | | |
| |bits are reserved for each of | | |
| |the four sub-partitions in | | |
| |row-major order. The 2-bit | | |
| |values are as per the inter | | |
| |macro-block major direction | | |
| |values. If the | | |
| |skip_block_partition_type | | |
| |indicates a 16x16 shape, then | | |
| |only the 1st 2 bits contains | | |
| |the direction, and other bits | | |
| |must be zeroed. | | |
+------------------------------------+-------------------------------+ | |
Payload type conversion functions | |
+++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to convert IME payload to MCE payloads and vice- | |
versa. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_mce_payload_t |Convert the SIC payload to a | | |
|intel_sub_group_avc_sic_convert_ |generic MCE payload. | | |
|to_mce_payload( | | | |
| intel_sub_group_avc_sic_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Convert the generic MCE payload| | |
|intel_sub_group_avc_ |to a SIC payload. | | |
|mce_convert_to_sic_payload( | | | |
| intel_sub_group_avc_mce_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
Multi-reference cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable multi-reference image costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_ | | |
|inter_base_multi_reference_penalty( |set_inter_base_multi_reference_| | |
| uchar reference_base_penalty, |penalty(..) the payload | | |
| intel_sub_group_avc_sic_payload_t|conversions with to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
Inter shape and direction cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable shape costing. | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_shape_penalty( |inter_shape_penalty(..) with | | |
| ulong packed_shape_cost, |the payload conversions | | |
| intel_sub_group_avc_sic_payload_t |to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+-------------------------------------+------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|inter_direction_penalty( |inter_direction_penalty(..) | | |
| uchar direction_cost, |with the payload conversions | | |
| intel_sub_group_avc_sic_payload_t|to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+-------------------------------------+------------------------------+ | |
Inter motion vector cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable motion vector costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|motion_vector_cost_function( |motion_vector_cost_function(..)| | |
| ulong packed_cost_center_delta, |with the payload conversions | | |
| uint2 packed_cost_table, |to/from MCE types. See MCE | | |
| uchar cost_precision, |version for description. | | |
| intel_sub_group_avc_sic_payload_t| | | |
| payload ) | | | |
+------------------------------------+-------------------------------+ | |
Intra shape cost configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable intra shape costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Set the luma shape cost penalty| | |
|intel_sub_group_avc_sic_set_ |for intra motion estimation. | | |
|intra_luma_shape_penalty( | | | |
| uint packed_shape_cost, |The value of packed_shape_cost | | |
| intel_sub_group_avc_sic_payload_t|is an integer value with the | | |
| payload ) |following bits specifying the | | |
| |shape cost in U4U4 format as | | |
| |follows: | | |
| | | | |
| |7:0 : must be zero | | |
| |15:8 : 16x16 cost | | |
| |23:16 : 8x8 cost | | |
| |31:24 : 4x4 cost | | |
| | | | |
| |The U4U4 decoded integer values| | |
| |must bit fit within 12 bits. | | |
+------------------------------------+-------------------------------+ | |
Intra mode cost configuration phase functions | |
+++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable intra mode costing. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |The value of luma_mode_penalty | | |
|intel_sub_group_avc_sic_set_ |specifies the penalty to be | | |
|intra_luma_mode_cost_function( |applied to the estimated luma | | |
| uchar luma_mode_penalty, |mode if it differs from its | | |
| uint luma_packed_neighbor_modes, |predicted luma mode (based on | | |
| uint luma_packed_non_dc_penalty, |its neighbor intra modes). It | | |
| intel_sub_group_avc_sic_payload_t|is specified in U4U4 format and| | |
| payload ) |must bit in 10 bits. | | |
| | | | |
| |The value of | | |
| |luma_packed_neighbor_modes | | |
| |specifies the values of the | | |
| |already computed top and left | | |
| |neighbor modes for the | | |
| |bordering 4x4 blocks, with the | | |
| |4x4 blocks numbered in the | | |
| |traditional Z-order as shown | | |
| |below. | | |
| | | | |
| | 0 1 4 5 | | |
| | 2 3 6 7 | | |
| | 8 9 C D | | |
| | A B E F | | |
| | | | |
| |the following bits specify the | | |
| |neighbor modes. | | |
| | | | |
| |3:0 => Left neighbor block 5 | | |
| |7:4 => Left neighbor block 7 | | |
| |11:8 => Left neighbor block D | | |
| |15:12 => Left neighbor block F | | |
| |19:16 => Top neighbor block A | | |
| |23:20 => Top neighbor block B | | |
| |27:24 => Top neighbor block E | | |
| |31:28 => Top neighbor block F | | |
| | | | |
| |The value of | | |
| |luma_packed_non_dc_penalty | | |
| |specifies the penalty to be | | |
| |applied for any computed non-DC| | |
| |luma mode for each of the | | |
| |16x16, 8x8, and 4x4 shapes, | | |
| |with the following bits | | |
| |specifying the penalties. | | |
| | | | |
| |7:0 => Intra16x16 non-dc | | |
| | penalty | | |
| |15:8 => Intra8x8 non-dc | | |
| | penalty | | |
| |23:16 => Intra4x4 non-dc | | |
| | penalty | | |
| |31:24 => Must be zero | | |
| | | | |
| |The component values are | | |
| |specified in 8-bit integer | | |
| |format. | | |
| | | | |
| |The intra distortion for each | | |
| |intra luma block can be | | |
| |described by the following | | |
| |formulas: | | |
| | | | |
| |Intra_4x4 SAD (or Haar) + | | |
| |luma_shape_penalty_4x4 + | | |
| |luma_non_dc_4x4_penalty (if not| | |
| |DC) + | | |
| |luma_mode_penalty (if computed | | |
| |mode is not the same predicted | | |
| |mode from neighbor modes) | | |
| | | | |
| |Intra_8x8 SAD (or Haar) + | | |
| |luma_shape_penalty_8x8 + | | |
| |luma_non_dc_8x8_penalty (if not| | |
| |DC) + | | |
| |luma_mode_penalty (if computed | | |
| |mode is not the same predicted | | |
| |mode from neighbor modes ) | | |
| | | | |
| |Intra_16x16 SAD (or Haar) + | | |
| |luma_shape_penalty_16x16 + | | |
| |luma_non_dc_4x4_penalty (if not| | |
| |DC) | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |Set the intra chroma mode cost | | |
|intel_sub_group_avc_sic_set_ |function by specifying the | | |
|intra_chroma_mode_cost_function( |penalty to be applied to the | | |
| uchar chroma_mode_penalty, |computed chroma mode. | | |
| intel_sub_group_avc_sic_payload_t| | | |
| payload ) |The chroma_mode_base_penalty is| | |
| |the base penalty to be applied | | |
| |to the computed intra chroma | | |
| |modes. This penalty is in U4U4 | | |
| |format. | | |
| | | | |
| |The U4U4 decoded integer value | | |
| |must fit in 12 bits. | | |
| | | | |
| |The base penalty is scaled | | |
| |based on the computed mode as | | |
| |defined below. | | |
| | | | |
| | DC : 0x | | |
| | HORZ : 1x | | |
| | VERT : 1x | | |
| | PLANE: 2x | | |
| | | | |
| |The intra distortion for each | | |
| |intra 8x8 chroma block can be | | |
| |described by the following | | |
| |formulas: | | |
| | | | |
| |distortion = | | |
| |SAD (or Haar) + | | |
| |chroma_mode_penalty (scaled | | |
| |based on computed mode) | | |
+------------------------------------+-------------------------------+ | |
Miscellaneous property configuration phase functions | |
++++++++++++++++++++++++++++++++++++++++++++++++++++ | |
These are optional built-in functions that may be called following the | |
configuration phase to enable miscellaneous properties setting in the | |
payload. | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|source_interlaced_field_polarity( |source_interlaced_field_ | | |
| uchar src_field_polarity, |polarity(..) with the result | | |
| intel_sub_group_avc_sic_payload_t|conversions to/from MCE | | |
| payload ) |types. See MCE version for | | |
| |description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|single_reference_interlaced_ |single_reference_ | | |
|field_polarity( |interlaced_field_polarity(..) | | |
| uchar ref_field_polarity, |with the result conversions | | |
| intel_sub_group_avc_sic_payload_t|to/from MCE types. See MCE | | |
| payload ) |version for description. | | |
+------------------------------------+-------------------------------+ | |
|intel_sub_group_avc_sic_payload_t |This is a wrapper for | | |
|intel_sub_group_avc_sic_set_ |intel_sub_group_avc_mce_set_ | | |
|dual_reference_interlaced_ |dual_reference_interlaced_ | | |
|field_polarities( |field_polarities(..) with the | | |
| uchar fwd_ref_field_polarity, |result conversions to/from MCE | | |
| uchar |