1 /*
2 * Copyright (c) 2022 Paul B Mahol
3 *
4 * This file is part of FFmpeg.
5 *
6 * FFmpeg is free software; you can redistribute it and/or
7 * modify it under the terms of the GNU Lesser General Public
8 * License as published by the Free Software Foundation; either
9 * version 2.1 of the License, or (at your option) any later version.
10 *
11 * FFmpeg is distributed in the hope that it will be useful,
12 * but WITHOUT ANY WARRANTY; without even the implied warranty of
13 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
14 * Lesser General Public License for more details.
15 *
16 * You should have received a copy of the GNU Lesser General Public
17 * License along with FFmpeg; if not, write to the Free Software
18 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19 */
20
21 #include "libavutil/colorspace.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/pixdesc.h"
24 #include "libavutil/opt.h"
25 #include "avfilter.h"
26 #include "drawutils.h"
27 #include "formats.h"
28 #include "framesync.h"
29 #include "internal.h"
30 #include "opencl.h"
31 #include "opencl_source.h"
32 #include "video.h"
33
34 typedef struct RemapOpenCLContext {
35 OpenCLFilterContext ocf;
36
37 int nb_planes;
38 int interp;
39 uint8_t fill_rgba[4];
40 cl_float4 cl_fill_color;
41
42 int initialised;
43 cl_kernel kernel;
44 cl_command_queue command_queue;
45
46 FFFrameSync fs;
47 } RemapOpenCLContext;
48
49 #define OFFSET(x) offsetof(RemapOpenCLContext, x)
50 #define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
51
52 static const AVOption remap_opencl_options[] = {
53 { "interp", "set interpolation method", OFFSET(interp), AV_OPT_TYPE_INT, {.i64=1}, 0, 1, FLAGS, "interp" },
54 { "near", NULL, 0, AV_OPT_TYPE_CONST, {.i64=0}, 0, 0, FLAGS, "interp" },
55 { "linear", NULL, 0, AV_OPT_TYPE_CONST, {.i64=1}, 0, 0, FLAGS, "interp" },
56 { "fill", "set the color of the unmapped pixels", OFFSET(fill_rgba), AV_OPT_TYPE_COLOR, {.str="black"}, .flags = FLAGS },
57 { NULL }
58 };
59
60 AVFILTER_DEFINE_CLASS(remap_opencl);
61
remap_opencl_init(AVFilterContext * avctx)62 static av_cold int remap_opencl_init(AVFilterContext *avctx)
63 {
64 return ff_opencl_filter_init(avctx);
65 }
66
67 static const char *kernels[] = { "remap_near", "remap_linear" };
68
remap_opencl_load(AVFilterContext * avctx,enum AVPixelFormat main_format,enum AVPixelFormat xmap_format,enum AVPixelFormat ymap_format)69 static int remap_opencl_load(AVFilterContext *avctx,
70 enum AVPixelFormat main_format,
71 enum AVPixelFormat xmap_format,
72 enum AVPixelFormat ymap_format)
73 {
74 RemapOpenCLContext *ctx = avctx->priv;
75 cl_int cle;
76 const char *source = ff_opencl_source_remap;
77 const char *kernel = kernels[ctx->interp];
78 const AVPixFmtDescriptor *main_desc;
79 int err, main_planes;
80 const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(main_format);
81 int is_rgb = !!(desc->flags & AV_PIX_FMT_FLAG_RGB);
82 const float scale = 1.f / 255.f;
83 uint8_t rgba_map[4];
84
85 ff_fill_rgba_map(rgba_map, main_format);
86
87 if (is_rgb) {
88 ctx->cl_fill_color.s[rgba_map[0]] = ctx->fill_rgba[0] * scale;
89 ctx->cl_fill_color.s[rgba_map[1]] = ctx->fill_rgba[1] * scale;
90 ctx->cl_fill_color.s[rgba_map[2]] = ctx->fill_rgba[2] * scale;
91 ctx->cl_fill_color.s[rgba_map[3]] = ctx->fill_rgba[3] * scale;
92 } else {
93 ctx->cl_fill_color.s[0] = RGB_TO_Y_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2]) * scale;
94 ctx->cl_fill_color.s[1] = RGB_TO_U_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
95 ctx->cl_fill_color.s[2] = RGB_TO_V_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
96 ctx->cl_fill_color.s[3] = ctx->fill_rgba[3] * scale;
97 }
98
99 main_desc = av_pix_fmt_desc_get(main_format);
100
101 main_planes = 0;
102 for (int i = 0; i < main_desc->nb_components; i++)
103 main_planes = FFMAX(main_planes,
104 main_desc->comp[i].plane + 1);
105
106 ctx->nb_planes = main_planes;
107
108 err = ff_opencl_filter_load_program(avctx, &source, 1);
109 if (err < 0)
110 goto fail;
111
112 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
113 ctx->ocf.hwctx->device_id,
114 0, &cle);
115 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
116 "command queue %d.\n", cle);
117
118 ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
119 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
120
121 ctx->initialised = 1;
122 return 0;
123
124 fail:
125 if (ctx->command_queue)
126 clReleaseCommandQueue(ctx->command_queue);
127 if (ctx->kernel)
128 clReleaseKernel(ctx->kernel);
129 return err;
130 }
131
remap_opencl_process_frame(FFFrameSync * fs)132 static int remap_opencl_process_frame(FFFrameSync *fs)
133 {
134 AVFilterContext *avctx = fs->parent;
135 AVFilterLink *outlink = avctx->outputs[0];
136 RemapOpenCLContext *ctx = avctx->priv;
137 AVFrame *input_main, *input_xmap, *input_ymap;
138 AVFrame *output;
139 cl_mem mem;
140 cl_int cle;
141 size_t global_work[2];
142 int kernel_arg = 0;
143 int err, plane;
144
145 err = ff_framesync_get_frame(fs, 0, &input_main, 0);
146 if (err < 0)
147 return err;
148 err = ff_framesync_get_frame(fs, 1, &input_xmap, 0);
149 if (err < 0)
150 return err;
151 err = ff_framesync_get_frame(fs, 2, &input_ymap, 0);
152 if (err < 0)
153 return err;
154
155 if (!ctx->initialised) {
156 AVHWFramesContext *main_fc =
157 (AVHWFramesContext*)input_main->hw_frames_ctx->data;
158 AVHWFramesContext *xmap_fc =
159 (AVHWFramesContext*)input_xmap->hw_frames_ctx->data;
160 AVHWFramesContext *ymap_fc =
161 (AVHWFramesContext*)input_ymap->hw_frames_ctx->data;
162
163 err = remap_opencl_load(avctx, main_fc->sw_format,
164 xmap_fc->sw_format,
165 ymap_fc->sw_format);
166 if (err < 0)
167 return err;
168 }
169
170 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
171 if (!output) {
172 err = AVERROR(ENOMEM);
173 goto fail;
174 }
175
176 for (plane = 0; plane < ctx->nb_planes; plane++) {
177 cl_float4 cl_fill_color;
178 kernel_arg = 0;
179
180 if (ctx->nb_planes == 1)
181 cl_fill_color = ctx->cl_fill_color;
182 else
183 cl_fill_color.s[0] = ctx->cl_fill_color.s[plane];
184
185 mem = (cl_mem)output->data[plane];
186 CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
187 kernel_arg++;
188
189 mem = (cl_mem)input_main->data[plane];
190 CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
191 kernel_arg++;
192
193 mem = (cl_mem)input_xmap->data[0];
194 CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
195 kernel_arg++;
196
197 mem = (cl_mem)input_ymap->data[0];
198 CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
199 kernel_arg++;
200
201 CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float4, &cl_fill_color);
202 kernel_arg++;
203
204 err = ff_opencl_filter_work_size_from_image(avctx, global_work,
205 output, plane, 0);
206 if (err < 0)
207 goto fail;
208
209 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
210 global_work, NULL, 0, NULL, NULL);
211 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue remap kernel "
212 "for plane %d: %d.\n", plane, cle);
213 }
214
215 cle = clFinish(ctx->command_queue);
216 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
217
218 err = av_frame_copy_props(output, input_main);
219
220 av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
221 av_get_pix_fmt_name(output->format),
222 output->width, output->height, output->pts);
223
224 return ff_filter_frame(outlink, output);
225
226 fail:
227 av_frame_free(&output);
228 return err;
229 }
230
config_output(AVFilterLink * outlink)231 static int config_output(AVFilterLink *outlink)
232 {
233 AVFilterContext *ctx = outlink->src;
234 RemapOpenCLContext *s = ctx->priv;
235 AVFilterLink *srclink = ctx->inputs[0];
236 AVFilterLink *xlink = ctx->inputs[1];
237 AVFilterLink *ylink = ctx->inputs[2];
238 FFFrameSyncIn *in;
239 int ret;
240
241 if (xlink->w != ylink->w || xlink->h != ylink->h) {
242 av_log(ctx, AV_LOG_ERROR, "Second input link %s parameters "
243 "(size %dx%d) do not match the corresponding "
244 "third input link %s parameters (%dx%d)\n",
245 ctx->input_pads[1].name, xlink->w, xlink->h,
246 ctx->input_pads[2].name, ylink->w, ylink->h);
247 return AVERROR(EINVAL);
248 }
249
250 outlink->w = xlink->w;
251 outlink->h = xlink->h;
252 outlink->sample_aspect_ratio = srclink->sample_aspect_ratio;
253 outlink->frame_rate = srclink->frame_rate;
254
255 ret = ff_framesync_init(&s->fs, ctx, 3);
256 if (ret < 0)
257 return ret;
258
259 in = s->fs.in;
260 in[0].time_base = srclink->time_base;
261 in[1].time_base = xlink->time_base;
262 in[2].time_base = ylink->time_base;
263 in[0].sync = 2;
264 in[0].before = EXT_STOP;
265 in[0].after = EXT_STOP;
266 in[1].sync = 1;
267 in[1].before = EXT_NULL;
268 in[1].after = EXT_INFINITY;
269 in[2].sync = 1;
270 in[2].before = EXT_NULL;
271 in[2].after = EXT_INFINITY;
272 s->fs.opaque = s;
273 s->fs.on_event = remap_opencl_process_frame;
274
275 ret = ff_framesync_configure(&s->fs);
276 outlink->time_base = s->fs.time_base;
277 if (ret < 0)
278 return ret;
279
280 s->ocf.output_width = outlink->w;
281 s->ocf.output_height = outlink->h;
282
283 return ff_opencl_filter_config_output(outlink);
284 }
285
activate(AVFilterContext * ctx)286 static int activate(AVFilterContext *ctx)
287 {
288 RemapOpenCLContext *s = ctx->priv;
289 return ff_framesync_activate(&s->fs);
290 }
291
remap_opencl_uninit(AVFilterContext * avctx)292 static av_cold void remap_opencl_uninit(AVFilterContext *avctx)
293 {
294 RemapOpenCLContext *ctx = avctx->priv;
295 cl_int cle;
296
297 if (ctx->kernel) {
298 cle = clReleaseKernel(ctx->kernel);
299 if (cle != CL_SUCCESS)
300 av_log(avctx, AV_LOG_ERROR, "Failed to release "
301 "kernel: %d.\n", cle);
302 }
303
304 if (ctx->command_queue) {
305 cle = clReleaseCommandQueue(ctx->command_queue);
306 if (cle != CL_SUCCESS)
307 av_log(avctx, AV_LOG_ERROR, "Failed to release "
308 "command queue: %d.\n", cle);
309 }
310
311 ff_opencl_filter_uninit(avctx);
312
313 ff_framesync_uninit(&ctx->fs);
314 }
315
316 static const AVFilterPad remap_opencl_inputs[] = {
317 {
318 .name = "source",
319 .type = AVMEDIA_TYPE_VIDEO,
320 .config_props = &ff_opencl_filter_config_input,
321 },
322 {
323 .name = "xmap",
324 .type = AVMEDIA_TYPE_VIDEO,
325 .config_props = &ff_opencl_filter_config_input,
326 },
327 {
328 .name = "ymap",
329 .type = AVMEDIA_TYPE_VIDEO,
330 .config_props = &ff_opencl_filter_config_input,
331 },
332 };
333
334 static const AVFilterPad remap_opencl_outputs[] = {
335 {
336 .name = "default",
337 .type = AVMEDIA_TYPE_VIDEO,
338 .config_props = config_output,
339 },
340 };
341
342 const AVFilter ff_vf_remap_opencl = {
343 .name = "remap_opencl",
344 .description = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."),
345 .priv_size = sizeof(RemapOpenCLContext),
346 .init = remap_opencl_init,
347 .uninit = remap_opencl_uninit,
348 .activate = activate,
349 FILTER_INPUTS(remap_opencl_inputs),
350 FILTER_OUTPUTS(remap_opencl_outputs),
351 FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
352 .priv_class = &remap_opencl_class,
353 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
354 };
355