Add example to show libva buffer sharing with extension clCreateImageFromLibvaIntel.

This example reads a source nv12 file to a VASurface, and creates a
target
VASurface. Then creates corresponding cl image objects from them. After
using ocl to do mirror effect post-processing on source VASurface,
target
VASurface is shown on screen by default.

Code of loading nv12 file to VASurface are referenced from
libva/test/encode/avcenc.c.

v2:
Delete 1920x1080.nv12 and 640x480.nv12 because of large size, add
256_128.nv12 as default test image.

v3:
1. Re-org files: add libva as a submodule then use display related
   files.
2. Show result on screen by default instead of saving as a file.
3. Fix warnings.

v4: Fix whitespace format warnings.

v5:
1. Modify upload_nv12_to_surface to read a nv12 file and then upload
   it to an NV12 VASurface. Also modify store_surface_to_nv12.
2. Change the cl post-processing kernel from gray effect to mirror
   effect, which make demo cooler.
3. Minor fix of other problems.

v6: Remove unnecessary OUTPUT_NV12_DEFAULT related code.

Signed-off-by: Chuanbo Weng <chuanbo.weng@intel.com>
Signed-off-by: Zhao Yakui <yakui.zhao@intel.com>
Reviewed-by: "Guo, Yejun" <yejun.guo@intel.com>

1 files changed, 24 insertions, 0 deletions

diff --git a/kernels/runtime_mirror_effect.cl b/kernels/runtime_mirror_effect.cl
new file mode 100644
index 00000000..a9787920
--- /dev/null
+++ b/kernels/runtime_mirror_effect.cl
@@ -0,0 +1,24 @@
+__kernel void
+runtime_mirror_effect(__read_only image2d_t src_y,
+                      __read_only image2d_t src_uv,
+                      __write_only image2d_t dst_y,
+                      __write_only image2d_t dst_uv,
+                      int src_height)
+{
+  const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST;
+  int2 loc = (int2)(get_global_id(0), get_global_id(1));
+  uint4 color_y, color_uv;
+
+  if (loc.y < src_height/2) {
+    color_y = read_imageui(src_y, sampler, loc);
+    color_uv = read_imageui(src_uv, sampler, loc/2);
+  } else {
+    int2 newloc = (int2)(loc.x, src_height - loc.y);
+    color_y = read_imageui(src_y, sampler, newloc);
+    color_uv = read_imageui(src_uv, sampler, newloc/2);
+  }
+
+  write_imageui(dst_y, loc, color_y);
+  write_imageui(dst_uv, loc/2, color_uv);
+
+}