[FFmpeg-devel] [PATCH 1/2] libavutil/libavfilter: opencl wrapper based on comments on 20130328

Michael Niedermayer michaelni at gmx.at
Fri Mar 29 01:08:02 CET 2013


On Thu, Mar 28, 2013 at 08:57:44PM +0800, Wei Gao wrote:
> 

>  configure          |    4 
>  libavutil/Makefile |    3 
>  libavutil/opencl.c |  658 +++++++++++++++++++++++++++++++++++++++++++++++++++++
>  libavutil/opencl.h |  198 +++++++++++++++
>  4 files changed, 863 insertions(+)
> b6faea13ffbbe062ca802507e91e2da2c4e5e5ef  0001-opencl-wrapper-based-on-comments-on-20130328.patch
> From 38aa9bb65c63d44cd1586383f37197bee93c61ab Mon Sep 17 00:00:00 2001
> From: highgod0401 <highgod0401 at gmail.com>
> Date: Thu, 28 Mar 2013 20:42:21 +0800
> Subject: [PATCH 1/2] opencl wrapper based on comments on 20130328
> 
> ---
>  configure          |   4 +
>  libavutil/Makefile |   3 +
>  libavutil/opencl.c | 658 +++++++++++++++++++++++++++++++++++++++++++++++++++++
>  libavutil/opencl.h | 198 ++++++++++++++++
>  4 files changed, 863 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..d04433a
> --- /dev/null
> +++ b/libavutil/opencl.c
> @@ -0,0 +1,658 @@
> +/*
> + * 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_MEMORYBARRIER && !HAVE_SYNC_VAL_COMPARE_AND_SWAP && !HAVE_MACHINE_RW_BARRIER
> +#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
> +#else
> +#define LOCK_OPENCL
> +#define UNLOCK_OPENCL
> +#endif

the #ifs are quite misleading as pthreads are a strict dependancy
currently


[...]
> +int av_opencl_register_kernel_function(const char *kernel_name, av_opencl_kernel_function function)
> +{
> +    int i;
> +    for (i = 0; i < gpu_env.kernel_count; i++) {
> +        if (av_strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
> +            LOCK_OPENCL
> +            gpu_env.kernel_functions[i] = function;
> +            gpu_env.runtime_kernel_count++;
> +            UNLOCK_OPENCL
> +            return 0;
> +        }
> +    }

kernel_count access needs to be inside the mutex


[...]
> +int av_opencl_run_kernel(const char *kernel_name, void **userdata)
> +{
> +    av_opencl_kernel_function function = NULL;
> +    int i;

> +    for (i = 0; i < gpu_env.kernel_count; i++) {
> +        if (av_strcasecmp(kernel_name, gpu_env.kernel_names[i]) == 0) {
> +            function = gpu_env.kernel_functions[i];
> +            break;
> +        }
> +    }

kernel_count is not a constant, thus this needs a mutex


> +    if (!function) {
> +        av_log(&openclutils, AV_LOG_ERROR, "Could not find kernel: %s\n", kernel_name);
> +        return AVERROR(EINVAL);
> +    }
> +    return(function(userdata));
> +}
> +

> +int av_opencl_init(AVDictionary *options, AVOpenCLExternalInfo *ext_opencl_info)
> +{
> +    int ret = 0;
> +    LOCK_OPENCL
> +    if (!gpu_env.opencl_is_inited) {
> +        /*initialize devices, context, command_queue*/
> +        AVDictionaryEntry *opt_entry = av_dict_get(options, "build_option", NULL, 0);
> +        ret = init_opencl_env(&gpu_env, ext_opencl_info);
> +        if (ret < 0)
> +            goto end;
> +        /*initialize program, kernel_name, kernel_count*/
> +        ret = compile_kernel_file(&gpu_env, opt_entry->value);
> +        if (ret < 0)
> +            goto end;
> +        av_assert1(gpu_env.kernel_count > 0);
> +        gpu_env.opencl_is_inited = 1;
> +    }
> +end:
> +    UNLOCK_OPENCL
> +    return ret;
> +}

Registration of things like codecs formats and filters is not
guranteed to end before the first are opened and used.
The API makes no such gurantees and user applications can easily
register all codecs and start decoding and only later register
all avfilters and build a filter graph.
User applications can also register their own codecs

The code does not look like it would work if any registration would
happen after the first use

[...]

-- 
Michael     GnuPG fingerprint: 9FF2128B147EF6730BADF133611EC787040B0FAB

Concerning the gods, I have no means of knowing whether they exist or not
or of what sort they may be, because of the obscurity of the subject, and
the brevity of human life -- Protagoras
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 198 bytes
Desc: Digital signature
URL: <http://ffmpeg.org/pipermail/ffmpeg-devel/attachments/20130329/8cd42005/attachment.asc>


More information about the ffmpeg-devel mailing list