* [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
@ 2022-06-29 19:36 Paul B Mahol
2022-06-30 6:59 ` Paul B Mahol
0 siblings, 1 reply; 8+ messages in thread
From: Paul B Mahol @ 2022-06-29 19:36 UTC (permalink / raw)
To: FFmpeg development discussions and patches
[-- Attachment #1: Type: text/plain, Size: 26 bytes --]
Hello,
patches attached.
[-- Attachment #2: 0002-avutil-hwcontext_opencl-add-support-fro-FLOAT-format.patch --]
[-- Type: text/x-patch, Size: 1426 bytes --]
From 87841c44678f4ea205e4f25213596ef3eb0744b3 Mon Sep 17 00:00:00 2001
From: Paul B Mahol <onemda@gmail.com>
Date: Wed, 29 Jun 2022 21:32:42 +0200
Subject: [PATCH 2/2] avutil/hwcontext_opencl: add support fro FLOAT formats
Signed-off-by: Paul B Mahol <onemda@gmail.com>
---
libavutil/hwcontext_opencl.c | 5 ++++-
1 file changed, 4 insertions(+), 1 deletion(-)
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);
}
--
2.36.1
[-- Attachment #3: 0001-avfilter-add-remap-opencl-filter.patch --]
[-- Type: text/x-patch, Size: 15092 bytes --]
From 011ec1b924adad0a46ff036ebed13d24bca034d9 Mon Sep 17 00:00:00 2001
From: Paul B Mahol <onemda@gmail.com>
Date: Wed, 29 Jun 2022 19:12:24 +0200
Subject: [PATCH 1/2] avfilter: add remap opencl filter
Signed-off-by: Paul B Mahol <onemda@gmail.com>
---
libavfilter/Makefile | 2 +
libavfilter/allfilters.c | 1 +
libavfilter/opencl/remap.cl | 39 ++++
libavfilter/opencl_source.h | 1 +
libavfilter/vf_remap_opencl.c | 329 ++++++++++++++++++++++++++++++++++
5 files changed, 372 insertions(+)
create mode 100644 libavfilter/opencl/remap.cl
create mode 100644 libavfilter/vf_remap_opencl.c
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index b9ce1a715b..367eb92063 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -421,6 +421,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 0152acbb81..05f0fa85db 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -400,6 +400,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..8851cdc429
--- /dev/null
+++ b/libavfilter/opencl/remap.cl
@@ -0,0 +1,39 @@
+/*
+ * 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(__write_only image2d_t dst,
+ __read_only image2d_t src,
+ __read_only image2d_t xmapi,
+ __read_only image2d_t ymapi)
+{
+ int2 p = (int2)(get_global_id(0), get_global_id(1));
+
+ 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;
+ float4 val = read_imagef(src, linear_sampler, pos);
+
+ 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/vf_remap_opencl.c b/libavfilter/vf_remap_opencl.c
new file mode 100644
index 0000000000..0282b6b4d0
--- /dev/null
+++ b/libavfilter/vf_remap_opencl.c
@@ -0,0 +1,329 @@
+/*
+ * 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 nb_components;
+ uint8_t fill_rgba[4];
+ int fill_color[4];
+
+ 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[] = {
+ { "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 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 = "remap";
+ const AVPixFmtDescriptor *main_desc, *xmap_desc, *ymap_desc;
+ int err, main_planes, xmap_planes, ymap_planes;
+
+ main_desc = av_pix_fmt_desc_get(main_format);
+ xmap_desc = av_pix_fmt_desc_get(xmap_format);
+ ymap_desc = av_pix_fmt_desc_get(ymap_format);
+
+ main_planes = xmap_planes = ymap_planes = 0;
+ for (int i = 0; i < main_desc->nb_components; i++)
+ main_planes = FFMAX(main_planes,
+ main_desc->comp[i].plane + 1);
+ for (int i = 0; i < xmap_desc->nb_components; i++)
+ xmap_planes = FFMAX(xmap_planes,
+ xmap_desc->comp[i].plane + 1);
+ for (int i = 0; i < ymap_desc->nb_components; i++)
+ ymap_planes = FFMAX(ymap_planes,
+ ymap_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++) {
+ kernel_arg = 0;
+
+ 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++;
+
+ 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;
+
+ ret = ff_opencl_filter_config_output(outlink);
+ if (ret < 0)
+ return 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;
+
+ return ret;
+}
+
+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,
+};
--
2.36.1
[-- Attachment #4: Type: text/plain, Size: 251 bytes --]
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-06-29 19:36 [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter Paul B Mahol
@ 2022-06-30 6:59 ` Paul B Mahol
2022-06-30 8:34 ` Anton Khirnov
2022-07-01 21:04 ` Michael Niedermayer
0 siblings, 2 replies; 8+ messages in thread
From: Paul B Mahol @ 2022-06-30 6:59 UTC (permalink / raw)
To: FFmpeg development discussions and patches
[-- Attachment #1: Type: text/plain, Size: 9 bytes --]
Updated.
[-- Attachment #2: 0001-avfilter-add-remap-opencl-filter.patch --]
[-- Type: text/x-patch, Size: 19083 bytes --]
From 2b6b653bebec01d9d523102c1bfe2ce4b1be93dd Mon Sep 17 00:00:00 2001
From: Paul B Mahol <onemda@gmail.com>
Date: Wed, 29 Jun 2022 19:12:24 +0200
Subject: [PATCH] avfilter: add remap opencl filter
Signed-off-by: Paul B Mahol <onemda@gmail.com>
---
libavfilter/Makefile | 2 +
libavfilter/allfilters.c | 1 +
libavfilter/opencl/remap.cl | 69 +++++++
libavfilter/opencl_source.h | 1 +
libavfilter/vf_remap_opencl.c | 362 ++++++++++++++++++++++++++++++++++
libavutil/hwcontext_opencl.c | 5 +-
6 files changed, 439 insertions(+), 1 deletion(-)
create mode 100644 libavfilter/opencl/remap.cl
create mode 100644 libavfilter/vf_remap_opencl.c
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index b9ce1a715b..367eb92063 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -421,6 +421,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 0152acbb81..05f0fa85db 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -400,6 +400,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..4252bce355
--- /dev/null
+++ b/libavfilter/opencl/remap.cl
@@ -0,0 +1,69 @@
+/*
+ * 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;
+ 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;
+
+ m = ((pos.x >= 0) * (pos.y >= 0) * (pos.x < dimf.x) * (pos.y < dimf.y) * (p.x <= dimi.x) * (p.y <= dimi.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;
+ 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;
+
+ m = ((pos.x >= 0) * (pos.y >= 0) * (pos.x < dimf.x) * (pos.y < dimf.y) * (p.x <= dimi.x) * (p.y <= dimi.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/vf_remap_opencl.c b/libavfilter/vf_remap_opencl.c
new file mode 100644
index 0000000000..ce7df1c0d5
--- /dev/null
+++ b/libavfilter/vf_remap_opencl.c
@@ -0,0 +1,362 @@
+/*
+ * 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, *xmap_desc, *ymap_desc;
+ int err, main_planes, xmap_planes, ymap_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);
+ xmap_desc = av_pix_fmt_desc_get(xmap_format);
+ ymap_desc = av_pix_fmt_desc_get(ymap_format);
+
+ main_planes = xmap_planes = ymap_planes = 0;
+ for (int i = 0; i < main_desc->nb_components; i++)
+ main_planes = FFMAX(main_planes,
+ main_desc->comp[i].plane + 1);
+ for (int i = 0; i < xmap_desc->nb_components; i++)
+ xmap_planes = FFMAX(xmap_planes,
+ xmap_desc->comp[i].plane + 1);
+ for (int i = 0; i < ymap_desc->nb_components; i++)
+ ymap_planes = FFMAX(ymap_planes,
+ ymap_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);
}
--
2.36.1
[-- Attachment #3: Type: text/plain, Size: 251 bytes --]
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-06-30 6:59 ` Paul B Mahol
@ 2022-06-30 8:34 ` Anton Khirnov
2022-06-30 8:47 ` Paul B Mahol
2022-07-01 21:04 ` Michael Niedermayer
1 sibling, 1 reply; 8+ messages in thread
From: Anton Khirnov @ 2022-06-30 8:34 UTC (permalink / raw)
To: FFmpeg development discussions and patches
Quoting Paul B Mahol (2022-06-30 08:59:53)
> Updated.
This comment is entirely useless. What was updated and why?
--
Anton Khirnov
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-06-30 8:34 ` Anton Khirnov
@ 2022-06-30 8:47 ` Paul B Mahol
0 siblings, 0 replies; 8+ messages in thread
From: Paul B Mahol @ 2022-06-30 8:47 UTC (permalink / raw)
To: FFmpeg development discussions and patches
On Thu, Jun 30, 2022 at 10:34 AM Anton Khirnov <anton@khirnov.net> wrote:
> Quoting Paul B Mahol (2022-06-30 08:59:53)
> > Updated.
>
> This comment is entirely useless. What was updated and why?
>
Lot of stuff, so do you not need to improve it.
>
> --
> Anton Khirnov
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-06-30 6:59 ` Paul B Mahol
2022-06-30 8:34 ` Anton Khirnov
@ 2022-07-01 21:04 ` Michael Niedermayer
2022-07-01 21:39 ` Paul B Mahol
1 sibling, 1 reply; 8+ messages in thread
From: Michael Niedermayer @ 2022-07-01 21:04 UTC (permalink / raw)
To: FFmpeg development discussions and patches
[-- Attachment #1.1: Type: text/plain, Size: 1876 bytes --]
On Thu, Jun 30, 2022 at 08:59:53AM +0200, Paul B Mahol wrote:
> Updated.
> libavfilter/Makefile | 2
> libavfilter/allfilters.c | 1
> libavfilter/opencl/remap.cl | 69 ++++++++
> libavfilter/opencl_source.h | 1
> libavfilter/vf_remap_opencl.c | 362 ++++++++++++++++++++++++++++++++++++++++++
> libavutil/hwcontext_opencl.c | 5
> 6 files changed, 439 insertions(+), 1 deletion(-)
> d6ec16b957d70b99edb4afe997ede6c683749ede 0001-avfilter-add-remap-opencl-filter.patch
> From 2b6b653bebec01d9d523102c1bfe2ce4b1be93dd Mon Sep 17 00:00:00 2001
> From: Paul B Mahol <onemda@gmail.com>
> Date: Wed, 29 Jun 2022 19:12:24 +0200
> Subject: [PATCH] avfilter: add remap opencl filter
>
> Signed-off-by: Paul B Mahol <onemda@gmail.com>
> ---
> libavfilter/Makefile | 2 +
> libavfilter/allfilters.c | 1 +
> libavfilter/opencl/remap.cl | 69 +++++++
> libavfilter/opencl_source.h | 1 +
> libavfilter/vf_remap_opencl.c | 362 ++++++++++++++++++++++++++++++++++
> libavutil/hwcontext_opencl.c | 5 +-
> 6 files changed, 439 insertions(+), 1 deletion(-)
> create mode 100644 libavfilter/opencl/remap.cl
> create mode 100644 libavfilter/vf_remap_opencl.c
breaks build here (ubuntu x86-64)
make
CC libavfilter/opencl.o
In file included from libavfilter/opencl.h:31:0,
from libavfilter/opencl.c:26:
./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory
#include <CL/cl.h>
^~~~~~~~~
compilation terminated.
ffbuild/common.mak:81: recipe for target 'libavfilter/opencl.o' failed
make: *** [libavfilter/opencl.o] Error 1
thx
[...]
--
Michael GnuPG fingerprint: 9FF2128B147EF6730BADF133611EC787040B0FAB
It is dangerous to be right in matters on which the established authorities
are wrong. -- Voltaire
[-- Attachment #1.2: signature.asc --]
[-- Type: application/pgp-signature, Size: 195 bytes --]
[-- Attachment #2: Type: text/plain, Size: 251 bytes --]
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-07-01 21:04 ` Michael Niedermayer
@ 2022-07-01 21:39 ` Paul B Mahol
2022-07-02 6:37 ` Paul B Mahol
0 siblings, 1 reply; 8+ messages in thread
From: Paul B Mahol @ 2022-07-01 21:39 UTC (permalink / raw)
To: FFmpeg development discussions and patches
On Fri, Jul 1, 2022 at 11:04 PM Michael Niedermayer <michael@niedermayer.cc>
wrote:
> On Thu, Jun 30, 2022 at 08:59:53AM +0200, Paul B Mahol wrote:
> > Updated.
>
> > libavfilter/Makefile | 2
> > libavfilter/allfilters.c | 1
> > libavfilter/opencl/remap.cl | 69 ++++++++
> > libavfilter/opencl_source.h | 1
> > libavfilter/vf_remap_opencl.c | 362
> ++++++++++++++++++++++++++++++++++++++++++
> > libavutil/hwcontext_opencl.c | 5
> > 6 files changed, 439 insertions(+), 1 deletion(-)
> > d6ec16b957d70b99edb4afe997ede6c683749ede
> 0001-avfilter-add-remap-opencl-filter.patch
> > From 2b6b653bebec01d9d523102c1bfe2ce4b1be93dd Mon Sep 17 00:00:00 2001
> > From: Paul B Mahol <onemda@gmail.com>
> > Date: Wed, 29 Jun 2022 19:12:24 +0200
> > Subject: [PATCH] avfilter: add remap opencl filter
> >
> > Signed-off-by: Paul B Mahol <onemda@gmail.com>
> > ---
> > libavfilter/Makefile | 2 +
> > libavfilter/allfilters.c | 1 +
> > libavfilter/opencl/remap.cl | 69 +++++++
> > libavfilter/opencl_source.h | 1 +
> > libavfilter/vf_remap_opencl.c | 362 ++++++++++++++++++++++++++++++++++
> > libavutil/hwcontext_opencl.c | 5 +-
> > 6 files changed, 439 insertions(+), 1 deletion(-)
> > create mode 100644 libavfilter/opencl/remap.cl
> > create mode 100644 libavfilter/vf_remap_opencl.c
>
> breaks build here (ubuntu x86-64)
>
> make
> CC libavfilter/opencl.o
> In file included from libavfilter/opencl.h:31:0,
> from libavfilter/opencl.c:26:
> ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file
> or directory
> #include <CL/cl.h>
> ^~~~~~~~~
> compilation terminated.
> ffbuild/common.mak:81: recipe for target 'libavfilter/opencl.o' failed
> make: *** [libavfilter/opencl.o] Error 1
>
> thx
>
Already spotted and fixed days ago.
>
> [...]
> --
> Michael GnuPG fingerprint: 9FF2128B147EF6730BADF133611EC787040B0FAB
>
> It is dangerous to be right in matters on which the established authorities
> are wrong. -- Voltaire
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-07-01 21:39 ` Paul B Mahol
@ 2022-07-02 6:37 ` Paul B Mahol
2022-07-07 13:18 ` Paul B Mahol
0 siblings, 1 reply; 8+ messages in thread
From: Paul B Mahol @ 2022-07-02 6:37 UTC (permalink / raw)
To: FFmpeg development discussions and patches
[-- Attachment #1: Type: text/plain, Size: 55 bytes --]
New version:
Added docs, and more cleanups and fixes.
[-- Attachment #2: 0001-avfilter-add-remap-opencl-filter.patch --]
[-- Type: text/x-patch, Size: 20343 bytes --]
From 22716c9f853d33769c1e63f7e010a5540a2e455b Mon Sep 17 00:00:00 2001
From: Paul B Mahol <onemda@gmail.com>
Date: Wed, 29 Jun 2022 19:12:24 +0200
Subject: [PATCH] avfilter: add remap opencl filter
Signed-off-by: Paul B Mahol <onemda@gmail.com>
---
configure | 1 +
doc/filters.texi | 24 +++
libavfilter/Makefile | 2 +
libavfilter/allfilters.c | 1 +
libavfilter/opencl/remap.cl | 73 +++++++
libavfilter/opencl_source.h | 1 +
libavfilter/vf_remap_opencl.c | 354 ++++++++++++++++++++++++++++++++++
libavutil/hwcontext_opencl.c | 5 +-
8 files changed, 460 insertions(+), 1 deletion(-)
create mode 100644 libavfilter/opencl/remap.cl
create mode 100644 libavfilter/vf_remap_opencl.c
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 e525e87b3c..7d383715da 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -25683,6 +25683,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/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);
}
--
2.36.1
[-- Attachment #3: Type: text/plain, Size: 251 bytes --]
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
* Re: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter
2022-07-02 6:37 ` Paul B Mahol
@ 2022-07-07 13:18 ` Paul B Mahol
0 siblings, 0 replies; 8+ messages in thread
From: Paul B Mahol @ 2022-07-07 13:18 UTC (permalink / raw)
To: FFmpeg development discussions and patches
Will apply soon.
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 8+ messages in thread
end of thread, other threads:[~2022-07-07 13:16 UTC | newest]
Thread overview: 8+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-06-29 19:36 [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter Paul B Mahol
2022-06-30 6:59 ` Paul B Mahol
2022-06-30 8:34 ` Anton Khirnov
2022-06-30 8:47 ` Paul B Mahol
2022-07-01 21:04 ` Michael Niedermayer
2022-07-01 21:39 ` Paul B Mahol
2022-07-02 6:37 ` Paul B Mahol
2022-07-07 13:18 ` Paul B Mahol
Git Inbox Mirror of the ffmpeg-devel mailing list - see https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
This inbox may be cloned and mirrored by anyone:
git clone --mirror https://master.gitmailbox.com/ffmpegdev/0 ffmpegdev/git/0.git
# If you have public-inbox 1.1+ installed, you may
# initialize and index your mirror using the following commands:
public-inbox-init -V2 ffmpegdev ffmpegdev/ https://master.gitmailbox.com/ffmpegdev \
ffmpegdev@gitmailbox.com
public-inbox-index ffmpegdev
Example config snippet for mirrors.
AGPL code for this site: git clone https://public-inbox.org/public-inbox.git