[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