avfilter: add remap opencl filter

10 files changed, 462 insertions, 2 deletions

diff --git a/Changelog b/Changelog
index ba679aa64e..1a7c84b7f8 100644
--- a/Changelog
+++ b/Changelog
@@ -23,6 +23,7 @@ version 5.1:
 - virtualbass audio filter
 - VDPAU AV1 hwaccel
 - PHM image format support
+- remap_opencl filter
 
 
 version 5.0:
diff --git a/configure b/configure
index fea512e8ef..7d5c4900bf 100755
--- a/configure
+++ b/configure
@@ -3706,6 +3706,7 @@ prewitt_opencl_filter_deps="opencl"
 procamp_vaapi_filter_deps="vaapi"
 program_opencl_filter_deps="opencl"
 pullup_filter_deps="gpl"
+remap_opencl_filter_deps="opencl"
 removelogo_filter_deps="avcodec avformat swscale"
 repeatfields_filter_deps="gpl"
 roberts_opencl_filter_deps="opencl"
diff --git a/doc/filters.texi b/doc/filters.texi
index 5a889895c6..296b0693ae 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -25701,6 +25701,30 @@ __kernel void blend_images(__write_only image2d_t dst,
 
 @end itemize
 
+@section remap_opencl
+
+Remap pixels using 2nd: Xmap and 3rd: Ymap input video stream.
+
+Destination pixel at position (X, Y) will be picked from source (x, y) position
+where x = Xmap(X, Y) and y = Ymap(X, Y). If mapping values are out of range, zero
+value for pixel will be used for destination pixel.
+
+Xmap and Ymap input video streams must be of same dimensions. Output video stream
+will have Xmap/Ymap video stream dimensions.
+Xmap and Ymap input video streams are 32bit float pixel format, single channel.
+
+@table @option
+@item interp
+Specify interpolation used for remapping of pixels.
+Allowed values are @code{near} and @code{linear}.
+Default value is @code{linear}.
+
+@item fill
+Specify the color of the unmapped pixels. For the syntax of this option,
+check the @ref{color syntax,,"Color" section in the ffmpeg-utils
+manual,ffmpeg-utils}. Default color is @code{black}.
+@end table
+
 @section roberts_opencl
 Apply the Roberts cross operator (@url{https://en.wikipedia.org/wiki/Roberts_cross}) to input video stream.
 
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index 22b0a0ca15..139f7cb751 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -420,6 +420,8 @@ OBJS-$(CONFIG_READEIA608_FILTER)             += vf_readeia608.o
 OBJS-$(CONFIG_READVITC_FILTER)               += vf_readvitc.o
 OBJS-$(CONFIG_REALTIME_FILTER)               += f_realtime.o
 OBJS-$(CONFIG_REMAP_FILTER)                  += vf_remap.o framesync.o
+OBJS-$(CONFIG_REMAP_OPENCL_FILTER)           += vf_remap_opencl.o framesync.o opencl.o \
+                                                opencl/remap.o
 OBJS-$(CONFIG_REMOVEGRAIN_FILTER)            += vf_removegrain.o
 OBJS-$(CONFIG_REMOVELOGO_FILTER)             += bbox.o lswsutils.o lavfutils.o vf_removelogo.o
 OBJS-$(CONFIG_REPEATFIELDS_FILTER)           += vf_repeatfields.o
diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
index ec70feef11..3018850b4b 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -399,6 +399,7 @@ extern const AVFilter ff_vf_readeia608;
 extern const AVFilter ff_vf_readvitc;
 extern const AVFilter ff_vf_realtime;
 extern const AVFilter ff_vf_remap;
+extern const AVFilter ff_vf_remap_opencl;
 extern const AVFilter ff_vf_removegrain;
 extern const AVFilter ff_vf_removelogo;
 extern const AVFilter ff_vf_repeatfields;
diff --git a/libavfilter/opencl/remap.cl b/libavfilter/opencl/remap.cl
new file mode 100644
index 0000000000..fba82d134e
--- /dev/null
+++ b/libavfilter/opencl/remap.cl
@@ -0,0 +1,73 @@
+/*
+ * 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
+ */
+
+const sampler_t linear_sampler = (CLK_NORMALIZED_COORDS_FALSE |
+                                  CLK_FILTER_LINEAR);
+
+const sampler_t nearest_sampler = (CLK_NORMALIZED_COORDS_FALSE |
+                                   CLK_FILTER_NEAREST);
+
+__kernel void remap_near(__write_only image2d_t dst,
+                         __read_only  image2d_t src,
+                         __read_only  image2d_t xmapi,
+                         __read_only  image2d_t ymapi,
+                         float4 fill_color)
+{
+    int2 p = (int2)(get_global_id(0), get_global_id(1));
+    int2 dimi = get_image_dim(src);
+    float2 dimf = (float2)(dimi.x, dimi.y);
+    float4 val;
+    int2 mi;
+    float m;
+
+    float4 xmap = read_imagef(xmapi, nearest_sampler, p);
+    float4 ymap = read_imagef(ymapi, nearest_sampler, p);
+    float2 pos  = (float2)(xmap.x, ymap.x);
+    pos.xy = pos.xy * 65535.f;
+
+    mi = ((pos >= (float2)(0.f, 0.f)) * (pos < dimf) * (p <= dimi));
+    m = mi.x && mi.y;
+    val = mix(fill_color, read_imagef(src, nearest_sampler, pos), m);
+
+    write_imagef(dst, p, val);
+}
+
+__kernel void remap_linear(__write_only image2d_t dst,
+                           __read_only  image2d_t src,
+                           __read_only  image2d_t xmapi,
+                           __read_only  image2d_t ymapi,
+                           float4 fill_color)
+{
+    int2 p = (int2)(get_global_id(0), get_global_id(1));
+    int2 dimi = get_image_dim(src);
+    float2 dimf = (float2)(dimi.x, dimi.y);
+    float4 val;
+    int2 mi;
+    float m;
+
+    float4 xmap = read_imagef(xmapi, nearest_sampler, p);
+    float4 ymap = read_imagef(ymapi, nearest_sampler, p);
+    float2 pos  = (float2)(xmap.x, ymap.x);
+    pos.xy = pos.xy * 65535.f;
+
+    mi = ((pos >= (float2)(0.f, 0.f)) * (pos < dimf) * (p <= dimi));
+    m = mi.x && mi.y;
+    val = mix(fill_color, read_imagef(src, linear_sampler, pos), m);
+
+    write_imagef(dst, p, val);
+}
diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
index 7e8133090e..9eac2dc516 100644
--- a/libavfilter/opencl_source.h
+++ b/libavfilter/opencl_source.h
@@ -28,6 +28,7 @@ extern const char *ff_opencl_source_neighbor;
 extern const char *ff_opencl_source_nlmeans;
 extern const char *ff_opencl_source_overlay;
 extern const char *ff_opencl_source_pad;
+extern const char *ff_opencl_source_remap;
 extern const char *ff_opencl_source_tonemap;
 extern const char *ff_opencl_source_transpose;
 extern const char *ff_opencl_source_unsharp;
diff --git a/libavfilter/version.h b/libavfilter/version.h
index 814ab071da..37f03b44a6 100644
--- a/libavfilter/version.h
+++ b/libavfilter/version.h
@@ -31,7 +31,7 @@
 
 #include "version_major.h"
 
-#define LIBAVFILTER_VERSION_MINOR  41
+#define LIBAVFILTER_VERSION_MINOR  42
 #define LIBAVFILTER_VERSION_MICRO 100
 
 
diff --git a/libavfilter/vf_remap_opencl.c b/libavfilter/vf_remap_opencl.c
new file mode 100644
index 0000000000..f3f84bde64
--- /dev/null
+++ b/libavfilter/vf_remap_opencl.c
@@ -0,0 +1,354 @@
+/*
+ * Copyright (c) 2022 Paul B Mahol
+ *
+ * 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/colorspace.h"
+#include "libavutil/imgutils.h"
+#include "libavutil/pixdesc.h"
+#include "libavutil/opt.h"
+#include "avfilter.h"
+#include "drawutils.h"
+#include "formats.h"
+#include "framesync.h"
+#include "internal.h"
+#include "opencl.h"
+#include "opencl_source.h"
+#include "video.h"
+
+typedef struct RemapOpenCLContext {
+    OpenCLFilterContext ocf;
+
+    int nb_planes;
+    int interp;
+    uint8_t fill_rgba[4];
+    cl_float4 cl_fill_color;
+
+    int              initialised;
+    cl_kernel        kernel;
+    cl_command_queue command_queue;
+
+    FFFrameSync fs;
+} RemapOpenCLContext;
+
+#define OFFSET(x) offsetof(RemapOpenCLContext, x)
+#define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
+
+static const AVOption remap_opencl_options[] = {
+    { "interp", "set interpolation method", OFFSET(interp), AV_OPT_TYPE_INT,   {.i64=1}, 0, 1, FLAGS, "interp" },
+    {  "near",   NULL, 0, AV_OPT_TYPE_CONST, {.i64=0}, 0, 0, FLAGS, "interp" },
+    {  "linear", NULL, 0, AV_OPT_TYPE_CONST, {.i64=1}, 0, 0, FLAGS, "interp" },
+    { "fill", "set the color of the unmapped pixels", OFFSET(fill_rgba), AV_OPT_TYPE_COLOR, {.str="black"}, .flags = FLAGS },
+    { NULL }
+};
+
+AVFILTER_DEFINE_CLASS(remap_opencl);
+
+static av_cold int remap_opencl_init(AVFilterContext *avctx)
+{
+    return ff_opencl_filter_init(avctx);
+}
+
+static const char *kernels[] = { "remap_near", "remap_linear" };
+
+static int remap_opencl_load(AVFilterContext *avctx,
+                             enum AVPixelFormat main_format,
+                             enum AVPixelFormat xmap_format,
+                             enum AVPixelFormat ymap_format)
+{
+    RemapOpenCLContext *ctx = avctx->priv;
+    cl_int cle;
+    const char *source = ff_opencl_source_remap;
+    const char *kernel = kernels[ctx->interp];
+    const AVPixFmtDescriptor *main_desc;
+    int err, main_planes;
+    const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(main_format);
+    int is_rgb = !!(desc->flags & AV_PIX_FMT_FLAG_RGB);
+    const float scale = 1.f / 255.f;
+    uint8_t rgba_map[4];
+
+    ff_fill_rgba_map(rgba_map, main_format);
+
+    if (is_rgb) {
+        ctx->cl_fill_color.s[rgba_map[0]] = ctx->fill_rgba[0] * scale;
+        ctx->cl_fill_color.s[rgba_map[1]] = ctx->fill_rgba[1] * scale;
+        ctx->cl_fill_color.s[rgba_map[2]] = ctx->fill_rgba[2] * scale;
+        ctx->cl_fill_color.s[rgba_map[3]] = ctx->fill_rgba[3] * scale;
+    } else {
+        ctx->cl_fill_color.s[0] = RGB_TO_Y_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2]) * scale;
+        ctx->cl_fill_color.s[1] = RGB_TO_U_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
+        ctx->cl_fill_color.s[2] = RGB_TO_V_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
+        ctx->cl_fill_color.s[3] = ctx->fill_rgba[3] * scale;
+    }
+
+    main_desc = av_pix_fmt_desc_get(main_format);
+
+    main_planes = 0;
+    for (int i = 0; i < main_desc->nb_components; i++)
+        main_planes = FFMAX(main_planes,
+                            main_desc->comp[i].plane + 1);
+
+    ctx->nb_planes = main_planes;
+
+    err = ff_opencl_filter_load_program(avctx, &source, 1);
+    if (err < 0)
+        goto fail;
+
+    ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
+                                              ctx->ocf.hwctx->device_id,
+                                              0, &cle);
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
+
+    ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
+
+    ctx->initialised = 1;
+    return 0;
+
+fail:
+    if (ctx->command_queue)
+        clReleaseCommandQueue(ctx->command_queue);
+    if (ctx->kernel)
+        clReleaseKernel(ctx->kernel);
+    return err;
+}
+
+static int remap_opencl_process_frame(FFFrameSync *fs)
+{
+    AVFilterContext *avctx = fs->parent;
+    AVFilterLink *outlink = avctx->outputs[0];
+    RemapOpenCLContext *ctx = avctx->priv;
+    AVFrame *input_main, *input_xmap, *input_ymap;
+    AVFrame *output;
+    cl_mem mem;
+    cl_int cle;
+    size_t global_work[2];
+    int kernel_arg = 0;
+    int err, plane;
+
+    err = ff_framesync_get_frame(fs, 0, &input_main, 0);
+    if (err < 0)
+        return err;
+    err = ff_framesync_get_frame(fs, 1, &input_xmap, 0);
+    if (err < 0)
+        return err;
+    err = ff_framesync_get_frame(fs, 2, &input_ymap, 0);
+    if (err < 0)
+        return err;
+
+    if (!ctx->initialised) {
+        AVHWFramesContext *main_fc =
+           (AVHWFramesContext*)input_main->hw_frames_ctx->data;
+        AVHWFramesContext *xmap_fc =
+            (AVHWFramesContext*)input_xmap->hw_frames_ctx->data;
+        AVHWFramesContext *ymap_fc =
+            (AVHWFramesContext*)input_ymap->hw_frames_ctx->data;
+
+        err = remap_opencl_load(avctx, main_fc->sw_format,
+                                xmap_fc->sw_format,
+                                ymap_fc->sw_format);
+        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 < ctx->nb_planes; plane++) {
+        cl_float4 cl_fill_color;
+        kernel_arg = 0;
+
+        if (ctx->nb_planes == 1)
+            cl_fill_color = ctx->cl_fill_color;
+        else
+            cl_fill_color.s[0] = ctx->cl_fill_color.s[plane];
+
+        mem = (cl_mem)output->data[plane];
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
+
+        mem = (cl_mem)input_main->data[plane];
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
+
+        mem = (cl_mem)input_xmap->data[0];
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
+
+        mem = (cl_mem)input_ymap->data[0];
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+        kernel_arg++;
+
+        CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float4, &cl_fill_color);
+        kernel_arg++;
+
+        err = ff_opencl_filter_work_size_from_image(avctx, global_work,
+                                                    output, plane, 0);
+        if (err < 0)
+            goto fail;
+
+        cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+                                     global_work, NULL, 0, NULL, NULL);
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue remap kernel "
+                         "for plane %d: %d.\n", plane, cle);
+    }
+
+    cle = clFinish(ctx->command_queue);
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
+
+    err = av_frame_copy_props(output, input_main);
+
+    av_log(avctx, 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:
+    av_frame_free(&output);
+    return err;
+}
+
+static int config_output(AVFilterLink *outlink)
+{
+    AVFilterContext *ctx = outlink->src;
+    RemapOpenCLContext *s = ctx->priv;
+    AVFilterLink *srclink = ctx->inputs[0];
+    AVFilterLink *xlink = ctx->inputs[1];
+    AVFilterLink *ylink = ctx->inputs[2];
+    FFFrameSyncIn *in;
+    int ret;
+
+    if (xlink->w != ylink->w || xlink->h != ylink->h) {
+        av_log(ctx, AV_LOG_ERROR, "Second input link %s parameters "
+               "(size %dx%d) do not match the corresponding "
+               "third input link %s parameters (%dx%d)\n",
+               ctx->input_pads[1].name, xlink->w, xlink->h,
+               ctx->input_pads[2].name, ylink->w, ylink->h);
+        return AVERROR(EINVAL);
+    }
+
+    outlink->w = xlink->w;
+    outlink->h = xlink->h;
+    outlink->sample_aspect_ratio = srclink->sample_aspect_ratio;
+    outlink->frame_rate = srclink->frame_rate;
+
+    ret = ff_framesync_init(&s->fs, ctx, 3);
+    if (ret < 0)
+        return ret;
+
+    in = s->fs.in;
+    in[0].time_base = srclink->time_base;
+    in[1].time_base = xlink->time_base;
+    in[2].time_base = ylink->time_base;
+    in[0].sync   = 2;
+    in[0].before = EXT_STOP;
+    in[0].after  = EXT_STOP;
+    in[1].sync   = 1;
+    in[1].before = EXT_NULL;
+    in[1].after  = EXT_INFINITY;
+    in[2].sync   = 1;
+    in[2].before = EXT_NULL;
+    in[2].after  = EXT_INFINITY;
+    s->fs.opaque   = s;
+    s->fs.on_event = remap_opencl_process_frame;
+
+    ret = ff_framesync_configure(&s->fs);
+    outlink->time_base = s->fs.time_base;
+    if (ret < 0)
+        return ret;
+
+    s->ocf.output_width  = outlink->w;
+    s->ocf.output_height = outlink->h;
+
+    return ff_opencl_filter_config_output(outlink);
+}
+
+static int activate(AVFilterContext *ctx)
+{
+    RemapOpenCLContext *s = ctx->priv;
+    return ff_framesync_activate(&s->fs);
+}
+
+static av_cold void remap_opencl_uninit(AVFilterContext *avctx)
+{
+    RemapOpenCLContext *ctx = avctx->priv;
+    cl_int cle;
+
+    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);
+
+    ff_framesync_uninit(&ctx->fs);
+}
+
+static const AVFilterPad remap_opencl_inputs[] = {
+    {
+        .name         = "source",
+        .type         = AVMEDIA_TYPE_VIDEO,
+        .config_props = &ff_opencl_filter_config_input,
+    },
+    {
+        .name         = "xmap",
+        .type         = AVMEDIA_TYPE_VIDEO,
+        .config_props = &ff_opencl_filter_config_input,
+    },
+    {
+        .name         = "ymap",
+        .type         = AVMEDIA_TYPE_VIDEO,
+        .config_props = &ff_opencl_filter_config_input,
+    },
+};
+
+static const AVFilterPad remap_opencl_outputs[] = {
+    {
+        .name          = "default",
+        .type          = AVMEDIA_TYPE_VIDEO,
+        .config_props  = config_output,
+    },
+};
+
+const AVFilter ff_vf_remap_opencl = {
+    .name          = "remap_opencl",
+    .description   = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."),
+    .priv_size     = sizeof(RemapOpenCLContext),
+    .init          = remap_opencl_init,
+    .uninit        = remap_opencl_uninit,
+    .activate      = activate,
+    FILTER_INPUTS(remap_opencl_inputs),
+    FILTER_OUTPUTS(remap_opencl_outputs),
+    FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
+    .priv_class    = &remap_opencl_class,
+    .flags_internal  = FF_FILTER_FLAG_HWFRAME_AWARE,
+};
diff --git a/libavutil/hwcontext_opencl.c b/libavutil/hwcontext_opencl.c
index 4f4bd13405..32aec7d228 100644
--- a/libavutil/hwcontext_opencl.c
+++ b/libavutil/hwcontext_opencl.c
@@ -1411,7 +1411,8 @@ static int opencl_get_plane_format(enum AVPixelFormat pixfmt,
         // The bits in each component must be packed in the
         // most-significant-bits of the relevant bytes.
         if (comp->shift + comp->depth != 8 &&
-            comp->shift + comp->depth != 16)
+            comp->shift + comp->depth != 16 &&
+            comp->shift + comp->depth != 32)
             return AVERROR(EINVAL);
         // The depth must not vary between components.
         if (depth && comp->depth != depth)
@@ -1455,6 +1456,8 @@ static int opencl_get_plane_format(enum AVPixelFormat pixfmt,
     } else {
         if (depth <= 16)
             image_format->image_channel_data_type = CL_UNORM_INT16;
+        else if (depth == 32)
+            image_format->image_channel_data_type = CL_FLOAT;
         else
             return AVERROR(EINVAL);
     }