[FFmpeg-devel] [PATCH] libavutil/libavfilter:add opencl warpper and opencl deshake filter
Stefano Sabatini
stefasab at gmail.com
Fri Mar 8 18:13:44 CET 2013
On date Friday 2013-03-08 10:29:49 +0800, Wei Gao encoded:
> Hi,
>
> Thank you for your comments, I have some questions as follow.
>
> Thanks
> Best regards
>
> 2013/3/8 Stefano Sabatini <stefasab at gmail.com>
>
> > Sorry again for the slow reply.
> >
> > On date Friday 2013-02-22 15:31:09 +0800, Wei Gao encoded:
> > >
> > > From 4824a3cb8108afe95756f46cfe5a185d3a6308dd Mon Sep 17 00:00:00 2001
> > > From: highgod0401 <highgod0401 at gmail.com>
> > > Date: Fri, 22 Feb 2013 15:25:48 +0800
> > > Subject: [PATCH] add opencl warpper and opencl deshake filter
> > >
> > > ---
> > > configure | 5 +
> > > libavfilter/Makefile | 3 +
> > > libavfilter/all_filter_kernels.c | 18 +
> > > libavfilter/all_filter_kernels.h | 9 +
> > > libavfilter/allfilters.c | 8 +-
> > > libavfilter/deshake_kernel.h | 182 +++++++++
> > > libavfilter/transform_opencl.c | 157 ++++++++
> > > libavfilter/transform_opencl.h | 38 ++
> > > libavfilter/vf_deshake.c | 213 ++++++++++-
> >
> > > libavutil/Makefile | 4 +
> > > libavutil/openclwrapper.c | 808
> > +++++++++++++++++++++++++++++++++++++++
> > > libavutil/openclwrapper.h | 198 ++++++++++
[...]
> > > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
> > > index 47158f9..e92cf8b 100644
> > > --- a/libavfilter/allfilters.c
> > > +++ b/libavfilter/allfilters.c
> > > @@ -21,7 +21,9 @@
> > >
> > > #include "avfilter.h"
> > > #include "config.h"
> > > -
> > > +#if CONFIG_OPENCL
> > > +#include "all_filter_kernels.h"
> > > +#endif
> >
> > avoid indirections, just insert the code here.
> >
> The function " ff_opencl_regist_kernel" is defined in openclwrapper, and
> only can be compiled if config CONFIG_OPENCL,so should I code here and
> still use indirections?
Yes.
[...]
> > > diff --git a/libavfilter/deshake_kernel.h b/libavfilter/deshake_kernel.h
> > > new file mode 100644
[...]
> > > +
> > > +kernel void avfilter_transform(global unsigned char *src,
> > > + global unsigned char *dst,
> > > + global float *matrix,
> > > + global float *matrix2,
> >
> > > + int interpolate,
> > > + int fillmethod,
> >
> > you may use some enum here, they greatly help debugging and
> > readability
> >
> Sorry, I don't get what you mean, it is a OpenCL kernel function,and each
> parameter is set from c code using AV_OPENCLCHECK_SET_KERNEL_ARG, and they
> should be set one by one.So I think it may be easy to write like this.
OK if you can't import the enum definitions from transform.h.
[...]
> > > diff --git a/libavfilter/transform_opencl.c
> > b/libavfilter/transform_opencl.c
> > > new file mode 100644
> > > index 0000000..4acec2b
> > > --- /dev/null
> > > +++ b/libavfilter/transform_opencl.c
> > > @@ -0,0 +1,157 @@
> > > +/*
> > > + * Copyright (C) 2013 Wei Gao <weigao at multicorewareinc.com>
> > > + *
> > > + *
> >
> > Nit: an empty line is enough
> >
> > > + * This file is part of FFmpeg.
> > > + *
> > > + * FFmpeg is free software; you can redistribute it and/or
> > > + * modify it under the terms of the GNU Lesser General Public
> > > + * License as published by the Free Software Foundation; either
> > > + * version 2.1 of the License, or (at your option) any later version.
> > > + *
> > > + * FFmpeg is distributed in the hope that it will be useful,
> > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > > + * Lesser General Public License for more details.
> > > + *
> > > + * You should have received a copy of the GNU Lesser General Public
> > > + * License along with FFmpeg; if not, write to the Free Software
> > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> > 02110-1301 USA
> > > + */
> > > +
> > > +/**
> > > + * @file
> > > + * transform input video
> > > + */
> > > +
> > > +#include "libavutil/common.h"
> > > +#include "libavutil/avassert.h"
> > > +#include "libavutil/avstring.h"
> > > +#include "libavutil/openclwrapper.h"
> > > +#include "transform.h"
> > > +#include "transform_opencl.h"
> > > +
> > > +
> > > +
> > > +static int ff_filter_transform_func(void **userdata, AV_KernelEnv *kenv)
> > > +{
> > > + cl_mem src = (cl_mem)userdata[0];
> > > + cl_mem dst = (cl_mem)userdata[1];
> > > + int src_stride_lu = (int)userdata[2];
> > > + int dst_stride_lu = (int)userdata[3];
> > > + int src_stride_ch = (int)userdata[4];
> > > + int dst_stride_ch = (int)userdata[5];
> > > + int width = (int)userdata[6];
> > > + int height = (int)userdata[7];
> > > + int cw = (int)userdata[8];
> > > + int ch = (int)userdata[9];
> > > + float *matrix = (float *)userdata[10];
> > > + float *matrix2 = (float *)userdata[11];
> > > + int interpolate = (int)userdata[12];
> > > + int fillmethod = (int)userdata[13];
> > > + cl_mem matrix_buf = (cl_mem)userdata[14];
> > > + cl_mem matrix_buf2 = (cl_mem)userdata[15];
> > > + AV_KernelEnv *env = (AV_KernelEnv *)userdata[16];
> > > + cl_uint status;
> > > + void *mapped;
> > > + const size_t global_work_size = width * height + 2 * ch * cw;
> > > + int m_size = 6;
> > > + cl_kernel kernel;
> > > + int arg_no;
> > > +
> > > +
> > > + mapped = clEnqueueMapBuffer(kenv->command_queue, matrix_buf,
> > CL_TRUE, CL_MAP_WRITE, 0, m_size*sizeof(cl_float),0,NULL, NULL, NULL);
> > > + memcpy(mapped,matrix,m_size*sizeof(cl_float));
> > > + clEnqueueUnmapMemObject(kenv->command_queue, matrix_buf, mapped, 0,
> > NULL, NULL);
> > > +
> > > + mapped = clEnqueueMapBuffer(kenv->command_queue, matrix_buf2,
> > CL_TRUE, CL_MAP_WRITE, 0, m_size*sizeof(cl_float),0,NULL, NULL, NULL);
> > > + memcpy(mapped,matrix2,m_size*sizeof(cl_float));
> > > + clEnqueueUnmapMemObject(kenv->command_queue, matrix_buf2, mapped,
> > 0, NULL, NULL);
> > > +
> > > + if (!env->kernel) {
> > > + status = av_opencl_create_kernel("avfilter_transform", kenv);
> > > + if (status) {
> > > + av_log(NULL,AV_LOG_ERROR,"clCreateKernel Error
> > %s\n","avfilter_transform");
> > > + return 0;
> > > + }
> > > + env->command_queue = kenv->command_queue;
> > > + env->context = kenv->context;
> > > + env->kernel = kenv->kernel;
> > > + av_strlcpy(env->kernel_name,kenv->kernel_name,150);
> > > + env->program = kenv->program;
> > > + }
> > > + kernel = env->kernel;
> > > + arg_no = 0;
> >
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&src);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&dst);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&matrix_buf);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&matrix_buf2);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&interpolate);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&fillmethod);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&src_stride_lu);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&dst_stride_lu);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&src_stride_ch);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&dst_stride_ch);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&height);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&width);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&ch);
> > > + AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&cw);
> > > +
> > > + AV_OPENCLCHECK( clEnqueueNDRangeKernel, env->command_queue,
> > env->kernel, 1, NULL,
> > > + &global_work_size, NULL, 0, NULL, NULL);
> >
> > since this is the only user, move the macro definition locally to this
> > file (rule: avoid indirections unless strictly necessary).
> >
>
> This is because we only submit deshake filter recently,other OpenCL
> filters we will submit later will use the macro
> AV_OPENCLCHECK_SET_KERNEL_ARG and AV_OPENCLCHECK to easy the code write.So
> we add it in opencl wraper.If not, every OpenCL filter will add the two
> macro.
OK, but please use names:
AV_OPENCL_SET_KERNEL_ARG
AV_OPENCL_CHECK
also you could probably move sizeof() within the macro.
[...]
> > > diff --git a/libavfilter/transform_opencl.h
> > b/libavfilter/transform_opencl.h
> > > new file mode 100644
> > > index 0000000..aa1b4e0
> > > --- /dev/null
> > > +++ b/libavfilter/transform_opencl.h
> > > @@ -0,0 +1,38 @@
> > > +/*
> > > + * Copyright (C) 2013 Wei Gao <weigao at multicorewareinc.com>
> > > + *
> > > + * This file is part of FFmpeg.
> > > + *
> > > + * FFmpeg is free software; you can redistribute it and/or
> > > + * modify it under the terms of the GNU Lesser General Public
> > > + * License as published by the Free Software Foundation; either
> > > + * version 2.1 of the License, or (at your option) any later version.
> > > + *
> > > + * FFmpeg is distributed in the hope that it will be useful,
> > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > > + * Lesser General Public License for more details.
> > > + *
> > > + * You should have received a copy of the GNU Lesser General Public
> > > + * License along with FFmpeg; if not, write to the Free Software
> > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> > 02110-1301 USA
> > > + */
> > > +
> > > +#ifndef AVFILTER_TRANSFORM_OPENCL_H
> > > +#define AVFILTER_TRANSFORM_OPENCL_H
> > > +
> > > +#include <stdint.h>
> > > +
> > > +void avfilter_transform_cl( void *src, void *dst,
> > > + int src_stride_lu, int dst_stride_lu,
> > > + int src_stride_ch, int dst_stride_ch,
> > > + int width, int height, int cw, int ch,
> > > + const float *matrix, const float *matrix2,
> > > + const void *matrix_cl, const void *matrix2_cl,
> >
> > > + enum InterpolateMethod interpolate,
> > > + enum FillMethod fill, AV_KernelEnv *env);
> >
> > where are these enums defined?
> >
> They are defined in "transform.h", I move the header file here.
>
> >
> > > +int ff_init_transform(void);
> >
> > Is this file supposed to be public? If not you should avoid to use
> > avfilter_ prefix (reserved for public functions).
> >
> > I suggest something along the lines of:
> >
> > ff_opencl_transform()
> > ff_opencl_transform_init()
> >
> This file,I referenced "transform.h",this file will be used in
> deshake_opencl filter.
Thus it shouldn't need to be made public, and the ff_prefix should be
used.
[...]
> > > +AVFilter avfilter_vf_deshake_opencl = {
> > > + .name = "deshake_opencl",
> > > + .description = NULL_IF_CONFIG_SMALL("Stabilize shaky video using
> > OpenCL."),
> > > + .priv_size = sizeof(DeshakeContext),
> > > + .init = init_opencl,
> > > + .uninit = uninit_opencl,
> > > + .query_formats = query_formats,
> > > + .inputs = deshake_opencl_inputs,
> > > + .outputs = deshake_opencl_outputs,
> > > +};
> > > +#endif
> >
> > General question: is this filter supposed to be binary compatible with
> > deshake? Or it just shares the same syntax?
> >
>
> Sorry, I don't get what your mean.Our idea is to make the deshake_opencl
> filter can be used both in ffmpeg.exe and a filter in libavfilter.lib witch
> can be used in other application program.
Is the output created by the opencl deshake filter the same as the one
produced by deshake?
You can test using:
ffmpeg -f lavfi -i testsrc=d=10,deshake -f md5 - -nostats
ffmpeg -f lavfi -i testsrc=d=10,deshake_opencl -f md5 - -nostats
and comparing the output, or using the showinfo filter.
[...]
> > > +int av_opencl_create_kernel(const char * kernelname, AV_KernelEnv * env)
> > > +{
> > > + int status;
> > > + env->kernel = clCreateKernel(gpu_env.programs[0], kernelname,
> > &status);
> > > + env->context = gpu_env.context;
> > > + env->command_queue = gpu_env.command_queue;
> >
> > > + if (status != CL_SUCCESS) {
> > > + av_log(&openclwrapperutils,AV_LOG_ERROR,"status = %d\n",status);
> >
> > Error messages should be intelligible, reporting just a status is not
> > enough.
> >
>
> There are a lot of statuses,and you can find the status in CL.h(this file
> is the OpenCL SDK),so I think just return the status then you can find the
> error.
> Also if you think it should return the error exactly as string, I
> will write a function to get the error string.
Yes. The point is that software should do its best to explain what's
going wrong. Returning a number is not particularly useful to the
user, a string containing an error description should be already better:
Something like this should be good:
if (status != CL_SUCCESS) {
av_log(&openclwrapperutils, AV_LOG_ERROR, "Error creating OpenCL kernel: %s", opencl_errstr(status));
[...]
> > (av_opencl_read_cl_buffer(cl_inBuf,gpu_env.temp_buffer,sizeof(uint8_t)*(linesize0
> > + linesize1)*height)) {
> > > + memcpy(Ybuf,gpu_env.temp_buffer,linesize0 * height);
> > > + memcpy(Ubuf,gpu_env.temp_buffer + linesize0 * height,linesize1
> > *chrH);
> > > + memcpy(Vbuf,gpu_env.temp_buffer + linesize0 * height +
> > linesize1 * chrH,linesize2 * chrH);
> > > + }
> > > + return 1;
> > > +}
> > > +void av_opencl_save_buffer(const char *filtername,void *cl_inbuf)
> > > +{
> > > + int i = 0;
> > > + while (strlen(filter_buffer[i].filter_name)) {
> > > + i++;
> > > + }
> > > + if (i > (MAX_FILTER_NUM - 1)) {
> > > + av_log(&openclwrapperutils,AV_LOG_ERROR,"filter num is too
> > large\n");
> > > + return;
> > > + }
> > > + if(strlen(filtername) > MAX_FILTER_NAME_LEN) {
> > > + av_log(&openclwrapperutils,AV_LOG_ERROR,"filter name is too
> > long\n");
> > > + return;
> > > + }
> > > +
> > av_strlcpy(filter_buffer[i].filter_name,filtername,MAX_FILTER_NAME_LEN+1);
> > > + filter_buffer[i].cl_inbuf = cl_inbuf;
> > > + return;
> >
> > aren't you supposed to return an error in case of too long buffers?
> >
>
> This functions is used to save the OpenCL memory pointer and if the follow
> filter is OpenCL filter, it can get the OpenCL buffer as input, it can
> avoid copy data between GPU and CPU.So the "cl_inbuf" is the buffer
> per-OpenCL filter has created,this function just save the pointer.
Yes but what happens if the buffer are too short to contain the copied
data?
[...]
> > > diff --git a/libavutil/openclwrapper.h b/libavutil/openclwrapper.h
Keep the name short, opencl.h should be enough (assuming we are going
to keep this file in libavutil).
> > > new file mode 100644
> > > index 0000000..a3ac8c2
> > > --- /dev/null
> > > +++ b/libavutil/openclwrapper.h
> > > @@ -0,0 +1,198 @@
> > > +/*
> > > + * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
> > > + * Copyright (C) 2012 Li Cao <li at multicorewareinc.com>
> > > + * Copyright (C) 2012 Wei Gao <weigao at multicorewareinc.com>
> > > + *
> > > + * This file is part of FFmpeg.
> > > + *
> > > + * FFmpeg is free software; you can redistribute it and/or
> > > + * modify it under the terms of the GNU Lesser General Public
> > > + * License as published by the Free Software Foundation; either
> > > + * version 2.1 of the License, or (at your option) any later version.
> > > + *
> > > + * FFmpeg is distributed in the hope that it will be useful,
> > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > > + * Lesser General Public License for more details.
> > > + *
> > > + * You should have received a copy of the GNU Lesser General Public
> > > + * License along with FFmpeg; if not, write to the Free Software
> > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> > 02110-1301 USA
> > > + */
> > > +
> > > +#include "config.h"
> > > +
> > > +#ifndef LIBAVUTIL_OPENCLWRAPPER_H
> > > +#define LIBAVUTIL_OPENCLWRAPPER_H
> > > +
> > > +#include <CL/cl.h>
> > > +
> >
> > > +#define AV_OPENCLCHECK(method, ...)\
> > > + status = method( __VA_ARGS__ ); if( status != CL_SUCCESS ) {\
> > > + av_log(NULL,AV_LOG_ERROR, " error %s %d\n", # method, status );
> > return status; }
> > > +
> > > +#define AV_OPENCLCHECK_SET_KERNEL_ARG(arg_size,arg_ptr) AV_OPENCLCHECK(
> > clSetKernelArg,kernel,(arg_no++),(arg_size),(void*)(arg_ptr))
> >
> > avoid nested macros
> >
> > > +
> > > +#define FF_OPENCL_KERNEL( ... )# __VA_ARGS__
> > > +
> > > +typedef struct AV_KernelEnv {
> > > + cl_context context;
> > > + cl_command_queue command_queue;
> > > + cl_program program;
> > > + cl_kernel kernel;
> > > + char kernel_name[150];
> > > +}AV_KernelEnv;
> >
> > AVKernelEnv?
> >
> > > +
> > > +typedef struct AV_ExtOpenCLInfo {
> > > + cl_platform_id platform;
> >
> > > + cl_device_type devide_type;
> >
> > typo?
> >
> > > + cl_context context;
> > > + cl_device_id *devices;
> >
> > devices_id?
> >
> > > + cl_device_id dev;
> > > + cl_command_queue command_queue;
> > > + char *platformName;
> >
> > camelCase
> >
> > > +}AV_ExtOpenCLInfo;
> >
> > AVExtOpenCLInfo
> >
> > also what's "Ext" for?
> >
>
> Let me explain it. There are two cases to use OpenCL filters, the first one
> is just as ffmpeg.exe,the application program doesn't create OpenCL
> environment but it want to use OpenCL filter in ffmpeg or
> libavfilter.lib,under this condition,libavfilter(ffmpeg) should create the
> OpenCL filter it self,so it doesn't use AVExtOpenCLInfo.The other one is
> that the application program has created the OpenCL environment, under this
> condition, libavfilter(ffmpeg) may use the OpenCL environment wtich
> application program created,so the application program just set the
> AVExtOpenCLInfo and pass it two the API "av_opencl_init_run_env",
> libavfilter will use it to init opencl filter.
Why two distinct interfaces if just one will be used?
Also my comment was more about what "Ext" stands for, I suppose it
stands for "External". In each case the name should be something like
AVOpenCLExtInfo. The usual convention here is library prefix ("AV"),
followed by module name ("OpenCL"), followed by a possibly descriptive
name.
[...]
--
FFmpeg = Fierce & Faithless Mean Power Enlightening Guru
More information about the ffmpeg-devel
mailing list