[FFmpeg-devel] [PATCH] lavfi/deshake: modify seting opencl parameters

Stefano Sabatini stefasab at gmail.com
Sat Apr 27 19:04:57 CEST 2013


On date Saturday 2013-04-27 10:22:25 +0800, Wei Gao encoded:
> 

> From 50a76e1e72cf773fd1da387026a043e10ba8a420 Mon Sep 17 00:00:00 2001
> From: highgod0401 <highgod0401 at gmail.com>
> Date: Sat, 27 Apr 2013 10:19:49 +0800
> Subject: [PATCH] lavfi/deshake: modify seting opencl parameters
> 
> ---
>  libavfilter/deshake_opencl.c | 65 ++++++++++++++++++++------------------------
>  1 file changed, 30 insertions(+), 35 deletions(-)
> 
> diff --git a/libavfilter/deshake_opencl.c b/libavfilter/deshake_opencl.c
> index 0f6dcc4..51bf515 100644
> --- a/libavfilter/deshake_opencl.c
> +++ b/libavfilter/deshake_opencl.c
> @@ -27,64 +27,59 @@
>  #include "libavutil/dict.h"
>  #include "libavutil/pixdesc.h"
>  #include "deshake_opencl.h"
> +#include "libavutil/opencl_internal.h"
>  
>  #define MATRIX_SIZE 6
>  #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 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 arg_no, ret = 0;
> +    int ret = 0;
>      const size_t global_work_size = width * height + 2 * ch * cw;
> -    cl_kernel kernel;
>      cl_int status;
>      DeshakeContext *deshake = ctx->priv;
> +    FFOpenclParam opencl_param = {0};
>      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)
>          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));
>      if (ret < 0)
>          return ret;
> -    kernel = deshake->opencl_ctx.kernel_env.kernel;
> -    arg_no = 0;
> -
>      if ((unsigned int)interpolate > INTERPOLATE_BIQUADRATIC) {
>          av_log(ctx, AV_LOG_ERROR, "Selected interpolate method is invalid\n");
>          return AVERROR(EINVAL);
>      }
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_inbuf);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_outbuf);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_y);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(deshake->opencl_ctx.cl_matrix_uv);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(interpolate);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(fill);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[0]);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[0]);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(in->linesize[1]);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(out->linesize[1]);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(height);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(width);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(ch);
> -    TRANSFORM_OPENCL_SET_KERNEL_ARG(cw);
> -    TRANSFORM_OPENCL_CHECK(clEnqueueNDRangeKernel, deshake->opencl_ctx.kernel_env.command_queue, deshake->opencl_ctx.kernel_env.kernel, 1, NULL,
> -                           &global_work_size, NULL, 0, NULL, NULL);
> +    opencl_param.ctx = ctx;
> +    opencl_param.kernel = deshake->opencl_ctx.kernel_env.kernel;
> +    ret = ff_opencl_set_parameter(&opencl_param,
> +                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_inbuf),
> +                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_outbuf),
> +                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_y),
> +                                  FF_OPENCL_PARAM_INFO(deshake->opencl_ctx.cl_matrix_uv),
> +                                  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.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: %d\n", status);

you may make use of av_opencl_strerror()

> +        return AVERROR_EXTERNAL;
> +    }
>      clFinish(deshake->opencl_ctx.kernel_env.command_queue);
>      ret = av_opencl_buffer_read_image(out->data, deshake->opencl_ctx.out_plane_size,
>                                        deshake->opencl_ctx.plane_num, deshake->opencl_ctx.cl_outbuf,

LGTM otherwise, thanks.
-- 
FFmpeg = Free Funny Miracolous Ponderous Elfic Guru


More information about the ffmpeg-devel mailing list