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                "(%zux%zu).\n", plane, global_work[0], global_work[1]);
 148
 149         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
 150                                      global_work, NULL, 0, NULL, NULL);
 151         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
 152     }
 153
 154     cle = clFinish(ctx->command_queue);
 155     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 156
 157     if (ctx->nb_inputs > 0) {
 158         err = av_frame_copy_props(output, ctx->frames[0]);
 159         if (err < 0)
 160             goto fail;
 161     } else {
 162         output->pts = ctx->index;
 163     }
 164     ++ctx->index;
 165
 166     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
 167            av_get_pix_fmt_name(output->format),
 168            output->width, output->height, output->pts);
 169
 170     return ff_filter_frame(outlink, output);
 171
 172 fail:
 173     clFinish(ctx->command_queue);
 174     av_frame_free(&output);
 175     return err;
 176 }
 177
 178 static int program_opencl_request_frame(AVFilterLink *outlink)
 179 {
 180     AVFilterContext *avctx = outlink->src;
 181
 182     return program_opencl_run(avctx);
 183 }
 184
 185 static int program_opencl_filter(FFFrameSync *fs)
 186 {
 187     AVFilterContext    *avctx = fs->parent;
 188     ProgramOpenCLContext *ctx = avctx->priv;
 189     int err, i;
 190
 191     for (i = 0; i < ctx->nb_inputs; i++) {
 192         err = ff_framesync_get_frame(&ctx->fs, i, &ctx->frames[i], 0);
 193         if (err < 0)
 194             return err;
 195     }
 196
 197     return program_opencl_run(avctx);
 198 }
 199
 200 static int program_opencl_activate(AVFilterContext *avctx)
 201 {
 202     ProgramOpenCLContext *ctx = avctx->priv;
 203
 204     av_assert0(ctx->nb_inputs > 0);
 205
 206     return ff_framesync_activate(&ctx->fs);
 207 }
 208
 209 static int program_opencl_config_output(AVFilterLink *outlink)
 210 {
 211     AVFilterContext    *avctx = outlink->src;
 212     ProgramOpenCLContext *ctx = avctx->priv;
 213     int err;
 214
 215     err = ff_opencl_filter_config_output(outlink);
 216     if (err < 0)
 217         return err;
 218
 219     if (ctx->nb_inputs > 0) {
 220         FFFrameSyncIn *in;
 221         int i;
 222
 223         err = ff_framesync_init(&ctx->fs, avctx, ctx->nb_inputs);
 224         if (err < 0)
 225             return err;
 226
 227         ctx->fs.opaque = ctx;
 228         ctx->fs.on_event = &program_opencl_filter;
 229
 230         in = ctx->fs.in;
 231         for (i = 0; i < ctx->nb_inputs; i++) {
 232             const AVFilterLink *inlink = avctx->inputs[i];
 233
 234             in[i].time_base = inlink->time_base;
 235             in[i].sync      = 1;
 236             in[i].before    = EXT_STOP;
 237             in[i].after     = EXT_INFINITY;
 238         }
 239
 240         err = ff_framesync_configure(&ctx->fs);
 241         if (err < 0)
 242             return err;
 243
 244     } else {
 245         outlink->time_base = av_inv_q(ctx->source_rate);
 246     }
 247
 248     return 0;
 249 }
 250
 251 static av_cold int program_opencl_init(AVFilterContext *avctx)
 252 {
 253     ProgramOpenCLContext *ctx = avctx->priv;
 254     int err;
 255
 256     ff_opencl_filter_init(avctx);
 257
 258     ctx->ocf.output_width  = ctx->width;
 259     ctx->ocf.output_height = ctx->height;
 260
 261     if (!strcmp(avctx->filter->name, "openclsrc")) {
 262         if (!ctx->ocf.output_width || !ctx->ocf.output_height) {
 263             av_log(avctx, AV_LOG_ERROR, "OpenCL source requires output "
 264                    "dimensions to be specified.\n");
 265             return AVERROR(EINVAL);
 266         }
 267
 268         ctx->nb_inputs = 0;
 269         ctx->ocf.output_format = ctx->source_format;
 270     } else {
 271         int i;
 272
 273         ctx->frames = av_mallocz_array(ctx->nb_inputs,
 274                                        sizeof(*ctx->frames));
 275         if (!ctx->frames)
 276             return AVERROR(ENOMEM);
 277
 278         for (i = 0; i < ctx->nb_inputs; i++) {
 279             AVFilterPad input;
 280             memset(&input, 0, sizeof(input));
 281
 282             input.type = AVMEDIA_TYPE_VIDEO;
 283             input.name = av_asprintf("input%d", i);
 284             if (!input.name)
 285                 return AVERROR(ENOMEM);
 286
 287             input.config_props = &ff_opencl_filter_config_input;
 288
 289             err = ff_insert_inpad(avctx, i, &input);
 290             if (err < 0) {
 291                 av_freep(&input.name);
 292                 return err;
 293             }
 294         }
 295     }
 296
 297     return 0;
 298 }
 299
 300 static av_cold void program_opencl_uninit(AVFilterContext *avctx)
 301 {
 302     ProgramOpenCLContext *ctx = avctx->priv;
 303     cl_int cle;
 304     int i;
 305
 306     if (ctx->nb_inputs > 0) {
 307         ff_framesync_uninit(&ctx->fs);
 308
 309         av_freep(&ctx->frames);
 310         for (i = 0; i < avctx->nb_inputs; i++)
 311             av_freep(&avctx->input_pads[i].name);
 312     }
 313
 314     if (ctx->kernel) {
 315         cle = clReleaseKernel(ctx->kernel);
 316         if (cle != CL_SUCCESS)
 317             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 318                    "kernel: %d.\n", cle);
 319     }
 320
 321     if (ctx->command_queue) {
 322         cle = clReleaseCommandQueue(ctx->command_queue);
 323         if (cle != CL_SUCCESS)
 324             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 325                    "command queue: %d.\n", cle);
 326     }
 327
 328     ff_opencl_filter_uninit(avctx);
 329 }
 330
 331 #define OFFSET(x) offsetof(ProgramOpenCLContext, x)
 332 #define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM)
 333
 334 #if CONFIG_PROGRAM_OPENCL_FILTER
 335
 336 static const AVOption program_opencl_options[] = {
 337     { "source", "OpenCL program source file", OFFSET(source_file),
 338       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 339     { "kernel", "Kernel name in program",     OFFSET(kernel_name),
 340       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 341
 342     { "inputs", "Number of inputs", OFFSET(nb_inputs),
 343       AV_OPT_TYPE_INT,              { .i64 = 1 }, 1, INT_MAX, FLAGS },
 344
 345     { "size",   "Video size",       OFFSET(width),
 346       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 347     { "s",      "Video size",       OFFSET(width),
 348       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 349
 350     { NULL },
 351 };
 352
 353 FRAMESYNC_DEFINE_CLASS(program_opencl, ProgramOpenCLContext, fs);
 354
 355 static const AVFilterPad program_opencl_outputs[] = {
 356     {
 357         .name         = "default",
 358         .type         = AVMEDIA_TYPE_VIDEO,
 359         .config_props = &program_opencl_config_output,
 360     },
 361     { NULL }
 362 };
 363
 364 AVFilter ff_vf_program_opencl = {
 365     .name           = "program_opencl",
 366     .description    = NULL_IF_CONFIG_SMALL("Filter video using an OpenCL program"),
 367     .priv_size      = sizeof(ProgramOpenCLContext),
 368     .priv_class     = &program_opencl_class,
 369     .preinit        = &program_opencl_framesync_preinit,
 370     .init           = &program_opencl_init,
 371     .uninit         = &program_opencl_uninit,
 372     .query_formats  = &ff_opencl_filter_query_formats,
 373     .activate       = &program_opencl_activate,
 374     .inputs         = NULL,
 375     .outputs        = program_opencl_outputs,
 376     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 377 };
 378
 379 #endif
 380
 381 #if CONFIG_OPENCLSRC_FILTER
 382
 383 static const AVOption openclsrc_options[] = {
 384     { "source", "OpenCL program source file", OFFSET(source_file),
 385       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 386     { "kernel", "Kernel name in program",     OFFSET(kernel_name),
 387       AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
 388
 389     { "size",   "Video size",       OFFSET(width),
 390       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 391     { "s",      "Video size",       OFFSET(width),
 392       AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
 393
 394     { "format", "Video format",     OFFSET(source_format),
 395       AV_OPT_TYPE_PIXEL_FMT,        { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS },
 396
 397     { "rate",   "Video frame rate", OFFSET(source_rate),
 398       AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
 399     { "r",      "Video frame rate", OFFSET(source_rate),
 400       AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
 401
 402     { NULL },
 403 };
 404
 405 AVFILTER_DEFINE_CLASS(openclsrc);
 406
 407 static const AVFilterPad openclsrc_outputs[] = {
 408     {
 409         .name          = "default",
 410         .type          = AVMEDIA_TYPE_VIDEO,
 411         .config_props  = &program_opencl_config_output,
 412         .request_frame = &program_opencl_request_frame,
 413     },
 414     { NULL }
 415 };
 416
 417 AVFilter ff_vsrc_openclsrc = {
 418     .name           = "openclsrc",
 419     .description    = NULL_IF_CONFIG_SMALL("Generate video using an OpenCL program"),
 420     .priv_size      = sizeof(ProgramOpenCLContext),
 421     .priv_class     = &openclsrc_class,
 422     .init           = &program_opencl_init,
 423     .uninit         = &program_opencl_uninit,
 424     .query_formats  = &ff_opencl_filter_query_formats,
 425     .inputs         = NULL,
 426     .outputs        = openclsrc_outputs,
 427     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
 428 };
 429
 430 #endif