git.sesse.net Git - ffmpeg/blob - libavfilter/vf_program_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/avstring.h"
  20 #include "libavutil/log.h"
  21 #include "libavutil/mem.h"
  22 #include "libavutil/opt.h"
  23 #include "libavutil/pixdesc.h"
  24
  25 #include "avfilter.h"
  26 #include "framesync.h"
  27 #include "internal.h"
  28 #include "opencl.h"
  29 #include "video.h"
  30
  31 typedef struct ProgramOpenCLContext {
  32     OpenCLFilterContext ocf;
  33
  34     int                 loaded;
  35     cl_uint             index;
  36     cl_kernel           kernel;
  37     cl_command_queue    command_queue;
  38
  39     FFFrameSync         fs;
  40     AVFrame           **frames;
  41
  42     const char         *source_file;
  43     const char         *kernel_name;
  44     int                 nb_inputs;
  45     int                 width, height;
  46     enum AVPixelFormat  source_format;
  47     AVRational          source_rate;
  48 } ProgramOpenCLContext;
  49
  50 static int program_opencl_load(AVFilterContext *avctx)
  51 {
  52     ProgramOpenCLContext *ctx = avctx->priv;
  53     cl_int cle;
  54     int err;
  55
  56     err = ff_opencl_filter_load_program_from_file(avctx, ctx->source_file);
  57     if (err < 0)
  58         return err;
  59
  60     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
  61                                               ctx->ocf.hwctx->device_id,
  62                                               0, &cle);
  63     if (!ctx->command_queue) {
  64         av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
  65                "command queue: %d.\n", cle);
  66         return AVERROR(EIO);
  67     }
  68
  69     ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->kernel_name, &cle);
  70     if (!ctx->kernel) {
  71         if (cle == CL_INVALID_KERNEL_NAME) {
  72             av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in "
  73                    "program.\n", ctx->kernel_name);
  74         } else {
  75             av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
  76         }
  77         return AVERROR(EIO);
  78     }
  79
  80     ctx->loaded = 1;
  81     return 0;
  82 }
  83
  84 static int program_opencl_run(AVFilterContext *avctx)
  85 {
  86     AVFilterLink     *outlink = avctx->outputs[0];
  87     ProgramOpenCLContext *ctx = avctx->priv;
  88     AVFrame *output = NULL;
  89     cl_int cle;
  90     size_t global_work[2];
  91     cl_mem src, dst;
  92     int err, input, plane;
  93
  94     if (!ctx->loaded) {
  95         err = program_opencl_load(avctx);
  96         if (err < 0)
  97             return err;
  98     }
  99
 100     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
 101     if (!output) {
 102         err = AVERROR(ENOMEM);
 103         goto fail;
 104     }
 105
 106     for (plane = 0; plane < FF_ARRAY_ELEMS(output->data); plane++) {
 107         dst = (cl_mem)output->data[plane];
 108         if (!dst)
 109             break;
 110
 111         cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
 112         if (cle != CL_SUCCESS) {
 113             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 114                    "destination image argument: %d.\n", cle);
 115             err = AVERROR_UNKNOWN;
 116             goto fail;
 117         }
 118         cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_uint), &ctx->index);
 119         if (cle != CL_SUCCESS) {
 120             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 121                    "index argument: %d.\n", cle);
 122             err = AVERROR_UNKNOWN;
 123             goto fail;
 124         }
 125
 126         for (input = 0; input < ctx->nb_inputs; input++) {
 127             av_assert0(ctx->frames[input]);
 128
 129             src = (cl_mem)ctx->frames[input]->data[plane];
 130             av_assert0(src);
 131
 132             cle = clSetKernelArg(ctx->kernel, 2 + input, sizeof(cl_mem), &src);
 133             if (cle != CL_SUCCESS) {
 134                 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
 135                        "source image argument %d: %d.\n", input, cle);
 136                 err = AVERROR_UNKNOWN;
 137                 goto fail;
 138             }
 139         }
 140
 141         err = ff_opencl_filter_work_size_from_image(avctx, global_work,
 142                                                     output, plane, 0);
 143         if (err < 0)
 144             goto fail;
 145
 146         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
 147                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
 148                plane, global_work[0], global_work[1]);
 149
 150         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
 151                                      global_work, NULL, 0, NULL, NULL);
 152         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
 153     }
 154
 155     cle = clFinish(ctx->command_queue);
 156     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 157
 158     if (ctx->nb_inputs > 0) {
 159         err = av_frame_copy_props(output, ctx->frames[0]);
 160         if (err < 0)
 161             goto fail;
 162     } else {
 163         output->pts = ctx->index;
 164     }
 165     ++ctx->index;
 166
 167     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
 168            av_get_pix_fmt_name(output->format),
 169            output->width, output->height, output->pts);
 170
 171     return ff_filter_frame(outlink, output);
 172
 173 fail:
 174     clFinish(ctx->command_queue);
 175     av_frame_free(&output);
 176     return err;
 177 }
 178
 179 static int program_opencl_request_frame(AVFilterLink *outlink)
 180 {
 181     AVFilterContext *avctx = outlink->src;
 182
 183     return program_opencl_run(avctx);
 184 }
 185
 186 static int program_opencl_filter(FFFrameSync *fs)
 187 {
 188     AVFilterContext    *avctx = fs->parent;
 189     ProgramOpenCLContext *ctx = avctx->priv;
 190     int err, i;
 191
 192     for (i = 0; i < ctx->nb_inputs; i++) {
 193         err = ff_framesync_get_frame(&ctx->fs, i, &ctx->frames[i], 0);
 194         if (err < 0)
 195             return err;
 196     }
 197
 198     return program_opencl_run(avctx);
 199 }
 200
 201 static int program_opencl_activate(AVFilterContext *avctx)
 202 {
 203     ProgramOpenCLContext *ctx = avctx->priv;
 204
 205     av_assert0(ctx->nb_inputs > 0);
 206
 207     return ff_framesync_activate(&ctx->fs);
 208 }
 209
 210 static int program_opencl_config_output(AVFilterLink *outlink)
 211 {
 212     AVFilterContext    *avctx = outlink->src;
 213     ProgramOpenCLContext *ctx = avctx->priv;
 214     int err;
 215
 216     err = ff_opencl_filter_config_output(outlink);
 217     if (err < 0)
 218         return err;
 219
 220     if (ctx->nb_inputs > 0) {
 221         FFFrameSyncIn *in;
 222         int i;
 223
 224         err = ff_framesync_init(&ctx->fs, avctx, ctx->nb_inputs);
 225         if (err < 0)
 226             return err;
 227
 228         ctx->fs.opaque = ctx;
 229         ctx->fs.on_event = &program_opencl_filter;
 230
 231         in = ctx->fs.in;
 232         for (i = 0; i < ctx->nb_inputs; i++) {
 233             const AVFilterLink *inlink = avctx->inputs[i];
 234
 235             in[i].time_base = inlink->time_base;
 236             in[i].sync      = 1;
 237             in[i].before    = EXT_STOP;
 238             in[i].after     = EXT_INFINITY;
 239         }
 240
 241         err = ff_framesync_configure(&ctx->fs);
 242         if (err < 0)
 243             return err;
 244
 245     } else {
 246         outlink->time_base = av_inv_q(ctx->source_rate);
 247     }
 248
 249     return 0;
 250 }
 251
 252 static av_cold int program_opencl_init(AVFilterContext *avctx)
 253 {
 254     ProgramOpenCLContext *ctx = avctx->priv;
 255     int err;
 256
 257     ff_opencl_filter_init(avctx);
 258
 259     ctx->ocf.output_width  = ctx->width;
 260     ctx->ocf.output_height = ctx->height;
 261
 262     if (!strcmp(avctx->filter->name, "openclsrc")) {
 263         if (!ctx->ocf.output_width || !ctx->ocf.output_height) {
 264             av_log(avctx, AV_LOG_ERROR, "OpenCL source requires output "
 265                    "dimensions to be specified.\n");
 266             return AVERROR(EINVAL);
 267         }
 268
 269         ctx->nb_inputs = 0;
 270         ctx->ocf.output_format = ctx->source_format;
 271     } else {
 272         int i;
 273
 274         ctx->frames = av_mallocz_array(ctx->nb_inputs,
 275                                        sizeof(*ctx->frames));
 276         if (!ctx->frames)
 277             return AVERROR(ENOMEM);
 278
 279         for (i = 0; i < ctx->nb_inputs; i++) {
 280             AVFilterPad input;
 281             memset(&input, 0, sizeof(input));
 282
 283             input.type = AVMEDIA_TYPE_VIDEO;
 284             input.name = av_asprintf("input%d", i);
 285             if (!input.name)
 286                 return AVERROR(ENOMEM);
 287
 288             input.config_props = &ff_opencl_filter_config_input;
 289
 290             err = ff_insert_inpad(avctx, i, &input);
 291             if (err < 0) {
 292                 av_freep(&input.name);
 293                 return err;
 294             }
 295         }
 296     }
 297
 298     return 0;
 299 }
 300
 301 static av_cold void program_opencl_uninit(AVFilterContext *avctx)
 302 {
 303     ProgramOpenCLContext *ctx = avctx->priv;
 304     cl_int cle;
 305     int i;
 306
 307     if (ctx->nb_inputs > 0) {
 308         ff_framesync_uninit(&ctx->fs);
 309
 310         av_freep(&ctx->frames);
 311         for (i = 0; i < avctx->nb_inputs; i++)
 312             av_freep(&avctx->input_pads[i].name);
 313     }
 314
 315     if (ctx->kernel) {
 316         cle = clReleaseKernel(ctx->kernel);
 317         if (cle != CL_SUCCESS)
 318             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 319                    "kernel: %d.\n", cle);
 320     }
 321
 322     if (ctx->command_queue) {
 323         cle = clReleaseCommandQueue(ctx->command_queue);
 324         if (cle != CL_SUCCESS)
 325             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 326                    "command queue: %d.\n", cle);
 327     }
 328
 329     ff_opencl_filter_uninit(avctx);
 330 }
 331
 332 #define OFFSET(x) offsetof(ProgramOpenCLContext, x)
 333 #define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM)
 334
 335 #if CONFIG_PROGRAM_OPENCL_FILTER
 336
 337 static const AVOption program_opencl_options[] = {
 338     { "source", "OpenCL program source file", OFFSET(source_file),
 339       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 340     { "kernel", "Kernel name in program",     OFFSET(kernel_name),
 341       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 342
 343     { "inputs", "Number of inputs", OFFSET(nb_inputs),
 344       AV_OPT_TYPE_INT,              { .i64 = 1 }, 1, INT_MAX, FLAGS },
 345
 346     { "size",   "Video size",       OFFSET(width),
 347       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 348     { "s",      "Video size",       OFFSET(width),
 349       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 350
 351     { NULL },
 352 };
 353
 354 FRAMESYNC_DEFINE_CLASS(program_opencl, ProgramOpenCLContext, fs);
 355
 356 static const AVFilterPad program_opencl_outputs[] = {
 357     {
 358         .name         = "default",
 359         .type         = AVMEDIA_TYPE_VIDEO,
 360         .config_props = &program_opencl_config_output,
 361     },
 362     { NULL }
 363 };
 364
 365 AVFilter ff_vf_program_opencl = {
 366     .name           = "program_opencl",
 367     .description    = NULL_IF_CONFIG_SMALL("Filter video using an OpenCL program"),
 368     .priv_size      = sizeof(ProgramOpenCLContext),
 369     .priv_class     = &program_opencl_class,
 370     .preinit        = &program_opencl_framesync_preinit,
 371     .init           = &program_opencl_init,
 372     .uninit         = &program_opencl_uninit,
 373     .query_formats  = &ff_opencl_filter_query_formats,
 374     .activate       = &program_opencl_activate,
 375     .inputs         = NULL,
 376     .outputs        = program_opencl_outputs,
 377     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 378 };
 379
 380 #endif
 381
 382 #if CONFIG_OPENCLSRC_FILTER
 383
 384 static const AVOption openclsrc_options[] = {
 385     { "source", "OpenCL program source file", OFFSET(source_file),
 386       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 387     { "kernel", "Kernel name in program",     OFFSET(kernel_name),
 388       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 389
 390     { "size",   "Video size",       OFFSET(width),
 391       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 392     { "s",      "Video size",       OFFSET(width),
 393       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 394
 395     { "format", "Video format",     OFFSET(source_format),
 396       AV_OPT_TYPE_PIXEL_FMT,        { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS },
 397
 398     { "rate",   "Video frame rate", OFFSET(source_rate),
 399       AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
 400     { "r",      "Video frame rate", OFFSET(source_rate),
 401       AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
 402
 403     { NULL },
 404 };
 405
 406 AVFILTER_DEFINE_CLASS(openclsrc);
 407
 408 static const AVFilterPad openclsrc_outputs[] = {
 409     {
 410         .name          = "default",
 411         .type          = AVMEDIA_TYPE_VIDEO,
 412         .config_props  = &program_opencl_config_output,
 413         .request_frame = &program_opencl_request_frame,
 414     },
 415     { NULL }
 416 };
 417
 418 AVFilter ff_vsrc_openclsrc = {
 419     .name           = "openclsrc",
 420     .description    = NULL_IF_CONFIG_SMALL("Generate video using an OpenCL program"),
 421     .priv_size      = sizeof(ProgramOpenCLContext),
 422     .priv_class     = &openclsrc_class,
 423     .init           = &program_opencl_init,
 424     .uninit         = &program_opencl_uninit,
 425     .query_formats  = &ff_opencl_filter_query_formats,
 426     .inputs         = NULL,
 427     .outputs        = openclsrc_outputs,
 428     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 429 };
 430
 431 #endif