path: root/libavfilter/vf_program_opencl.c

/*
 * This file is part of FFmpeg.
 *
 * FFmpeg is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Lesser General Public
 * License as published by the Free Software Foundation; either
 * version 2.1 of the License, or (at your option) any later version.
 *
 * FFmpeg is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Lesser General Public License for more details.
 *
 * You should have received a copy of the GNU Lesser General Public
 * License along with FFmpeg; if not, write to the Free Software
 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
 */

#include "libavutil/avstring.h"
#include "libavutil/log.h"
#include "libavutil/mem.h"
#include "libavutil/opt.h"
#include "libavutil/pixdesc.h"

#include "avfilter.h"
#include "framesync.h"
#include "internal.h"
#include "opencl.h"
#include "video.h"

typedef struct ProgramOpenCLContext {
    OpenCLFilterContext ocf;

    int                 loaded;
    cl_uint             index;
    cl_kernel           kernel;
    cl_command_queue    command_queue;

    FFFrameSync         fs;
    AVFrame           **frames;

    const char         *source_file;
    const char         *kernel_name;
    int                 nb_inputs;
    int                 width, height;
    enum AVPixelFormat  source_format;
    AVRational          source_rate;
} ProgramOpenCLContext;

static int program_opencl_load(AVFilterContext *avctx)
{
    ProgramOpenCLContext *ctx = avctx->priv;
    cl_int cle;
    int err;

    err = ff_opencl_filter_load_program_from_file(avctx, ctx->source_file);
    if (err < 0)
        return err;

    ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                              ctx->ocf.hwctx->device_id,
                                              0, &cle);
    if (!ctx->command_queue) {
        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
               "command queue: %d.\n", cle);
        return AVERROR(EIO);
    }

    ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->kernel_name, &cle);
    if (!ctx->kernel) {
        if (cle == CL_INVALID_KERNEL_NAME) {
            av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in "
                   "program.\n", ctx->kernel_name);
        } else {
            av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
        }
        return AVERROR(EIO);
    }

    ctx->loaded = 1;
    return 0;
}

static int program_opencl_run(AVFilterContext *avctx)
{
    AVFilterLink     *outlink = avctx->outputs[0];
    ProgramOpenCLContext *ctx = avctx->priv;
    AVFrame *output = NULL;
    cl_int cle;
    size_t global_work[2];
    cl_mem src, dst;
    int err, input, plane;

    if (!ctx->loaded) {
        err = program_opencl_load(avctx);
        if (err < 0)
            return err;
    }

    output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
    if (!output) {
        err = AVERROR(ENOMEM);
        goto fail;
    }

    for (plane = 0; plane < FF_ARRAY_ELEMS(output->data); plane++) {
        dst = (cl_mem)output->data[plane];
        if (!dst)
            break;

        cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
        if (cle != CL_SUCCESS) {
            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                   "destination image argument: %d.\n", cle);
            err = AVERROR_UNKNOWN;
            goto fail;
        }
        cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_uint), &ctx->index);
        if (cle != CL_SUCCESS) {
            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                   "index argument: %d.\n", cle);
            err = AVERROR_UNKNOWN;
            goto fail;
        }

        for (input = 0; input < ctx->nb_inputs; input++) {
            av_assert0(ctx->frames[input]);

            src = (cl_mem)ctx->frames[input]->data[plane];
            av_assert0(src);

            cle = clSetKernelArg(ctx->kernel, 2 + input, sizeof(cl_mem), &src);
            if (cle != CL_SUCCESS) {
                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
                       "source image argument %d: %d.\n", input, cle);
                err = AVERROR_UNKNOWN;
                goto fail;
            }
        }

        err = ff_opencl_filter_work_size_from_image(avctx, global_work,
                                                    output, plane, 0);
        if (err < 0)
            goto fail;

        av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
               "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
               plane, global_work[0], global_work[1]);

        cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
                                     global_work, NULL, 0, NULL, NULL);
        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
    }

    cle = clFinish(ctx->command_queue);
    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);

    if (ctx->nb_inputs > 0) {
        err = av_frame_copy_props(output, ctx->frames[0]);
        if (err < 0)
            goto fail;
    } else {
        output->pts = ctx->index;
    }
    ++ctx->index;

    av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
           av_get_pix_fmt_name(output->format),
           output->width, output->height, output->pts);

    return ff_filter_frame(outlink, output);

fail:
    clFinish(ctx->command_queue);
    av_frame_free(&output);
    return err;
}

static int program_opencl_request_frame(AVFilterLink *outlink)
{
    AVFilterContext *avctx = outlink->src;

    return program_opencl_run(avctx);
}

static int program_opencl_filter(FFFrameSync *fs)
{
    AVFilterContext    *avctx = fs->parent;
    ProgramOpenCLContext *ctx = avctx->priv;
    int err, i;

    for (i = 0; i < ctx->nb_inputs; i++) {
        err = ff_framesync_get_frame(&ctx->fs, i, &ctx->frames[i], 0);
        if (err < 0)
            return err;
    }

    return program_opencl_run(avctx);
}

static int program_opencl_activate(AVFilterContext *avctx)
{
    ProgramOpenCLContext *ctx = avctx->priv;

    av_assert0(ctx->nb_inputs > 0);

    return ff_framesync_activate(&ctx->fs);
}

static int program_opencl_config_output(AVFilterLink *outlink)
{
    AVFilterContext    *avctx = outlink->src;
    ProgramOpenCLContext *ctx = avctx->priv;
    int err;

    err = ff_opencl_filter_config_output(outlink);
    if (err < 0)
        return err;

    if (ctx->nb_inputs > 0) {
        FFFrameSyncIn *in;
        int i;

        err = ff_framesync_init(&ctx->fs, avctx, ctx->nb_inputs);
        if (err < 0)
            return err;

        ctx->fs.opaque = ctx;
        ctx->fs.on_event = &program_opencl_filter;

        in = ctx->fs.in;
        for (i = 0; i < ctx->nb_inputs; i++) {
            const AVFilterLink *inlink = avctx->inputs[i];

            in[i].time_base = inlink->time_base;
            in[i].sync      = 1;
            in[i].before    = EXT_STOP;
            in[i].after     = EXT_INFINITY;
        }

        err = ff_framesync_configure(&ctx->fs);
        if (err < 0)
            return err;

    } else {
        outlink->time_base = av_inv_q(ctx->source_rate);
    }

    return 0;
}

static av_cold int program_opencl_init(AVFilterContext *avctx)
{
    ProgramOpenCLContext *ctx = avctx->priv;
    int err;

    ff_opencl_filter_init(avctx);

    ctx->ocf.output_width  = ctx->width;
    ctx->ocf.output_height = ctx->height;

    if (!strcmp(avctx->filter->name, "openclsrc")) {
        if (!ctx->ocf.output_width || !ctx->ocf.output_height) {
            av_log(avctx, AV_LOG_ERROR, "OpenCL source requires output "
                   "dimensions to be specified.\n");
            return AVERROR(EINVAL);
        }

        ctx->nb_inputs = 0;
        ctx->ocf.output_format = ctx->source_format;
    } else {
        int i;

        ctx->frames = av_calloc(ctx->nb_inputs, sizeof(*ctx->frames));
        if (!ctx->frames)
            return AVERROR(ENOMEM);

        for (i = 0; i < ctx->nb_inputs; i++) {
            AVFilterPad input;
            memset(&input, 0, sizeof(input));

            input.type = AVMEDIA_TYPE_VIDEO;
            input.name = av_asprintf("input%d", i);
            if (!input.name)
                return AVERROR(ENOMEM);

            input.config_props = &ff_opencl_filter_config_input;

            err = ff_append_inpad_free_name(avctx, &input);
            if (err < 0)
                return err;
        }
    }

    return 0;
}

static av_cold void program_opencl_uninit(AVFilterContext *avctx)
{
    ProgramOpenCLContext *ctx = avctx->priv;
    cl_int cle;

    if (ctx->nb_inputs > 0) {
        ff_framesync_uninit(&ctx->fs);

        av_freep(&ctx->frames);
    }

    if (ctx->kernel) {
        cle = clReleaseKernel(ctx->kernel);
        if (cle != CL_SUCCESS)
            av_log(avctx, AV_LOG_ERROR, "Failed to release "
                   "kernel: %d.\n", cle);
    }

    if (ctx->command_queue) {
        cle = clReleaseCommandQueue(ctx->command_queue);
        if (cle != CL_SUCCESS)
            av_log(avctx, AV_LOG_ERROR, "Failed to release "
                   "command queue: %d.\n", cle);
    }

    ff_opencl_filter_uninit(avctx);
}

#define OFFSET(x) offsetof(ProgramOpenCLContext, x)
#define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM)

#if CONFIG_PROGRAM_OPENCL_FILTER

static const AVOption program_opencl_options[] = {
    { "source", "OpenCL program source file", OFFSET(source_file),
      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
    { "kernel", "Kernel name in program",     OFFSET(kernel_name),
      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },

    { "inputs", "Number of inputs", OFFSET(nb_inputs),
      AV_OPT_TYPE_INT,              { .i64 = 1 }, 1, INT_MAX, FLAGS },

    { "size",   "Video size",       OFFSET(width),
      AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
    { "s",      "Video size",       OFFSET(width),
      AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },

    { NULL },
};

FRAMESYNC_DEFINE_CLASS(program_opencl, ProgramOpenCLContext, fs);

static const AVFilterPad program_opencl_outputs[] = {
    {
        .name         = "default",
        .type         = AVMEDIA_TYPE_VIDEO,
        .config_props = &program_opencl_config_output,
    },
};

const AVFilter ff_vf_program_opencl = {
    .name           = "program_opencl",
    .description    = NULL_IF_CONFIG_SMALL("Filter video using an OpenCL program"),
    .priv_size      = sizeof(ProgramOpenCLContext),
    .priv_class     = &program_opencl_class,
    .flags          = AVFILTER_FLAG_DYNAMIC_INPUTS,
    .preinit        = &program_opencl_framesync_preinit,
    .init           = &program_opencl_init,
    .uninit         = &program_opencl_uninit,
    .query_formats  = &ff_opencl_filter_query_formats,
    .activate       = &program_opencl_activate,
    .inputs         = NULL,
    FILTER_OUTPUTS(program_opencl_outputs),
    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
};

#endif

#if CONFIG_OPENCLSRC_FILTER

static const AVOption openclsrc_options[] = {
    { "source", "OpenCL program source file", OFFSET(source_file),
      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
    { "kernel", "Kernel name in program",     OFFSET(kernel_name),
      AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },

    { "size",   "Video size",       OFFSET(width),
      AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },
    { "s",      "Video size",       OFFSET(width),
      AV_OPT_TYPE_IMAGE_SIZE,       { .str = NULL }, 0, 0, FLAGS },

    { "format", "Video format",     OFFSET(source_format),
      AV_OPT_TYPE_PIXEL_FMT,        { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS },

    { "rate",   "Video frame rate", OFFSET(source_rate),
      AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },
    { "r",      "Video frame rate", OFFSET(source_rate),
      AV_OPT_TYPE_VIDEO_RATE,       { .str = "25" }, 0, INT_MAX, FLAGS },

    { NULL },
};

AVFILTER_DEFINE_CLASS(openclsrc);

static const AVFilterPad openclsrc_outputs[] = {
    {
        .name          = "default",
        .type          = AVMEDIA_TYPE_VIDEO,
        .config_props  = &program_opencl_config_output,
        .request_frame = &program_opencl_request_frame,
    },
};

const AVFilter ff_vsrc_openclsrc = {
    .name           = "openclsrc",
    .description    = NULL_IF_CONFIG_SMALL("Generate video using an OpenCL program"),
    .priv_size      = sizeof(ProgramOpenCLContext),
    .priv_class     = &openclsrc_class,
    .init           = &program_opencl_init,
    .uninit         = &program_opencl_uninit,
    .query_formats  = &ff_opencl_filter_query_formats,
    .inputs         = NULL,
    FILTER_OUTPUTS(openclsrc_outputs),
    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
};

#endif