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