[Beignet] [PATCH v3 1/2] Add example to show v4l2 buffer sharing with extension clGetMemObjectFdIntel.
Chuanbo Weng
chuanbo.weng at intel.com
Tue Apr 7 23:50:47 PDT 2015
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.
Signed-off-by: Chuanbo Weng <chuanbo.weng at 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 904f259..ab31fe7 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)
EXEC_PROGRAM(ls ARGS "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva" OUTPUT_VARIABLE LS_OUTPUT)
IF(NOT LS_OUTPUT)
EXEC_PROGRAM(git "${CMAKE_CURRENT_SOURCE_DIR}/.." ARGS "submodule init")
@@ -5,17 +11,13 @@ EXEC_PROGRAM(git "${CMAKE_CURRENT_SOURCE_DIR}/.." ARGS "submodule update")
EXEC_PROGRAM(git "${CMAKE_CURRENT_SOURCE_DIR}/thirdparty/libva" ARGS "checkout master")
ENDIF(NOT LS_OUTPUT)
-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..60fc61c
--- /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
More information about the Beignet
mailing list