[v4,1/2] Add example to show v4l2 buffer sharing with extension clGetMemObjectFdIntel.

Submitted by Chuanbo Weng on June 18, 2015, 8:29 a.m.

Details

Message ID 1434616179-30006-1-git-send-email-chuanbo.weng@intel.com
State New
Headers show

Not browsing as part of any series.

Commit Message

Chuanbo Weng June 18, 2015, 8:29 a.m.
This example captures yuy2 frame directly to cl buffer object by the way
of dma, processed by OpenCL kernel, then convert to nv12 format and
shown by libva.

v2:
Close cl buffer's fd by clCloseMemObjectFdIntel instead of close
function.
v3:
Just use close function, no need of clCloseMemObjectFdIntel.
v4:
Some modifcation of examples/CMakeLists.txt after code rebase.

Signed-off-by: Chuanbo Weng <chuanbo.weng@intel.com>
---
 CMakeLists.txt                                     |  35 +-
 examples/CMakeLists.txt                            |  29 +-
 .../v4l2_buffer_sharing/v4l2_buffer_sharing.cpp    | 590 +++++++++++++++++++++
 kernels/runtime_yuy2_processing.cl                 |  15 +
 4 files changed, 645 insertions(+), 24 deletions(-)
 create mode 100644 examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp
 create mode 100644 kernels/runtime_yuy2_processing.cl

Patch hide | download patch | download mbox

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 5474447..4f627cf 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -216,23 +216,30 @@  IF(BUILD_EXAMPLES)
 IF(NOT X11_FOUND)
   MESSAGE(FATAL_ERROR "XLib is necessary for examples - not found")
 ENDIF(NOT X11_FOUND)
-# libva
-pkg_check_modules(LIBVA REQUIRED libva>=0.36.0)
-IF(LIBVA_FOUND)
+# libva & libva-x11
+#pkg_check_modules(LIBVA REQUIRED libva>=0.36.0)
+pkg_check_modules(LIBVA REQUIRED libva)
+pkg_check_modules(LIBVA-X11 REQUIRED libva-x11)
+set(LIBVA_BUF_SH_DEP false)
+set(V4L2_BUF_SH_DEP false)
+IF(LIBVA_FOUND AND LIBVA-X11_FOUND)
   MESSAGE(STATUS "Looking for LIBVA - found at ${LIBVA_PREFIX} ${LIBVA_VERSION}")
-  INCLUDE_DIRECTORIES(${LIBVA_INCLUDE_DIRS})
-ELSE(LIBVA_FOUND)
-  MESSAGE(STATUS "Looking for LIBVA (>= 0.36.0) - not found")
-ENDIF(LIBVA_FOUND)
-
-# libva-x11
-pkg_check_modules(LIBVA-X11 REQUIRED libva-x11>=0.36.0)
-IF(LIBVA-X11_FOUND)
   MESSAGE(STATUS "Looking for LIBVA-X11 - found at ${LIBVA-X11_PREFIX} ${LIBVA-X11_VERSION}")
+  INCLUDE_DIRECTORIES(${LIBVA_INCLUDE_DIRS})
   INCLUDE_DIRECTORIES(${LIBVA-X11_INCLUDE_DIRS})
-ELSE(LIBVA-X11_FOUND)
-  MESSAGE(STATUS "Looking for LIBVA-X11 (>= 0.36.0) - not found")
-ENDIF(LIBVA-X11_FOUND)
+  set(V4L2_BUF_SH_DEP true)
+  IF(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION VERSION_LESS "0.36.0")
+    IF(LIBVA_VERSION VERSION_LESS "0.36.0")
+      MESSAGE(STATUS "Looking for LIBVA (>= 0.36.0) - not found")
+    ENDIF(LIBVA_VERSION VERSION_LESS "0.36.0")
+    IF(LIBVA-X11_VERSION VERSION_LESS "0.36.0")
+      MESSAGE(STATUS "Looking for LIBVA-X11 (>= 0.36.0) - not found")
+    ENDIF(LIBVA-X11_VERSION VERSION_LESS "0.36.0")
+    MESSAGE(STATUS "Example libva_buffer_sharing will not be built")
+  ELSE(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION VERSION_LESS "0.36.0")
+    set(LIBVA_BUF_SH_DEP true)
+  ENDIF(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION VERSION_LESS "0.36.0")
+ENDIF(LIBVA_FOUND AND LIBVA-X11_FOUND)
 ENDIF(BUILD_EXAMPLES)
 
 ADD_SUBDIRECTORY(include)
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
index fe4e5f6..850b3d9 100644
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@@ -1,3 +1,9 @@ 
+INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}
+                    ${CMAKE_CURRENT_SOURCE_DIR}/../utests
+                    ${CMAKE_CURRENT_SOURCE_DIR}/../include
+                    ${X11_INCLUDE_DIR})
+
+IF(LIBVA_BUF_SH_DEP OR V4L2_BUF_SH_DEP)
 EXECUTE_PROCESS(COMMAND ls "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva" OUTPUT_VARIABLE LS_RESULT)
 IF ("LS_RESULT" STREQUAL "")
 EXECUTE_PROCESS(COMMAND git submodule init WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/..)
@@ -5,17 +11,13 @@  EXECUTE_PROCESS(COMMAND git submodule update WORKING_DIRECTORY ${CMAKE_CURRENT_S
 EXECUTE_PROCESS(COMMAND git checkout master WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva)
 ENDIF ("LS_RESULT" STREQUAL "")
 
-INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}
-                    ${CMAKE_CURRENT_SOURCE_DIR}/../utests
-                    ${CMAKE_CURRENT_SOURCE_DIR}/../include
-                    ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/va
-                    ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/test/common
-                    ${X11_INCLUDE_DIR})
+INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/va
+                    ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/test/common)
 
 link_directories (${LIBVA_LIBDIR}
                   ${LIBVA-X11_LIBDIR})
 
-set (examples_sources
+set (va_ocl_basic_sources
   ../utests/utest_error.c
   ../utests/utest_assert.cpp
   ../utests/utest_file_map.cpp
@@ -23,13 +25,20 @@  set (examples_sources
   ./thirdparty/libva/test/common/va_display.c
   ./thirdparty/libva/test/common/va_display_x11.c)
 
-
 ADD_DEFINITIONS(-DHAVE_VA_X11)
-ADD_DEFINITIONS(-DINPUT_NV12_DEFAULT="${CMAKE_CURRENT_SOURCE_DIR}/libva_buffer_sharing/256_128.nv12")
 
-ADD_LIBRARY(va_ocl_basic SHARED ${examples_sources})
+ADD_LIBRARY(va_ocl_basic SHARED ${va_ocl_basic_sources})
 
 TARGET_LINK_LIBRARIES(va_ocl_basic cl m va va-x11 ${X11_X11_LIB})
 
+IF(LIBVA_BUF_SH_DEP)
+ADD_DEFINITIONS(-DINPUT_NV12_DEFAULT="${CMAKE_CURRENT_SOURCE_DIR}/libva_buffer_sharing/256_128.nv12")
 ADD_EXECUTABLE(example-libva_buffer_sharing ./libva_buffer_sharing/libva_buffer_sharing.cpp)
 TARGET_LINK_LIBRARIES(example-libva_buffer_sharing va_ocl_basic)
+ENDIF(LIBVA_BUF_SH_DEP)
+
+IF(V4L2_BUF_SH_DEP)
+ADD_EXECUTABLE(example-v4l2_buffer_sharing ./v4l2_buffer_sharing/v4l2_buffer_sharing.cpp)
+TARGET_LINK_LIBRARIES(example-v4l2_buffer_sharing va_ocl_basic)
+ENDIF(V4L2_BUF_SH_DEP)
+ENDIF(LIBVA_BUF_SH_DEP OR V4L2_BUF_SH_DEP)
diff --git a/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp b/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp
new file mode 100644
index 0000000..42ab642
--- /dev/null
+++ b/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp
@@ -0,0 +1,590 @@ 
+/*
+ ** Copyright (c) 2012, 2015 Intel Corporation. All Rights Reserved.
+ **
+ ** Permission is hereby granted, free of charge, to any person obtaining a
+ ** copy of this software and associated documentation files (the
+ ** "Software"), to deal in the Software without restriction, including
+ ** without limitation the rights to use, copy, modify, merge, publish,
+ ** distribute, sub license, and/or sell copies of the Software, and to
+ ** permit persons to whom the Software is furnished to do so, subject to
+ ** the following conditions:
+ **
+ ** The above copyright notice and this permission notice (including the
+ ** next paragraph) shall be included in all copies or substantial portions
+ ** of the Software.
+ **
+ ** THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ ** OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ ** MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT.
+ ** IN NO EVENT SHALL PRECISION INSIGHT AND/OR ITS SUPPLIERS BE LIABLE FOR
+ ** ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ ** TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ ** SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ **/
+
+#include <getopt.h>
+#include <errno.h>
+#include <assert.h>
+#include <fcntl.h>
+#include <linux/videodev2.h>
+#include <poll.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+#include <string.h>
+#include <sys/ioctl.h>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <unistd.h>
+#include <sys/time.h>
+#include <time.h>
+
+#include <inttypes.h>
+#include <ctype.h>
+
+#include <va/va.h>
+#include <va/va_drmcommon.h>
+
+#include "va_display.h"
+#include "utest_helper.hpp"
+
+using namespace std;
+
+#define BUFFER_NUM_DEFAULT 5
+#define VIDEO_NODE_DEFAULT "/dev/video0"
+#define WIDTH_DEFAULT 640
+#define HEIGHT_DEFAULT 480
+
+#define CHECK_VASTATUS(va_status,func)                                  \
+  if (va_status != VA_STATUS_SUCCESS) {                                   \
+    fprintf(stderr, "status = %d, %s: %s(line %d) failed, exit\n",va_status, __func__, func, __LINE__); \
+    exit(1);                                                            \
+  }
+
+#define CHECK_CLSTATUS(status,func)                                  \
+  if (status != CL_SUCCESS) {                                   \
+    fprintf(stderr, "status = %d, %s: %s(line %d) failed, exit\n", status, __func__, func, __LINE__); \
+    exit(1);                                                            \
+  }
+
+#define CHECK_V4L2ERROR(ret, STR)                               \
+  if (ret){                             \
+    fprintf(stderr, STR);            \
+    perror(" ");                            \
+    fprintf(stderr, "ret = %d, %s: %s(line %d) failed, exit\n", ret, __func__, STR, __LINE__);      \
+    exit(1);                                  \
+  }
+
+VADisplay	va_dpy;
+cl_int cl_status;
+VAStatus va_status;
+VASurfaceID nv12_surface_id;
+VAImage nv12_image;
+
+int dev_fd;
+uint64_t image_size;
+unsigned int pitch;
+cl_mem *import_buf = NULL;
+typedef cl_int (OCLGETMEMOBJECTFD)(cl_context, cl_mem, int *);
+OCLGETMEMOBJECTFD *oclGetMemObjectFd = NULL;
+
+int frame_count = 0;
+struct v4l2_options{
+  const char *dev_name;
+  unsigned int width, height;
+  unsigned int spec_res;
+  unsigned int buffer_num;
+  unsigned int do_list;
+} vo;
+int *import_buf_fd = NULL;
+
+static const char short_options[] = "d:r:b:lh";
+
+static const struct option
+long_options[] = {
+  { "device", required_argument, NULL, 'd' },
+  { "help",   no_argument,       NULL, 'h' },
+  { "resolution", required_argument,       NULL, 'r' },
+  { "buffer_num",  required_argument, NULL, 'b' },
+  { "list",  no_argument, NULL, 'l' },
+  { 0, 0, 0, 0 }
+};
+
+static void usage(FILE *fp, int argc, char **argv)
+{
+  fprintf(fp,
+      "This example aims to demostrate the usage of DMABUF buffer sharing between v4l2 and Beignet.\n"
+      "For more details, please read docs/howto/v4l2-buffer-sharing-howto.mdwn.\n"
+      "Usage: %s [options]\n\n"
+      "Options:\n"
+      "-d | --device=<dev>  Specify device by <dev> instead of /dev/video0\n"
+      "-h | --help          Print this message\n"
+      "-r | --resolution=<width,height>    Set image resolution\n"
+      "-b | --buffer_num=<num>  Set number of buffers\n"
+      "-l | --list  List available resolution of format 'V4L2_PIX_FMT_YUYV'\n"
+      "",
+      argv[0]);
+}
+
+static void list_resolution(){
+  int ret;
+  struct v4l2_capability cap;
+  struct v4l2_frmsizeenum frm_sz;
+
+  dev_fd = open(vo.dev_name, O_RDWR | O_NONBLOCK, 0);
+  if (dev_fd < 0) {
+    fprintf(stderr, "Can not open %s: %s\n",
+        vo.dev_name, strerror(errno));
+    exit(1);
+  }
+
+  memset(&cap, 0, sizeof(cap));
+  ret = ioctl(dev_fd, VIDIOC_QUERYCAP, &cap);
+  CHECK_V4L2ERROR(ret, "VIDIOC_QUERYCAP");
+
+  if(!(cap.capabilities & V4L2_CAP_VIDEO_CAPTURE)){
+    fprintf(stderr, "The device is not video capture device\n");
+    exit(1);
+  }
+  if(!(cap.capabilities & V4L2_CAP_STREAMING)){
+    fprintf(stderr, "The device does not support streaming i/o\n");
+    exit(1);
+  }
+
+  printf("Supported resolution under pixel format 'V4L2_PIX_FMT_YUYV':\n");
+  frm_sz.pixel_format = V4L2_PIX_FMT_YUYV;
+  frm_sz.index = 0;
+  bool extra_info = true;
+  while (ioctl(dev_fd, VIDIOC_ENUM_FRAMESIZES, &frm_sz) == 0) {
+    if (frm_sz.type == V4L2_FRMSIZE_TYPE_DISCRETE) {
+      if(extra_info){
+        printf("(width, height) = \n");
+        extra_info = false;
+      }
+      printf("(%d, %d)", frm_sz.discrete.width, frm_sz.discrete.height);
+      printf("\n");
+    }
+    else if (frm_sz.type == V4L2_FRMSIZE_TYPE_STEPWISE) {
+      printf("(width, height) from (%d, %d) to (%d, %d) with step (%d, %d)",
+          frm_sz.stepwise.min_width,
+          frm_sz.stepwise.min_height,
+          frm_sz.stepwise.max_width,
+          frm_sz.stepwise.max_height,
+          frm_sz.stepwise.step_width,
+          frm_sz.stepwise.step_height);
+      continue;
+    }
+    frm_sz.index++;
+  }
+
+  ret = close(dev_fd);
+  if (ret) {
+    fprintf(stderr, "Failed to close %s: %s\n",
+        vo.dev_name, strerror(errno));
+    exit(1);
+  }
+}
+
+static void analyse_args(int argc, char *argv[])
+{
+  vo.dev_name = NULL;
+  vo.width = 0;
+  vo.height = 0;
+  vo.spec_res = 0;
+  vo.buffer_num = BUFFER_NUM_DEFAULT;
+  vo.do_list = 0;
+
+  int c, idx;
+  for (;;) {
+
+    c = getopt_long(argc, argv,
+        short_options, long_options, &idx);
+
+    if (-1 == c)
+      break;
+
+    switch (c) {
+      case 0:
+        break;
+
+      case 'd':
+        vo.dev_name = optarg;
+        break;
+
+      case '?':
+      case 'h':
+        usage(stdout, argc, argv);
+        exit(0);
+
+      case 'r':
+        sscanf(optarg, "%d,%d", &vo.width, &vo.height);
+        vo.spec_res = 1;
+        break;
+
+      case 'b':
+        vo.buffer_num = strtoul(optarg, NULL, 0);
+        break;
+
+      case 'l':
+        vo.do_list = 1;
+        break;
+
+      default:
+        usage(stderr, argc, argv);
+        exit(1);
+    }
+  }
+
+  if(!vo.dev_name){
+    printf("Haven't specified device, use default device: %s\n",
+        VIDEO_NODE_DEFAULT);
+  }
+  if(!vo.dev_name)
+    vo.dev_name = VIDEO_NODE_DEFAULT;
+  if(vo.do_list){
+    list_resolution();
+    exit(0);
+  }
+  if(!vo.spec_res){
+    printf("Haven't specified resolution, use default resolution: (width,height) = (%d, %d)\n",
+        WIDTH_DEFAULT, HEIGHT_DEFAULT);
+    vo.width = WIDTH_DEFAULT;
+    vo.height = HEIGHT_DEFAULT;
+  }
+  return;
+}
+
+static void initialize_va_ocl(){
+  int major_ver, minor_ver;
+
+  printf("\n***********************libva info: ***********************\n");
+  fflush(stdout);
+  va_dpy = va_open_display();
+  va_status = vaInitialize(va_dpy, &major_ver, &minor_ver);
+  CHECK_VASTATUS(va_status, "vaInitialize");
+
+  VASurfaceAttrib forcc;
+  forcc.type =VASurfaceAttribPixelFormat;
+  forcc.flags=VA_SURFACE_ATTRIB_SETTABLE;
+  forcc.value.type=VAGenericValueTypeInteger;
+  forcc.value.value.i = VA_FOURCC_NV12;
+  va_status = vaCreateSurfaces(va_dpy, VA_RT_FORMAT_YUV420,
+                               vo.width, vo.height,
+                               &nv12_surface_id, 1, &forcc, 1);
+  CHECK_VASTATUS(va_status, "vaCreateSurfaces");
+
+  VAImageFormat image_fmt;
+  image_fmt.fourcc = VA_FOURCC_NV12;
+  image_fmt.byte_order = VA_LSB_FIRST;
+  image_fmt.bits_per_pixel = 12;
+  va_status = vaCreateImage(va_dpy, &image_fmt, vo.width, vo.height, &nv12_image);
+  CHECK_VASTATUS(va_status, "vaCreateImage");
+
+  //ocl initialization: basic & create kernel & get extension
+  printf("\n***********************OpenCL info: ***********************\n");
+  if ((cl_status = cl_test_init("runtime_yuy2_processing.cl", "runtime_yuy2_processing", SOURCE)) != 0){
+    fprintf(stderr, "cl_test_init error\n");
+    exit(1);
+  }
+
+#ifdef CL_VERSION_1_2
+  oclGetMemObjectFd = (OCLGETMEMOBJECTFD *)clGetExtensionFunctionAddressForPlatform(platform, "clGetMemObjectFdIntel");
+#else
+  oclGetMemObjectFd = (OCLGETMEMOBJECTFD *)clGetExtensionFunctionAddress("clGetMemObjectFdIntel");
+#endif
+  if(!oclGetMemObjectFd){
+    fprintf(stderr, "Failed to get extension clGetMemObjectFdIntel\n");
+    exit(1);
+  }
+  printf("\n***********************************************************\n");
+}
+
+static void create_dmasharing_buffers()
+{
+  if(import_buf_fd == NULL)
+    import_buf_fd = (int *)malloc(sizeof(int) * vo.buffer_num);
+  if(import_buf == NULL){
+    import_buf = (cl_mem *)malloc(sizeof(cl_mem) * vo.buffer_num);
+  }
+
+  for (unsigned int i = 0; i < vo.buffer_num; ++i){
+    import_buf[i] = clCreateBuffer(ctx, CL_MEM_READ_WRITE, image_size, NULL, &cl_status);
+    CHECK_CLSTATUS(cl_status, "clCreateBuffer");
+
+    //get cl buffer object's fd
+    cl_status = oclGetMemObjectFd(ctx, import_buf[i], &import_buf_fd[i]);
+    CHECK_CLSTATUS(cl_status, "clGetMemObjectFdIntel");
+  }
+}
+
+static void release_va_ocl(){
+  va_status = vaDestroySurfaces(va_dpy,&nv12_surface_id,1);
+  CHECK_VASTATUS(va_status, "vaDestroySurfaces");
+  va_status = vaDestroyImage(va_dpy, nv12_image.image_id);
+  CHECK_VASTATUS(va_status, "vaDestroyImage");
+  va_status = vaTerminate(va_dpy);
+  CHECK_VASTATUS(va_status, "vaTerminate");
+  va_close_display(va_dpy);
+
+  int ret;
+  for (unsigned int i = 0; i < vo.buffer_num; ++i) {
+    ret = close(import_buf_fd[i]);
+    if (ret) {
+      fprintf(stderr, "Failed to close import_buf[%u]'s fd: %s\n", i, strerror(errno));
+    }
+    cl_status = clReleaseMemObject(import_buf[i]);
+    CHECK_CLSTATUS(cl_status, "clReleaseMemObject");
+  }
+}
+
+static void process_show_frame(int index)
+{
+  //process import_buf[index] by ocl
+  size_t global_size[2];
+  global_size[0] = vo.width * 2 / 4;
+  global_size[1] = vo.height;
+  cl_status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &import_buf[index]);
+  CHECK_CLSTATUS(cl_status, "clSetKernelArg");
+  cl_status = clSetKernelArg(kernel, 1, sizeof(int), &vo.height);
+  CHECK_CLSTATUS(cl_status, "clSetKernelArg");
+  cl_status = clSetKernelArg(kernel, 2, sizeof(int), &pitch);
+  CHECK_CLSTATUS(cl_status, "clSetKernelArg");
+  cl_status = clEnqueueNDRangeKernel(queue, kernel, 2, NULL,
+                                     global_size, NULL, 0, NULL, NULL);
+  CHECK_CLSTATUS(cl_status, "clEnqueueNDRangeKernel");
+  cl_status = clFinish(queue);
+  CHECK_CLSTATUS(cl_status, "clFinish");
+
+  //create corresponding VASurface
+  VASurfaceID yuy2_surface_id;
+  VASurfaceAttrib sa[2];
+  sa[0].type = VASurfaceAttribMemoryType;
+  sa[0].flags = VA_SURFACE_ATTRIB_SETTABLE;
+  sa[0].value.type = VAGenericValueTypeInteger;
+  sa[0].value.value.i = VA_SURFACE_ATTRIB_MEM_TYPE_DRM_PRIME;
+  sa[1].type = VASurfaceAttribExternalBufferDescriptor;
+  sa[1].flags = VA_SURFACE_ATTRIB_SETTABLE;
+  sa[1].value.type = VAGenericValueTypePointer;
+  VASurfaceAttribExternalBuffers sa_eb;
+  sa_eb.pixel_format = VA_FOURCC_YUY2;
+  sa_eb.width = vo.width;
+  sa_eb.height = vo.height;
+  sa_eb.data_size = image_size;
+  sa_eb.num_planes = 1;
+  sa_eb.pitches[0] = pitch;
+  sa_eb.offsets[0] = 0;
+  sa_eb.num_buffers = 1;
+  sa_eb.buffers = (unsigned long *)malloc(sizeof(unsigned long) * sa_eb.num_buffers);
+  sa_eb.buffers[0] = import_buf_fd[index];
+  sa_eb.flags = 0;
+  sa[1].value.value.p = &sa_eb;
+  va_status = vaCreateSurfaces(va_dpy, VA_RT_FORMAT_YUV422,
+                               vo.width, vo.height,
+                               &yuy2_surface_id, 1, sa, 2);
+  CHECK_VASTATUS(va_status, "vaCreateSurfaces");
+
+  //convert to NV12 format
+  va_status = vaGetImage (va_dpy, yuy2_surface_id, 0, 0,
+                          vo.width, vo.height, nv12_image.image_id);
+  CHECK_VASTATUS(va_status, "vaGetImage");
+  va_status = vaPutImage(va_dpy, nv12_surface_id, nv12_image.image_id,
+                         0, 0, vo.width, vo.height, 0, 0,
+                         vo.width, vo.height);
+  CHECK_VASTATUS(va_status, "vaPutImage");
+
+  //show by vaPutsurface
+  VARectangle src_rect, dst_rect;
+  src_rect.x      = 0;
+  src_rect.y      = 0;
+  src_rect.width  = vo.width;
+  src_rect.height = vo.height;
+  dst_rect        = src_rect;
+  va_status = va_put_surface(va_dpy, nv12_surface_id, &src_rect, &dst_rect);
+  CHECK_VASTATUS(va_status, "vaPutSurface");
+
+  vaDestroySurfaces(va_dpy,&yuy2_surface_id,1);
+  CHECK_VASTATUS(va_status, "vaDestroySurfaces");
+  free(sa_eb.buffers);
+  return;
+}
+
+static void init_dmabuf(void){
+  int ret;
+  struct v4l2_requestbuffers reqbuf;
+
+  memset(&reqbuf, 0, sizeof(reqbuf));
+  reqbuf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+  reqbuf.memory = V4L2_MEMORY_DMABUF;
+  reqbuf.count = vo.buffer_num;
+
+  ret = ioctl(dev_fd, VIDIOC_REQBUFS, &reqbuf);
+  if(ret == -1 && errno == EINVAL){
+    fprintf(stderr, "Video capturing or DMABUF streaming is not supported\n");
+    exit(1);
+  }
+  else
+    CHECK_V4L2ERROR(ret, "VIDIOC_REQBUFS");
+
+  create_dmasharing_buffers();
+  printf("Succeed to create %d dma buffers \n", vo.buffer_num);
+
+}
+
+static void init_device(void){
+
+  int ret;
+  struct v4l2_capability cap;
+  struct v4l2_format format;
+
+  dev_fd = open(vo.dev_name, O_RDWR | O_NONBLOCK, 0);
+  if (dev_fd < 0) {
+    fprintf(stderr, "Can not open %s: %s\n",
+        vo.dev_name, strerror(errno));
+    exit(1);
+  }
+
+  memset(&cap, 0, sizeof(cap));
+  ret = ioctl(dev_fd, VIDIOC_QUERYCAP, &cap);
+  CHECK_V4L2ERROR(ret, "VIDIOC_QUERYCAP");
+  if(!(cap.capabilities & V4L2_CAP_STREAMING)){
+    fprintf(stderr, "The device does not support streaming i/o\n");
+    exit(1);
+  }
+
+  memset(&format, 0, sizeof(format));
+  format.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+  format.fmt.pix.width = vo.width;
+  format.fmt.pix.height = vo.height;
+  format.fmt.pix.pixelformat = V4L2_PIX_FMT_YUYV;
+  format.fmt.pix.field = V4L2_FIELD_ANY;
+
+  ret = ioctl(dev_fd, VIDIOC_S_FMT, &format);
+  CHECK_V4L2ERROR(ret, "VIDIOC_S_FMT");
+
+  ret = ioctl(dev_fd, VIDIOC_G_FMT, &format);
+  CHECK_V4L2ERROR(ret, "VIDIOC_G_FMT");
+  if(format.fmt.pix.pixelformat != V4L2_PIX_FMT_YUYV){
+    fprintf(stderr, "V4L2_PIX_FMT_YUYV format is not supported by %s\n", vo.dev_name);
+    exit(1);
+  }
+  if(format.fmt.pix.width != vo.width  || format.fmt.pix.height != vo.height){
+    fprintf(stderr, "This resolution is not supported, please go through supported resolution by command './main -l'\n");
+    exit(1);
+  }
+  printf("Input image format: (width, height) = (%u, %u), pixel format = %.4s\n",
+      format.fmt.pix.width, format.fmt.pix.height, (char*)&format.fmt.pix.pixelformat);
+  image_size = format.fmt.pix.sizeimage;
+	pitch = format.fmt.pix.bytesperline;
+}
+
+static void start_capturing(void){
+  int ret;
+  for (unsigned int i = 0; i < vo.buffer_num; ++i) {
+    struct v4l2_buffer buf;
+
+    memset(&buf, 0, sizeof(buf));
+    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+    buf.memory = V4L2_MEMORY_DMABUF;
+    buf.index = i;
+    buf.m.fd = import_buf_fd[i];
+    ret = ioctl(dev_fd, VIDIOC_QBUF, &buf);
+    CHECK_V4L2ERROR(ret, "VIDIOC_QBUF");
+  }
+
+  int type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+  ret = ioctl(dev_fd, VIDIOC_STREAMON, &type);
+  CHECK_V4L2ERROR(ret, "VIDIOC_STREAMON");
+}
+
+static void mainloop(void){
+  int ret;
+  struct v4l2_buffer buf;
+  int index;
+
+  while (1) {
+    frame_count++;
+    printf("******************Frame %d\n", frame_count);
+    fd_set fds;
+    struct timeval tv;
+    int r;
+
+    FD_ZERO(&fds);
+    FD_SET(dev_fd, &fds);
+
+    /* Timeout. */
+    tv.tv_sec = 2;
+    tv.tv_usec = 0;
+
+
+    r = select(dev_fd + 1, &fds, NULL, NULL, &tv);
+
+    if (-1 == r) {
+      if (EINTR == errno)
+        continue;
+      perror("select");
+    }
+
+    if(r == 0){
+      fprintf(stderr, "Select timeout\n");
+      exit(1);
+    }
+
+    memset(&buf, 0, sizeof(buf));
+    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+    buf.memory = V4L2_MEMORY_DMABUF;
+    ret = ioctl(dev_fd, VIDIOC_DQBUF, &buf);
+    CHECK_V4L2ERROR(ret, "VIDIOC_DQBUF");
+    index = buf.index;
+
+    //process by ocl and show on screen by libva
+    process_show_frame(index);
+
+    //Then queue this buffer(buf.index) by QBUF
+    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+    buf.memory = V4L2_MEMORY_DMABUF;
+    buf.m.fd = import_buf_fd[index];
+    buf.index = index;
+
+    ret = ioctl(dev_fd, VIDIOC_QBUF, &buf);
+    CHECK_V4L2ERROR(ret, "VIDIOC_QBUF");
+  }
+}
+
+static void stop_capturing(void)
+{
+  int ret;
+  int type = V4L2_BUF_TYPE_VIDEO_CAPTURE;
+
+  ret = ioctl(dev_fd, VIDIOC_STREAMOFF, &type);
+  CHECK_V4L2ERROR(ret, "VIDIOC_STREAMOFF");
+}
+
+static void uninit_device(void){
+  free(import_buf_fd);
+  free(import_buf);
+  int ret = close(dev_fd);
+  if (ret) {
+    fprintf(stderr, "Failed to close %s: %s\n",
+        vo.dev_name, strerror(errno));
+    exit(1);
+  }
+}
+
+int main(int argc, char *argv[])
+{
+  analyse_args(argc, argv);
+
+  init_device();
+  initialize_va_ocl();
+  init_dmabuf();
+
+  start_capturing();
+  mainloop();
+
+  stop_capturing();
+  release_va_ocl();
+  uninit_device();
+
+  return 0;
+}
diff --git a/kernels/runtime_yuy2_processing.cl b/kernels/runtime_yuy2_processing.cl
new file mode 100644
index 0000000..1478e65
--- /dev/null
+++ b/kernels/runtime_yuy2_processing.cl
@@ -0,0 +1,15 @@ 
+__kernel void
+runtime_yuy2_processing(__global uchar *src,
+                        int image_height,
+                        int image_pitch)
+{
+  int gx = get_global_id(0);
+  int gy = get_global_id(1);
+
+  int src_y = image_height / 2 + gy;
+  int mirror_y = image_height - src_y;
+
+  uchar4 mirror_val = *(__global uchar4*)(src + mirror_y*image_pitch + gx*4);
+  *(__global uchar4*)(src + src_y*image_pitch + gx*4) = mirror_val;
+
+}

Comments

Pushed, thanks.

> -----Original Message-----

> From: Beignet [mailto:beignet-bounces@lists.freedesktop.org] On Behalf Of

> Chuanbo Weng

> Sent: Thursday, June 18, 2015 16:30

> To: beignet@lists.freedesktop.org

> Cc: Weng, Chuanbo

> Subject: [Beignet] [PATCH v4 1/2] Add example to show v4l2 buffer sharing

> with extension clGetMemObjectFdIntel.

> 

> This example captures yuy2 frame directly to cl buffer object by the way of

> dma, processed by OpenCL kernel, then convert to nv12 format and shown

> by libva.

> 

> v2:

> Close cl buffer's fd by clCloseMemObjectFdIntel instead of close function.

> v3:

> Just use close function, no need of clCloseMemObjectFdIntel.

> v4:

> Some modifcation of examples/CMakeLists.txt after code rebase.

> 

> Signed-off-by: Chuanbo Weng <chuanbo.weng@intel.com>

> ---

>  CMakeLists.txt                                     |  35 +-

>  examples/CMakeLists.txt                            |  29 +-

>  .../v4l2_buffer_sharing/v4l2_buffer_sharing.cpp    | 590

> +++++++++++++++++++++

>  kernels/runtime_yuy2_processing.cl                 |  15 +

>  4 files changed, 645 insertions(+), 24 deletions(-)  create mode 100644

> examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp

>  create mode 100644 kernels/runtime_yuy2_processing.cl

> 

> diff --git a/CMakeLists.txt b/CMakeLists.txt index 5474447..4f627cf 100644

> --- a/CMakeLists.txt

> +++ b/CMakeLists.txt

> @@ -216,23 +216,30 @@ IF(BUILD_EXAMPLES)  IF(NOT X11_FOUND)

>    MESSAGE(FATAL_ERROR "XLib is necessary for examples - not found")

> ENDIF(NOT X11_FOUND) -# libva -pkg_check_modules(LIBVA REQUIRED

> libva>=0.36.0)

> -IF(LIBVA_FOUND)

> +# libva & libva-x11

> +#pkg_check_modules(LIBVA REQUIRED libva>=0.36.0)

> +pkg_check_modules(LIBVA REQUIRED libva)

> +pkg_check_modules(LIBVA-X11 REQUIRED libva-x11)

> set(LIBVA_BUF_SH_DEP

> +false) set(V4L2_BUF_SH_DEP false) IF(LIBVA_FOUND AND LIBVA-

> X11_FOUND)

>    MESSAGE(STATUS "Looking for LIBVA - found at ${LIBVA_PREFIX}

> ${LIBVA_VERSION}")

> -  INCLUDE_DIRECTORIES(${LIBVA_INCLUDE_DIRS})

> -ELSE(LIBVA_FOUND)

> -  MESSAGE(STATUS "Looking for LIBVA (>= 0.36.0) - not found")

> -ENDIF(LIBVA_FOUND)

> -

> -# libva-x11

> -pkg_check_modules(LIBVA-X11 REQUIRED libva-x11>=0.36.0)

> -IF(LIBVA-X11_FOUND)

>    MESSAGE(STATUS "Looking for LIBVA-X11 - found at ${LIBVA-X11_PREFIX}

> ${LIBVA-X11_VERSION}")

> +  INCLUDE_DIRECTORIES(${LIBVA_INCLUDE_DIRS})

>    INCLUDE_DIRECTORIES(${LIBVA-X11_INCLUDE_DIRS})

> -ELSE(LIBVA-X11_FOUND)

> -  MESSAGE(STATUS "Looking for LIBVA-X11 (>= 0.36.0) - not found")

> -ENDIF(LIBVA-X11_FOUND)

> +  set(V4L2_BUF_SH_DEP true)

> +  IF(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION

> VERSION_LESS "0.36.0")

> +    IF(LIBVA_VERSION VERSION_LESS "0.36.0")

> +      MESSAGE(STATUS "Looking for LIBVA (>= 0.36.0) - not found")

> +    ENDIF(LIBVA_VERSION VERSION_LESS "0.36.0")

> +    IF(LIBVA-X11_VERSION VERSION_LESS "0.36.0")

> +      MESSAGE(STATUS "Looking for LIBVA-X11 (>= 0.36.0) - not found")

> +    ENDIF(LIBVA-X11_VERSION VERSION_LESS "0.36.0")

> +    MESSAGE(STATUS "Example libva_buffer_sharing will not be built")

> +  ELSE(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION

> VERSION_LESS "0.36.0")

> +    set(LIBVA_BUF_SH_DEP true)

> +  ENDIF(LIBVA_VERSION VERSION_LESS "0.36.0" OR LIBVA-X11_VERSION

> +VERSION_LESS "0.36.0") ENDIF(LIBVA_FOUND AND LIBVA-X11_FOUND)

>  ENDIF(BUILD_EXAMPLES)

> 

>  ADD_SUBDIRECTORY(include)

> diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index

> fe4e5f6..850b3d9 100644

> --- a/examples/CMakeLists.txt

> +++ b/examples/CMakeLists.txt

> @@ -1,3 +1,9 @@

> +INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}

> +                    ${CMAKE_CURRENT_SOURCE_DIR}/../utests

> +                    ${CMAKE_CURRENT_SOURCE_DIR}/../include

> +                    ${X11_INCLUDE_DIR})

> +

> +IF(LIBVA_BUF_SH_DEP OR V4L2_BUF_SH_DEP)

>  EXECUTE_PROCESS(COMMAND ls

> "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva" OUTPUT_VARIABLE

> LS_RESULT)  IF ("LS_RESULT" STREQUAL "")  EXECUTE_PROCESS(COMMAND

> git submodule init WORKING_DIRECTORY

> ${CMAKE_CURRENT_SOURCE_DIR}/..) @@ -5,17 +11,13 @@

> EXECUTE_PROCESS(COMMAND git submodule update

> WORKING_DIRECTORY ${CMAKE_CURRENT_S

> EXECUTE_PROCESS(COMMAND git checkout master WORKING_DIRECTORY

> ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva)

>  ENDIF ("LS_RESULT" STREQUAL "")

> 

> -INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}

> -                    ${CMAKE_CURRENT_SOURCE_DIR}/../utests

> -                    ${CMAKE_CURRENT_SOURCE_DIR}/../include

> -                    ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/va

> -

> ${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/test/common

> -                    ${X11_INCLUDE_DIR})

> +INCLUDE_DIRECTORIES(${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva

> /va

> +

> +${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva/test/common)

> 

>  link_directories (${LIBVA_LIBDIR}

>                    ${LIBVA-X11_LIBDIR})

> 

> -set (examples_sources

> +set (va_ocl_basic_sources

>    ../utests/utest_error.c

>    ../utests/utest_assert.cpp

>    ../utests/utest_file_map.cpp

> @@ -23,13 +25,20 @@ set (examples_sources

>    ./thirdparty/libva/test/common/va_display.c

>    ./thirdparty/libva/test/common/va_display_x11.c)

> 

> -

>  ADD_DEFINITIONS(-DHAVE_VA_X11)

> -ADD_DEFINITIONS(-

> DINPUT_NV12_DEFAULT="${CMAKE_CURRENT_SOURCE_DIR}/libva_buffer_

> sharing/256_128.nv12")

> 

> -ADD_LIBRARY(va_ocl_basic SHARED ${examples_sources})

> +ADD_LIBRARY(va_ocl_basic SHARED ${va_ocl_basic_sources})

> 

>  TARGET_LINK_LIBRARIES(va_ocl_basic cl m va va-x11 ${X11_X11_LIB})

> 

> +IF(LIBVA_BUF_SH_DEP)

> +ADD_DEFINITIONS(-

> DINPUT_NV12_DEFAULT="${CMAKE_CURRENT_SOURCE_DIR}/libva

> +_buffer_sharing/256_128.nv12")

>  ADD_EXECUTABLE(example-

> libva_buffer_sharing ./libva_buffer_sharing/libva_buffer_sharing.cpp)

>  TARGET_LINK_LIBRARIES(example-libva_buffer_sharing va_ocl_basic)

> +ENDIF(LIBVA_BUF_SH_DEP)

> +

> +IF(V4L2_BUF_SH_DEP)

> +ADD_EXECUTABLE(example-v4l2_buffer_sharing

> +./v4l2_buffer_sharing/v4l2_buffer_sharing.cpp)

> +TARGET_LINK_LIBRARIES(example-v4l2_buffer_sharing va_ocl_basic)

> +ENDIF(V4L2_BUF_SH_DEP)

> +ENDIF(LIBVA_BUF_SH_DEP OR V4L2_BUF_SH_DEP)

> diff --git a/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp

> b/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp

> new file mode 100644

> index 0000000..42ab642

> --- /dev/null

> +++ b/examples/v4l2_buffer_sharing/v4l2_buffer_sharing.cpp

> @@ -0,0 +1,590 @@

> +/*

> + ** Copyright (c) 2012, 2015 Intel Corporation. All Rights Reserved.

> + **

> + ** Permission is hereby granted, free of charge, to any person

> +obtaining a

> + ** copy of this software and associated documentation files (the

> + ** "Software"), to deal in the Software without restriction, including

> + ** without limitation the rights to use, copy, modify, merge, publish,

> + ** distribute, sub license, and/or sell copies of the Software, and to

> + ** permit persons to whom the Software is furnished to do so, subject

> +to

> + ** the following conditions:

> + **

> + ** The above copyright notice and this permission notice (including

> +the

> + ** next paragraph) shall be included in all copies or substantial

> +portions

> + ** of the Software.

> + **

> + ** THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY

> KIND,

> +EXPRESS

> + ** OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF

> + ** MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NON-

> INFRINGEMENT.

> + ** IN NO EVENT SHALL PRECISION INSIGHT AND/OR ITS SUPPLIERS BE

> LIABLE

> +FOR

> + ** ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF

> +CONTRACT,

> + ** TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION

> WITH THE

> + ** SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.

> + **/

> +

> +#include <getopt.h>

> +#include <errno.h>

> +#include <assert.h>

> +#include <fcntl.h>

> +#include <linux/videodev2.h>

> +#include <poll.h>

> +#include <stdio.h>

> +#include <stdlib.h>

> +#include <stdint.h>

> +#include <string.h>

> +#include <sys/ioctl.h>

> +#include <sys/mman.h>

> +#include <sys/stat.h>

> +#include <sys/types.h>

> +#include <unistd.h>

> +#include <sys/time.h>

> +#include <time.h>

> +

> +#include <inttypes.h>

> +#include <ctype.h>

> +

> +#include <va/va.h>

> +#include <va/va_drmcommon.h>

> +

> +#include "va_display.h"

> +#include "utest_helper.hpp"

> +

> +using namespace std;

> +

> +#define BUFFER_NUM_DEFAULT 5

> +#define VIDEO_NODE_DEFAULT "/dev/video0"

> +#define WIDTH_DEFAULT 640

> +#define HEIGHT_DEFAULT 480

> +

> +#define CHECK_VASTATUS(va_status,func)                                  \

> +  if (va_status != VA_STATUS_SUCCESS) {                                   \

> +    fprintf(stderr, "status = %d, %s: %s(line %d) failed, exit\n",va_status,

> __func__, func, __LINE__); \

> +    exit(1);                                                            \

> +  }

> +

> +#define CHECK_CLSTATUS(status,func)                                  \

> +  if (status != CL_SUCCESS) {                                   \

> +    fprintf(stderr, "status = %d, %s: %s(line %d) failed, exit\n", status,

> __func__, func, __LINE__); \

> +    exit(1);                                                            \

> +  }

> +

> +#define CHECK_V4L2ERROR(ret, STR)                               \

> +  if (ret){                             \

> +    fprintf(stderr, STR);            \

> +    perror(" ");                            \

> +    fprintf(stderr, "ret = %d, %s: %s(line %d) failed, exit\n", ret, __func__,

> STR, __LINE__);      \

> +    exit(1);                                  \

> +  }

> +

> +VADisplay	va_dpy;

> +cl_int cl_status;

> +VAStatus va_status;

> +VASurfaceID nv12_surface_id;

> +VAImage nv12_image;

> +

> +int dev_fd;

> +uint64_t image_size;

> +unsigned int pitch;

> +cl_mem *import_buf = NULL;

> +typedef cl_int (OCLGETMEMOBJECTFD)(cl_context, cl_mem, int *);

> +OCLGETMEMOBJECTFD *oclGetMemObjectFd = NULL;

> +

> +int frame_count = 0;

> +struct v4l2_options{

> +  const char *dev_name;

> +  unsigned int width, height;

> +  unsigned int spec_res;

> +  unsigned int buffer_num;

> +  unsigned int do_list;

> +} vo;

> +int *import_buf_fd = NULL;

> +

> +static const char short_options[] = "d:r:b:lh";

> +

> +static const struct option

> +long_options[] = {

> +  { "device", required_argument, NULL, 'd' },

> +  { "help",   no_argument,       NULL, 'h' },

> +  { "resolution", required_argument,       NULL, 'r' },

> +  { "buffer_num",  required_argument, NULL, 'b' },

> +  { "list",  no_argument, NULL, 'l' },

> +  { 0, 0, 0, 0 }

> +};

> +

> +static void usage(FILE *fp, int argc, char **argv) {

> +  fprintf(fp,

> +      "This example aims to demostrate the usage of DMABUF buffer sharing

> between v4l2 and Beignet.\n"

> +      "For more details, please read docs/howto/v4l2-buffer-sharing-

> howto.mdwn.\n"

> +      "Usage: %s [options]\n\n"

> +      "Options:\n"

> +      "-d | --device=<dev>  Specify device by <dev> instead of

> /dev/video0\n"

> +      "-h | --help          Print this message\n"

> +      "-r | --resolution=<width,height>    Set image resolution\n"

> +      "-b | --buffer_num=<num>  Set number of buffers\n"

> +      "-l | --list  List available resolution of format 'V4L2_PIX_FMT_YUYV'\n"

> +      "",

> +      argv[0]);

> +}

> +

> +static void list_resolution(){

> +  int ret;

> +  struct v4l2_capability cap;

> +  struct v4l2_frmsizeenum frm_sz;

> +

> +  dev_fd = open(vo.dev_name, O_RDWR | O_NONBLOCK, 0);  if (dev_fd < 0)

> + {

> +    fprintf(stderr, "Can not open %s: %s\n",

> +        vo.dev_name, strerror(errno));

> +    exit(1);

> +  }

> +

> +  memset(&cap, 0, sizeof(cap));

> +  ret = ioctl(dev_fd, VIDIOC_QUERYCAP, &cap);  CHECK_V4L2ERROR(ret,

> + "VIDIOC_QUERYCAP");

> +

> +  if(!(cap.capabilities & V4L2_CAP_VIDEO_CAPTURE)){

> +    fprintf(stderr, "The device is not video capture device\n");

> +    exit(1);

> +  }

> +  if(!(cap.capabilities & V4L2_CAP_STREAMING)){

> +    fprintf(stderr, "The device does not support streaming i/o\n");

> +    exit(1);

> +  }

> +

> +  printf("Supported resolution under pixel format

> + 'V4L2_PIX_FMT_YUYV':\n");  frm_sz.pixel_format = V4L2_PIX_FMT_YUYV;

> + frm_sz.index = 0;  bool extra_info = true;  while (ioctl(dev_fd,

> + VIDIOC_ENUM_FRAMESIZES, &frm_sz) == 0) {

> +    if (frm_sz.type == V4L2_FRMSIZE_TYPE_DISCRETE) {

> +      if(extra_info){

> +        printf("(width, height) = \n");

> +        extra_info = false;

> +      }

> +      printf("(%d, %d)", frm_sz.discrete.width, frm_sz.discrete.height);

> +      printf("\n");

> +    }

> +    else if (frm_sz.type == V4L2_FRMSIZE_TYPE_STEPWISE) {

> +      printf("(width, height) from (%d, %d) to (%d, %d) with step (%d, %d)",

> +          frm_sz.stepwise.min_width,

> +          frm_sz.stepwise.min_height,

> +          frm_sz.stepwise.max_width,

> +          frm_sz.stepwise.max_height,

> +          frm_sz.stepwise.step_width,

> +          frm_sz.stepwise.step_height);

> +      continue;

> +    }

> +    frm_sz.index++;

> +  }

> +

> +  ret = close(dev_fd);

> +  if (ret) {

> +    fprintf(stderr, "Failed to close %s: %s\n",

> +        vo.dev_name, strerror(errno));

> +    exit(1);

> +  }

> +}

> +

> +static void analyse_args(int argc, char *argv[]) {

> +  vo.dev_name = NULL;

> +  vo.width = 0;

> +  vo.height = 0;

> +  vo.spec_res = 0;

> +  vo.buffer_num = BUFFER_NUM_DEFAULT;

> +  vo.do_list = 0;

> +

> +  int c, idx;

> +  for (;;) {

> +

> +    c = getopt_long(argc, argv,

> +        short_options, long_options, &idx);

> +

> +    if (-1 == c)

> +      break;

> +

> +    switch (c) {

> +      case 0:

> +        break;

> +

> +      case 'd':

> +        vo.dev_name = optarg;

> +        break;

> +

> +      case '?':

> +      case 'h':

> +        usage(stdout, argc, argv);

> +        exit(0);

> +

> +      case 'r':

> +        sscanf(optarg, "%d,%d", &vo.width, &vo.height);

> +        vo.spec_res = 1;

> +        break;

> +

> +      case 'b':

> +        vo.buffer_num = strtoul(optarg, NULL, 0);

> +        break;

> +

> +      case 'l':

> +        vo.do_list = 1;

> +        break;

> +

> +      default:

> +        usage(stderr, argc, argv);

> +        exit(1);

> +    }

> +  }

> +

> +  if(!vo.dev_name){

> +    printf("Haven't specified device, use default device: %s\n",

> +        VIDEO_NODE_DEFAULT);

> +  }

> +  if(!vo.dev_name)

> +    vo.dev_name = VIDEO_NODE_DEFAULT;

> +  if(vo.do_list){

> +    list_resolution();

> +    exit(0);

> +  }

> +  if(!vo.spec_res){

> +    printf("Haven't specified resolution, use default resolution: (width,height)

> = (%d, %d)\n",

> +        WIDTH_DEFAULT, HEIGHT_DEFAULT);

> +    vo.width = WIDTH_DEFAULT;

> +    vo.height = HEIGHT_DEFAULT;

> +  }

> +  return;

> +}

> +

> +static void initialize_va_ocl(){

> +  int major_ver, minor_ver;

> +

> +  printf("\n***********************libva info:

> + ***********************\n");  fflush(stdout);  va_dpy =

> + va_open_display();  va_status = vaInitialize(va_dpy, &major_ver,

> + &minor_ver);  CHECK_VASTATUS(va_status, "vaInitialize");

> +

> +  VASurfaceAttrib forcc;

> +  forcc.type =VASurfaceAttribPixelFormat;

> + forcc.flags=VA_SURFACE_ATTRIB_SETTABLE;

> +  forcc.value.type=VAGenericValueTypeInteger;

> +  forcc.value.value.i = VA_FOURCC_NV12;  va_status =

> + vaCreateSurfaces(va_dpy, VA_RT_FORMAT_YUV420,

> +                               vo.width, vo.height,

> +                               &nv12_surface_id, 1, &forcc, 1);

> + CHECK_VASTATUS(va_status, "vaCreateSurfaces");

> +

> +  VAImageFormat image_fmt;

> +  image_fmt.fourcc = VA_FOURCC_NV12;

> +  image_fmt.byte_order = VA_LSB_FIRST;

> +  image_fmt.bits_per_pixel = 12;

> +  va_status = vaCreateImage(va_dpy, &image_fmt, vo.width, vo.height,

> + &nv12_image);  CHECK_VASTATUS(va_status, "vaCreateImage");

> +

> +  //ocl initialization: basic & create kernel & get extension

> + printf("\n***********************OpenCL info:

> + ***********************\n");  if ((cl_status =

> cl_test_init("runtime_yuy2_processing.cl", "runtime_yuy2_processing",

> SOURCE)) != 0){

> +    fprintf(stderr, "cl_test_init error\n");

> +    exit(1);

> +  }

> +

> +#ifdef CL_VERSION_1_2

> +  oclGetMemObjectFd = (OCLGETMEMOBJECTFD

> +*)clGetExtensionFunctionAddressForPlatform(platform,

> +"clGetMemObjectFdIntel"); #else

> +  oclGetMemObjectFd = (OCLGETMEMOBJECTFD

> +*)clGetExtensionFunctionAddress("clGetMemObjectFdIntel");

> +#endif

> +  if(!oclGetMemObjectFd){

> +    fprintf(stderr, "Failed to get extension clGetMemObjectFdIntel\n");

> +    exit(1);

> +  }

> +

> +printf("\n**************************************************

> *********\n

> +");

> +}

> +

> +static void create_dmasharing_buffers() {

> +  if(import_buf_fd == NULL)

> +    import_buf_fd = (int *)malloc(sizeof(int) * vo.buffer_num);

> +  if(import_buf == NULL){

> +    import_buf = (cl_mem *)malloc(sizeof(cl_mem) * vo.buffer_num);

> +  }

> +

> +  for (unsigned int i = 0; i < vo.buffer_num; ++i){

> +    import_buf[i] = clCreateBuffer(ctx, CL_MEM_READ_WRITE, image_size,

> NULL, &cl_status);

> +    CHECK_CLSTATUS(cl_status, "clCreateBuffer");

> +

> +    //get cl buffer object's fd

> +    cl_status = oclGetMemObjectFd(ctx, import_buf[i], &import_buf_fd[i]);

> +    CHECK_CLSTATUS(cl_status, "clGetMemObjectFdIntel");

> +  }

> +}

> +

> +static void release_va_ocl(){

> +  va_status = vaDestroySurfaces(va_dpy,&nv12_surface_id,1);

> +  CHECK_VASTATUS(va_status, "vaDestroySurfaces");

> +  va_status = vaDestroyImage(va_dpy, nv12_image.image_id);

> +  CHECK_VASTATUS(va_status, "vaDestroyImage");

> +  va_status = vaTerminate(va_dpy);

> +  CHECK_VASTATUS(va_status, "vaTerminate");

> +  va_close_display(va_dpy);

> +

> +  int ret;

> +  for (unsigned int i = 0; i < vo.buffer_num; ++i) {

> +    ret = close(import_buf_fd[i]);

> +    if (ret) {

> +      fprintf(stderr, "Failed to close import_buf[%u]'s fd: %s\n", i,

> strerror(errno));

> +    }

> +    cl_status = clReleaseMemObject(import_buf[i]);

> +    CHECK_CLSTATUS(cl_status, "clReleaseMemObject");

> +  }

> +}

> +

> +static void process_show_frame(int index) {

> +  //process import_buf[index] by ocl

> +  size_t global_size[2];

> +  global_size[0] = vo.width * 2 / 4;

> +  global_size[1] = vo.height;

> +  cl_status = clSetKernelArg(kernel, 0, sizeof(cl_mem),

> +&import_buf[index]);

> +  CHECK_CLSTATUS(cl_status, "clSetKernelArg");

> +  cl_status = clSetKernelArg(kernel, 1, sizeof(int), &vo.height);

> +  CHECK_CLSTATUS(cl_status, "clSetKernelArg");

> +  cl_status = clSetKernelArg(kernel, 2, sizeof(int), &pitch);

> +  CHECK_CLSTATUS(cl_status, "clSetKernelArg");

> +  cl_status = clEnqueueNDRangeKernel(queue, kernel, 2, NULL,

> +                                     global_size, NULL, 0, NULL, NULL);

> +  CHECK_CLSTATUS(cl_status, "clEnqueueNDRangeKernel");

> +  cl_status = clFinish(queue);

> +  CHECK_CLSTATUS(cl_status, "clFinish");

> +

> +  //create corresponding VASurface

> +  VASurfaceID yuy2_surface_id;

> +  VASurfaceAttrib sa[2];

> +  sa[0].type = VASurfaceAttribMemoryType;  sa[0].flags =

> + VA_SURFACE_ATTRIB_SETTABLE;  sa[0].value.type =

> + VAGenericValueTypeInteger;  sa[0].value.value.i =

> + VA_SURFACE_ATTRIB_MEM_TYPE_DRM_PRIME;

> +  sa[1].type = VASurfaceAttribExternalBufferDescriptor;

> +  sa[1].flags = VA_SURFACE_ATTRIB_SETTABLE;  sa[1].value.type =

> + VAGenericValueTypePointer;  VASurfaceAttribExternalBuffers sa_eb;

> + sa_eb.pixel_format = VA_FOURCC_YUY2;  sa_eb.width = vo.width;

> + sa_eb.height = vo.height;  sa_eb.data_size = image_size;

> + sa_eb.num_planes = 1;  sa_eb.pitches[0] = pitch;  sa_eb.offsets[0] =

> + 0;  sa_eb.num_buffers = 1;  sa_eb.buffers = (unsigned long

> + *)malloc(sizeof(unsigned long) * sa_eb.num_buffers);  sa_eb.buffers[0]

> + = import_buf_fd[index];  sa_eb.flags = 0;  sa[1].value.value.p =

> + &sa_eb;  va_status = vaCreateSurfaces(va_dpy, VA_RT_FORMAT_YUV422,

> +                               vo.width, vo.height,

> +                               &yuy2_surface_id, 1, sa, 2);

> + CHECK_VASTATUS(va_status, "vaCreateSurfaces");

> +

> +  //convert to NV12 format

> +  va_status = vaGetImage (va_dpy, yuy2_surface_id, 0, 0,

> +                          vo.width, vo.height, nv12_image.image_id);

> + CHECK_VASTATUS(va_status, "vaGetImage");  va_status =

> + vaPutImage(va_dpy, nv12_surface_id, nv12_image.image_id,

> +                         0, 0, vo.width, vo.height, 0, 0,

> +                         vo.width, vo.height);

> + CHECK_VASTATUS(va_status, "vaPutImage");

> +

> +  //show by vaPutsurface

> +  VARectangle src_rect, dst_rect;

> +  src_rect.x      = 0;

> +  src_rect.y      = 0;

> +  src_rect.width  = vo.width;

> +  src_rect.height = vo.height;

> +  dst_rect        = src_rect;

> +  va_status = va_put_surface(va_dpy, nv12_surface_id, &src_rect,

> + &dst_rect);  CHECK_VASTATUS(va_status, "vaPutSurface");

> +

> +  vaDestroySurfaces(va_dpy,&yuy2_surface_id,1);

> +  CHECK_VASTATUS(va_status, "vaDestroySurfaces");

> +  free(sa_eb.buffers);

> +  return;

> +}

> +

> +static void init_dmabuf(void){

> +  int ret;

> +  struct v4l2_requestbuffers reqbuf;

> +

> +  memset(&reqbuf, 0, sizeof(reqbuf));

> +  reqbuf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;  reqbuf.memory =

> + V4L2_MEMORY_DMABUF;  reqbuf.count = vo.buffer_num;

> +

> +  ret = ioctl(dev_fd, VIDIOC_REQBUFS, &reqbuf);  if(ret == -1 && errno

> + == EINVAL){

> +    fprintf(stderr, "Video capturing or DMABUF streaming is not

> supported\n");

> +    exit(1);

> +  }

> +  else

> +    CHECK_V4L2ERROR(ret, "VIDIOC_REQBUFS");

> +

> +  create_dmasharing_buffers();

> +  printf("Succeed to create %d dma buffers \n", vo.buffer_num);

> +

> +}

> +

> +static void init_device(void){

> +

> +  int ret;

> +  struct v4l2_capability cap;

> +  struct v4l2_format format;

> +

> +  dev_fd = open(vo.dev_name, O_RDWR | O_NONBLOCK, 0);  if (dev_fd < 0)

> + {

> +    fprintf(stderr, "Can not open %s: %s\n",

> +        vo.dev_name, strerror(errno));

> +    exit(1);

> +  }

> +

> +  memset(&cap, 0, sizeof(cap));

> +  ret = ioctl(dev_fd, VIDIOC_QUERYCAP, &cap);  CHECK_V4L2ERROR(ret,

> + "VIDIOC_QUERYCAP");  if(!(cap.capabilities & V4L2_CAP_STREAMING)){

> +    fprintf(stderr, "The device does not support streaming i/o\n");

> +    exit(1);

> +  }

> +

> +  memset(&format, 0, sizeof(format));

> +  format.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;  format.fmt.pix.width =

> + vo.width;  format.fmt.pix.height = vo.height;

> + format.fmt.pix.pixelformat = V4L2_PIX_FMT_YUYV;  format.fmt.pix.field

> + = V4L2_FIELD_ANY;

> +

> +  ret = ioctl(dev_fd, VIDIOC_S_FMT, &format);  CHECK_V4L2ERROR(ret,

> + "VIDIOC_S_FMT");

> +

> +  ret = ioctl(dev_fd, VIDIOC_G_FMT, &format);

> +  CHECK_V4L2ERROR(ret, "VIDIOC_G_FMT");

> +  if(format.fmt.pix.pixelformat != V4L2_PIX_FMT_YUYV){

> +    fprintf(stderr, "V4L2_PIX_FMT_YUYV format is not supported by %s\n",

> vo.dev_name);

> +    exit(1);

> +  }

> +  if(format.fmt.pix.width != vo.width  || format.fmt.pix.height !=

> vo.height){

> +    fprintf(stderr, "This resolution is not supported, please go through

> supported resolution by command './main -l'\n");

> +    exit(1);

> +  }

> +  printf("Input image format: (width, height) = (%u, %u), pixel format

> = %.4s\n",

> +      format.fmt.pix.width, format.fmt.pix.height,

> +(char*)&format.fmt.pix.pixelformat);

> +  image_size = format.fmt.pix.sizeimage;

> +	pitch = format.fmt.pix.bytesperline;

> +}

> +

> +static void start_capturing(void){

> +  int ret;

> +  for (unsigned int i = 0; i < vo.buffer_num; ++i) {

> +    struct v4l2_buffer buf;

> +

> +    memset(&buf, 0, sizeof(buf));

> +    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;

> +    buf.memory = V4L2_MEMORY_DMABUF;

> +    buf.index = i;

> +    buf.m.fd = import_buf_fd[i];

> +    ret = ioctl(dev_fd, VIDIOC_QBUF, &buf);

> +    CHECK_V4L2ERROR(ret, "VIDIOC_QBUF");  }

> +

> +  int type = V4L2_BUF_TYPE_VIDEO_CAPTURE;

> +  ret = ioctl(dev_fd, VIDIOC_STREAMON, &type);

> +  CHECK_V4L2ERROR(ret, "VIDIOC_STREAMON"); }

> +

> +static void mainloop(void){

> +  int ret;

> +  struct v4l2_buffer buf;

> +  int index;

> +

> +  while (1) {

> +    frame_count++;

> +    printf("******************Frame %d\n", frame_count);

> +    fd_set fds;

> +    struct timeval tv;

> +    int r;

> +

> +    FD_ZERO(&fds);

> +    FD_SET(dev_fd, &fds);

> +

> +    /* Timeout. */

> +    tv.tv_sec = 2;

> +    tv.tv_usec = 0;

> +

> +

> +    r = select(dev_fd + 1, &fds, NULL, NULL, &tv);

> +

> +    if (-1 == r) {

> +      if (EINTR == errno)

> +        continue;

> +      perror("select");

> +    }

> +

> +    if(r == 0){

> +      fprintf(stderr, "Select timeout\n");

> +      exit(1);

> +    }

> +

> +    memset(&buf, 0, sizeof(buf));

> +    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;

> +    buf.memory = V4L2_MEMORY_DMABUF;

> +    ret = ioctl(dev_fd, VIDIOC_DQBUF, &buf);

> +    CHECK_V4L2ERROR(ret, "VIDIOC_DQBUF");

> +    index = buf.index;

> +

> +    //process by ocl and show on screen by libva

> +    process_show_frame(index);

> +

> +    //Then queue this buffer(buf.index) by QBUF

> +    buf.type = V4L2_BUF_TYPE_VIDEO_CAPTURE;

> +    buf.memory = V4L2_MEMORY_DMABUF;

> +    buf.m.fd = import_buf_fd[index];

> +    buf.index = index;

> +

> +    ret = ioctl(dev_fd, VIDIOC_QBUF, &buf);

> +    CHECK_V4L2ERROR(ret, "VIDIOC_QBUF");

> +  }

> +}

> +

> +static void stop_capturing(void)

> +{

> +  int ret;

> +  int type = V4L2_BUF_TYPE_VIDEO_CAPTURE;

> +

> +  ret = ioctl(dev_fd, VIDIOC_STREAMOFF, &type);

> +  CHECK_V4L2ERROR(ret, "VIDIOC_STREAMOFF"); }

> +

> +static void uninit_device(void){

> +  free(import_buf_fd);

> +  free(import_buf);

> +  int ret = close(dev_fd);

> +  if (ret) {

> +    fprintf(stderr, "Failed to close %s: %s\n",

> +        vo.dev_name, strerror(errno));

> +    exit(1);

> +  }

> +}

> +

> +int main(int argc, char *argv[])

> +{

> +  analyse_args(argc, argv);

> +

> +  init_device();

> +  initialize_va_ocl();

> +  init_dmabuf();

> +

> +  start_capturing();

> +  mainloop();

> +

> +  stop_capturing();

> +  release_va_ocl();

> +  uninit_device();

> +

> +  return 0;

> +}

> diff --git a/kernels/runtime_yuy2_processing.cl

> b/kernels/runtime_yuy2_processing.cl

> new file mode 100644

> index 0000000..1478e65

> --- /dev/null

> +++ b/kernels/runtime_yuy2_processing.cl

> @@ -0,0 +1,15 @@

> +__kernel void

> +runtime_yuy2_processing(__global uchar *src,

> +                        int image_height,

> +                        int image_pitch) {

> +  int gx = get_global_id(0);

> +  int gy = get_global_id(1);

> +

> +  int src_y = image_height / 2 + gy;

> +  int mirror_y = image_height - src_y;

> +

> +  uchar4 mirror_val = *(__global uchar4*)(src + mirror_y*image_pitch +

> + gx*4);  *(__global uchar4*)(src + src_y*image_pitch + gx*4) =

> + mirror_val;

> +

> +}

> --

> 1.9.1

> 

> _______________________________________________

> Beignet mailing list

> Beignet@lists.freedesktop.org

> http://lists.freedesktop.org/mailman/listinfo/beignet