generic-library/ffmpeg
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

lavfi/deshake_opencl: use ff_opencl_set_parameter

Browse Source 
Signed-off-by: Michael Niedermayer <michaelni@gmx.at>
This commit is contained in:
highgod0401 2013-05-06 10:02:50 +08:00 committed by Michael Niedermayer
parent 0b5f4fdc38
commit b02f073ad4
1 changed files with 31 additions and 34 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

65
libavfilter/deshake_opencl.c
View File

@ -27,64 +27,61 @@
#include "libavutil/dict.h" #include "libavutil/dict.h"
#include "libavutil/pixdesc.h" #include "libavutil/pixdesc.h"
#include "deshake_opencl.h" #include "deshake_opencl.h"
#include "libavutil/opencl_internal.h"
#define MATRIX_SIZE 6 #define MATRIX_SIZE 6
#define PLANE_NUM 3 #define PLANE_NUM 3
#define TRANSFORM_OPENCL_CHECK(method, ...) \
status = method(__VA_ARGS__); \
if (status != CL_SUCCESS) { \
av_log(ctx, AV_LOG_ERROR, "error %s %d\n", # method, status); \
return AVERROR_EXTERNAL; \
}
#define TRANSFORM_OPENCL_SET_KERNEL_ARG(arg_ptr) \
status = clSetKernelArg((kernel),(arg_no++),(sizeof(arg_ptr)),(void*)(&(arg_ptr))); \
if (status != CL_SUCCESS) { \
av_log(ctx, AV_LOG_ERROR, "cannot set kernel argument: %d\n", status ); \
return AVERROR_EXTERNAL; \
}
int ff_opencl_transform(AVFilterContext *ctx, int ff_opencl_transform(AVFilterContext *ctx,
int width, int height, int cw, int ch, int width, int height, int cw, int ch,
const float *matrix_y, const float *matrix_uv, const float *matrix_y, const float *matrix_uv,
enum InterpolateMethod interpolate, enum InterpolateMethod interpolate,
enum FillMethod fill, AVFrame *in, AVFrame *out) enum FillMethod fill, AVFrame *in, AVFrame *out)
{ {
int arg_no, ret = 0; int ret = 0;
const size_t global_work_size = width * height + 2 * ch * cw; const size_t global_work_size = width * height + 2 * ch * cw;
cl_kernel kernel;
cl_int status; cl_int status;
DeshakeContext *deshake = ctx->priv; DeshakeContext *deshake = ctx->priv;
FFOpenclParam opencl_param = {0};
opencl_param.ctx = ctx;
opencl_param.kernel = deshake->opencl_ctx.kernel_env.kernel;
ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float)); ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_y, (uint8_t *)matrix_y, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
if (ret < 0) if (ret < 0)
return ret; return ret;
ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_uv, (uint8_t *)matrix_uv, deshake->opencl_ctx.matrix_size * sizeof(cl_float)); ret = av_opencl_buffer_write(deshake->opencl_ctx.cl_matrix_uv, (uint8_t *)matrix_uv, deshake->opencl_ctx.matrix_size * sizeof(cl_float));
if (ret < 0) if (ret < 0)
return ret; return ret;
kernel = deshake->opencl_ctx.kernel_env.kernel;
arg_no = 0;
if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) { if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n"); av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
return AVERROR(EINVAL); return AVERROR(EINVAL);
} }
TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_inbuf); ret = ff_opencl_set_parameter(&opencl_param,
TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_outbuf); FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_y); FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_uv); FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_y),
TRANSFORM_OPENCL_SET_KERNEL_ARG(interpolate); FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_uv),
TRANSFORM_OPENCL_SET_KERNEL_ARG(fill); FF_OPENCL_PARAM_INFO(interpolate),
TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[0]); FF_OPENCL_PARAM_INFO(fill),
TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[0]); FF_OPENCL_PARAM_INFO(in->linesize[0]),
TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[1]); FF_OPENCL_PARAM_INFO(out->linesize[0]),
TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[1]); FF_OPENCL_PARAM_INFO(in->linesize[1]),
TRANSFORM_OPENCL_SET_KERNEL_ARG(height); FF_OPENCL_PARAM_INFO(out->linesize[1]),
TRANSFORM_OPENCL_SET_KERNEL_ARG(width); FF_OPENCL_PARAM_INFO(height),
TRANSFORM_OPENCL_SET_KERNEL_ARG(ch); FF_OPENCL_PARAM_INFO(width),
TRANSFORM_OPENCL_SET_KERNEL_ARG(cw); FF_OPENCL_PARAM_INFO(ch),
TRANSFORM_OPENCL_CHECK(clEnqueueNDRangeKernel, deshake->opencl_ctx.kernel_env.command_queue, deshake->opencl_ctx.kernel_env.kernel, 1, NULL, FF_OPENCL_PARAM_INFO(cw),
&global_work_size, NULL, 0, NULL, NULL); NULL);
if (ret < 0)
return ret;
status = clEnqueueNDRangeKernel(deshake->opencl_ctx.kernel_env.command_queue,
deshake->opencl_ctx.kernel_env.kernel, 1, NULL,
&global_work_size, 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;
}
clFinish(deshake->opencl_ctx.kernel_env.command_queue); clFinish(deshake->opencl_ctx.kernel_env.command_queue);
ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size, 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.plane_num, deshake->opencl_ctx.cl_outbuf,