Chiebot-Mirror
/
FFmpeg
mirror of https://github.com/FFmpeg/FFmpeg.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

lavu,lavfi,ffmpeg: Remove experimental OpenCL API

Browse Source 
This was added in early 2013 and abandoned several months later; as far as
I can tell, there are no external users.  Future OpenCL use will be via
hwcontext, which requires neither special OpenCL-only API nor global state
in libavutil.

All internal users are also deleted - this is just the unsharp filter
(replaced by unsharp_opencl, which is more flexible) and the deshake filter
(no replacement).
pull/272/head
Mark Thompson 7 years ago
parent 0f93cef2d6
commit 3650cb2dfa
26 changed files with 6 additions and 3054 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 2
      configure
  2. 3
      doc/APIchanges
  3. 8
      doc/filters.texi
  4. 30
      doc/utils.texi
  5. 1
      fftools/Makefile
  6. 18
      fftools/cmdutils.h
  7. 283
      fftools/cmdutils_opencl.c
  8. 3
      libavfilter/Makefile
  9. 2
      libavfilter/allfilters.c
  10. 24
      libavfilter/deshake.h
  11. 198
      libavfilter/deshake_opencl.c
  12. 45
      libavfilter/deshake_opencl.h
  13. 225
      libavfilter/deshake_opencl_kernel.h
  14. 41
      libavfilter/opencl_allkernels.c
  15. 29
      libavfilter/opencl_allkernels.h
  16. 31
      libavfilter/unsharp.h
  17. 422
      libavfilter/unsharp_opencl.c
  18. 34
      libavfilter/unsharp_opencl.h
  19. 342
      libavfilter/unsharp_opencl_kernel.h
  20. 23
      libavfilter/vf_deshake.c
  21. 26
      libavfilter/vf_unsharp.c
  22. 4
      libavutil/Makefile
  23. 875
      libavutil/opencl.c
  24. 292
      libavutil/opencl.h
  25. 59
      libavutil/opencl_internal.c
  26. 40
      libavutil/opencl_internal.h

2
configure vendored
Unescape View File

@ -3209,7 +3209,6 @@ deinterlace_qsv_filter_deps="libmfx"
deinterlace_vaapi_filter_deps="vaapi" deinterlace_vaapi_filter_deps="vaapi"
delogo_filter_deps="gpl" delogo_filter_deps="gpl"
deshake_filter_select="pixelutils" deshake_filter_select="pixelutils"
deshake_filter_suggest="opencl"
drawtext_filter_deps="libfreetype" drawtext_filter_deps="libfreetype"
drawtext_filter_suggest="libfontconfig libfribidi" drawtext_filter_suggest="libfontconfig libfribidi"
elbg_filter_deps="avcodec" elbg_filter_deps="avcodec"
@ -3285,7 +3284,6 @@ tinterlace_pad_test_deps="tinterlace_filter"
tonemap_filter_deps="const_nan" tonemap_filter_deps="const_nan"
unsharp_opencl_filter_deps="opencl" unsharp_opencl_filter_deps="opencl"
uspp_filter_deps="gpl avcodec" uspp_filter_deps="gpl avcodec"
unsharp_filter_suggest="opencl"
vaguedenoiser_filter_deps="gpl" vaguedenoiser_filter_deps="gpl"
vidstabdetect_filter_deps="libvidstab" vidstabdetect_filter_deps="libvidstab"
vidstabtransform_filter_deps="libvidstab" vidstabtransform_filter_deps="libvidstab"

3
doc/APIchanges
Unescape View File

@ -15,6 +15,9 @@ libavutil: 2017-10-21
API changes, most recent first: API changes, most recent first:
2017-11-xx - xxxxxxx - lavu 55.3.0 - opencl.h
Remove experiental OpenCL API (av_opencl_*).
2017-11-xx - xxxxxxx - lavu 55.2.0 - hwcontext.h hwcontext_opencl.h 2017-11-xx - xxxxxxx - lavu 55.2.0 - hwcontext.h hwcontext_opencl.h
Add AV_HWDEVICE_TYPE_OPENCL and a new installed header with Add AV_HWDEVICE_TYPE_OPENCL and a new installed header with
OpenCL-specific hwcontext definitions. OpenCL-specific hwcontext definitions.

8
doc/filters.texi
Unescape View File

@ -6913,10 +6913,6 @@ Default value is @samp{exhaustive}.
If set then a detailed log of the motion search is written to the If set then a detailed log of the motion search is written to the
specified file. specified file.
@item opencl
If set to 1, specify using OpenCL capabilities, only available if
FFmpeg was configured with @code{--enable-opencl}. Default value is 0.
@end table @end table
@section despill @section despill
@ -15048,10 +15044,6 @@ sharpen it, a value of zero will disable the effect.
Default value is 0.0. Default value is 0.0.
@item opencl
If set to 1, specify using OpenCL capabilities, only available if
FFmpeg was configured with @code{--enable-opencl}. Default value is 0.
@end table @end table
All parameters are optional and default to the equivalent of the All parameters are optional and default to the equivalent of the

30
doc/utils.texi
Unescape View File

@ -1057,33 +1057,3 @@ indication of the corresponding powers of 10 and of 2.
@end table @end table
@c man end EXPRESSION EVALUATION @c man end EXPRESSION EVALUATION
@chapter OpenCL Options
@c man begin OPENCL OPTIONS
When FFmpeg is configured with @code{--enable-opencl}, it is possible
to set the options for the global OpenCL context.
The list of supported options follows:
@table @option
@item build_options
Set build options used to compile the registered kernels.
See reference "OpenCL Specification Version: 1.2 chapter 5.6.4".
@item platform_idx
Select the index of the platform to run OpenCL code.
The specified index must be one of the indexes in the device list
which can be obtained with @code{ffmpeg -opencl_bench} or @code{av_opencl_get_device_list()}.
@item device_idx
Select the index of the device used to run OpenCL code.
The specified index must be one of the indexes in the device list which
can be obtained with @code{ffmpeg -opencl_bench} or @code{av_opencl_get_device_list()}.
@end table
@c man end OPENCL OPTIONS

1
fftools/Makefile
Unescape View File

@ -29,7 +29,6 @@ $(1)$(PROGSSUF)_g$(EXESUF): FF_EXTRALIBS += $(EXTRALIBS-$(1))
-include $$(OBJS-$(1):.o=.d) -include $$(OBJS-$(1):.o=.d)
endef endef
$(foreach P,$(AVPROGS-yes),$(eval OBJS-$(P)-$(CONFIG_OPENCL) += fftools/cmdutils_opencl.o))
$(foreach P,$(AVPROGS-yes),$(eval $(call DOFFTOOL,$(P)))) $(foreach P,$(AVPROGS-yes),$(eval $(call DOFFTOOL,$(P))))
all: $(AVPROGS) all: $(AVPROGS)

18
fftools/cmdutils.h
Unescape View File

@ -105,12 +105,6 @@ int opt_max_alloc(void *optctx, const char *opt, const char *arg);
int opt_codec_debug(void *optctx, const char *opt, const char *arg); int opt_codec_debug(void *optctx, const char *opt, const char *arg);
#if CONFIG_OPENCL
int opt_opencl(void *optctx, const char *opt, const char *arg);
int opt_opencl_bench(void *optctx, const char *opt, const char *arg);
#endif
/** /**
* Limit the execution time. * Limit the execution time.
*/ */
@ -207,17 +201,6 @@ typedef struct OptionDef {
void show_help_options(const OptionDef *options, const char *msg, int req_flags, void show_help_options(const OptionDef *options, const char *msg, int req_flags,
int rej_flags, int alt_flags); int rej_flags, int alt_flags);
#if CONFIG_OPENCL
#define CMDUTILS_COMMON_OPTIONS_OPENCL \
{ "opencl_bench", OPT_EXIT, {.func_arg = opt_opencl_bench}, \
"run benchmark on all OpenCL devices and show results" }, \
{ "opencl_options", HAS_ARG, {.func_arg = opt_opencl}, \
"set OpenCL environment options" }, \
#else
#define CMDUTILS_COMMON_OPTIONS_OPENCL
#endif
#if CONFIG_AVDEVICE #if CONFIG_AVDEVICE
#define CMDUTILS_COMMON_OPTIONS_AVDEVICE \ #define CMDUTILS_COMMON_OPTIONS_AVDEVICE \
{ "sources" , OPT_EXIT | HAS_ARG, { .func_arg = show_sources }, \ { "sources" , OPT_EXIT | HAS_ARG, { .func_arg = show_sources }, \
@ -257,7 +240,6 @@ void show_help_options(const OptionDef *options, const char *msg, int req_flags,
{ "max_alloc", HAS_ARG, { .func_arg = opt_max_alloc }, "set maximum size of a single allocated block", "bytes" }, \ { "max_alloc", HAS_ARG, { .func_arg = opt_max_alloc }, "set maximum size of a single allocated block", "bytes" }, \
{ "cpuflags", HAS_ARG | OPT_EXPERT, { .func_arg = opt_cpuflags }, "force specific cpu flags", "flags" }, \ { "cpuflags", HAS_ARG | OPT_EXPERT, { .func_arg = opt_cpuflags }, "force specific cpu flags", "flags" }, \
{ "hide_banner", OPT_BOOL | OPT_EXPERT, {&hide_banner}, "do not show program banner", "hide_banner" }, \ { "hide_banner", OPT_BOOL | OPT_EXPERT, {&hide_banner}, "do not show program banner", "hide_banner" }, \
CMDUTILS_COMMON_OPTIONS_OPENCL \
CMDUTILS_COMMON_OPTIONS_AVDEVICE \ CMDUTILS_COMMON_OPTIONS_AVDEVICE \
/** /**

283
fftools/cmdutils_opencl.c
Unescape View File

@ -1,283 +0,0 @@
/*
* Copyright (C) 2013 Lenny Wang
*
* 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 "libavutil/opt.h"
#include "libavutil/time.h"
#include "libavutil/log.h"
#include "libavutil/opencl.h"
#include "libavutil/avstring.h"
#include "cmdutils.h"
typedef struct {
int platform_idx;
int device_idx;
char device_name[64];
int64_t runtime;
} OpenCLDeviceBenchmark;
const char *ocl_bench_source = AV_OPENCL_KERNEL(
inline unsigned char clip_uint8(int a)
{
if (a & (~0xFF))
return (-a)>>31;
else
return a;
}
kernel void unsharp_bench(
global unsigned char *src,
global unsigned char *dst,
global int *mask,
int width,
int height)
{
int i, j, local_idx, lc_idx, sum = 0;
int2 thread_idx, block_idx, global_idx, lm_idx;
thread_idx.x = get_local_id(0);
thread_idx.y = get_local_id(1);
block_idx.x = get_group_id(0);
block_idx.y = get_group_id(1);
global_idx.x = get_global_id(0);
global_idx.y = get_global_id(1);
local uchar data[32][32];
local int lc[128];
for (i = 0; i <= 1; i++) {
lm_idx.y = -8 + (block_idx.y + i) * 16 + thread_idx.y;
lm_idx.y = lm_idx.y < 0 ? 0 : lm_idx.y;
lm_idx.y = lm_idx.y >= height ? height - 1: lm_idx.y;
for (j = 0; j <= 1; j++) {
lm_idx.x = -8 + (block_idx.x + j) * 16 + thread_idx.x;
lm_idx.x = lm_idx.x < 0 ? 0 : lm_idx.x;
lm_idx.x = lm_idx.x >= width ? width - 1: lm_idx.x;
data[i*16 + thread_idx.y][j*16 + thread_idx.x] = src[lm_idx.y*width + lm_idx.x];
}
}
local_idx = thread_idx.y*16 + thread_idx.x;
if (local_idx < 128)
lc[local_idx] = mask[local_idx];
barrier(CLK_LOCAL_MEM_FENCE);
\n#pragma unroll\n
for (i = -4; i <= 4; i++) {
lm_idx.y = 8 + i + thread_idx.y;
\n#pragma unroll\n
for (j = -4; j <= 4; j++) {
lm_idx.x = 8 + j + thread_idx.x;
lc_idx = (i + 4)*8 + j + 4;
sum += (int)data[lm_idx.y][lm_idx.x] * lc[lc_idx];
}
}
int temp = (int)data[thread_idx.y + 8][thread_idx.x + 8];
int res = temp + (((temp - (int)((sum + 1<<15) >> 16))) >> 16);
if (global_idx.x < width && global_idx.y < height)
dst[global_idx.x + global_idx.y*width] = clip_uint8(res);
}
);
#define OCLCHECK(method, ... ) \
do { \
status = method(__VA_ARGS__); \
if (status != CL_SUCCESS) { \
av_log(NULL, AV_LOG_ERROR, # method " error '%s'\n", \
av_opencl_errstr(status)); \
ret = AVERROR_EXTERNAL; \
goto end; \
} \
} while (0)
#define CREATEBUF(out, flags, size) \
do { \
out = clCreateBuffer(ext_opencl_env->context, flags, size, NULL, &status); \
if (status != CL_SUCCESS) { \
av_log(NULL, AV_LOG_ERROR, "Could not create OpenCL buffer\n"); \
ret = AVERROR_EXTERNAL; \
goto end; \
} \
} while (0)
static void fill_rand_int(int *data, int n)
{
int i;
srand(av_gettime());
for (i = 0; i < n; i++)
data[i] = rand();
}
#define OPENCL_NB_ITER 5
static int64_t run_opencl_bench(AVOpenCLExternalEnv *ext_opencl_env)
{
int i, arg = 0, width = 1920, height = 1088;
int64_t start, ret = 0;
cl_int status;
size_t kernel_len;
char *inbuf;
int *mask = NULL;
int buf_size = width * height * sizeof(char);
int mask_size = sizeof(uint32_t) * 128;
cl_mem cl_mask = NULL, cl_inbuf = NULL, cl_outbuf = NULL;
cl_kernel kernel = NULL;
cl_program program = NULL;
size_t local_work_size_2d[2] = {16, 16};
size_t global_work_size_2d[2] = {(size_t)width, (size_t)height};
if (!(inbuf = av_malloc(buf_size)) || !(mask = av_malloc(mask_size))) {
av_log(NULL, AV_LOG_ERROR, "Out of memory\n");
ret = AVERROR(ENOMEM);
goto end;
}
fill_rand_int((int*)inbuf, buf_size/4);
fill_rand_int(mask, mask_size/4);
CREATEBUF(cl_mask, CL_MEM_READ_ONLY, mask_size);
CREATEBUF(cl_inbuf, CL_MEM_READ_ONLY, buf_size);
CREATEBUF(cl_outbuf, CL_MEM_READ_WRITE, buf_size);
kernel_len = strlen(ocl_bench_source);
program = clCreateProgramWithSource(ext_opencl_env->context, 1, &ocl_bench_source,
&kernel_len, &status);
if (status != CL_SUCCESS || !program) {
av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark program\n");
ret = AVERROR_EXTERNAL;
goto end;
}
status = clBuildProgram(program, 1, &(ext_opencl_env->device_id), NULL, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(NULL, AV_LOG_ERROR, "OpenCL unable to build benchmark program\n");
ret = AVERROR_EXTERNAL;
goto end;
}
kernel = clCreateKernel(program, "unsharp_bench", &status);
if (status != CL_SUCCESS) {
av_log(NULL, AV_LOG_ERROR, "OpenCL unable to create benchmark kernel\n");
ret = AVERROR_EXTERNAL;
goto end;
}
OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_inbuf, CL_TRUE, 0,
buf_size, inbuf, 0, NULL, NULL);
OCLCHECK(clEnqueueWriteBuffer, ext_opencl_env->command_queue, cl_mask, CL_TRUE, 0,
mask_size, mask, 0, NULL, NULL);
OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_inbuf);
OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_outbuf);
OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_mem), &cl_mask);
OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &width);
OCLCHECK(clSetKernelArg, kernel, arg++, sizeof(cl_int), &height);
start = av_gettime_relative();
for (i = 0; i < OPENCL_NB_ITER; i++)
OCLCHECK(clEnqueueNDRangeKernel, ext_opencl_env->command_queue, kernel, 2, NULL,
global_work_size_2d, local_work_size_2d, 0, NULL, NULL);
clFinish(ext_opencl_env->command_queue);
ret = (av_gettime_relative() - start)/OPENCL_NB_ITER;
end:
if (kernel)
clReleaseKernel(kernel);
if (program)
clReleaseProgram(program);
if (cl_inbuf)
clReleaseMemObject(cl_inbuf);
if (cl_outbuf)
clReleaseMemObject(cl_outbuf);
if (cl_mask)
clReleaseMemObject(cl_mask);
av_free(inbuf);
av_free(mask);
return ret;
}
static int compare_ocl_device_desc(const void *a, const void *b)
{
const OpenCLDeviceBenchmark* va = (const OpenCLDeviceBenchmark*)a;
const OpenCLDeviceBenchmark* vb = (const OpenCLDeviceBenchmark*)b;
return FFDIFFSIGN(va->runtime , vb->runtime);
}
int opt_opencl_bench(void *optctx, const char *opt, const char *arg)
{
int i, j, nb_devices = 0, count = 0, ret = 0;
int64_t score = 0;
AVOpenCLDeviceList *device_list;
AVOpenCLDeviceNode *device_node = NULL;
OpenCLDeviceBenchmark *devices = NULL;
cl_platform_id platform;
ret = av_opencl_get_device_list(&device_list);
if (ret < 0) {
return ret;
}
for (i = 0; i < device_list->platform_num; i++)
nb_devices += device_list->platform_node[i]->device_num;
if (!nb_devices) {
av_log(NULL, AV_LOG_ERROR, "No OpenCL device detected!\n");
av_opencl_free_device_list(&device_list);
return AVERROR(EINVAL);
}
if (!(devices = av_malloc_array(nb_devices, sizeof(OpenCLDeviceBenchmark)))) {
av_log(NULL, AV_LOG_ERROR, "Could not allocate buffer\n");
av_opencl_free_device_list(&device_list);
return AVERROR(ENOMEM);
}
for (i = 0; i < device_list->platform_num; i++) {
for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
device_node = device_list->platform_node[i]->device_node[j];
platform = device_list->platform_node[i]->platform_id;
score = av_opencl_benchmark(device_node, platform, run_opencl_bench);
if (score > 0) {
devices[count].platform_idx = i;
devices[count].device_idx = j;
devices[count].runtime = score;
av_strlcpy(devices[count].device_name, device_node->device_name,
sizeof(devices[count].device_name));
count++;
}
}
}
qsort(devices, count, sizeof(OpenCLDeviceBenchmark), compare_ocl_device_desc);
fprintf(stderr, "platform_idx\tdevice_idx\tdevice_name\truntime\n");
for (i = 0; i < count; i++)
fprintf(stdout, "%d\t%d\t%s\t%"PRId64"\n",
devices[i].platform_idx, devices[i].device_idx,
devices[i].device_name, devices[i].runtime);
av_opencl_free_device_list(&device_list);
av_free(devices);
return 0;
}
int opt_opencl(void *optctx, const char *opt, const char *arg)
{
char *key, *value;
const char *opts = arg;
int ret = 0;
while (*opts) {
ret = av_opt_get_key_value(&opts, "=", ":", 0, &key, &value);
if (ret < 0)
return ret;
ret = av_opencl_set_option(key, value);
if (ret < 0)
return ret;
if (*opts)
opts++;
}
return ret;
}

3
libavfilter/Makefile
Unescape View File

@ -19,7 +19,6 @@ OBJS = allfilters.o \
framequeue.o \ framequeue.o \
graphdump.o \ graphdump.o \
graphparser.o \ graphparser.o \
opencl_allkernels.o \
transform.o \ transform.o \
video.o \ video.o \
@ -250,7 +249,6 @@ OBJS-$(CONFIG_NOISE_FILTER) += vf_noise.o
OBJS-$(CONFIG_NULL_FILTER) += vf_null.o OBJS-$(CONFIG_NULL_FILTER) += vf_null.o
OBJS-$(CONFIG_OCR_FILTER) += vf_ocr.o OBJS-$(CONFIG_OCR_FILTER) += vf_ocr.o
OBJS-$(CONFIG_OCV_FILTER) += vf_libopencv.o OBJS-$(CONFIG_OCV_FILTER) += vf_libopencv.o
OBJS-$(CONFIG_OPENCL) += deshake_opencl.o unsharp_opencl.o
OBJS-$(CONFIG_OSCILLOSCOPE_FILTER) += vf_datascope.o OBJS-$(CONFIG_OSCILLOSCOPE_FILTER) += vf_datascope.o
OBJS-$(CONFIG_OVERLAY_FILTER) += vf_overlay.o framesync.o OBJS-$(CONFIG_OVERLAY_FILTER) += vf_overlay.o framesync.o
OBJS-$(CONFIG_OVERLAY_OPENCL_FILTER) += vf_overlay_opencl.o opencl.o \ OBJS-$(CONFIG_OVERLAY_OPENCL_FILTER) += vf_overlay_opencl.o opencl.o \
@ -396,7 +394,6 @@ OBJS-$(CONFIG_MOVIE_FILTER) += src_movie.o
SLIBOBJS-$(HAVE_GNU_WINDRES) += avfilterres.o SLIBOBJS-$(HAVE_GNU_WINDRES) += avfilterres.o
SKIPHEADERS-$(CONFIG_LIBVIDSTAB) += vidstabutils.h SKIPHEADERS-$(CONFIG_LIBVIDSTAB) += vidstabutils.h
SKIPHEADERS-$(CONFIG_OPENCL) += opencl_internal.h deshake_opencl_kernel.h unsharp_opencl_kernel.h
OBJS-$(CONFIG_SHARED) += log2_tab.o OBJS-$(CONFIG_SHARED) += log2_tab.o

2
libavfilter/allfilters.c
Unescape View File

@ -22,7 +22,6 @@
#include "libavutil/thread.h" #include "libavutil/thread.h"
#include "avfilter.h" #include "avfilter.h"
#include "config.h" #include "config.h"
#include "opencl_allkernels.h"
#define REGISTER_FILTER(X, x, y) \ #define REGISTER_FILTER(X, x, y) \
@ -407,7 +406,6 @@ static void register_all(void)
REGISTER_FILTER_UNCONDITIONAL(vsink_buffer); REGISTER_FILTER_UNCONDITIONAL(vsink_buffer);
REGISTER_FILTER_UNCONDITIONAL(af_afifo); REGISTER_FILTER_UNCONDITIONAL(af_afifo);
REGISTER_FILTER_UNCONDITIONAL(vf_fifo); REGISTER_FILTER_UNCONDITIONAL(vf_fifo);
ff_opencl_register_filter_kernel_code_all();
} }
void avfilter_register_all(void) void avfilter_register_all(void)

24
libavfilter/deshake.h
Unescape View File

@ -26,9 +26,6 @@
#include "avfilter.h" #include "avfilter.h"
#include "transform.h" #include "transform.h"
#include "libavutil/pixelutils.h" #include "libavutil/pixelutils.h"
#if CONFIG_OPENCL
#include "libavutil/opencl.h"
#endif
enum SearchMethod { enum SearchMethod {
@ -53,24 +50,6 @@ typedef struct Transform {
double zoom; ///< Zoom percentage double zoom; ///< Zoom percentage
} Transform; } Transform;
#if CONFIG_OPENCL
typedef struct DeshakeOpenclContext {
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel_luma;
cl_kernel kernel_chroma;
int in_plane_size[8];
int out_plane_size[8];
int plane_num;
cl_mem cl_inbuf;
size_t cl_inbuf_size;
cl_mem cl_outbuf;
size_t cl_outbuf_size;
} DeshakeOpenclContext;
#endif
#define MAX_R 64 #define MAX_R 64
typedef struct DeshakeContext { typedef struct DeshakeContext {
@ -96,9 +75,6 @@ typedef struct DeshakeContext {
int cy; int cy;
char *filename; ///< Motion search detailed log filename char *filename; ///< Motion search detailed log filename
int opencl; int opencl;
#if CONFIG_OPENCL
DeshakeOpenclContext opencl_ctx;
#endif
int (* transform)(AVFilterContext *ctx, int width, int height, int cw, int ch, int (* transform)(AVFilterContext *ctx, int width, int height, int cw, int ch,
const float *matrix_y, const float *matrix_uv, enum InterpolateMethod interpolate, const float *matrix_y, const float *matrix_uv, enum InterpolateMethod interpolate,
enum FillMethod fill, AVFrame *in, AVFrame *out); enum FillMethod fill, AVFrame *in, AVFrame *out);

198
libavfilter/deshake_opencl.c
Unescape View File

@ -1,198 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang
*
* 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/dict.h"
#include "libavutil/pixdesc.h"
#include "deshake_opencl.h"
#include "libavutil/opencl_internal.h"
#define PLANE_NUM 3
#define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
int ff_opencl_transform(AVFilterContext *ctx,
int width, int height, int cw, int ch,
const float *matrix_y, const float *matrix_uv,
enum InterpolateMethod interpolate,
enum FillMethod fill, AVFrame *in, AVFrame *out)
{
int ret = 0;
cl_int status;
DeshakeContext *deshake = ctx->priv;
float4 packed_matrix_lu = {matrix_y[0], matrix_y[1], matrix_y[2], matrix_y[5]};
float4 packed_matrix_ch = {matrix_uv[0], matrix_uv[1], matrix_uv[2], matrix_uv[5]};
size_t global_worksize_lu[2] = {(size_t)ROUND_TO_16(width), (size_t)ROUND_TO_16(height)};
size_t global_worksize_ch[2] = {(size_t)ROUND_TO_16(cw), (size_t)(2*ROUND_TO_16(ch))};
size_t local_worksize[2] = {16, 16};
FFOpenclParam param_lu = {0};
FFOpenclParam param_ch = {0};
param_lu.ctx = param_ch.ctx = ctx;
param_lu.kernel = deshake->opencl_ctx.kernel_luma;
param_ch.kernel = deshake->opencl_ctx.kernel_chroma;
if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
return AVERROR(EINVAL);
}
ret = avpriv_opencl_set_parameter(&param_lu,
FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
FF_OPENCL_PARAM_INFO(packed_matrix_lu),
FF_OPENCL_PARAM_INFO(interpolate),
FF_OPENCL_PARAM_INFO(fill),
FF_OPENCL_PARAM_INFO(in->linesize[0]),
FF_OPENCL_PARAM_INFO(out->linesize[0]),
FF_OPENCL_PARAM_INFO(height),
FF_OPENCL_PARAM_INFO(width),
NULL);
if (ret < 0)
return ret;
ret = avpriv_opencl_set_parameter(&param_ch,
FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
FF_OPENCL_PARAM_INFO(packed_matrix_ch),
FF_OPENCL_PARAM_INFO(interpolate),
FF_OPENCL_PARAM_INFO(fill),
FF_OPENCL_PARAM_INFO(in->linesize[0]),
FF_OPENCL_PARAM_INFO(out->linesize[0]),
FF_OPENCL_PARAM_INFO(in->linesize[1]),
FF_OPENCL_PARAM_INFO(out->linesize[1]),
FF_OPENCL_PARAM_INFO(height),
FF_OPENCL_PARAM_INFO(width),
FF_OPENCL_PARAM_INFO(ch),
FF_OPENCL_PARAM_INFO(cw),
NULL);
if (ret < 0)
return ret;
status = clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
deshake->opencl_ctx.kernel_luma, 2, NULL,
global_worksize_lu, local_worksize, 0, NULL, NULL);
status |= clEnqueueNDRangeKernel(deshake->opencl_ctx.command_queue,
deshake->opencl_ctx.kernel_chroma, 2, NULL,
global_worksize_ch, local_worksize, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,
deshake->opencl_ctx.cl_outbuf_size);
if (ret < 0)
return ret;
return ret;
}
int ff_opencl_deshake_init(AVFilterContext *ctx)
{
int ret = 0;
DeshakeContext *deshake = ctx->priv;
ret = av_opencl_init(NULL);
if (ret < 0)
return ret;
deshake->opencl_ctx.plane_num = PLANE_NUM;
deshake->opencl_ctx.command_queue = av_opencl_get_command_queue();
if (!deshake->opencl_ctx.command_queue) {
av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'deshake'\n");
return AVERROR(EINVAL);
}
deshake->opencl_ctx.program = av_opencl_compile("avfilter_transform", NULL);
if (!deshake->opencl_ctx.program) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'avfilter_transform'\n");
return AVERROR(EINVAL);
}
if (!deshake->opencl_ctx.kernel_luma) {
deshake->opencl_ctx.kernel_luma = clCreateKernel(deshake->opencl_ctx.program,
"avfilter_transform_luma", &ret);
if (ret != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform_luma'\n");
return AVERROR(EINVAL);
}
}
if (!deshake->opencl_ctx.kernel_chroma) {
deshake->opencl_ctx.kernel_chroma = clCreateKernel(deshake->opencl_ctx.program,
"avfilter_transform_chroma", &ret);
if (ret != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'avfilter_transform_chroma'\n");
return AVERROR(EINVAL);
}
}
return ret;
}
void ff_opencl_deshake_uninit(AVFilterContext *ctx)
{
DeshakeContext *deshake = ctx->priv;
av_opencl_buffer_release(&deshake->opencl_ctx.cl_inbuf);
av_opencl_buffer_release(&deshake->opencl_ctx.cl_outbuf);
clReleaseKernel(deshake->opencl_ctx.kernel_luma);
clReleaseKernel(deshake->opencl_ctx.kernel_chroma);
clReleaseProgram(deshake->opencl_ctx.program);
deshake->opencl_ctx.command_queue = NULL;
av_opencl_uninit();
}
int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
{
int ret = 0;
AVFilterLink *link = ctx->inputs[0];
DeshakeContext *deshake = ctx->priv;
const int hshift = av_pix_fmt_desc_get(link->format)->log2_chroma_h;
int chroma_height = AV_CEIL_RSHIFT(link->h, hshift);
if ((!deshake->opencl_ctx.cl_inbuf) || (!deshake->opencl_ctx.cl_outbuf)) {
deshake->opencl_ctx.in_plane_size[0] = (in->linesize[0] * in->height);
deshake->opencl_ctx.in_plane_size[1] = (in->linesize[1] * chroma_height);
deshake->opencl_ctx.in_plane_size[2] = (in->linesize[2] * chroma_height);
deshake->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
deshake->opencl_ctx.out_plane_size[1] = (out->linesize[1] * chroma_height);
deshake->opencl_ctx.out_plane_size[2] = (out->linesize[2] * chroma_height);
deshake->opencl_ctx.cl_inbuf_size = deshake->opencl_ctx.in_plane_size[0] +
deshake->opencl_ctx.in_plane_size[1] +
deshake->opencl_ctx.in_plane_size[2];
deshake->opencl_ctx.cl_outbuf_size = deshake->opencl_ctx.out_plane_size[0] +
deshake->opencl_ctx.out_plane_size[1] +
deshake->opencl_ctx.out_plane_size[2];
if (!deshake->opencl_ctx.cl_inbuf) {
ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_inbuf,
deshake->opencl_ctx.cl_inbuf_size,
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
}
if (!deshake->opencl_ctx.cl_outbuf) {
ret = av_opencl_buffer_create(&deshake->opencl_ctx.cl_outbuf,
deshake->opencl_ctx.cl_outbuf_size,
CL_MEM_READ_WRITE, NULL);
if (ret < 0)
return ret;
}
}
ret = av_opencl_buffer_write_image(deshake->opencl_ctx.cl_inbuf,
deshake->opencl_ctx.cl_inbuf_size,
0, in->data,deshake->opencl_ctx.in_plane_size,
deshake->opencl_ctx.plane_num);
return ret;
}

45
libavfilter/deshake_opencl.h
Unescape View File

@ -1,45 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@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_DESHAKE_OPENCL_H
#define AVFILTER_DESHAKE_OPENCL_H
#include "deshake.h"
typedef struct float4 {
float x;
float y;
float z;
float w;
} float4;
int ff_opencl_deshake_init(AVFilterContext *ctx);
void ff_opencl_deshake_uninit(AVFilterContext *ctx);
int ff_opencl_deshake_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
int ff_opencl_transform(AVFilterContext *ctx,
int width, int height, int cw, int ch,
const float *matrix_y, const float *matrix_uv,
enum InterpolateMethod interpolate,
enum FillMethod fill, AVFrame *in, AVFrame *out);
#endif /* AVFILTER_DESHAKE_OPENCL_H */

225
libavfilter/deshake_opencl_kernel.h
Unescape View File

@ -1,225 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang
*
*
* 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_DESHAKE_OPENCL_KERNEL_H
#define AVFILTER_DESHAKE_OPENCL_KERNEL_H
#include "libavutil/opencl.h"
const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
inline unsigned char pixel(global const unsigned char *src, int x, int y,
int w, int h,int stride, unsigned char def)
{
return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[x + y * stride];
}
unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{
return pixel(src, (int)(x + 0.5f), (int)(y + 0.5f), width, height, stride, def);
}
unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{
int x_c, x_f, y_c, y_f;
int v1, v2, v3, v4;
x_f = (int)x;
y_f = (int)y;
x_c = x_f + 1;
y_c = y_f + 1;
if (x_f < -1 || x_f > width || y_f < -1 || y_f > height) {
return def;
} else {
v4 = pixel(src, x_f, y_f, width, height, stride, def);
v2 = pixel(src, x_c, y_f, width, height, stride, def);
v3 = pixel(src, x_f, y_c, width, height, stride, def);
v1 = pixel(src, x_c, y_c, width, height, stride, def);
return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
}
}
unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
int width, int height, int stride, unsigned char def)
{
int x_c, x_f, y_c, y_f;
unsigned char v1, v2, v3, v4;
float f1, f2, f3, f4;
x_f = (int)x;
y_f = (int)y;
x_c = x_f + 1;
y_c = y_f + 1;
if (x_f < - 1 || x_f > width || y_f < -1 || y_f > height)
return def;
else {
v4 = pixel(src, x_f, y_f, width, height, stride, def);
v2 = pixel(src, x_c, y_f, width, height, stride, def);
v3 = pixel(src, x_f, y_c, width, height, stride, def);
v1 = pixel(src, x_c, y_c, width, height, stride, def);
f1 = 1 - sqrt((x_c - x) * (y_c - y));
f2 = 1 - sqrt((x_c - x) * (y - y_f));
f3 = 1 - sqrt((x - x_f) * (y_c - y));
f4 = 1 - sqrt((x - x_f) * (y - y_f));
return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
}
}
inline const float clipf(float a, float amin, float amax)
{
if (a < amin) return amin;
else if (a > amax) return amax;
else return a;
}
inline int mirror(int v, int m)
{
while ((unsigned)v > (unsigned)m) {
v = -v;
if (v < 0)
v += 2 * m;
}
return v;
}
kernel void avfilter_transform_luma(global unsigned char *src,
global unsigned char *dst,
float4 matrix,
int interpolate,
int fill,
int src_stride_lu,
int dst_stride_lu,
int height,
int width)
{
int x = get_global_id(0);
int y = get_global_id(1);
int idx_dst = y * dst_stride_lu + x;
unsigned char def = 0;
float x_s = x * matrix.x + y * matrix.y + matrix.z;
float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
if (x < width && y < height) {
switch (fill) {
case 0: //FILL_BLANK
def = 0;
break;
case 1: //FILL_ORIGINAL
def = src[y*src_stride_lu + x];
break;
case 2: //FILL_CLAMP
y_s = clipf(y_s, 0, height - 1);
x_s = clipf(x_s, 0, width - 1);
def = src[(int)y_s * src_stride_lu + (int)x_s];
break;
case 3: //FILL_MIRROR
y_s = mirror(y_s, height - 1);
x_s = mirror(x_s, width - 1);
def = src[(int)y_s * src_stride_lu + (int)x_s];
break;
}
switch (interpolate) {
case 0: //INTERPOLATE_NEAREST
dst[idx_dst] = interpolate_nearest(x_s, y_s, src, width, height, src_stride_lu, def);
break;
case 1: //INTERPOLATE_BILINEAR
dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, width, height, src_stride_lu, def);
break;
case 2: //INTERPOLATE_BIQUADRATIC
dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, width, height, src_stride_lu, def);
break;
default:
return;
}
}
}
kernel void avfilter_transform_chroma(global unsigned char *src,
global unsigned char *dst,
float4 matrix,
int interpolate,
int fill,
int src_stride_lu,
int dst_stride_lu,
int src_stride_ch,
int dst_stride_ch,
int height,
int width,
int ch,
int cw)
{
int x = get_global_id(0);
int y = get_global_id(1);
int pad_ch = get_global_size(1)>>1;
global unsigned char *dst_u = dst + height * dst_stride_lu;
global unsigned char *src_u = src + height * src_stride_lu;
global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
global unsigned char *src_v = src_u + ch * src_stride_ch;
src = y < pad_ch ? src_u : src_v;
dst = y < pad_ch ? dst_u : dst_v;
y = select(y - pad_ch, y, y < pad_ch);
float x_s = x * matrix.x + y * matrix.y + matrix.z;
float y_s = x * (-matrix.y) + y * matrix.x + matrix.w;
int idx_dst = y * dst_stride_ch + x;
unsigned char def;
if (x < cw && y < ch) {
switch (fill) {
case 0: //FILL_BLANK
def = 0;
break;
case 1: //FILL_ORIGINAL
def = src[y*src_stride_ch + x];
break;
case 2: //FILL_CLAMP
y_s = clipf(y_s, 0, ch - 1);
x_s = clipf(x_s, 0, cw - 1);
def = src[(int)y_s * src_stride_ch + (int)x_s];
break;
case 3: //FILL_MIRROR
y_s = mirror(y_s, ch - 1);
x_s = mirror(x_s, cw - 1);
def = src[(int)y_s * src_stride_ch + (int)x_s];
break;
}
switch (interpolate) {
case 0: //INTERPOLATE_NEAREST
dst[idx_dst] = interpolate_nearest(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
case 1: //INTERPOLATE_BILINEAR
dst[idx_dst] = interpolate_bilinear(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
case 2: //INTERPOLATE_BIQUADRATIC
dst[idx_dst] = interpolate_biquadratic(x_s, y_s, src, cw, ch, src_stride_ch, def);
break;
default:
return;
}
}
}
);
#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */

41
libavfilter/opencl_allkernels.c
Unescape View File

@ -1,41 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@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_allkernels.h"
#if CONFIG_OPENCL
#include "libavutil/opencl.h"
#include "deshake_opencl_kernel.h"
#include "unsharp_opencl_kernel.h"
#endif
#define OPENCL_REGISTER_KERNEL_CODE(X, x) \
{ \
if (CONFIG_##X##_FILTER) { \
av_opencl_register_kernel_code(ff_kernel_##x##_opencl); \
} \
}
void ff_opencl_register_filter_kernel_code_all(void)
{
#if CONFIG_OPENCL
OPENCL_REGISTER_KERNEL_CODE(DESHAKE, deshake);
OPENCL_REGISTER_KERNEL_CODE(UNSHARP, unsharp);
#endif
}

29
libavfilter/opencl_allkernels.h
Unescape View File

@ -1,29 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@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_OPENCL_ALLKERNELS_H
#define AVFILTER_OPENCL_ALLKERNELS_H
#include "avfilter.h"
#include "config.h"
void ff_opencl_register_filter_kernel_code_all(void);
#endif /* AVFILTER_OPENCL_ALLKERNELS_H */

31
libavfilter/unsharp.h
Unescape View File

@ -24,38 +24,10 @@
#include "config.h" #include "config.h"
#include "avfilter.h" #include "avfilter.h"
#if CONFIG_OPENCL
#include "libavutil/opencl.h"
#endif
#define MIN_MATRIX_SIZE 3 #define MIN_MATRIX_SIZE 3
#define MAX_MATRIX_SIZE 63 #define MAX_MATRIX_SIZE 63
#if CONFIG_OPENCL
typedef struct UnsharpOpenclContext {
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel_default;
cl_kernel kernel_luma;
cl_kernel kernel_chroma;
cl_mem cl_luma_mask;
cl_mem cl_chroma_mask;
cl_mem cl_luma_mask_x;
cl_mem cl_chroma_mask_x;
cl_mem cl_luma_mask_y;
cl_mem cl_chroma_mask_y;
int in_plane_size[8];
int out_plane_size[8];
int plane_num;
cl_mem cl_inbuf;
size_t cl_inbuf_size;
cl_mem cl_outbuf;
size_t cl_outbuf_size;
int use_fast_kernels;
} UnsharpOpenclContext;
#endif
typedef struct UnsharpFilterParam { typedef struct UnsharpFilterParam {
int msize_x; ///< matrix width int msize_x; ///< matrix width
@ -76,9 +48,6 @@ typedef struct UnsharpContext {
UnsharpFilterParam chroma; ///< chroma parameters (width, height, amount) UnsharpFilterParam chroma; ///< chroma parameters (width, height, amount)
int hsub, vsub; int hsub, vsub;
int opencl; int opencl;
#if CONFIG_OPENCL
UnsharpOpenclContext opencl_ctx;
#endif
int (* apply_unsharp)(AVFilterContext *ctx, AVFrame *in, AVFrame *out); int (* apply_unsharp)(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
} UnsharpContext; } UnsharpContext;

422
libavfilter/unsharp_opencl.c
Unescape View File

@ -1,422 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang
*
* 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
* unsharp input video
*/
#include "unsharp_opencl.h"
#include "libavutil/common.h"
#include "libavutil/opencl_internal.h"
#define PLANE_NUM 3
#define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
static inline void add_mask_counter(uint32_t *dst, uint32_t *counter1, uint32_t *counter2, int len)
{
int i;
for (i = 0; i < len; i++) {
dst[i] = counter1[i] + counter2[i];
}
}
static int compute_mask(int step, uint32_t *mask)
{
int i, z, ret = 0;
int counter_size = sizeof(uint32_t) * (2 * step + 1);
uint32_t *temp1_counter, *temp2_counter, **counter = NULL;
temp1_counter = av_mallocz(counter_size);
if (!temp1_counter) {
ret = AVERROR(ENOMEM);
goto end;
}
temp2_counter = av_mallocz(counter_size);
if (!temp2_counter) {
ret = AVERROR(ENOMEM);
goto end;
}
counter = av_mallocz_array(2 * step + 1, sizeof(uint32_t *));
if (!counter) {
ret = AVERROR(ENOMEM);
goto end;
}
for (i = 0; i < 2 * step + 1; i++) {
counter[i] = av_mallocz(counter_size);
if (!counter[i]) {
ret = AVERROR(ENOMEM);
goto end;
}
}
for (i = 0; i < 2 * step + 1; i++) {
memset(temp1_counter, 0, counter_size);
temp1_counter[i] = 1;
for (z = 0; z < step * 2; z += 2) {
add_mask_counter(temp2_counter, counter[z], temp1_counter, step * 2);
memcpy(counter[z], temp1_counter, counter_size);
add_mask_counter(temp1_counter, counter[z + 1], temp2_counter, step * 2);
memcpy(counter[z + 1], temp2_counter, counter_size);
}
}
memcpy(mask, temp1_counter, counter_size);
end:
av_freep(&temp1_counter);
av_freep(&temp2_counter);
for (i = 0; counter && i < 2 * step + 1; i++) {
av_freep(&counter[i]);
}
av_freep(&counter);
return ret;
}
static int copy_separable_masks(cl_mem cl_mask_x, cl_mem cl_mask_y, int step_x, int step_y)
{
int ret = 0;
uint32_t *mask_x, *mask_y;
size_t size_mask_x = sizeof(uint32_t) * (2 * step_x + 1);
size_t size_mask_y = sizeof(uint32_t) * (2 * step_y + 1);
mask_x = av_mallocz_array(2 * step_x + 1, sizeof(uint32_t));
if (!mask_x) {
ret = AVERROR(ENOMEM);
goto end;
}
mask_y = av_mallocz_array(2 * step_y + 1, sizeof(uint32_t));
if (!mask_y) {
ret = AVERROR(ENOMEM);
goto end;
}
ret = compute_mask(step_x, mask_x);
if (ret < 0)
goto end;
ret = compute_mask(step_y, mask_y);
if (ret < 0)
goto end;
ret = av_opencl_buffer_write(cl_mask_x, (uint8_t *)mask_x, size_mask_x);
ret = av_opencl_buffer_write(cl_mask_y, (uint8_t *)mask_y, size_mask_y);
end:
av_freep(&mask_x);
av_freep(&mask_y);
return ret;
}
static int generate_mask(AVFilterContext *ctx)
{
cl_mem masks[4];
cl_mem mask_matrix[2];
int i, ret = 0, step_x[2], step_y[2];
UnsharpContext *unsharp = ctx->priv;
mask_matrix[0] = unsharp->opencl_ctx.cl_luma_mask;
mask_matrix[1] = unsharp->opencl_ctx.cl_chroma_mask;
masks[0] = unsharp->opencl_ctx.cl_luma_mask_x;
masks[1] = unsharp->opencl_ctx.cl_luma_mask_y;
masks[2] = unsharp->opencl_ctx.cl_chroma_mask_x;
masks[3] = unsharp->opencl_ctx.cl_chroma_mask_y;
step_x[0] = unsharp->luma.steps_x;
step_x[1] = unsharp->chroma.steps_x;
step_y[0] = unsharp->luma.steps_y;
step_y[1] = unsharp->chroma.steps_y;
/* use default kernel if any matrix dim larger than 8 due to limited local mem size */
if (step_x[0]>8 || step_x[1]>8 || step_y[0]>8 || step_y[1]>8)
unsharp->opencl_ctx.use_fast_kernels = 0;
else
unsharp->opencl_ctx.use_fast_kernels = 1;
if (!masks[0] || !masks[1] || !masks[2] || !masks[3]) {
av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");
return AVERROR(EINVAL);
}
if (!mask_matrix[0] || !mask_matrix[1]) {
av_log(ctx, AV_LOG_ERROR, "Luma mask and chroma mask should not be NULL\n");
return AVERROR(EINVAL);
}
for (i = 0; i < 2; i++) {
ret = copy_separable_masks(masks[2*i], masks[2*i+1], step_x[i], step_y[i]);
if (ret < 0)
return ret;
}
return ret;
}
int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
{
int ret;
AVFilterLink *link = ctx->inputs[0];
UnsharpContext *unsharp = ctx->priv;
cl_int status;
FFOpenclParam kernel1 = {0};
FFOpenclParam kernel2 = {0};
int width = link->w;
int height = link->h;
int cw = AV_CEIL_RSHIFT(link->w, unsharp->hsub);
int ch = AV_CEIL_RSHIFT(link->h, unsharp->vsub);
size_t globalWorkSize1d = width * height + 2 * ch * cw;
size_t globalWorkSize2dLuma[2];
size_t globalWorkSize2dChroma[2];
size_t localWorkSize2d[2] = {16, 16};
if (unsharp->opencl_ctx.use_fast_kernels) {
globalWorkSize2dLuma[0] = (size_t)ROUND_TO_16(width);
globalWorkSize2dLuma[1] = (size_t)ROUND_TO_16(height);
globalWorkSize2dChroma[0] = (size_t)ROUND_TO_16(cw);
globalWorkSize2dChroma[1] = (size_t)(2*ROUND_TO_16(ch));
kernel1.ctx = ctx;
kernel1.kernel = unsharp->opencl_ctx.kernel_luma;
ret = avpriv_opencl_set_parameter(&kernel1,
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_x),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask_y),
FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
FF_OPENCL_PARAM_INFO(in->linesize[0]),
FF_OPENCL_PARAM_INFO(out->linesize[0]),
FF_OPENCL_PARAM_INFO(width),
FF_OPENCL_PARAM_INFO(height),
NULL);
if (ret < 0)
return ret;
kernel2.ctx = ctx;
kernel2.kernel = unsharp->opencl_ctx.kernel_chroma;
ret = avpriv_opencl_set_parameter(&kernel2,
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_x),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask_y),
FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
FF_OPENCL_PARAM_INFO(in->linesize[0]),
FF_OPENCL_PARAM_INFO(in->linesize[1]),
FF_OPENCL_PARAM_INFO(out->linesize[0]),
FF_OPENCL_PARAM_INFO(out->linesize[1]),
FF_OPENCL_PARAM_INFO(link->w),
FF_OPENCL_PARAM_INFO(link->h),
FF_OPENCL_PARAM_INFO(cw),
FF_OPENCL_PARAM_INFO(ch),
NULL);
if (ret < 0)
return ret;
status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
unsharp->opencl_ctx.kernel_luma, 2, NULL,
globalWorkSize2dLuma, localWorkSize2d, 0, NULL, NULL);
status |=clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
unsharp->opencl_ctx.kernel_chroma, 2, NULL,
globalWorkSize2dChroma, localWorkSize2d, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
} else { /* use default kernel */
kernel1.ctx = ctx;
kernel1.kernel = unsharp->opencl_ctx.kernel_default;
ret = avpriv_opencl_set_parameter(&kernel1,
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_inbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_outbuf),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_luma_mask),
FF_OPENCL_PARAM_INFO(unsharp->opencl_ctx.cl_chroma_mask),
FF_OPENCL_PARAM_INFO(unsharp->luma.amount),
FF_OPENCL_PARAM_INFO(unsharp->chroma.amount),
FF_OPENCL_PARAM_INFO(unsharp->luma.steps_x),
FF_OPENCL_PARAM_INFO(unsharp->luma.steps_y),
FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_x),
FF_OPENCL_PARAM_INFO(unsharp->chroma.steps_y),
FF_OPENCL_PARAM_INFO(unsharp->luma.scalebits),
FF_OPENCL_PARAM_INFO(unsharp->chroma.scalebits),
FF_OPENCL_PARAM_INFO(unsharp->luma.halfscale),
FF_OPENCL_PARAM_INFO(unsharp->chroma.halfscale),
FF_OPENCL_PARAM_INFO(in->linesize[0]),
FF_OPENCL_PARAM_INFO(in->linesize[1]),
FF_OPENCL_PARAM_INFO(out->linesize[0]),
FF_OPENCL_PARAM_INFO(out->linesize[1]),
FF_OPENCL_PARAM_INFO(link->h),
FF_OPENCL_PARAM_INFO(link->w),
FF_OPENCL_PARAM_INFO(ch),
FF_OPENCL_PARAM_INFO(cw),
NULL);
if (ret < 0)
return ret;
status = clEnqueueNDRangeKernel(unsharp->opencl_ctx.command_queue,
unsharp->opencl_ctx.kernel_default, 1, NULL,
&globalWorkSize1d, NULL, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
}
//blocking map is suffficient, no need for clFinish
//clFinish(unsharp->opencl_ctx.command_queue);
return av_opencl_buffer_read_image(out->data, unsharp->opencl_ctx.out_plane_size,
unsharp->opencl_ctx.plane_num, unsharp->opencl_ctx.cl_outbuf,
unsharp->opencl_ctx.cl_outbuf_size);
}
int ff_opencl_unsharp_init(AVFilterContext *ctx)
{
int ret = 0;
char build_opts[96];
UnsharpContext *unsharp = ctx->priv;
ret = av_opencl_init(NULL);
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask,
sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1) * (2 * unsharp->luma.steps_y + 1),
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
CL_MEM_READ_ONLY, NULL);
// separable filters
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_x,
sizeof(uint32_t) * (2 * unsharp->luma.steps_x + 1),
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_luma_mask_y,
sizeof(uint32_t) * (2 * unsharp->luma.steps_y + 1),
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_x,
sizeof(uint32_t) * (2 * unsharp->chroma.steps_x + 1),
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask_y,
sizeof(uint32_t) * (2 * unsharp->chroma.steps_y + 1),
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
ret = generate_mask(ctx);
if (ret < 0)
return ret;
unsharp->opencl_ctx.plane_num = PLANE_NUM;
unsharp->opencl_ctx.command_queue = av_opencl_get_command_queue();
if (!unsharp->opencl_ctx.command_queue) {
av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue in filter 'unsharp'\n");
return AVERROR(EINVAL);
}
snprintf(build_opts, 96, "-D LU_RADIUS_X=%d -D LU_RADIUS_Y=%d -D CH_RADIUS_X=%d -D CH_RADIUS_Y=%d",
2*unsharp->luma.steps_x+1, 2*unsharp->luma.steps_y+1, 2*unsharp->chroma.steps_x+1, 2*unsharp->chroma.steps_y+1);
unsharp->opencl_ctx.program = av_opencl_compile("unsharp", build_opts);
if (!unsharp->opencl_ctx.program) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program 'unsharp'\n");
return AVERROR(EINVAL);
}
if (unsharp->opencl_ctx.use_fast_kernels) {
if (!unsharp->opencl_ctx.kernel_luma) {
unsharp->opencl_ctx.kernel_luma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_luma", &ret);
if (ret != CL_SUCCESS) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_luma'\n");
return ret;
}
}
if (!unsharp->opencl_ctx.kernel_chroma) {
unsharp->opencl_ctx.kernel_chroma = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_chroma", &ret);
if (ret < 0) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_chroma'\n");
return ret;
}
}
}
else {
if (!unsharp->opencl_ctx.kernel_default) {
unsharp->opencl_ctx.kernel_default = clCreateKernel(unsharp->opencl_ctx.program, "unsharp_default", &ret);
if (ret < 0) {
av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel 'unsharp_default'\n");
return ret;
}
}
}
return ret;
}
void ff_opencl_unsharp_uninit(AVFilterContext *ctx)
{
UnsharpContext *unsharp = ctx->priv;
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_inbuf);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_outbuf);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_x);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_x);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_luma_mask_y);
av_opencl_buffer_release(&unsharp->opencl_ctx.cl_chroma_mask_y);
clReleaseKernel(unsharp->opencl_ctx.kernel_default);
clReleaseKernel(unsharp->opencl_ctx.kernel_luma);
clReleaseKernel(unsharp->opencl_ctx.kernel_chroma);
clReleaseProgram(unsharp->opencl_ctx.program);
unsharp->opencl_ctx.command_queue = NULL;
av_opencl_uninit();
}
int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out)
{
int ret = 0;
AVFilterLink *link = ctx->inputs[0];
UnsharpContext *unsharp = ctx->priv;
int ch = AV_CEIL_RSHIFT(link->h, unsharp->vsub);
if ((!unsharp->opencl_ctx.cl_inbuf) || (!unsharp->opencl_ctx.cl_outbuf)) {
unsharp->opencl_ctx.in_plane_size[0] = (in->linesize[0] * in->height);
unsharp->opencl_ctx.in_plane_size[1] = (in->linesize[1] * ch);
unsharp->opencl_ctx.in_plane_size[2] = (in->linesize[2] * ch);
unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] * out->height);
unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);
unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);
unsharp->opencl_ctx.cl_inbuf_size = unsharp->opencl_ctx.in_plane_size[0] +
unsharp->opencl_ctx.in_plane_size[1] +
unsharp->opencl_ctx.in_plane_size[2];
unsharp->opencl_ctx.cl_outbuf_size = unsharp->opencl_ctx.out_plane_size[0] +
unsharp->opencl_ctx.out_plane_size[1] +
unsharp->opencl_ctx.out_plane_size[2];
if (!unsharp->opencl_ctx.cl_inbuf) {
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_inbuf,
unsharp->opencl_ctx.cl_inbuf_size,
CL_MEM_READ_ONLY, NULL);
if (ret < 0)
return ret;
}
if (!unsharp->opencl_ctx.cl_outbuf) {
ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_outbuf,
unsharp->opencl_ctx.cl_outbuf_size,
CL_MEM_READ_WRITE, NULL);
if (ret < 0)
return ret;
}
}
return av_opencl_buffer_write_image(unsharp->opencl_ctx.cl_inbuf,
unsharp->opencl_ctx.cl_inbuf_size,
0, in->data, unsharp->opencl_ctx.in_plane_size,
unsharp->opencl_ctx.plane_num);
}

34
libavfilter/unsharp_opencl.h
Unescape View File

@ -1,34 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@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_UNSHARP_OPENCL_H
#define AVFILTER_UNSHARP_OPENCL_H
#include "unsharp.h"
int ff_opencl_unsharp_init(AVFilterContext *ctx);
void ff_opencl_unsharp_uninit(AVFilterContext *ctx);
int ff_opencl_unsharp_process_inout_buf(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
int ff_opencl_apply_unsharp(AVFilterContext *ctx, AVFrame *in, AVFrame *out);
#endif /* AVFILTER_UNSHARP_OPENCL_H */

342
libavfilter/unsharp_opencl_kernel.h
Unescape View File

@ -1,342 +0,0 @@
/*
* Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang
*
* 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_UNSHARP_OPENCL_KERNEL_H
#define AVFILTER_UNSHARP_OPENCL_KERNEL_H
#include "libavutil/opencl.h"
const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
inline unsigned char clip_uint8(int a)
{
if (a & (~0xFF))
return (-a)>>31;
else
return a;
}
kernel void unsharp_luma(
global unsigned char *src,
global unsigned char *dst,
global int *mask_x,
global int *mask_y,
int amount,
int scalebits,
int halfscale,
int src_stride,
int dst_stride,
int width,
int height)
{
int2 threadIdx, blockIdx, globalIdx;
threadIdx.x = get_local_id(0);
threadIdx.y = get_local_id(1);
blockIdx.x = get_group_id(0);
blockIdx.y = get_group_id(1);
globalIdx.x = get_global_id(0);
globalIdx.y = get_global_id(1);
if (!amount) {
if (globalIdx.x < width && globalIdx.y < height)
dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
return;
}
local unsigned int l[32][32];
local unsigned int lcx[LU_RADIUS_X];
local unsigned int lcy[LU_RADIUS_Y];
int indexIx, indexIy, i, j;
//load up tile: actual workspace + halo of 8 points in x and y \n
for(i = 0; i <= 1; i++) {
indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
indexIy = indexIy < 0 ? 0 : indexIy;
indexIy = indexIy >= height ? height - 1: indexIy;
for(j = 0; j <= 1; j++) {
indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
indexIx = indexIx < 0 ? 0 : indexIx;
indexIx = indexIx >= width ? width - 1: indexIx;
l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
}
}
int indexL = threadIdx.y*16 + threadIdx.x;
if (indexL < LU_RADIUS_X)
lcx[indexL] = mask_x[indexL];
if (indexL < LU_RADIUS_Y)
lcy[indexL] = mask_y[indexL];
barrier(CLK_LOCAL_MEM_FENCE);
//needed for unsharp mask application in the end \n
int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
int idx, idy, maskIndex;
int temp[2] = {0};
int steps_x = (LU_RADIUS_X-1)/2;
int steps_y = (LU_RADIUS_Y-1)/2;
// compute the actual workspace + left&right halos \n
\n#pragma unroll\n
for (j = 0; j <=1; j++) {
//extra work to cover left and right halos \n
idx = 16*j + threadIdx.x;
\n#pragma unroll\n
for (i = -steps_y; i <= steps_y; i++) {
idy = 8 + i + threadIdx.y;
maskIndex = (i + steps_y);
temp[j] += (int)l[idy][idx] * lcy[maskIndex];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
//save results from the vertical filter in local memory \n
idy = 8 + threadIdx.y;
\n#pragma unroll\n
for (j = 0; j <=1; j++) {
idx = 16*j + threadIdx.x;
l[idy][idx] = temp[j];
}
barrier(CLK_LOCAL_MEM_FENCE);
//compute results with the horizontal filter \n
int sum = 0;
idy = 8 + threadIdx.y;
\n#pragma unroll\n
for (j = -steps_x; j <= steps_x; j++) {
idx = 8 + j + threadIdx.x;
maskIndex = j + steps_x;
sum += (int)l[idy][idx] * lcx[maskIndex];
}
int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
if (globalIdx.x < width && globalIdx.y < height)
dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
}
kernel void unsharp_chroma(
global unsigned char *src_y,
global unsigned char *dst_y,
global int *mask_x,
global int *mask_y,
int amount,
int scalebits,
int halfscale,
int src_stride_lu,
int src_stride_ch,
int dst_stride_lu,
int dst_stride_ch,
int width,
int height,
int cw,
int ch)
{
global unsigned char *dst_u = dst_y + height * dst_stride_lu;
global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
global unsigned char *src_u = src_y + height * src_stride_lu;
global unsigned char *src_v = src_u + ch * src_stride_ch;
int2 threadIdx, blockIdx, globalIdx;
threadIdx.x = get_local_id(0);
threadIdx.y = get_local_id(1);
blockIdx.x = get_group_id(0);
blockIdx.y = get_group_id(1);
globalIdx.x = get_global_id(0);
globalIdx.y = get_global_id(1);
int padch = get_global_size(1)/2;
global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;
blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;
if (!amount) {
if (globalIdx.x < cw && globalIdx.y < ch)
dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
return;
}
local unsigned int l[32][32];
local unsigned int lcx[CH_RADIUS_X];
local unsigned int lcy[CH_RADIUS_Y];
int indexIx, indexIy, i, j;
for(i = 0; i <= 1; i++) {
indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
indexIy = indexIy < 0 ? 0 : indexIy;
indexIy = indexIy >= ch ? ch - 1: indexIy;
for(j = 0; j <= 1; j++) {
indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
indexIx = indexIx < 0 ? 0 : indexIx;
indexIx = indexIx >= cw ? cw - 1: indexIx;
l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
}
}
int indexL = threadIdx.y*16 + threadIdx.x;
if (indexL < CH_RADIUS_X)
lcx[indexL] = mask_x[indexL];
if (indexL < CH_RADIUS_Y)
lcy[indexL] = mask_y[indexL];
barrier(CLK_LOCAL_MEM_FENCE);
int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
int idx, idy, maskIndex;
int steps_x = CH_RADIUS_X/2;
int steps_y = CH_RADIUS_Y/2;
int temp[2] = {0,0};
\n#pragma unroll\n
for (j = 0; j <= 1; j++) {
idx = 16*j + threadIdx.x;
\n#pragma unroll\n
for (i = -steps_y; i <= steps_y; i++) {
idy = 8 + i + threadIdx.y;
maskIndex = i + steps_y;
temp[j] += (int)l[idy][idx] * lcy[maskIndex];
}
}
barrier(CLK_LOCAL_MEM_FENCE);
idy = 8 + threadIdx.y;
\n#pragma unroll\n
for (j = 0; j <= 1; j++) {
idx = 16*j + threadIdx.x;
l[idy][idx] = temp[j];
}
barrier(CLK_LOCAL_MEM_FENCE);
//compute results with the horizontal filter \n
int sum = 0;
idy = 8 + threadIdx.y;
\n#pragma unroll\n
for (j = -steps_x; j <= steps_x; j++) {
idx = 8 + j + threadIdx.x;
maskIndex = j + steps_x;
sum += (int)l[idy][idx] * lcx[maskIndex];
}
int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
if (globalIdx.x < cw && globalIdx.y < ch)
dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
}
kernel void unsharp_default(global unsigned char *src,
global unsigned char *dst,
const global unsigned int *mask_lu,
const global unsigned int *mask_ch,
int amount_lu,
int amount_ch,
int step_x_lu,
int step_y_lu,
int step_x_ch,
int step_y_ch,
int scalebits_lu,
int scalebits_ch,
int halfscale_lu,
int halfscale_ch,
int src_stride_lu,
int src_stride_ch,
int dst_stride_lu,
int dst_stride_ch,
int height,
int width,
int ch,
int cw)
{
global unsigned char *dst_y = dst;
global unsigned char *dst_u = dst_y + height * dst_stride_lu;
global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
global unsigned char *src_y = src;
global unsigned char *src_u = src_y + height * src_stride_lu;
global unsigned char *src_v = src_u + ch * src_stride_ch;
global unsigned char *temp_dst;
global unsigned char *temp_src;
const global unsigned int *temp_mask;
int global_id = get_global_id(0);
int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
if (global_id < width * height) {
y = global_id / width;
x = global_id % width;
temp_dst = dst_y;
temp_src = src_y;
temp_src_stride = src_stride_lu;
temp_dst_stride = dst_stride_lu;
temp_height = height;
temp_width = width;
temp_steps_x = step_x_lu;
temp_steps_y = step_y_lu;
temp_mask = mask_lu;
temp_amount = amount_lu;
temp_scalebits = scalebits_lu;
temp_halfscale = halfscale_lu;
} else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
y = (global_id - width * height) / cw;
x = (global_id - width * height) % cw;
temp_dst = dst_u;
temp_src = src_u;
temp_src_stride = src_stride_ch;
temp_dst_stride = dst_stride_ch;
temp_height = ch;
temp_width = cw;
temp_steps_x = step_x_ch;
temp_steps_y = step_y_ch;
temp_mask = mask_ch;
temp_amount = amount_ch;
temp_scalebits = scalebits_ch;
temp_halfscale = halfscale_ch;
} else {
y = (global_id - width * height - ch * cw) / cw;
x = (global_id - width * height - ch * cw) % cw;
temp_dst = dst_v;
temp_src = src_v;
temp_src_stride = src_stride_ch;
temp_dst_stride = dst_stride_ch;
temp_height = ch;
temp_width = cw;
temp_steps_x = step_x_ch;
temp_steps_y = step_y_ch;
temp_mask = mask_ch;
temp_amount = amount_ch;
temp_scalebits = scalebits_ch;
temp_halfscale = halfscale_ch;
}
if (temp_amount) {
sum = 0;
for (j = 0; j <= 2 * temp_steps_y; j++) {
idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
for (i = 0; i <= 2 * temp_steps_x; i++) {
idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
}
}
temp = (int)temp_src[x + y * temp_src_stride];
res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
} else {
temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
}
}
);
#endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */

23
libavfilter/vf_deshake.c
Unescape View File

@ -60,7 +60,6 @@
#include "libavutil/qsort.h" #include "libavutil/qsort.h"
#include "deshake.h" #include "deshake.h"
#include "deshake_opencl.h"
#define OFFSET(x) offsetof(DeshakeContext, x) #define OFFSET(x) offsetof(DeshakeContext, x)
#define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM
@ -83,7 +82,7 @@ static const AVOption deshake_options[] = {
{ "exhaustive", "exhaustive search", 0, AV_OPT_TYPE_CONST, {.i64=EXHAUSTIVE}, INT_MIN, INT_MAX, FLAGS, "smode" }, { "exhaustive", "exhaustive search", 0, AV_OPT_TYPE_CONST, {.i64=EXHAUSTIVE}, INT_MIN, INT_MAX, FLAGS, "smode" },
{ "less", "less exhaustive search", 0, AV_OPT_TYPE_CONST, {.i64=SMART_EXHAUSTIVE}, INT_MIN, INT_MAX, FLAGS, "smode" }, { "less", "less exhaustive search", 0, AV_OPT_TYPE_CONST, {.i64=SMART_EXHAUSTIVE}, INT_MIN, INT_MAX, FLAGS, "smode" },
{ "filename", "set motion search detailed log file name", OFFSET(filename), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, { "filename", "set motion search detailed log file name", OFFSET(filename), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS },
{ "opencl", "use OpenCL filtering capabilities", OFFSET(opencl), AV_OPT_TYPE_BOOL, {.i64=0}, 0, 1, .flags = FLAGS }, { "opencl", "ignored", OFFSET(opencl), AV_OPT_TYPE_BOOL, {.i64=0}, 0, 1, .flags = FLAGS },
{ NULL } { NULL }
}; };
@ -341,7 +340,6 @@ static int deshake_transform_c(AVFilterContext *ctx,
static av_cold int init(AVFilterContext *ctx) static av_cold int init(AVFilterContext *ctx)
{ {
int ret;
DeshakeContext *deshake = ctx->priv; DeshakeContext *deshake = ctx->priv;
deshake->sad = av_pixelutils_get_sad_fn(4, 4, 1, deshake); // 16x16, 2nd source unaligned deshake->sad = av_pixelutils_get_sad_fn(4, 4, 1, deshake); // 16x16, 2nd source unaligned
@ -369,17 +367,7 @@ static av_cold int init(AVFilterContext *ctx)
deshake->cx &= ~15; deshake->cx &= ~15;
} }
deshake->transform = deshake_transform_c; deshake->transform = deshake_transform_c;
if (!CONFIG_OPENCL && deshake->opencl) {
av_log(ctx, AV_LOG_ERROR, "OpenCL support was not enabled in this build, cannot be selected\n");
return AVERROR(EINVAL);
}
if (CONFIG_OPENCL && deshake->opencl) {
deshake->transform = ff_opencl_transform;
ret = ff_opencl_deshake_init(ctx);
if (ret < 0)
return ret;
}
av_log(ctx, AV_LOG_VERBOSE, "cx: %d, cy: %d, cw: %d, ch: %d, rx: %d, ry: %d, edge: %d blocksize: %d contrast: %d search: %d\n", av_log(ctx, AV_LOG_VERBOSE, "cx: %d, cy: %d, cw: %d, ch: %d, rx: %d, ry: %d, edge: %d blocksize: %d contrast: %d search: %d\n",
deshake->cx, deshake->cy, deshake->cw, deshake->ch, deshake->cx, deshake->cy, deshake->cw, deshake->ch,
deshake->rx, deshake->ry, deshake->edge, deshake->blocksize * 2, deshake->contrast, deshake->search); deshake->rx, deshake->ry, deshake->edge, deshake->blocksize * 2, deshake->contrast, deshake->search);
@ -416,9 +404,6 @@ static int config_props(AVFilterLink *link)
static av_cold void uninit(AVFilterContext *ctx) static av_cold void uninit(AVFilterContext *ctx)
{ {
DeshakeContext *deshake = ctx->priv; DeshakeContext *deshake = ctx->priv;
if (CONFIG_OPENCL && deshake->opencl) {
ff_opencl_deshake_uninit(ctx);
}
av_frame_free(&deshake->ref); av_frame_free(&deshake->ref);
av_freep(&deshake->angles); av_freep(&deshake->angles);
deshake->angles_size = 0; deshake->angles_size = 0;
@ -447,12 +432,6 @@ static int filter_frame(AVFilterLink *link, AVFrame *in)
} }
av_frame_copy_props(out, in); av_frame_copy_props(out, in);
if (CONFIG_OPENCL && deshake->opencl) {
ret = ff_opencl_deshake_process_inout_buf(link->dst,in, out);
if (ret < 0)
goto fail;
}
if (deshake->cx < 0 || deshake->cy < 0 || deshake->cw < 0 || deshake->ch < 0) { if (deshake->cx < 0 || deshake->cy < 0 || deshake->cw < 0 || deshake->ch < 0) {
// Find the most likely global motion for the current frame // Find the most likely global motion for the current frame
find_motion(deshake, (deshake->ref == NULL) ? in->data[0] : deshake->ref->data[0], in->data[0], link->w, link->h, in->linesize[0], &t); find_motion(deshake, (deshake->ref == NULL) ? in->data[0] : deshake->ref->data[0], in->data[0], link->w, link->h, in->linesize[0], &t);

26
libavfilter/vf_unsharp.c
Unescape View File

@ -46,7 +46,6 @@
#include "libavutil/opt.h" #include "libavutil/opt.h"
#include "libavutil/pixdesc.h" #include "libavutil/pixdesc.h"
#include "unsharp.h" #include "unsharp.h"
#include "unsharp_opencl.h"
static void apply_unsharp( uint8_t *dst, int dst_stride, static void apply_unsharp( uint8_t *dst, int dst_stride,
const uint8_t *src, int src_stride, const uint8_t *src, int src_stride,
@ -134,10 +133,8 @@ static void set_filter_param(UnsharpFilterParam *fp, int msize_x, int msize_y, f
static av_cold int init(AVFilterContext *ctx) static av_cold int init(AVFilterContext *ctx)
{ {
int ret = 0;
UnsharpContext *s = ctx->priv; UnsharpContext *s = ctx->priv;
set_filter_param(&s->luma, s->lmsize_x, s->lmsize_y, s->lamount); set_filter_param(&s->luma, s->lmsize_x, s->lmsize_y, s->lamount);
set_filter_param(&s->chroma, s->cmsize_x, s->cmsize_y, s->camount); set_filter_param(&s->chroma, s->cmsize_x, s->cmsize_y, s->camount);
@ -146,16 +143,6 @@ static av_cold int init(AVFilterContext *ctx)
return AVERROR(EINVAL); return AVERROR(EINVAL);
} }
s->apply_unsharp = apply_unsharp_c; s->apply_unsharp = apply_unsharp_c;
if (!CONFIG_OPENCL && s->opencl) {
av_log(ctx, AV_LOG_ERROR, "OpenCL support was not enabled in this build, cannot be selected\n");
return AVERROR(EINVAL);
}
if (CONFIG_OPENCL && s->opencl) {
s->apply_unsharp = ff_opencl_apply_unsharp;
ret = ff_opencl_unsharp_init(ctx);
if (ret < 0)
return ret;
}
return 0; return 0;
} }
@ -227,10 +214,6 @@ static av_cold void uninit(AVFilterContext *ctx)
{ {
UnsharpContext *s = ctx->priv; UnsharpContext *s = ctx->priv;
if (CONFIG_OPENCL && s->opencl) {
ff_opencl_unsharp_uninit(ctx);
}
free_filter_param(&s->luma); free_filter_param(&s->luma);
free_filter_param(&s->chroma); free_filter_param(&s->chroma);
} }
@ -248,14 +231,9 @@ static int filter_frame(AVFilterLink *link, AVFrame *in)
return AVERROR(ENOMEM); return AVERROR(ENOMEM);
} }
av_frame_copy_props(out, in); av_frame_copy_props(out, in);
if (CONFIG_OPENCL && s->opencl) {
ret = ff_opencl_unsharp_process_inout_buf(link->dst, in, out);
if (ret < 0)
goto end;
}
ret = s->apply_unsharp(link->dst, in, out); ret = s->apply_unsharp(link->dst, in, out);
end:
av_frame_free(&in); av_frame_free(&in);
if (ret < 0) { if (ret < 0) {
@ -282,7 +260,7 @@ static const AVOption unsharp_options[] = {
{ "cy", "set chroma matrix vertical size", OFFSET(cmsize_y), AV_OPT_TYPE_INT, { .i64 = 5 }, MIN_SIZE, MAX_SIZE, FLAGS }, { "cy", "set chroma matrix vertical size", OFFSET(cmsize_y), AV_OPT_TYPE_INT, { .i64 = 5 }, MIN_SIZE, MAX_SIZE, FLAGS },
{ "chroma_amount", "set chroma effect strength", OFFSET(camount), AV_OPT_TYPE_FLOAT, { .dbl = 0 }, -2, 5, FLAGS }, { "chroma_amount", "set chroma effect strength", OFFSET(camount), AV_OPT_TYPE_FLOAT, { .dbl = 0 }, -2, 5, FLAGS },
{ "ca", "set chroma effect strength", OFFSET(camount), AV_OPT_TYPE_FLOAT, { .dbl = 0 }, -2, 5, FLAGS }, { "ca", "set chroma effect strength", OFFSET(camount), AV_OPT_TYPE_FLOAT, { .dbl = 0 }, -2, 5, FLAGS },
{ "opencl", "use OpenCL filtering capabilities", OFFSET(opencl), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS }, { "opencl", "ignored", OFFSET(opencl), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
{ NULL } { NULL }
}; };

4
libavutil/Makefile
Unescape View File

@ -79,8 +79,6 @@ HEADERS = adler32.h \
HEADERS-$(CONFIG_LZO) += lzo.h HEADERS-$(CONFIG_LZO) += lzo.h
HEADERS-$(CONFIG_OPENCL) += opencl.h
ARCH_HEADERS = bswap.h \ ARCH_HEADERS = bswap.h \
intmath.h \ intmath.h \
intreadwrite.h \ intreadwrite.h \
@ -164,7 +162,6 @@ OBJS-$(CONFIG_DXVA2) += hwcontext_dxva2.o
OBJS-$(CONFIG_QSV) += hwcontext_qsv.o OBJS-$(CONFIG_QSV) += hwcontext_qsv.o
OBJS-$(CONFIG_LIBDRM) += hwcontext_drm.o OBJS-$(CONFIG_LIBDRM) += hwcontext_drm.o
OBJS-$(CONFIG_LZO) += lzo.o OBJS-$(CONFIG_LZO) += lzo.o
OBJS-$(CONFIG_OPENCL) += opencl.o opencl_internal.o
OBJS-$(CONFIG_OPENCL) += hwcontext_opencl.o OBJS-$(CONFIG_OPENCL) += hwcontext_opencl.o
OBJS-$(CONFIG_VAAPI) += hwcontext_vaapi.o OBJS-$(CONFIG_VAAPI) += hwcontext_vaapi.o
OBJS-$(CONFIG_VIDEOTOOLBOX) += hwcontext_videotoolbox.o OBJS-$(CONFIG_VIDEOTOOLBOX) += hwcontext_videotoolbox.o
@ -187,7 +184,6 @@ SKIPHEADERS-$(CONFIG_VDPAU) += hwcontext_vdpau.h
SKIPHEADERS-$(HAVE_ATOMICS_GCC) += atomic_gcc.h SKIPHEADERS-$(HAVE_ATOMICS_GCC) += atomic_gcc.h
SKIPHEADERS-$(HAVE_ATOMICS_SUNCC) += atomic_suncc.h SKIPHEADERS-$(HAVE_ATOMICS_SUNCC) += atomic_suncc.h
SKIPHEADERS-$(HAVE_ATOMICS_WIN32) += atomic_win32.h SKIPHEADERS-$(HAVE_ATOMICS_WIN32) += atomic_win32.h
SKIPHEADERS-$(CONFIG_OPENCL) += opencl.h
TESTPROGS = adler32 \ TESTPROGS = adler32 \
aes \ aes \

875
libavutil/opencl.c
Unescape View File

@ -1,875 +0,0 @@
/*
* Copyright (C) 2012 Peng Gao <peng@multicorewareinc.com>
* Copyright (C) 2012 Li Cao <li@multicorewareinc.com>
* Copyright (C) 2012 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang <lwanghpc@gmail.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"
#include "opt.h"
#if HAVE_THREADS
#include "thread.h"
#include "atomic.h"
static pthread_mutex_t * volatile atomic_opencl_lock = NULL;
#define LOCK_OPENCL pthread_mutex_lock(atomic_opencl_lock)
#define UNLOCK_OPENCL pthread_mutex_unlock(atomic_opencl_lock)
#else
#define LOCK_OPENCL
#define UNLOCK_OPENCL
#endif
#define MAX_KERNEL_CODE_NUM 200
typedef struct {
int is_compiled;
const char *kernel_string;
} KernelCode;
typedef struct {
const AVClass *class;
int log_offset;
void *log_ctx;
int init_count;
int opt_init_flag;
/**
* if set to 1, the OpenCL environment was created by the user and
* passed as AVOpenCLExternalEnv when initing ,0:created by opencl wrapper.
*/
int is_user_created;
int platform_idx;
int device_idx;
cl_platform_id platform_id;
cl_device_type device_type;
cl_context context;
cl_device_id device_id;
cl_command_queue command_queue;
int kernel_code_count;
KernelCode kernel_code[MAX_KERNEL_CODE_NUM];
AVOpenCLDeviceList device_list;
} OpenclContext;
#define OFFSET(x) offsetof(OpenclContext, x)
static const AVOption opencl_options[] = {
{ "platform_idx", "set platform index value", OFFSET(platform_idx), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX},
{ "device_idx", "set device index value", OFFSET(device_idx), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX},
{ NULL }
};
static const AVClass openclutils_class = {
.class_name = "opencl",
.option = opencl_options,
.item_name = av_default_item_name,
.version = LIBAVUTIL_VERSION_INT,
.log_level_offset_offset = offsetof(OpenclContext, log_offset),
.parent_log_context_offset = offsetof(OpenclContext, log_ctx),
};
static OpenclContext opencl_ctx = {&openclutils_class};
static const cl_device_type device_type[] = {CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_CPU};
typedef struct {
int err_code;
const char *err_str;
} OpenclErrorMsg;
static const OpenclErrorMsg opencl_err_msg[] = {
{CL_DEVICE_NOT_FOUND, "DEVICE NOT FOUND"},
{CL_DEVICE_NOT_AVAILABLE, "DEVICE NOT AVAILABLE"},
{CL_COMPILER_NOT_AVAILABLE, "COMPILER NOT AVAILABLE"},
{CL_MEM_OBJECT_ALLOCATION_FAILURE, "MEM OBJECT ALLOCATION FAILURE"},
{CL_OUT_OF_RESOURCES, "OUT OF RESOURCES"},
{CL_OUT_OF_HOST_MEMORY, "OUT OF HOST MEMORY"},
{CL_PROFILING_INFO_NOT_AVAILABLE, "PROFILING INFO NOT AVAILABLE"},
{CL_MEM_COPY_OVERLAP, "MEM COPY OVERLAP"},
{CL_IMAGE_FORMAT_MISMATCH, "IMAGE FORMAT MISMATCH"},
{CL_IMAGE_FORMAT_NOT_SUPPORTED, "IMAGE FORMAT NOT_SUPPORTED"},
{CL_BUILD_PROGRAM_FAILURE, "BUILD PROGRAM FAILURE"},
{CL_MAP_FAILURE, "MAP FAILURE"},
{CL_MISALIGNED_SUB_BUFFER_OFFSET, "MISALIGNED SUB BUFFER OFFSET"},
{CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST, "EXEC STATUS ERROR FOR EVENTS IN WAIT LIST"},
{CL_COMPILE_PROGRAM_FAILURE, "COMPILE PROGRAM FAILURE"},
{CL_LINKER_NOT_AVAILABLE, "LINKER NOT AVAILABLE"},
{CL_LINK_PROGRAM_FAILURE, "LINK PROGRAM FAILURE"},
{CL_DEVICE_PARTITION_FAILED, "DEVICE PARTITION FAILED"},
{CL_KERNEL_ARG_INFO_NOT_AVAILABLE, "KERNEL ARG INFO NOT AVAILABLE"},
{CL_INVALID_VALUE, "INVALID VALUE"},
{CL_INVALID_DEVICE_TYPE, "INVALID DEVICE TYPE"},
{CL_INVALID_PLATFORM, "INVALID PLATFORM"},
{CL_INVALID_DEVICE, "INVALID DEVICE"},
{CL_INVALID_CONTEXT, "INVALID CONTEXT"},
{CL_INVALID_QUEUE_PROPERTIES, "INVALID QUEUE PROPERTIES"},
{CL_INVALID_COMMAND_QUEUE, "INVALID COMMAND QUEUE"},
{CL_INVALID_HOST_PTR, "INVALID HOST PTR"},
{CL_INVALID_MEM_OBJECT, "INVALID MEM OBJECT"},
{CL_INVALID_IMAGE_FORMAT_DESCRIPTOR, "INVALID IMAGE FORMAT DESCRIPTOR"},
{CL_INVALID_IMAGE_SIZE, "INVALID IMAGE SIZE"},
{CL_INVALID_SAMPLER, "INVALID SAMPLER"},
{CL_INVALID_BINARY, "INVALID BINARY"},
{CL_INVALID_BUILD_OPTIONS, "INVALID BUILD OPTIONS"},
{CL_INVALID_PROGRAM, "INVALID PROGRAM"},
{CL_INVALID_PROGRAM_EXECUTABLE, "INVALID PROGRAM EXECUTABLE"},
{CL_INVALID_KERNEL_NAME, "INVALID KERNEL NAME"},
{CL_INVALID_KERNEL_DEFINITION, "INVALID KERNEL DEFINITION"},
{CL_INVALID_KERNEL, "INVALID KERNEL"},
{CL_INVALID_ARG_INDEX, "INVALID ARG INDEX"},
{CL_INVALID_ARG_VALUE, "INVALID ARG VALUE"},
{CL_INVALID_ARG_SIZE, "INVALID ARG_SIZE"},
{CL_INVALID_KERNEL_ARGS, "INVALID KERNEL ARGS"},
{CL_INVALID_WORK_DIMENSION, "INVALID WORK DIMENSION"},
{CL_INVALID_WORK_GROUP_SIZE, "INVALID WORK GROUP SIZE"},
{CL_INVALID_WORK_ITEM_SIZE, "INVALID WORK ITEM SIZE"},
{CL_INVALID_GLOBAL_OFFSET, "INVALID GLOBAL OFFSET"},
{CL_INVALID_EVENT_WAIT_LIST, "INVALID EVENT WAIT LIST"},
{CL_INVALID_EVENT, "INVALID EVENT"},
{CL_INVALID_OPERATION, "INVALID OPERATION"},
{CL_INVALID_GL_OBJECT, "INVALID GL OBJECT"},
{CL_INVALID_BUFFER_SIZE, "INVALID BUFFER SIZE"},
{CL_INVALID_MIP_LEVEL, "INVALID MIP LEVEL"},
{CL_INVALID_GLOBAL_WORK_SIZE, "INVALID GLOBAL WORK SIZE"},
{CL_INVALID_PROPERTY, "INVALID PROPERTY"},
{CL_INVALID_IMAGE_DESCRIPTOR, "INVALID IMAGE DESCRIPTOR"},
{CL_INVALID_COMPILER_OPTIONS, "INVALID COMPILER OPTIONS"},
{CL_INVALID_LINKER_OPTIONS, "INVALID LINKER OPTIONS"},
{CL_INVALID_DEVICE_PARTITION_COUNT, "INVALID DEVICE PARTITION COUNT"},
};
const char *av_opencl_errstr(cl_int status)
{
int i;
for (i = 0; i < FF_ARRAY_ELEMS(opencl_err_msg); i++) {
if (opencl_err_msg[i].err_code == status)
return opencl_err_msg[i].err_str;
}
return "unknown error";
}
static void free_device_list(AVOpenCLDeviceList *device_list)
{
int i, j;
if (!device_list || !device_list->platform_node)
return;
for (i = 0; i < device_list->platform_num; i++) {
if (!device_list->platform_node[i])
continue;
for (j = 0; j < device_list->platform_node[i]->device_num; j++) {
av_freep(&(device_list->platform_node[i]->device_node[j]->device_name));
av_freep(&(device_list->platform_node[i]->device_node[j]));
}
av_freep(&device_list->platform_node[i]->device_node);
av_freep(&(device_list->platform_node[i]->platform_name));
av_freep(&device_list->platform_node[i]);
}
av_freep(&device_list->platform_node);
device_list->platform_num = 0;
}
static int get_device_list(AVOpenCLDeviceList *device_list)
{
cl_int status;
int i, j, k, device_num, total_devices_num, ret = 0;
int *devices_num;
cl_platform_id *platform_ids = NULL;
cl_device_id *device_ids = NULL;
AVOpenCLDeviceNode *device_node = NULL;
size_t platform_name_size = 0;
size_t device_name_size = 0;
status = clGetPlatformIDs(0, NULL, &device_list->platform_num);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not get OpenCL platform ids: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
platform_ids = av_mallocz_array(device_list->platform_num, sizeof(cl_platform_id));
if (!platform_ids)
return AVERROR(ENOMEM);
status = clGetPlatformIDs(device_list->platform_num, platform_ids, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not get OpenCL platform ids: %s\n", av_opencl_errstr(status));
ret = AVERROR_EXTERNAL;
goto end;
}
device_list->platform_node = av_mallocz_array(device_list->platform_num, sizeof(AVOpenCLPlatformNode *));
if (!device_list->platform_node) {
ret = AVERROR(ENOMEM);
goto end;
}
devices_num = av_mallocz(sizeof(int) * FF_ARRAY_ELEMS(device_type));
if (!devices_num) {
ret = AVERROR(ENOMEM);
goto end;
}
for (i = 0; i < device_list->platform_num; i++) {
device_list->platform_node[i] = av_mallocz(sizeof(AVOpenCLPlatformNode));
if (!device_list->platform_node[i]) {
ret = AVERROR(ENOMEM);
goto end;
}
device_list->platform_node[i]->platform_id = platform_ids[i];
status = clGetPlatformInfo(platform_ids[i], CL_PLATFORM_VENDOR,
0, NULL, &platform_name_size);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not get size of platform name: %s\n", av_opencl_errstr(status));
} else {
device_list->platform_node[i]->platform_name = av_malloc(platform_name_size * sizeof(char));
if (!device_list->platform_node[i]->platform_name) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not allocate memory for device name: %s\n", av_opencl_errstr(status));
} else {
status = clGetPlatformInfo(platform_ids[i], CL_PLATFORM_VENDOR,
platform_name_size * sizeof(char),
device_list->platform_node[i]->platform_name, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not get platform name: %s\n", av_opencl_errstr(status));
}
}
}
total_devices_num = 0;
for (j = 0; j < FF_ARRAY_ELEMS(device_type); j++) {
status = clGetDeviceIDs(device_list->platform_node[i]->platform_id,
device_type[j], 0, NULL, &devices_num[j]);
total_devices_num += devices_num[j];
}
device_list->platform_node[i]->device_node = av_mallocz_array(total_devices_num, sizeof(AVOpenCLDeviceNode *));
if (!device_list->platform_node[i]->device_node) {
ret = AVERROR(ENOMEM);
goto end;
}
for (j = 0; j < FF_ARRAY_ELEMS(device_type); j++) {
if (devices_num[j]) {
device_ids = av_mallocz_array(devices_num[j], sizeof(cl_device_id));
if (!device_ids) {
ret = AVERROR(ENOMEM);
goto end;
}
status = clGetDeviceIDs(device_list->platform_node[i]->platform_id, device_type[j],
devices_num[j], device_ids, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not get device ID: %s:\n", av_opencl_errstr(status));
av_freep(&device_ids);
continue;
}
for (k = 0; k < devices_num[j]; k++) {
device_num = device_list->platform_node[i]->device_num;
device_list->platform_node[i]->device_node[device_num] = av_mallocz(sizeof(AVOpenCLDeviceNode));
if (!device_list->platform_node[i]->device_node[device_num]) {
ret = AVERROR(ENOMEM);
goto end;
}
device_node = device_list->platform_node[i]->device_node[device_num];
device_node->device_id = device_ids[k];
device_node->device_type = device_type[j];
status = clGetDeviceInfo(device_node->device_id, CL_DEVICE_NAME,
0, NULL, &device_name_size);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not get size of device name: %s\n", av_opencl_errstr(status));
continue;
}
device_node->device_name = av_malloc(device_name_size * sizeof(char));
if (!device_node->device_name) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not allocate memory for device name: %s\n", av_opencl_errstr(status));
continue;
}
status = clGetDeviceInfo(device_node->device_id, CL_DEVICE_NAME,
device_name_size * sizeof(char),
device_node->device_name, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Could not get device name: %s\n", av_opencl_errstr(status));
continue;
}
device_list->platform_node[i]->device_num++;
}
av_freep(&device_ids);
}
}
}
end:
av_freep(&platform_ids);
av_freep(&devices_num);
av_freep(&device_ids);
if (ret < 0)
free_device_list(device_list);
return ret;
}
int av_opencl_get_device_list(AVOpenCLDeviceList **device_list)
{
int ret = 0;
*device_list = av_mallocz(sizeof(AVOpenCLDeviceList));
if (!(*device_list)) {
av_log(&opencl_ctx, AV_LOG_ERROR, "Could not allocate opencl device list\n");
return AVERROR(ENOMEM);
}
ret = get_device_list(*device_list);
if (ret < 0) {
av_log(&opencl_ctx, AV_LOG_ERROR, "Could not get device list from environment\n");
free_device_list(*device_list);
av_freep(device_list);
return ret;
}
return ret;
}
void av_opencl_free_device_list(AVOpenCLDeviceList **device_list)
{
free_device_list(*device_list);
av_freep(device_list);
}
static inline int init_opencl_mtx(void)
{
#if HAVE_THREADS
if (!atomic_opencl_lock) {
int err;
pthread_mutex_t *tmp = av_malloc(sizeof(pthread_mutex_t));
if (!tmp)
return AVERROR(ENOMEM);
if ((err = pthread_mutex_init(tmp, NULL))) {
av_free(tmp);
return AVERROR(err);
}
if (avpriv_atomic_ptr_cas((void * volatile *)&atomic_opencl_lock, NULL, tmp)) {
pthread_mutex_destroy(tmp);
av_free(tmp);
}
}
#endif
return 0;
}
int av_opencl_set_option(const char *key, const char *val)
{
int ret = init_opencl_mtx( );
if (ret < 0)
return ret;
LOCK_OPENCL;
if (!opencl_ctx.opt_init_flag) {
av_opt_set_defaults(&opencl_ctx);
opencl_ctx.opt_init_flag = 1;
}
ret = av_opt_set(&opencl_ctx, key, val, 0);
UNLOCK_OPENCL;
return ret;
}
int av_opencl_get_option(const char *key, uint8_t **out_val)
{
int ret = 0;
LOCK_OPENCL;
ret = av_opt_get(&opencl_ctx, key, 0, out_val);
UNLOCK_OPENCL;
return ret;
}
void av_opencl_free_option(void)
{
/*FIXME: free openclutils context*/
LOCK_OPENCL;
av_opt_free(&opencl_ctx);
UNLOCK_OPENCL;
}
AVOpenCLExternalEnv *av_opencl_alloc_external_env(void)
{
AVOpenCLExternalEnv *ext = av_mallocz(sizeof(AVOpenCLExternalEnv));
if (!ext) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not malloc external opencl environment data space\n");
}
return ext;
}
void av_opencl_free_external_env(AVOpenCLExternalEnv **ext_opencl_env)
{
av_freep(ext_opencl_env);
}
int av_opencl_register_kernel_code(const char *kernel_code)
{
int i, ret = init_opencl_mtx( );
if (ret < 0)
return ret;
LOCK_OPENCL;
if (opencl_ctx.kernel_code_count >= MAX_KERNEL_CODE_NUM) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not register kernel code, maximum number of registered kernel code %d already reached\n",
MAX_KERNEL_CODE_NUM);
ret = AVERROR(EINVAL);
goto end;
}
for (i = 0; i < opencl_ctx.kernel_code_count; i++) {
if (opencl_ctx.kernel_code[i].kernel_string == kernel_code) {
av_log(&opencl_ctx, AV_LOG_WARNING, "Same kernel code has been registered\n");
goto end;
}
}
opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].kernel_string = kernel_code;
opencl_ctx.kernel_code[opencl_ctx.kernel_code_count].is_compiled = 0;
opencl_ctx.kernel_code_count++;
end:
UNLOCK_OPENCL;
return ret;
}
cl_program av_opencl_compile(const char *program_name, const char *build_opts)
{
int i;
cl_int status, build_status;
int kernel_code_idx = 0;
const char *kernel_source = NULL;
size_t kernel_code_len;
char* ptr = NULL;
cl_program program = NULL;
size_t log_size;
char *log = NULL;
LOCK_OPENCL;
for (i = 0; i < opencl_ctx.kernel_code_count; i++) {
// identify a program using a unique name within the kernel source
ptr = av_stristr(opencl_ctx.kernel_code[i].kernel_string, program_name);
if (ptr && !opencl_ctx.kernel_code[i].is_compiled) {
kernel_source = opencl_ctx.kernel_code[i].kernel_string;
kernel_code_len = strlen(opencl_ctx.kernel_code[i].kernel_string);
kernel_code_idx = i;
break;
}
}
if (!kernel_source) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Unable to find OpenCL kernel source '%s'\n", program_name);
goto end;
}
/* create a CL program from kernel source */
program = clCreateProgramWithSource(opencl_ctx.context, 1, &kernel_source, &kernel_code_len, &status);
if(status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Unable to create OpenCL program '%s': %s\n", program_name, av_opencl_errstr(status));
program = NULL;
goto end;
}
build_status = clBuildProgram(program, 1, &(opencl_ctx.device_id), build_opts, NULL, NULL);
status = clGetProgramBuildInfo(program, opencl_ctx.device_id,
CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Failed to get compilation log: %s\n",
av_opencl_errstr(status));
} else {
log = av_malloc(log_size);
if (log) {
status = clGetProgramBuildInfo(program, opencl_ctx.device_id,
CL_PROGRAM_BUILD_LOG, log_size,
log, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_WARNING,
"Failed to get compilation log: %s\n",
av_opencl_errstr(status));
} else {
int level = build_status == CL_SUCCESS ? AV_LOG_DEBUG :
AV_LOG_ERROR;
av_log(&opencl_ctx, level, "Compilation log:\n%s\n", log);
}
}
av_freep(&log);
}
if (build_status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Compilation failed with OpenCL program '%s': %s\n",
program_name, av_opencl_errstr(build_status));
program = NULL;
goto end;
}
opencl_ctx.kernel_code[kernel_code_idx].is_compiled = 1;
end:
UNLOCK_OPENCL;
return program;
}
cl_command_queue av_opencl_get_command_queue(void)
{
return opencl_ctx.command_queue;
}
static int init_opencl_env(OpenclContext *opencl_ctx, AVOpenCLExternalEnv *ext_opencl_env)
{
cl_int status;
cl_context_properties cps[3];
int i, ret = 0;
AVOpenCLDeviceNode *device_node = NULL;
if (ext_opencl_env) {
if (opencl_ctx->is_user_created)
return 0;
opencl_ctx->platform_id = ext_opencl_env->platform_id;
opencl_ctx->is_user_created = 1;
opencl_ctx->command_queue = ext_opencl_env->command_queue;
opencl_ctx->context = ext_opencl_env->context;
opencl_ctx->device_id = ext_opencl_env->device_id;
opencl_ctx->device_type = ext_opencl_env->device_type;
} else {
if (!opencl_ctx->is_user_created) {
if (!opencl_ctx->device_list.platform_num) {
ret = get_device_list(&opencl_ctx->device_list);
if (ret < 0) {
return ret;
}
}
if (opencl_ctx->platform_idx >= 0) {
if (opencl_ctx->device_list.platform_num < opencl_ctx->platform_idx + 1) {
av_log(opencl_ctx, AV_LOG_ERROR, "User set platform index not exist\n");
return AVERROR(EINVAL);
}
if (!opencl_ctx->device_list.platform_node[opencl_ctx->platform_idx]->device_num) {
av_log(opencl_ctx, AV_LOG_ERROR, "No devices in user specific platform with index %d\n",
opencl_ctx->platform_idx);
return AVERROR(EINVAL);
}
opencl_ctx->platform_id = opencl_ctx->device_list.platform_node[opencl_ctx->platform_idx]->platform_id;
} else {
/* get a usable platform by default*/
for (i = 0; i < opencl_ctx->device_list.platform_num; i++) {
if (opencl_ctx->device_list.platform_node[i]->device_num) {
opencl_ctx->platform_id = opencl_ctx->device_list.platform_node[i]->platform_id;
opencl_ctx->platform_idx = i;
break;
}
}
}
if (!opencl_ctx->platform_id) {
av_log(opencl_ctx, AV_LOG_ERROR, "Could not get OpenCL platforms\n");
return AVERROR_EXTERNAL;
}
/* get a usable device*/
if (opencl_ctx->device_idx >= 0) {
if (opencl_ctx->device_list.platform_node[opencl_ctx->platform_idx]->device_num < opencl_ctx->device_idx + 1) {
av_log(opencl_ctx, AV_LOG_ERROR,
"Could not get OpenCL device idx %d in the user set platform\n", opencl_ctx->platform_idx);
return AVERROR(EINVAL);
}
} else {
opencl_ctx->device_idx = 0;
}
device_node = opencl_ctx->device_list.platform_node[opencl_ctx->platform_idx]->device_node[opencl_ctx->device_idx];
opencl_ctx->device_id = device_node->device_id;
opencl_ctx->device_type = device_node->device_type;
/*
* Use available platform.
*/
av_log(opencl_ctx, AV_LOG_VERBOSE, "Platform Name: %s, Device Name: %s\n",
opencl_ctx->device_list.platform_node[opencl_ctx->platform_idx]->platform_name,
device_node->device_name);
cps[0] = CL_CONTEXT_PLATFORM;
cps[1] = (cl_context_properties)opencl_ctx->platform_id;
cps[2] = 0;
opencl_ctx->context = clCreateContextFromType(cps, opencl_ctx->device_type,
NULL, NULL, &status);
if (status != CL_SUCCESS) {
av_log(opencl_ctx, AV_LOG_ERROR,
"Could not get OpenCL context from device type: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
opencl_ctx->command_queue = clCreateCommandQueue(opencl_ctx->context, opencl_ctx->device_id,
0, &status);
if (status != CL_SUCCESS) {
av_log(opencl_ctx, AV_LOG_ERROR,
"Could not create OpenCL command queue: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
}
}
return ret;
}
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env)
{
int ret = init_opencl_mtx( );
if (ret < 0)
return ret;
LOCK_OPENCL;
if (!opencl_ctx.init_count) {
if (!opencl_ctx.opt_init_flag) {
av_opt_set_defaults(&opencl_ctx);
opencl_ctx.opt_init_flag = 1;
}
ret = init_opencl_env(&opencl_ctx, ext_opencl_env);
if (ret < 0)
goto end;
if (opencl_ctx.kernel_code_count <= 0) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"No kernel code is registered, compile kernel file failed\n");
ret = AVERROR(EINVAL);
goto end;
}
}
opencl_ctx.init_count++;
end:
UNLOCK_OPENCL;
return ret;
}
void av_opencl_uninit(void)
{
int i;
cl_int status;
LOCK_OPENCL;
opencl_ctx.init_count--;
if (opencl_ctx.is_user_created)
goto end;
if (opencl_ctx.init_count > 0)
goto end;
if (opencl_ctx.command_queue) {
status = clReleaseCommandQueue(opencl_ctx.command_queue);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not release OpenCL command queue: %s\n", av_opencl_errstr(status));
}
opencl_ctx.command_queue = NULL;
}
if (opencl_ctx.context) {
status = clReleaseContext(opencl_ctx.context);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not release OpenCL context: %s\n", av_opencl_errstr(status));
}
opencl_ctx.context = NULL;
}
for (i = 0; i < opencl_ctx.kernel_code_count; i++) {
opencl_ctx.kernel_code[i].is_compiled = 0;
}
free_device_list(&opencl_ctx.device_list);
end:
if (opencl_ctx.init_count <= 0)
av_opt_free(&opencl_ctx); //FIXME: free openclutils context
UNLOCK_OPENCL;
}
int av_opencl_buffer_create(cl_mem *cl_buf, size_t cl_buf_size, int flags, void *host_ptr)
{
cl_int status;
*cl_buf = clCreateBuffer(opencl_ctx.context, flags, cl_buf_size, host_ptr, &status);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR, "Could not create OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
return 0;
}
void av_opencl_buffer_release(cl_mem *cl_buf)
{
cl_int status = 0;
if (!cl_buf)
return;
status = clReleaseMemObject(*cl_buf);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not release OpenCL buffer: %s\n", av_opencl_errstr(status));
}
memset(cl_buf, 0, sizeof(*cl_buf));
}
int av_opencl_buffer_write(cl_mem dst_cl_buf, uint8_t *src_buf, size_t buf_size)
{
cl_int status;
void *mapped = clEnqueueMapBuffer(opencl_ctx.command_queue, dst_cl_buf,
CL_TRUE, CL_MAP_WRITE, 0, sizeof(uint8_t) * buf_size,
0, NULL, NULL, &status);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not map OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
memcpy(mapped, src_buf, buf_size);
status = clEnqueueUnmapMemObject(opencl_ctx.command_queue, dst_cl_buf, mapped, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not unmap OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
return 0;
}
int av_opencl_buffer_read(uint8_t *dst_buf, cl_mem src_cl_buf, size_t buf_size)
{
cl_int status;
void *mapped = clEnqueueMapBuffer(opencl_ctx.command_queue, src_cl_buf,
CL_TRUE, CL_MAP_READ, 0, buf_size,
0, NULL, NULL, &status);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not map OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
memcpy(dst_buf, mapped, buf_size);
status = clEnqueueUnmapMemObject(opencl_ctx.command_queue, src_cl_buf, mapped, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not unmap OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
return 0;
}
int av_opencl_buffer_write_image(cl_mem dst_cl_buf, size_t cl_buffer_size, int dst_cl_offset,
uint8_t **src_data, int *plane_size, int plane_num)
{
int i, buffer_size = 0;
uint8_t *temp;
cl_int status;
void *mapped;
if ((unsigned int)plane_num > 8) {
return AVERROR(EINVAL);
}
for (i = 0;i < plane_num;i++) {
buffer_size += plane_size[i];
}
if (buffer_size > cl_buffer_size) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Cannot write image to OpenCL buffer: buffer too small\n");
return AVERROR(EINVAL);
}
mapped = clEnqueueMapBuffer(opencl_ctx.command_queue, dst_cl_buf,
CL_TRUE, CL_MAP_WRITE, 0, buffer_size + dst_cl_offset,
0, NULL, NULL, &status);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not map OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
temp = mapped;
temp += dst_cl_offset;
for (i = 0; i < plane_num; i++) {
memcpy(temp, src_data[i], plane_size[i]);
temp += plane_size[i];
}
status = clEnqueueUnmapMemObject(opencl_ctx.command_queue, dst_cl_buf, mapped, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not unmap OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
return 0;
}
int av_opencl_buffer_read_image(uint8_t **dst_data, int *plane_size, int plane_num,
cl_mem src_cl_buf, size_t cl_buffer_size)
{
int i,buffer_size = 0,ret = 0;
uint8_t *temp;
void *mapped;
cl_int status;
if ((unsigned int)plane_num > 8) {
return AVERROR(EINVAL);
}
for (i = 0; i < plane_num; i++) {
buffer_size += plane_size[i];
}
if (buffer_size > cl_buffer_size) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Cannot write image to CPU buffer: OpenCL buffer too small\n");
return AVERROR(EINVAL);
}
mapped = clEnqueueMapBuffer(opencl_ctx.command_queue, src_cl_buf,
CL_TRUE, CL_MAP_READ, 0, buffer_size,
0, NULL, NULL, &status);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not map OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
temp = mapped;
if (ret >= 0) {
for (i = 0; i < plane_num; i++) {
memcpy(dst_data[i], temp, plane_size[i]);
temp += plane_size[i];
}
}
status = clEnqueueUnmapMemObject(opencl_ctx.command_queue, src_cl_buf, mapped, 0, NULL, NULL);
if (status != CL_SUCCESS) {
av_log(&opencl_ctx, AV_LOG_ERROR,
"Could not unmap OpenCL buffer: %s\n", av_opencl_errstr(status));
return AVERROR_EXTERNAL;
}
return 0;
}
int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device_node, cl_platform_id platform,
int64_t (*benchmark)(AVOpenCLExternalEnv *ext_opencl_env))
{
int64_t ret = 0;
cl_int status;
cl_context_properties cps[3];
AVOpenCLExternalEnv *ext_opencl_env = NULL;
ext_opencl_env = av_opencl_alloc_external_env();
ext_opencl_env->device_id = device_node->device_id;
ext_opencl_env->device_type = device_node->device_type;
av_log(&opencl_ctx, AV_LOG_VERBOSE, "Performing test on OpenCL device %s\n",
device_node->device_name);
cps[0] = CL_CONTEXT_PLATFORM;
cps[1] = (cl_context_properties)platform;
cps[2] = 0;
ext_opencl_env->context = clCreateContextFromType(cps, ext_opencl_env->device_type,
NULL, NULL, &status);
if (status != CL_SUCCESS || !ext_opencl_env->context) {
ret = AVERROR_EXTERNAL;
goto end;
}
ext_opencl_env->command_queue = clCreateCommandQueue(ext_opencl_env->context,
ext_opencl_env->device_id, 0, &status);
if (status != CL_SUCCESS || !ext_opencl_env->command_queue) {
ret = AVERROR_EXTERNAL;
goto end;
}
ret = benchmark(ext_opencl_env);
if (ret < 0)
av_log(&opencl_ctx, AV_LOG_ERROR, "Benchmark failed with OpenCL device %s\n",
device_node->device_name);
end:
if (ext_opencl_env->command_queue)
clReleaseCommandQueue(ext_opencl_env->command_queue);
if (ext_opencl_env->context)
clReleaseContext(ext_opencl_env->context);
av_opencl_free_external_env(&ext_opencl_env);
return ret;
}

292
libavutil/opencl.h
Unescape View File

@ -1,292 +0,0 @@
/*
* Copyright (C) 2012 Peng Gao <peng@multicorewareinc.com>
* Copyright (C) 2012 Li Cao <li@multicorewareinc.com>
* Copyright (C) 2012 Wei Gao <weigao@multicorewareinc.com>
* Copyright (C) 2013 Lenny Wang <lwanghpc@gmail.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
*/
/**
* @file
* OpenCL wrapper
*
* This interface is considered still experimental and its API and ABI may
* change without prior notice.
*/
#ifndef AVUTIL_OPENCL_H
#define AVUTIL_OPENCL_H
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS 1
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
#include <stdint.h>
#include "dict.h"
#include "libavutil/version.h"
#define AV_OPENCL_KERNEL( ... )# __VA_ARGS__
typedef struct {
int device_type;
char *device_name;
cl_device_id device_id;
} AVOpenCLDeviceNode;
typedef struct {
cl_platform_id platform_id;
char *platform_name;
int device_num;
AVOpenCLDeviceNode **device_node;
} AVOpenCLPlatformNode;
typedef struct {
int platform_num;
AVOpenCLPlatformNode **platform_node;
} AVOpenCLDeviceList;
typedef struct {
cl_platform_id platform_id;
cl_device_type device_type;
cl_context context;
cl_device_id device_id;
cl_command_queue command_queue;
char *platform_name;
} AVOpenCLExternalEnv;
/**
* Get OpenCL device list.
*
* It must be freed with av_opencl_free_device_list().
*
* @param device_list pointer to OpenCL environment device list,
* should be released by av_opencl_free_device_list()
*
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_get_device_list(AVOpenCLDeviceList **device_list);
/**
* Free OpenCL device list.
*
* @param device_list pointer to OpenCL environment device list
* created by av_opencl_get_device_list()
*/
void av_opencl_free_device_list(AVOpenCLDeviceList **device_list);
/**
* Set option in the global OpenCL context.
*
* This options affect the operation performed by the next
* av_opencl_init() operation.
*
* The currently accepted options are:
* - platform: set index of platform in device list
* - device: set index of device in device list
*
* See reference "OpenCL Specification Version: 1.2 chapter 5.6.4".
*
* @param key option key
* @param val option value
* @return >=0 on success, a negative error code in case of failure
* @see av_opencl_get_option()
*/
int av_opencl_set_option(const char *key, const char *val);
/**
* Get option value from the global OpenCL context.
*
* @param key option key
* @param out_val pointer to location where option value will be
* written, must be freed with av_freep()
* @return >=0 on success, a negative error code in case of failure
* @see av_opencl_set_option()
*/
int av_opencl_get_option(const char *key, uint8_t **out_val);
/**
* Free option values of the global OpenCL context.
*
*/
void av_opencl_free_option(void);
/**
* Allocate OpenCL external environment.
*
* It must be freed with av_opencl_free_external_env().
*
* @return pointer to allocated OpenCL external environment
*/
AVOpenCLExternalEnv *av_opencl_alloc_external_env(void);
/**
* Free OpenCL external environment.
*
* @param ext_opencl_env pointer to OpenCL external environment
* created by av_opencl_alloc_external_env()
*/
void av_opencl_free_external_env(AVOpenCLExternalEnv **ext_opencl_env);
/**
* Get OpenCL error string.
*
* @param status OpenCL error code
* @return OpenCL error string
*/
const char *av_opencl_errstr(cl_int status);
/**
* Register kernel code.
*
* The registered kernel code is stored in a global context, and compiled
* in the runtime environment when av_opencl_init() is called.
*
* @param kernel_code kernel code to be compiled in the OpenCL runtime environment
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_register_kernel_code(const char *kernel_code);
/**
* Initialize the run time OpenCL environment
*
* @param ext_opencl_env external OpenCL environment, created by an
* application program, ignored if set to NULL
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_init(AVOpenCLExternalEnv *ext_opencl_env);
/**
* compile specific OpenCL kernel source
*
* @param program_name pointer to a program name used for identification
* @param build_opts pointer to a string that describes the preprocessor
* build options to be used for building the program
* @return a cl_program object
*/
cl_program av_opencl_compile(const char *program_name, const char* build_opts);
/**
* get OpenCL command queue
*
* @return a cl_command_queue object
*/
cl_command_queue av_opencl_get_command_queue(void);
/**
* Create OpenCL buffer.
*
* The buffer is used to save the data used or created by an OpenCL
* kernel.
* The created buffer must be released with av_opencl_buffer_release().
*
* See clCreateBuffer() function reference for more information about
* the parameters.
*
* @param cl_buf pointer to OpenCL buffer
* @param cl_buf_size size in bytes of the OpenCL buffer to create
* @param flags flags used to control buffer attributes
* @param host_ptr host pointer of the OpenCL buffer
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_buffer_create(cl_mem *cl_buf, size_t cl_buf_size, int flags, void *host_ptr);
/**
* Write OpenCL buffer with data from src_buf.
*
* @param dst_cl_buf pointer to OpenCL destination buffer
* @param src_buf pointer to source buffer
* @param buf_size size in bytes of the source and destination buffers
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_buffer_write(cl_mem dst_cl_buf, uint8_t *src_buf, size_t buf_size);
/**
* Read data from OpenCL buffer to memory buffer.
*
* @param dst_buf pointer to destination buffer (CPU memory)
* @param src_cl_buf pointer to source OpenCL buffer
* @param buf_size size in bytes of the source and destination buffers
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_buffer_read(uint8_t *dst_buf, cl_mem src_cl_buf, size_t buf_size);
/**
* Write image data from memory to OpenCL buffer.
*
* The source must be an array of pointers to image plane buffers.
*
* @param dst_cl_buf pointer to destination OpenCL buffer
* @param dst_cl_buf_size size in bytes of OpenCL buffer
* @param dst_cl_buf_offset the offset of the OpenCL buffer start position
* @param src_data array of pointers to source plane buffers
* @param src_plane_sizes array of sizes in bytes of the source plane buffers
* @param src_plane_num number of source image planes
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_buffer_write_image(cl_mem dst_cl_buf, size_t cl_buffer_size, int dst_cl_offset,
uint8_t **src_data, int *plane_size, int plane_num);
/**
* Read image data from OpenCL buffer.
*
* @param dst_data array of pointers to destination plane buffers
* @param dst_plane_sizes array of pointers to destination plane buffers
* @param dst_plane_num number of destination image planes
* @param src_cl_buf pointer to source OpenCL buffer
* @param src_cl_buf_size size in bytes of OpenCL buffer
* @return >=0 on success, a negative error code in case of failure
*/
int av_opencl_buffer_read_image(uint8_t **dst_data, int *plane_size, int plane_num,
cl_mem src_cl_buf, size_t cl_buffer_size);
/**
* Release OpenCL buffer.
*
* @param cl_buf pointer to OpenCL buffer to release, which was
* previously filled with av_opencl_buffer_create()
*/
void av_opencl_buffer_release(cl_mem *cl_buf);
/**
* Release OpenCL environment.
*
* The OpenCL environment is effectively released only if all the created
* kernels had been released with av_opencl_release_kernel().
*/
void av_opencl_uninit(void);
/**
* Benchmark an OpenCL device with a user defined callback function. This function
* sets up an external OpenCL environment including context and command queue on
* the device then tears it down in the end. The callback function should perform
* the rest of the work.
*
* @param device pointer to the OpenCL device to be used
* @param platform cl_platform_id handle to which the device belongs to
* @param benchmark callback function to perform the benchmark, return a
* negative value in case of failure
* @return the score passed from the callback function, a negative error code in case
* of failure
*/
int64_t av_opencl_benchmark(AVOpenCLDeviceNode *device, cl_platform_id platform,
int64_t (*benchmark)(AVOpenCLExternalEnv *ext_opencl_env));
#endif /* AVUTIL_OPENCL_H */

59
libavutil/opencl_internal.c
Unescape View File

@ -1,59 +0,0 @@
/*
* Copyright (C) 2012 Peng Gao <peng@multicorewareinc.com>
* Copyright (C) 2012 Li Cao <li@multicorewareinc.com>
* Copyright (C) 2012 Wei Gao <weigao@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_internal.h"
#include "libavutil/log.h"
int avpriv_opencl_set_parameter(FFOpenclParam *opencl_param, ...)
{
int ret = 0;
va_list arg_ptr;
void *param;
size_t param_size;
cl_int status;
if (!opencl_param->kernel) {
av_log(opencl_param->ctx, AV_LOG_ERROR, "OpenCL kernel must be set\n");
return AVERROR(EINVAL);
}
va_start(arg_ptr, opencl_param);
do {
param = va_arg(arg_ptr, void *);
if (!param)
break;
param_size = va_arg(arg_ptr, size_t);
if (!param_size) {
av_log(opencl_param->ctx, AV_LOG_ERROR, "Parameter size must not be 0\n");
ret = AVERROR(EINVAL);
goto end;
}
status = clSetKernelArg(opencl_param->kernel, opencl_param->param_num, param_size, param);
if (status != CL_SUCCESS) {
av_log(opencl_param->ctx, AV_LOG_ERROR, "Cannot set kernel argument: %s\n", av_opencl_errstr(status));
ret = AVERROR_EXTERNAL;
goto end;
}
opencl_param->param_num++;
} while (param && param_size);
end:
va_end(arg_ptr);
return ret;
}

40
libavutil/opencl_internal.h
Unescape View File

@ -1,40 +0,0 @@
/*
* Copyright (C) 2012 Peng Gao <peng@multicorewareinc.com>
* Copyright (C) 2012 Li Cao <li@multicorewareinc.com>
* Copyright (C) 2012 Wei Gao <weigao@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 AVUTIL_OPENCL_INTERNAL_H
#define AVUTIL_OPENCL_INTERNAL_H
#include "attributes.h"
#include "opencl.h"
#define FF_OPENCL_PARAM_INFO(a) ((void*)(&(a))), (sizeof(a))
typedef struct {
cl_kernel kernel;
int param_num;
void *ctx;
} FFOpenclParam;
av_warn_unused_result
int avpriv_opencl_set_parameter(FFOpenclParam *opencl_param, ...);
#endif /* AVUTIL_OPENCL_INTERNAL_H */
Write Preview
Loading…
Cancel
Save