git.sesse.net Git - ffmpeg/blob - libavfilter/vf_pad_opencl.c

   1 /*
   2  * This file is part of FFmpeg.
   3  *
   4  * FFmpeg is free software; you can redistribute it and/or
   5  * modify it under the terms of the GNU Lesser General Public
   6  * License as published by the Free Software Foundation; either
   7  * version 2.1 of the License, or (at your option) any later version.
   8  *
   9  * FFmpeg is distributed in the hope that it will be useful,
  10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  12  * Lesser General Public License for more details.
  13  *
  14  * You should have received a copy of the GNU Lesser General Public
  15  * License along with FFmpeg; if not, write to the Free Software
  16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  17  */
  18
  19 #include "libavutil/colorspace.h"
  20 #include "libavutil/eval.h"
  21 #include "libavutil/opt.h"
  22 #include "libavutil/imgutils.h"
  23 #include "avfilter.h"
  24 #include "drawutils.h"
  25 #include "formats.h"
  26 #include "internal.h"
  27 #include "opencl.h"
  28 #include "opencl_source.h"
  29 #include "video.h"
  30
  31 static const char *const var_names[] = {
  32     "in_w",   "iw",
  33     "in_h",   "ih",
  34     "out_w",  "ow",
  35     "out_h",  "oh",
  36     "x",
  37     "y",
  38     "a",
  39     "sar",
  40     "dar",
  41     NULL
  42 };
  43
  44 enum var_name {
  45     VAR_IN_W,   VAR_IW,
  46     VAR_IN_H,   VAR_IH,
  47     VAR_OUT_W,  VAR_OW,
  48     VAR_OUT_H,  VAR_OH,
  49     VAR_X,
  50     VAR_Y,
  51     VAR_A,
  52     VAR_SAR,
  53     VAR_DAR,
  54     VARS_NB
  55 };
  56
  57 typedef struct PadOpenCLContext {
  58     OpenCLFilterContext ocf;
  59     int initialized;
  60     int is_rgb;
  61     int is_packed;
  62     int hsub, vsub;
  63
  64     char *w_expr;
  65     char *h_expr;
  66     char *x_expr;
  67     char *y_expr;
  68     AVRational aspect;
  69
  70     cl_command_queue command_queue;
  71     cl_kernel kernel_pad;
  72
  73     int w, h;
  74     int x, y;
  75     uint8_t pad_rgba[4];
  76     uint8_t pad_color[4];
  77     cl_float4 pad_color_float;
  78     cl_int2 pad_pos;
  79 } PadOpenCLContext;
  80
  81 static int pad_opencl_init(AVFilterContext *avctx, AVFrame *input_frame)
  82 {
  83     PadOpenCLContext *ctx = avctx->priv;
  84     AVHWFramesContext *input_frames_ctx = (AVHWFramesContext *)input_frame->hw_frames_ctx->data;
  85     const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(input_frames_ctx->sw_format);
  86     uint8_t rgba_map[4];
  87     cl_int cle;
  88     int err;
  89
  90     ff_fill_rgba_map(rgba_map, input_frames_ctx->sw_format);
  91     ctx->is_rgb = !!(desc->flags & AV_PIX_FMT_FLAG_RGB);
  92     ctx->is_packed = !(desc->flags & AV_PIX_FMT_FLAG_PLANAR);
  93     ctx->hsub = desc->log2_chroma_w;
  94     ctx->vsub = desc->log2_chroma_h;
  95
  96     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_pad, 1);
  97     if (err < 0)
  98         goto fail;
  99
 100     ctx->command_queue = clCreateCommandQueue(
 101         ctx->ocf.hwctx->context,
 102         ctx->ocf.hwctx->device_id,
 103         0,
 104         &cle
 105     );
 106
 107     if (ctx->is_rgb) {
 108         ctx->pad_color[rgba_map[0]] = ctx->pad_rgba[0];
 109         ctx->pad_color[rgba_map[1]] = ctx->pad_rgba[1];
 110         ctx->pad_color[rgba_map[2]] = ctx->pad_rgba[2];
 111         ctx->pad_color[rgba_map[3]] = ctx->pad_rgba[3];
 112     } else {
 113         ctx->pad_color[0] = RGB_TO_Y_BT709(ctx->pad_rgba[0], ctx->pad_rgba[1], ctx->pad_rgba[2]);
 114         ctx->pad_color[1] = RGB_TO_U_BT709(ctx->pad_rgba[0], ctx->pad_rgba[1], ctx->pad_rgba[2], 0);
 115         ctx->pad_color[2] = RGB_TO_V_BT709(ctx->pad_rgba[0], ctx->pad_rgba[1], ctx->pad_rgba[2], 0);
 116         ctx->pad_color[3] = ctx->pad_rgba[3];
 117     }
 118
 119     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL command queue %d.\n", cle);
 120
 121     ctx->kernel_pad = clCreateKernel(ctx->ocf.program, "pad", &cle);
 122     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create pad kernel: %d.\n", cle);
 123
 124     for (int i = 0; i < 4; ++i) {
 125         ctx->pad_color_float.s[i] = (float)ctx->pad_color[i] / 255.0;
 126     }
 127
 128     ctx->pad_pos.s[0] = ctx->x;
 129     ctx->pad_pos.s[1] = ctx->y;
 130
 131     ctx->initialized = 1;
 132     return 0;
 133
 134 fail:
 135     if (ctx->command_queue)
 136         clReleaseCommandQueue(ctx->command_queue);
 137     if (ctx->kernel_pad)
 138         clReleaseKernel(ctx->kernel_pad);
 139     return err;
 140 }
 141
 142 static int filter_frame(AVFilterLink *link, AVFrame *input_frame)
 143 {
 144     AVFilterContext *avctx = link->dst;
 145     AVFilterLink *outlink = avctx->outputs[0];
 146     PadOpenCLContext *pad_ctx = avctx->priv;
 147     AVFrame *output_frame = NULL;
 148     int err;
 149     cl_int cle;
 150     size_t global_work[2];
 151     cl_mem src, dst;
 152
 153     if (!input_frame->hw_frames_ctx)
 154         return AVERROR(EINVAL);
 155
 156     if (!pad_ctx->initialized) {
 157         err = pad_opencl_init(avctx, input_frame);
 158         if (err < 0)
 159             goto fail;
 160     }
 161
 162     output_frame = ff_get_video_buffer(outlink, outlink->w, outlink->h);
 163     if (!output_frame) {
 164         err = AVERROR(ENOMEM);
 165         goto fail;
 166     }
 167
 168     for (int p = 0; p < FF_ARRAY_ELEMS(output_frame->data); p++) {
 169         cl_float4 pad_color_float;
 170         cl_int2 pad_pos;
 171
 172         if (pad_ctx->is_packed) {
 173             pad_color_float = pad_ctx->pad_color_float;
 174         } else {
 175             pad_color_float.s[0] = pad_ctx->pad_color_float.s[p];
 176             pad_color_float.s[1] = pad_ctx->pad_color_float.s[2];
 177         }
 178
 179         if (p > 0 && p < 3) {
 180             pad_pos.s[0] = pad_ctx->pad_pos.s[0] >> pad_ctx->hsub;
 181             pad_pos.s[1] = pad_ctx->pad_pos.s[1] >> pad_ctx->vsub;
 182         } else {
 183             pad_pos.s[0] = pad_ctx->pad_pos.s[0];
 184             pad_pos.s[1] = pad_ctx->pad_pos.s[1];
 185         }
 186
 187         src = (cl_mem)input_frame->data[p];
 188         dst = (cl_mem)output_frame->data[p];
 189
 190         if (!dst)
 191             break;
 192
 193         CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 0, cl_mem, &src);
 194         CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 1, cl_mem, &dst);
 195         CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 2, cl_float4, &pad_color_float);
 196         CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 3, cl_int2, &pad_pos);
 197
 198         err = ff_opencl_filter_work_size_from_image(avctx, global_work, output_frame, p, 16);
 199         if (err < 0)
 200             goto fail;
 201
 202         cle = clEnqueueNDRangeKernel(pad_ctx->command_queue, pad_ctx->kernel_pad, 2, NULL,
 203                                      global_work, NULL, 0, NULL, NULL);
 204
 205         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue pad kernel: %d.\n", cle);
 206     }
 207
 208     // Run queued kernel
 209     cle = clFinish(pad_ctx->command_queue);
 210     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 211
 212     err = av_frame_copy_props(output_frame, input_frame);
 213     if (err < 0)
 214         goto fail;
 215
 216     av_frame_free(&input_frame);
 217
 218     return ff_filter_frame(outlink, output_frame);
 219
 220 fail:
 221     clFinish(pad_ctx->command_queue);
 222     av_frame_free(&input_frame);
 223     av_frame_free(&output_frame);
 224     return err;
 225 }
 226
 227 static av_cold void pad_opencl_uninit(AVFilterContext *avctx)
 228 {
 229     PadOpenCLContext *ctx = avctx->priv;
 230     cl_int cle;
 231
 232     if (ctx->kernel_pad) {
 233         cle = clReleaseKernel(ctx->kernel_pad);
 234         if (cle != CL_SUCCESS)
 235             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 236                    "kernel: %d.\n", cle);
 237     }
 238
 239     if (ctx->command_queue) {
 240         cle = clReleaseCommandQueue(ctx->command_queue);
 241         if (cle != CL_SUCCESS)
 242             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 243                    "command queue: %d.\n", cle);
 244     }
 245
 246     ff_opencl_filter_uninit(avctx);
 247 }
 248
 249 static int pad_opencl_config_output(AVFilterLink *outlink)
 250 {
 251     AVFilterContext *avctx = outlink->src;
 252     AVFilterLink *inlink = avctx->inputs[0];
 253     PadOpenCLContext *ctx = avctx->priv;
 254     AVRational adjusted_aspect = ctx->aspect;
 255     double var_values[VARS_NB], res;
 256     int err, ret;
 257     char *expr;
 258
 259     var_values[VAR_IN_W]  = var_values[VAR_IW] = inlink->w;
 260     var_values[VAR_IN_H]  = var_values[VAR_IH] = inlink->h;
 261     var_values[VAR_OUT_W] = var_values[VAR_OW] = NAN;
 262     var_values[VAR_OUT_H] = var_values[VAR_OH] = NAN;
 263     var_values[VAR_A]     = (double) inlink->w / inlink->h;
 264     var_values[VAR_SAR]   = inlink->sample_aspect_ratio.num ?
 265         (double) inlink->sample_aspect_ratio.num / inlink->sample_aspect_ratio.den : 1;
 266     var_values[VAR_DAR]   = var_values[VAR_A] * var_values[VAR_SAR];
 267
 268     av_expr_parse_and_eval(&res, (expr = ctx->w_expr),
 269                            var_names, var_values,
 270                            NULL, NULL, NULL, NULL, NULL, 0, ctx);
 271     ctx->w = var_values[VAR_OUT_W] = var_values[VAR_OW] = res;
 272     if ((ret = av_expr_parse_and_eval(&res, (expr = ctx->h_expr),
 273                                       var_names, var_values,
 274                                       NULL, NULL, NULL, NULL, NULL, 0, ctx)) < 0)
 275         return ret;
 276     ctx->h = var_values[VAR_OUT_H] = var_values[VAR_OH] = res;
 277     if (!ctx->h)
 278         var_values[VAR_OUT_H] = var_values[VAR_OH] = ctx->h = inlink->h;
 279
 280     /* evaluate the width again, as it may depend on the evaluated output height */
 281     if ((ret = av_expr_parse_and_eval(&res, (expr = ctx->w_expr),
 282                                       var_names, var_values,
 283                                       NULL, NULL, NULL, NULL, NULL, 0, ctx)) < 0)
 284         return ret;
 285     ctx->w = var_values[VAR_OUT_W] = var_values[VAR_OW] = res;
 286     if (!ctx->w)
 287         var_values[VAR_OUT_W] = var_values[VAR_OW] = ctx->w = inlink->w;
 288
 289     if (adjusted_aspect.num && adjusted_aspect.den) {
 290         adjusted_aspect = av_div_q(adjusted_aspect, inlink->sample_aspect_ratio);
 291         if (ctx->h < av_rescale(ctx->w, adjusted_aspect.den, adjusted_aspect.num)) {
 292             ctx->h = var_values[VAR_OUT_H] = var_values[VAR_OH] = av_rescale(ctx->w, adjusted_aspect.den, adjusted_aspect.num);
 293         } else {
 294             ctx->w = var_values[VAR_OUT_W] = var_values[VAR_OW] = av_rescale(ctx->h, adjusted_aspect.num, adjusted_aspect.den);
 295         }
 296     }
 297
 298     /* evaluate x and y */
 299     av_expr_parse_and_eval(&res, (expr = ctx->x_expr),
 300                            var_names, var_values,
 301                            NULL, NULL, NULL, NULL, NULL, 0, ctx);
 302     ctx->x = var_values[VAR_X] = res;
 303     if ((ret = av_expr_parse_and_eval(&res, (expr = ctx->y_expr),
 304                                       var_names, var_values,
 305                                       NULL, NULL, NULL, NULL, NULL, 0, ctx)) < 0)
 306         return ret;
 307     ctx->y = var_values[VAR_Y] = res;
 308     /* evaluate x again, as it may depend on the evaluated y value */
 309     if ((ret = av_expr_parse_and_eval(&res, (expr = ctx->x_expr),
 310                                       var_names, var_values,
 311                                       NULL, NULL, NULL, NULL, NULL, 0, ctx)) < 0)
 312         return ret;
 313     ctx->x = var_values[VAR_X] = res;
 314
 315     if (ctx->x < 0 || ctx->x + inlink->w > ctx->w)
 316         ctx->x = var_values[VAR_X] = (ctx->w - inlink->w) / 2;
 317     if (ctx->y < 0 || ctx->y + inlink->h > ctx->h)
 318         ctx->y = var_values[VAR_Y] = (ctx->h - inlink->h) / 2;
 319
 320     /* sanity check params */
 321     if (ctx->w < inlink->w || ctx->h < inlink->h) {
 322         av_log(ctx, AV_LOG_ERROR, "Padded dimensions cannot be smaller than input dimensions.\n");
 323         return AVERROR(EINVAL);
 324     }
 325
 326     if (ctx->w > avctx->inputs[0]->w) {
 327         ctx->ocf.output_width  = ctx->w;
 328     } else {
 329         ctx->ocf.output_width  = avctx->inputs[0]->w;
 330     }
 331
 332     if (ctx->h > avctx->inputs[0]->h) {
 333         ctx->ocf.output_height = ctx->h;
 334     } else {
 335         ctx->ocf.output_height = avctx->inputs[0]->h;
 336     }
 337
 338     if (ctx->x + avctx->inputs[0]->w > ctx->ocf.output_width ||
 339         ctx->y + avctx->inputs[0]->h > ctx->ocf.output_height) {
 340         return AVERROR(EINVAL);
 341     }
 342
 343     err = ff_opencl_filter_config_output(outlink);
 344     if (err < 0)
 345         return err;
 346
 347     return 0;
 348 }
 349
 350 static const AVFilterPad pad_opencl_inputs[] = {
 351     {
 352         .name = "default",
 353         .type = AVMEDIA_TYPE_VIDEO,
 354         .filter_frame = filter_frame,
 355         .config_props = &ff_opencl_filter_config_input,
 356     },
 357     { NULL }
 358 };
 359
 360 static const AVFilterPad pad_opencl_outputs[] = {
 361     {
 362         .name = "default",
 363         .type = AVMEDIA_TYPE_VIDEO,
 364         .config_props = &pad_opencl_config_output,
 365     },
 366     { NULL }
 367 };
 368
 369 #define OFFSET(x) offsetof(PadOpenCLContext, x)
 370 #define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
 371
 372 static const AVOption pad_opencl_options[] = {
 373     { "width",  "set the pad area width",       OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS },
 374     { "w",      "set the pad area width",       OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS },
 375     { "height", "set the pad area height",      OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS },
 376     { "h",      "set the pad area height",      OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS },
 377     { "x",      "set the x offset for the input image position", OFFSET(x_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, INT16_MAX, FLAGS },
 378     { "y",      "set the y offset for the input image position", OFFSET(y_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, INT16_MAX, FLAGS },
 379     { "color", "set the color of the padded area border", OFFSET(pad_rgba), AV_OPT_TYPE_COLOR, { .str = "black" }, 0, 0, FLAGS },
 380     { "aspect",  "pad to fit an aspect instead of a resolution", OFFSET(aspect), AV_OPT_TYPE_RATIONAL, {.dbl = 0}, 0, INT16_MAX, FLAGS },
 381     { NULL }
 382 };
 383
 384 AVFILTER_DEFINE_CLASS(pad_opencl);
 385
 386 const AVFilter ff_vf_pad_opencl = {
 387     .name           = "pad_opencl",
 388     .description    = NULL_IF_CONFIG_SMALL("Pad the input video."),
 389     .priv_size      = sizeof(PadOpenCLContext),
 390     .priv_class     = &pad_opencl_class,
 391     .init           = &ff_opencl_filter_init,
 392     .uninit         = &pad_opencl_uninit,
 393     .query_formats  = &ff_opencl_filter_query_formats,
 394     .inputs         = pad_opencl_inputs,
 395     .outputs        = pad_opencl_outputs,
 396     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE
 397 };