[FFmpeg-devel] [PATCH 1/2] libavutil/libavfilter: opencl wrapper based on comments on 20130326
Wei Gao
highgod0401 at gmail.com
Wed Mar 27 08:11:46 CET 2013
Hi,
Stefano Sabatini, thanks for your reply. some questions and explanations as
follows
Thanks
Best regards
2013/3/27 Stefano Sabatini <stefasab at gmail.com>
> On date Tuesday 2013-03-26 18:55:05 +0800, Wei Gao encoded:
> >
>
> > From f91df6a8315a1b7bdc7b69517831fc745fcbd4fd Mon Sep 17 00:00:00 2001
> > From: highgod0401 <highgod0401 at gmail.com>
> > Date: Tue, 26 Mar 2013 18:43:00 +0800
> > Subject: [PATCH 1/2] opencl wrapper based on comments on 20130326
> >
> > ---
> > configure | 4 +
> > libavutil/Makefile | 3 +
> > libavutil/opencl.c | 653
> +++++++++++++++++++++++++++++++++++++++++++++++++++++
> > libavutil/opencl.h | 219 ++++++++++++++++++
> > 4 files changed, 879 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..6375e10 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
> \
> > @@ -115,6 +117,7 @@ SKIPHEADERS-$(HAVE_MACHINE_RW_BARRIER) +=
> atomic_suncc.h
> > SKIPHEADERS-$(HAVE_MEMORYBARRIER) += atomic_win32.h
> > SKIPHEADERS-$(HAVE_SYNC_VAL_COMPARE_AND_SWAP) += atomic_gcc.h
> >
> > +OBJS-$(CONFIG_OPENCL) += opencl.o
> > TESTPROGS = adler32
> \
> > aes
> \
> > atomic
> \
> > diff --git a/libavutil/opencl.c b/libavutil/opencl.c
> > new file mode 100644
> > index 0000000..929aae7
> > --- /dev/null
> > +++ b/libavutil/opencl.c
> > @@ -0,0 +1,653 @@
> > +/*
> > + * 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"
> > +
> > +#define MAX_KERNEL_NAME_LEN 64
> > +#define MAX_KERNEL_NUM 200
> > +
> > +typedef struct GPUEnv {
> > + 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;
>
> > + char kernel_names[MAX_KERNEL_NUM][MAX_KERNEL_NAME_LEN+1];
> > + av_opencl_kernel_function kernel_functions[MAX_KERNEL_NUM];
> > + const char *kernel_code[MAX_KERNEL_NUM];
>
> Are these related to kernel *functions*, right?
>
> In this case I think:
> char kernel_code_name
> [MAX_KERNEL_FUNCTIONS_NUM][MAX_KERNEL_FUNCTION_NAME_LEN+1];
> const char *kernel_code
> [MAX_KERNEL_FUNCTIONS_NUM];
> av_opencl_kernel_function kernel_functions
> [MAX_KERNEL_FUNCTIONS_NUM];
>
> may be more clear.
>
> > + int kernel_count;
>
> This also seems related to kernel function names, so I guess:
>
> int kernel_function_count;
>
> could be better
>
this is the register kernel count, the actually opencl filter number, not
the function number.
>
> > + int runtime_kernel_count;
> > + int is_user_created; // 1: the opencl env is created by user and
> use AVOpenCLExternalInfo to pass to ffmpeg ,0:created by opencl wrapper
> > + uint8_t *temp_buffer;
> > + int temp_buffer_size;
> > +} GPUEnv;
> > +
>
> if (gpu_env.kernel_count >= MAX_KERNEL_NUM) {
> ...
> return err;
> }
>
> if (strlen(...) >= ...)
> ...
> return err;
> }
>
> register kernel;
> }
>
> Uhm so basically this function register a kernel name and code in the
> global environment, and this will be compiled when doing
> av_opencl_init(), right?
>
Yes.
>
> More about the overall design later.
> > + return 0;
> > +}
> > +static const char *opencl_errstr(int status)
> > +{
> > + int i;
> > + for (i = 0; i < sizeof(opencl_err_msg); i++) {
> > + if (opencl_err_msg[i].err_code == status)
> > + return opencl_err_msg[i].err_str;
> > + }
> > + return "unknown error";
> > +}
> > +
>
> > +int av_opencl_create_kernel(const char *kernel_name, AVOpenCLKernelEnv
> *env)
> > +{
> > + int status;
> > + if (!env->kernel) {
> > + env->kernel = clCreateKernel(gpu_env.program,
> kernel_name, &status);
> > + env->context = gpu_env.context;
> > + env->command_queue = gpu_env.command_queue;
> > + env->program = gpu_env.program;
> > + av_strlcpy(env->kernel_name, kernel_name,
> AV_OPENCL_MAX_KERNEL_NAME_SIZE);
>
> You could abort here in case the kernel name is too long.
>
> > + if (status != CL_SUCCESS) {
> > + av_log(&openclutils, AV_LOG_ERROR, "Could not create OpenCL
> kernel: %s\n", opencl_errstr(status));
> > + return AVERROR_EXTERNAL;
> > + }
> > + }
> > + return 0;
> > +}
>
> You should abort in case the kernel was already created.
>
> Also maybe:
> int av_opencl_create_kernel(AVOpenCLKernelEnv *env, const char
> *kernel_name);
>
> could make more sense.
>
> Also about the param "kernel_name", this is rather the name of the
> function defined in a kernel if I understand correctly, so maybe:
> "kernel_function_name" or "kernel_entry_point" could be less
> confusing.
>
" if (!env->kernel)" , this can check whether the kernel has been created.
>
> > +
>
>
> > +#include <CL/cl.h>
>
> This is my understanding of how this OpenCL API works.
>
> You register some code with a name (which are currently stored
> in the global environment), with av_opencl_register_kernel().
>
> Then the previously registered framments of code are compiled by
> OpenCL, when doing: av_opencl_init() -> compile_kernel_file()
>
> With av_opencl_init() you also specify some parameters (build_options)
> which are used when compiling the code of the specified functions.
>
> av_opencl_init() also creates the OpenCL program, which is stored in
> the global environment. The program is unique for all the kernels
> registered so far.
>
> At this point you need to create an entry point for each kernel, to
> run a specific *function* defined within it. This is done by creating
> a kernel, with av_opencl_create_kernel()
>
> av_opencl_create_kernel() is used to create a kernel (a sort of
> handler to communicate with the *compiled* kernel). The kernel is
> created specifying the name of the function to run *in the kernel
> code*. The kernel is set in the passed AVOpenCLEnv environment.
>
> In order to run a function specified in a kernel, you also need to
> provide some parameters/data to it.
>
> This is done through av_opencl_register_kernel_function(), which is
> used to register a function which is associated to one of the
> previously registered kernel in the global environment.
>
> so we have: kernel(global env) -> function(global env)
>
> To run the code of a kernel, av_opencl_run_kernel() must be called,
> with the name of the registered kernel on which the function is to be
> called.
>
> This function lookups the functions registered in the global
> environment, and executes the registered function with provided user
> data/parameters, which in particular must contain the opencl
> environment. The environment should contain the kernel handler created
> with av_opencl_create_kerne(), and is used to set the arguments for
> the function defined in the kernel code, and eventually run the code
> for it (see the deshake patch for an example of such usage).
>
> ...
>
> So basically this is the workflow:
>
> kernel code registration (done in the global environment) ->
> av_opencl_register_kernel()
> kernel code compilation/init (always done in the global environment) ->
> av_opencl_init()
>
> kernel function registration (can be eventually done *before* init) ->
> av_opencl_register_kernel_function()
> kernel object creation, which is required to run the code ->
> av_opencl_create_kernel
> kernel code execution with user data parameters ->
> av_opencl_run_kernel()
>
> Cleanup:
> kernel object (stored in an environment) ->
> av_opencl_release_kernel()
> global environment ->
> av_opencl_uninit()
>
> ...
>
> Can you confirm that this is an accurate description of the
> design/workflow?
>
yes,you are right
>
> The main problem with this design is that different threads and
> components can messup with the global environment.
>
> For example you may want to init a filter, this creates a global
> environment, then you create another filter/component which requires
> to build a different kernel etc., which can't be done since you're
> supposed to init the global environment just once.
>
all the opencl code an use the same, such as reuse context, command queue,
device, platform....... just init once is enough.
>
> Ideally we should have one OpenCL context per component, so we don't
> need to know everything (kernel code and functions) when we init the
> OpenCL system, and by using a global environment you are prevented
> from doing that.
>
> all kernels can reuse the opencl environment, no need to create for
each.also the init function accept the externel OpenCL environment
"AVOpenCLExternalInfo" from application, if the environment is created by
application, the library will not release the environment.
> In a similar way, when you uninit the OpenCL system you don't know if
> other components are actually using it, so the only safe way is to
> uninit() it when you close the *application*, which is not ideal for a
> library.
>
I set a counter to count whether all kernels are released
"gpu_env.runtime_kernel_count--; ", if all the kernel is released ,then
release the OpenCL environment(it should wait for all the kernels have been
released.).the unint function just release the environment created by the
library itself,"gpu_env.is_user_created" can indicate that whether the
OpenCL enviroment is created by application or the opencl library itself,
it only set to 1 if the OpenCL enviroment is created by application in
av_opencl_init.There are two paths, one is application create the opencl
library itself,then it release it. the other is created by application,
then application release it.
>
> > +
> > +#define AV_OPENCL_KERNEL( ... )# __VA_ARGS__
> > +
> > +#define AV_OPENCL_MAX_KERNEL_NAME_SIZE 150
> > +
> > +typedef struct AVOpenCLKernelEnv {
> > + cl_context context;
> > + cl_command_queue command_queue;
> > + cl_program program;
> > + cl_kernel kernel;
> > + char kernel_name[AV_OPENCL_MAX_KERNEL_NAME_SIZE];
> > +} AVOpenCLKernelEnv;
> > +
> > +typedef struct AVOpenCLExternalInfo {
> > + 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;
> > + char *platform_name;
> > +} AVOpenCLExternalInfo;
> > +
> > +/**
> > + * User defined function, used to set the input parameter in the kernel
> > + *environment. This function launches kernel and copies data from GPU to
> > + *CPU, or from CPU to GPU.
> > + */
> > +typedef int (* av_opencl_kernel_function) (void **userdata);
> > +
> > +/**
> > + * Register a function for running the kernel specified by the kernel
> name.
> > + *@param kernel_name this kernel name is used to find the kernel in
> OpenCL runtime environment.
> > + *@param function user defined function,should not be NULL, it
> is used to set the input parameter in the kernel environment
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_register_kernel_function(const char *kernel_name,
> av_opencl_kernel_function function);
> > +
> > +/**
> > + *Load OpenCL kernel.
> > + *
> > + *@param kernel_name this kernel name is used to find the kernel in
> OpenCL runtime environment.
> > + *@param userdata this userdata is the all parameters for
> running the kernel specified by kernel name
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_run_kernel(const char *kernel_name, void **userdata);
> > +
> > +/**
> > + * Init the run time OpenCL environment.
> > + *
> > + *This function must be called befor calling any function related to
> OpenCL.This function should be called by a single thread.
> > + *
> > + *
> > + *@param build_option option of compile the kernel in OpenCL
> runtime environment,reference "OpenCL Specification Version: 1.2 chapter
> 5.6.4"
> > + *@param ext_opencl_info this is the extern OpenCL environment which
> the application program has created
> > + *@return >=0 on success, a negative error value on failure
> > + */
> > +int av_opencl_init(const char *build_option, AVOpenCLExternalInfo
> *ext_opencl_info);
>
> It would be better to pass an AVDictionary here, so we can add more
> options (e.g. paths and security policy options) without breaking the
> API later. Or this could be done with a sort of per-component context,
> in a similar way with what is done for example with avcodec_open2()
> (where we pass all the options through an AVDictionary).
>
Sorry, I don't get what you mean.
> --
> FFmpeg = Fanciful and Fiendish MultiPurpose Exxagerate Genius
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
More information about the ffmpeg-devel
mailing list