[FFmpeg-devel] [PATCH] avfilter: add pad opencl filter

[Mark Thompson]

On 09/02/2020 19:33, Paul B Mahol wrote:
> On 2/9/20, Mark Thompson <sw at jkqxz.net> wrote:
>> On 06/02/2020 18:54, Paul B Mahol wrote:
>>> Signed-off-by: Paul B Mahol <onemda at gmail.com>
>>> ---
>>>  configure                   |   1 +
>>>  doc/filters.texi            |  29 ++++
>>>  libavfilter/Makefile        |   1 +
>>>  libavfilter/allfilters.c    |   1 +
>>>  libavfilter/opencl/pad.cl   |  34 +++++
>>>  libavfilter/opencl_source.h |   1 +
>>>  libavfilter/vf_pad_opencl.c | 289 ++++++++++++++++++++++++++++++++++++
>>>  7 files changed, 356 insertions(+)
>>>  create mode 100644 libavfilter/opencl/pad.cl
>>>  create mode 100644 libavfilter/vf_pad_opencl.c
>>>
>>> ...
>>> +
>>> +static int filter_frame(AVFilterLink *link, AVFrame *input_frame)
>>> +{
>>> +    AVFilterContext *avctx = link->dst;
>>> +    AVFilterLink *outlink = avctx->outputs[0];
>>> +    PadOpenCLContext *pad_ctx = avctx->priv;
>>> +    AVFrame *output_frame = NULL;
>>> +    int err;
>>> +    cl_int cle;
>>> +    size_t global_work[2];
>>> +    cl_mem src, dst;
>>> +
>>> +    if (!input_frame->hw_frames_ctx)
>>> +        return AVERROR(EINVAL);
>>> +
>>> +    if (!pad_ctx->initialized) {
>>> +        err = pad_opencl_init(avctx, input_frame);
>>> +        if (err < 0)
>>> +            goto fail;
>>> +    }
>>> +
>>> +    output_frame = ff_get_video_buffer(outlink, outlink->w, outlink->h);
>>> +    if (!output_frame) {
>>> +        err = AVERROR(ENOMEM);
>>> +        goto fail;
>>> +    }
>>> +
>>> +    for (int p = 0; p < FF_ARRAY_ELEMS(output_frame->data); p++) {
>>> +        cl_float4 pad_color_float;
>>> +        cl_int2 pad_pos;
>>> +
>>> +        if (pad_ctx->is_packed) {
>>> +            pad_color_float = pad_ctx->pad_color_float;
>>> +        } else {
>>> +            pad_color_float.s[0] = pad_ctx->pad_color_float.s[p];
>>> +        }
>>
>> This colour choice is missing some cases: it's off for GBRP (wrong order),
>> and for NV12/P010 (missing the second component on the chroma plane).
>>
>> (Check the format list that hwcontext_opencl logs on AV_LOG_DEBUG in
>> get_constraints() from hwupload.)
> 
> How to fix?

I think those two are the only interesting cases, so just apply them manually.  Something like:

If RGB and planar and p in 0..2 then use pad_color_float.s[0] = pad_ctx->pad_color_float[(p + 1) % 3].

If YUV and planar and p is 1 then also set pad_color_float.s[1] = pad_ctx->pad_color_float.s[p + 1].

>>> +
>>> +        if (p > 0 && p < 3) {
>>> +            pad_pos.s[0] = pad_ctx->pad_pos.s[0] >> pad_ctx->hsub;
>>> +            pad_pos.s[1] = pad_ctx->pad_pos.s[1] >> pad_ctx->vsub;
>>> +        } else {
>>> +            pad_pos.s[0] = pad_ctx->pad_pos.s[0];
>>> +            pad_pos.s[1] = pad_ctx->pad_pos.s[1];
>>> +        }
>>> +
>>> +        src = (cl_mem)input_frame->data[p];
>>> +        dst = (cl_mem)output_frame->data[p];
>>> +
>>> +        if (!dst)
>>> +            break;
>>> +
>>> +        CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 0, cl_mem, &src);
>>> +        CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 1, cl_mem, &dst);
>>> +        CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 2, cl_float4,
>>> &pad_color_float);
>>> +        CL_SET_KERNEL_ARG(pad_ctx->kernel_pad, 3, cl_int2, &pad_pos);
>>> +
>>> +        err = ff_opencl_filter_work_size_from_image(avctx, global_work,
>>> output_frame, p, 16);
>>> +        if (err < 0)
>>> +            goto fail;
>>> +
>>> +        cle = clEnqueueNDRangeKernel(pad_ctx->command_queue,
>>> pad_ctx->kernel_pad, 2, NULL,
>>> +                                     global_work, NULL, 0, NULL, NULL);
>>> +
>>> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue pad kernel:
>>> %d.\n", cle);
>>> +    }
>>> +
>>> +    // Run queued kernel
>>> +    cle = clFinish(pad_ctx->command_queue);
>>> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue:
>>> %d.\n", cle);
>>> +
>>> +    err = av_frame_copy_props(output_frame, input_frame);
>>> +    if (err < 0)
>>> +        goto fail;
>>> +
>>> +    av_frame_free(&input_frame);
>>> +
>>> +    return ff_filter_frame(outlink, output_frame);
>>> +
>>> +fail:
>>> +    clFinish(pad_ctx->command_queue);
>>> +    av_frame_free(&input_frame);
>>> +    av_frame_free(&output_frame);
>>> +    return err;
>>> +}
>>> +
>>> ...
>>> +
>>> +static int pad_opencl_config_output(AVFilterLink *outlink)
>>> +{
>>> +    AVFilterContext *avctx = outlink->src;
>>> +    PadOpenCLContext *ctx = avctx->priv;
>>> +    int err;
>>> +
>>> +    if (ctx->w < avctx->inputs[0]->w ||
>>> +        ctx->h < avctx->inputs[0]->h) {
>>> +        return AVERROR(EINVAL);
>>> +    }
>>> +
>>> +    if (ctx->w > avctx->inputs[0]->w && ctx->h > avctx->inputs[0]->h) {
>>> +        ctx->ocf.output_width  = ctx->w;
>>> +        ctx->ocf.output_height = ctx->h;
>>> +    } else {
>>> +        ctx->ocf.output_width  = avctx->inputs[0]->w;
>>> +        ctx->ocf.output_height = avctx->inputs[0]->h;
>>> +    }
>>
>> This goes wrong if you're only padding in one direction (e.g. to change
>> aspect ratio).
>>
>> Consider a 1080p input with args like h=1200:y=60.
> 
> I do not follow.

If I pad top and bottom only with h=1200:y=60 then the first half of the condition is not true, so it falls into the second branch and incorrectly uses the height of the input stream rather than the height I specified.

>>> +
>>> +    if (ctx->x + avctx->inputs[0]->w > ctx->ocf.output_width ||
>>> +        ctx->y + avctx->inputs[0]->h > ctx->ocf.output_height) {
>>> +        return AVERROR(EINVAL);
>>> +    }
>>> +
>>> +    err = ff_opencl_filter_config_output(outlink);
>>> +    if (err < 0)
>>> +        return err;
>>> +
>>> +    return 0;
>>> +}
>>> +
>>> ...