On 09/02/2020 19:33, Paul B Mahol wrote:
> On 2/9/20, Mark Thompson <s...@jkqxz.net> wrote:
>> On 06/02/2020 18:54, Paul B Mahol wrote:
>>> Signed-off-by: Paul B Mahol <one...@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;
>>> +}
>>> +
>>> ...
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-requ...@ffmpeg.org with subject "unsubscribe".
