[Beignet] [PATCH v4 1/2] Add example to show v4l2 buffer sharing with extension clGetMemObjectFdIntel.
Yang, Rong R
rong.r.yang at intel.com
Thu Jun 18 19:30:24 PDT 2015
Pushed, thanks.
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Chuanbo Weng
> Sent: Thursday, June 18, 2015 16:30
> To: beignet at 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 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
> 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 at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list