[FFmpeg-devel] [PATCH 1/2] libavutil/libavfilter: opencl wrapper based on comments on 20130329
Wei Gao
highgod0401 at gmail.com
Sat Mar 30 07:47:10 CET 2013
Hi
Thanks very much for your reviewing, some explanations as follows
Thanks
Best regards
2013/3/30 Stefano Sabatini <stefasab at gmail.com>
> On date Friday 2013-03-29 21:32:35 +0800, Wei Gao encoded:
> >
>
> > From dc939475b0ac425604e1293b816bbaeb0e651b7f Mon Sep 17 00:00:00 2001
> > From: highgod0401 <highgod0401 at gmail.com>
> > Date: Fri, 29 Mar 2013 20:13:55 +0800
> > Subject: [PATCH 1/2] opencl wrapper based on comments on 20130329
> >
> > ---
> > configure | 4 +
> > libavutil/Makefile | 3 +
> > libavutil/opencl.c | 708
> +++++++++++++++++++++++++++++++++++++++++++++++++++++
> > libavutil/opencl.h | 203 +++++++++++++++
> > 4 files changed, 918 insertions(+)
> > create mode 100644 libavutil/opencl.c
> > create mode 100644 libavutil/opencl.h
> >
> > diff --git a/configure b/configure
> > index 8443db4..9c42a85 100755
> > --- a/configure
> > +++ b/configure
> > @@ -233,6 +233,7 @@ External library support:
> > --enable-libxvid enable Xvid encoding via xvidcore,
> > native MPEG-4/Xvid encoder exists [no]
> > --enable-openal enable OpenAL 1.1 capture support [no]
> > + --enable-opencl enable OpenCL code
> > --enable-openssl enable openssl [no]
> > --enable-x11grab enable X11 grabbing [no]
> > --enable-zlib enable zlib [autodetect]
> > @@ -1178,6 +1179,7 @@ EXTERNAL_LIBRARY_LIST="
> > libxavs
> > libxvid
> > openal
> > + opencl
> > openssl
> > x11grab
> > zlib
> > @@ -3982,6 +3984,7 @@ enabled openal && { { for al_libs in
> "${OPENAL_LIBS}" "-lopenal" "-lOpenAL32
> > die "ERROR: openal not found"; } &&
> > { check_cpp_condition "AL/al.h"
> "defined(AL_VERSION_1_1)" ||
> > die "ERROR: openal must be installed and
> version must be 1.1 or compatible"; }
> > +enabled opencl && require2 opencl CL/cl.h clEnqueueNDRangeKernel
> -lOpenCL
> > enabled openssl && { check_lib openssl/ssl.h SSL_library_init -lssl
> -lcrypto ||
> > check_lib openssl/ssl.h SSL_library_init
> -lssl32 -leay32 ||
> > check_lib openssl/ssl.h SSL_library_init -lssl
> -lcrypto -lws2_32 -lgdi32 ||
> > @@ -4350,6 +4353,7 @@ echo "network support ${network-no}"
> > echo "threading support ${thread_type-no}"
> > echo "safe bitstream reader ${safe_bitstream_reader-no}"
> > echo "SDL support ${sdl-no}"
> > +echo "opencl enabled ${opencl-no}"
> > echo "texi2html enabled ${texi2html-no}"
> > echo "perl enabled ${perl-no}"
> > echo "pod2man enabled ${pod2man-no}"
> > diff --git a/libavutil/Makefile b/libavutil/Makefile
> > index 103ce5e..b520473 100644
> > --- a/libavutil/Makefile
> > +++ b/libavutil/Makefile
> > @@ -52,6 +52,8 @@ HEADERS = adler32.h
> \
> >
> > HEADERS-$(CONFIG_LZO) += lzo.h
> >
> > +HEADERS-$(CONFIG_OPENCL) += opencl.h
> > +
> > ARCH_HEADERS = bswap.h
> \
> > intmath.h
> \
> > intreadwrite.h
> \
> > @@ -106,6 +108,7 @@ OBJS = adler32.o
> \
> > xtea.o
> \
> >
> > OBJS-$(CONFIG_LZO) += lzo.o
> > +OBJS-$(CONFIG_OPENCL) += opencl.o
> >
> > OBJS += $(COMPAT_OBJS:%=../compat/%)
> >
> > diff --git a/libavutil/opencl.c b/libavutil/opencl.c
> > new file mode 100644
> > index 0000000..6281b32
> > --- /dev/null
> > +++ b/libavutil/opencl.c
> > @@ -0,0 +1,708 @@
> > +/*
> > + * 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 "opencl.h"
> > +#include "avstring.h"
> > +#include "log.h"
> > +#include "avassert.h"
> > +
> > +#if HAVE_PTHREADS
> > +
> > +#include <pthread.h>
> > +static pthread_mutex_t atomic_opencl_lock = PTHREAD_MUTEX_INITIALIZER;
> > +
> > +#define LOCK_OPENCL pthread_mutex_lock(&atomic_opencl_lock);
> > +#define UNLOCK_OPENCL pthread_mutex_unlock(&atomic_opencl_lock);
> > +
> > +#elif !HAVE_THREADS
> > +#define LOCK_OPENCL
> > +#define UNLOCK_OPENCL
> > +#endif
> > +
> > +
> > +#define MAX_KERNEL_NUM 500
>
> > +#define MAX_REG_NAME_LEN 64
> > +#define MAX_REG_NUM 200
> > +
> > +typedef struct KernelNode {
> > + char kernel_name[AV_OPENCL_MAX_KERNEL_NAME_SIZE];
> > + cl_kernel kernel;
> > + int kernel_ref;
> > +} KernelNode;
> > +
> > +typedef struct GPUEnv {
> > + int opencl_is_inited;
> > + cl_platform_id platform;
> > + cl_device_type device_type;
> > + cl_context context;
> > + cl_device_id *device_ids;
> > + cl_device_id device_id;
> > + cl_command_queue command_queue;
> > + cl_program program;
>
> > + int register_count;
> > + char register_names[MAX_REG_NUM][MAX_REG_NAME_LEN + 1];
> > + const char *kernel_code[MAX_REG_NUM];
> > + av_opencl_kernel_function kernel_functions[MAX_REG_NUM];
>
> "register_names" sounds pretty confusing, it reminds "CPU registers"
> to me, which have a completely different meaning in this context, and
> doesn't tell much about what it's actually representing.
>
will delete the register_names and av_opencl_kernel_function
kernel_functions[MAX_REG_NUM];
>
>
> > + env->context = gpu_env.context;
> > + env->command_queue = gpu_env.command_queue;
> > + env->program = gpu_env.program;
>
> What's the need to set these fields in env if you can always access the
> global environment?
>
the global opencl environment is static variable,static GPUEnv gpu_env;,
and it is better to use apis to access than access it directly I think.
>
> > + }
> > + gpu_env.kernel_node[i].kernel_ref--;
> > + }
> > + }
> > +end:
> > + UNLOCK_OPENCL
>
> I feel this kernel referencing code is a bit overdesigned and could be
> avoided if you allow the creation of more kernels associated to the
> same kernel function. On the other hand I don't know what are the
> drawbacks of this approach (performance loss?).
>
the kernel it released is the resource on GPU created by clCreateKernel, it
is not good to create many kernels with the same name(but different kernel
handle), it a waste of resource.
>
> > +}
> > +
> What about specifying the function when you run the kernel?
>
> I mean something like:
> int av_opencl_run_kernel(av_opencl_function fn, void **userdata);
>
> This way you avoid the registration process, simplify the API, avoid
> function look-up when running the kernel, and also increase
> flexibility since you can pass different functions to the kernel.
>
> Also I do not understand why do you need to use a register name (used
> to register a kernel code) to associate a function to a kernel.
>
delete this api
>
> > +end:
> > + UNLOCK_OPENCL
> > + return ret;
> > +}
> > +
>
>
> @param kernel_name name of a function defined in the OpenCL runtime
> environment
>
this name is the real kernel function name you can reference the
deshake_kernel.h, the function has the prefix "kernel"
"avfilter_transform" is the name
>
> > + * @return >=0 on success, a negative error code on failure
> > + */
> > +int av_opencl_create_kernel(AVOpenCLKernelEnv *env, const char
> *kernel_name);
> > +
>
> [...]
> --
> FFmpeg = Frightening and Faithless Multipurpose Practical Extravagant Gem
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
More information about the ffmpeg-devel
mailing list