summaryrefslogtreecommitdiff
path: root/Software/Beignet/howto/video-motion-estimation-howto.mdwn
blob: 6b99d54cefc5c2d0908c1bcade0a885ac0e98726 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
Video Motion Vector HowTo
==========================

Beignet now supports cl_intel_accelerator and cl_intel_motion_estimation, which are
Khronos official extensions. It provides a hardware acceleration of video motion
vector to users.

Supported hardware platform
---------------------------

Only 3rd Generation Intel Core Processors is supported for vme now. We will consider
to support more platforms if necessary.

Steps
-----

In order to use video motion estimation provided by Beignet in your program, please follow
the steps as below:

- Create a cl_accelerator_intel object using extension API clCreateAcceleratorINTEL, like
  this:
    
        _accelerator_type_intel accelerator_type = CL_ACCELERATOR_TYPE_MOTION_ESTIMATION_INTEL;
        cl_motion_estimation_desc_intel vmedesc = {CL_ME_MB_TYPE_16x16_INTEL,
                                                   CL_ME_SUBPIXEL_MODE_INTEGER_INTEL,
                                                   CL_ME_SAD_ADJUST_MODE_NONE_INTEL,
                                                   CL_ME_SEARCH_PATH_RADIUS_16_12_INTEL
                                                   };
    
- Invoke clCreateProgramWithBuiltInKernels to create a program object with built-in kernels
  information, and invoke clCreateKernel to create a kernel object whose kernel name is
  block_motion_estimate_intel.

- The prototype of built-in kernel block_motion_estimate_intel is as following:

        _kernel void
        block_motion_estimate_intel
        (
          accelerator_intel_t accelerator,
          __read_only  image2d_t src_image,
          __read_only  image2d_t ref_image,
          __global short2 * prediction_motion_vector_buffer,
          __global short2 * motion_vector_buffer,
          __global ushort * residuals
        );
  So you should create related objects and setup these kernel arguments by clSetKernelArg.
  Create source and reference image object, on which you want to do video motion estimation.
  The image_channel_order should be CL_R and image_channel_data_type should be CL_UNORM_INT8.
  Create a buffer object to get the motion vector result. This motion vector buffer representing
  a vector field of pixel block motion vectors, stored linearly in row-major order. The elements
  (pixels) of this image contain a motion vector for the corresponding pixel block, with its x/y
  components packed as two 16-bit integer values. Each component is encoded as a S13.2 fixed
  point value(two's complement).

- Use clEnqueueNDRangeKernel to enqueue this kernel. The only thing you need to setup is global_work_size:
  global_work_size[0] equal to width of source image, global_work_size[1] equal to height of source
  image.

- Use clEnqueueReadBuffer or clEnqueueMapBuffer to get motion vector result.


Sample code
-----------

We have developed an utest case of using video motion vector in utests/builtin_kernel_block_motion_estimate_intel.cpp.
Please go through it for details.

More references
---------------

<https://www.khronos.org/registry/cl/extensions/intel/cl_intel_accelerator.txt>  
<https://www.khronos.org/registry/cl/extensions/intel/cl_intel_motion_estimation.txt>  
<https://software.intel.com/en-us/articles/intro-to-motion-estimation-extension-for-opencl>