FFmpeg  4.4.4
vf_transpose_opencl.c
Go to the documentation of this file.
1 /*
2  * This file is part of FFmpeg.
3  *
4  * FFmpeg is free software; you can redistribute it and/or
5  * modify it under the terms of the GNU Lesser General Public
6  * License as published by the Free Software Foundation; either
7  * version 2.1 of the License, or (at your option) any later version.
8  *
9  * FFmpeg is distributed in the hope that it will be useful,
10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12  * Lesser General Public License for more details.
13  *
14  * You should have received a copy of the GNU Lesser General Public
15  * License along with FFmpeg; if not, write to the Free Software
16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
17  */
18 #include <float.h>
19 
20 #include "libavutil/avassert.h"
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/mem.h"
24 #include "libavutil/opt.h"
25 #include "libavutil/pixdesc.h"
26 
27 #include "avfilter.h"
28 #include "internal.h"
29 #include "opencl.h"
30 #include "opencl_source.h"
31 #include "video.h"
32 #include "transpose.h"
33 
34 typedef struct TransposeOpenCLContext {
37  int passthrough; ///< PassthroughType, landscape passthrough mode enabled
38  int dir; ///< TransposeDir
39  cl_kernel kernel;
40  cl_command_queue command_queue;
42 
44 {
46  cl_int cle;
47  int err;
48 
50  if (err < 0)
51  goto fail;
52 
53  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
54  ctx->ocf.hwctx->device_id,
55  0, &cle);
56  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
57  "command queue %d.\n", cle);
58 
59  ctx->kernel = clCreateKernel(ctx->ocf.program, "transpose", &cle);
60  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
61 
62 
63  ctx->initialised = 1;
64  return 0;
65 
66 fail:
67  if (ctx->command_queue)
68  clReleaseCommandQueue(ctx->command_queue);
69  if (ctx->kernel)
70  clReleaseKernel(ctx->kernel);
71  return err;
72 }
73 
75 {
76  AVFilterContext *avctx = outlink->src;
77  TransposeOpenCLContext *s = avctx->priv;
78  AVFilterLink *inlink = avctx->inputs[0];
79  const AVPixFmtDescriptor *desc_in = av_pix_fmt_desc_get(inlink->format);
80  int ret;
81 
82  if ((inlink->w >= inlink->h &&
83  s->passthrough == TRANSPOSE_PT_TYPE_LANDSCAPE) ||
84  (inlink->w <= inlink->h &&
85  s->passthrough == TRANSPOSE_PT_TYPE_PORTRAIT)) {
86  if (inlink->hw_frames_ctx) {
87  outlink->hw_frames_ctx = av_buffer_ref(inlink->hw_frames_ctx);
88  if (!outlink->hw_frames_ctx)
89  return AVERROR(ENOMEM);
90  }
91  av_log(avctx, AV_LOG_VERBOSE,
92  "w:%d h:%d -> w:%d h:%d (passthrough mode)\n",
93  inlink->w, inlink->h, inlink->w, inlink->h);
94 
95  return 0;
96  } else {
97  s->passthrough = TRANSPOSE_PT_TYPE_NONE;
98  }
99 
100  if (desc_in->log2_chroma_w != desc_in->log2_chroma_h) {
101  av_log(avctx, AV_LOG_ERROR, "Input format %s not supported.\n",
102  desc_in->name);
103  return AVERROR(EINVAL);
104  }
105 
106  s->ocf.output_width = inlink->h;
107  s->ocf.output_height = inlink->w;
108  ret = ff_opencl_filter_config_output(outlink);
109  if (ret < 0)
110  return ret;
111 
112  if (inlink->sample_aspect_ratio.num)
113  outlink->sample_aspect_ratio = av_div_q((AVRational) { 1, 1 },
114  inlink->sample_aspect_ratio);
115  else
116  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
117 
118  av_log(avctx, AV_LOG_VERBOSE,
119  "w:%d h:%d dir:%d -> w:%d h:%d rotation:%s vflip:%d\n",
120  inlink->w, inlink->h, s->dir, outlink->w, outlink->h,
121  s->dir == 1 || s->dir == 3 ? "clockwise" : "counterclockwise",
122  s->dir == 0 || s->dir == 3);
123  return 0;
124 }
125 
126 static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
127 {
128  TransposeOpenCLContext *s = inlink->dst->priv;
129 
130  return s->passthrough ?
131  ff_null_get_video_buffer (inlink, w, h) :
132  ff_default_get_video_buffer(inlink, w, h);
133 }
134 
136 {
137  AVFilterContext *avctx = inlink->dst;
138  AVFilterLink *outlink = avctx->outputs[0];
139  TransposeOpenCLContext *ctx = avctx->priv;
140  AVFrame *output = NULL;
141  size_t global_work[2];
142  cl_mem src, dst;
143  cl_int cle;
144  int err, p;
145 
146  av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
147  av_get_pix_fmt_name(input->format),
148  input->width, input->height, input->pts);
149 
150  if (!input->hw_frames_ctx)
151  return AVERROR(EINVAL);
152 
153  if (ctx->passthrough)
154  return ff_filter_frame(outlink, input);
155 
156  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
157  if (!output) {
158  err = AVERROR(ENOMEM);
159  goto fail;
160  }
161 
162  err = av_frame_copy_props(output, input);
163  if (err < 0)
164  goto fail;
165 
166  if (input->sample_aspect_ratio.num == 0) {
167  output->sample_aspect_ratio = input->sample_aspect_ratio;
168  } else {
171  }
172 
173  if (!ctx->initialised) {
174  err = transpose_opencl_init(avctx);
175  if (err < 0)
176  goto fail;
177  }
178 
179  for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
180  src = (cl_mem) input->data[p];
181  dst = (cl_mem) output->data[p];
182 
183  if (!dst)
184  break;
185  CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst);
186  CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src);
187  CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_int, &ctx->dir);
188 
189  err = ff_opencl_filter_work_size_from_image(avctx, global_work, output,
190  p, 16);
191 
192  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
193  global_work, NULL,
194  0, NULL, NULL);
195  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
196  }
197  cle = clFinish(ctx->command_queue);
198  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
199 
200  av_frame_free(&input);
201 
202  av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
203  av_get_pix_fmt_name(output->format),
204  output->width, output->height, output->pts);
205 
206  return ff_filter_frame(outlink, output);
207 
208 fail:
209  clFinish(ctx->command_queue);
210  av_frame_free(&input);
211  av_frame_free(&output);
212  return err;
213 }
214 
216 {
217  TransposeOpenCLContext *ctx = avctx->priv;
218  cl_int cle;
219 
220  if (ctx->kernel) {
221  cle = clReleaseKernel(ctx->kernel);
222  if (cle != CL_SUCCESS)
223  av_log(avctx, AV_LOG_ERROR, "Failed to release "
224  "kernel: %d.\n", cle);
225  }
226 
227  if (ctx->command_queue) {
228  cle = clReleaseCommandQueue(ctx->command_queue);
229  if (cle != CL_SUCCESS)
230  av_log(avctx, AV_LOG_ERROR, "Failed to release "
231  "command queue: %d.\n", cle);
232  }
233 
235 }
236 
237 #define OFFSET(x) offsetof(TransposeOpenCLContext, x)
238 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
240  { "dir", "set transpose direction", OFFSET(dir), AV_OPT_TYPE_INT, { .i64 = TRANSPOSE_CCLOCK_FLIP }, 0, 3, FLAGS, "dir" },
241  { "cclock_flip", "rotate counter-clockwise with vertical flip", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CCLOCK_FLIP }, .flags=FLAGS, .unit = "dir" },
242  { "clock", "rotate clockwise", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CLOCK }, .flags=FLAGS, .unit = "dir" },
243  { "cclock", "rotate counter-clockwise", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CCLOCK }, .flags=FLAGS, .unit = "dir" },
244  { "clock_flip", "rotate clockwise with vertical flip", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CLOCK_FLIP }, .flags=FLAGS, .unit = "dir" },
245 
246  { "passthrough", "do not apply transposition if the input matches the specified geometry",
247  OFFSET(passthrough), AV_OPT_TYPE_INT, {.i64=TRANSPOSE_PT_TYPE_NONE}, 0, INT_MAX, FLAGS, "passthrough" },
248  { "none", "always apply transposition", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_NONE}, INT_MIN, INT_MAX, FLAGS, "passthrough" },
249  { "portrait", "preserve portrait geometry", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_PORTRAIT}, INT_MIN, INT_MAX, FLAGS, "passthrough" },
250  { "landscape", "preserve landscape geometry", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_LANDSCAPE}, INT_MIN, INT_MAX, FLAGS, "passthrough" },
251 
252  { NULL }
253 };
254 
255 AVFILTER_DEFINE_CLASS(transpose_opencl);
256 
258  {
259  .name = "default",
260  .type = AVMEDIA_TYPE_VIDEO,
261  .get_video_buffer = get_video_buffer,
262  .filter_frame = &transpose_opencl_filter_frame,
263  .config_props = &ff_opencl_filter_config_input,
264  },
265  { NULL }
266 };
267 
269  {
270  .name = "default",
271  .type = AVMEDIA_TYPE_VIDEO,
272  .config_props = &transpose_opencl_config_output,
273  },
274  { NULL }
275 };
276 
278  .name = "transpose_opencl",
279  .description = NULL_IF_CONFIG_SMALL("Transpose input video"),
280  .priv_size = sizeof(TransposeOpenCLContext),
281  .priv_class = &transpose_opencl_class,
287  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
288 };
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:243
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define av_cold
Definition: attributes.h:88
simple assert() macros that are a bit more flexible than ISO C assert().
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1096
Main libavfilter public API header.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:31
#define s(width, name)
Definition: cbs_vp9.c:257
#define fail()
Definition: checkasm.h:133
common internal and external API header
#define NULL
Definition: coverity.c:32
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
@ AV_OPT_TYPE_INT
Definition: opt.h:225
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
#define AVERROR(e)
Definition: error.h:43
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:203
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:658
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:215
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:210
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:194
AVRational av_div_q(AVRational b, AVRational c)
Divide one rational by another.
Definition: rational.c:88
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
misc image utilities
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: internal.h:339
common internal API header
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:117
uint8_t w
Definition: llviddspenc.c:39
Memory handling functions.
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:156
int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings)
Load a new OpenCL program from strings in memory.
Definition: opencl.c:171
int ff_opencl_filter_config_input(AVFilterLink *inlink)
Check that the input link contains a suitable hardware frames context and extract the device from it.
Definition: opencl.c:60
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:147
int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment)
Find the work size needed needed for a given plane of an image.
Definition: opencl.c:278
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:96
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:28
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
AVOptions.
const char * av_get_pix_fmt_name(enum AVPixelFormat pix_fmt)
Return the short name for a pixel format, NULL in case pix_fmt is unknown.
Definition: pixdesc.c:2489
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2573
#define FF_ARRAY_ELEMS(a)
An instance of a filter.
Definition: avfilter.h:341
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:349
void * priv
private data for use by the filter
Definition: avfilter.h:356
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:353
A filter pad used for either input or output.
Definition: internal.h:54
const char * name
Pad name.
Definition: internal.h:60
Filter definition.
Definition: avfilter.h:145
const char * name
Filter name.
Definition: avfilter.h:149
This structure describes decoded (raw) audio or video data.
Definition: frame.h:318
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:411
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:332
int width
Definition: frame.h:376
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
Definition: frame.h:657
int height
Definition: frame.h:376
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
Definition: frame.h:406
int format
format of the frame, -1 if unknown or unset Values correspond to enum AVPixelFormat for video frames,...
Definition: frame.h:391
AVOption.
Definition: opt.h:248
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:81
const char * name
Definition: pixdesc.h:82
uint8_t log2_chroma_w
Amount to shift the luma width right to find the chroma width.
Definition: pixdesc.h:92
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:101
Rational number (pair of numerator and denominator).
Definition: rational.h:58
int num
Numerator.
Definition: rational.h:59
int den
Denominator.
Definition: rational.h:60
OpenCLFilterContext ocf
cl_command_queue command_queue
int passthrough
PassthroughType, landscape passthrough mode enabled.
#define av_log(a,...)
#define src
Definition: vp8dsp.c:255
AVFormatContext * ctx
Definition: movenc.c:48
const char * ff_opencl_source_transpose
Definition: transpose.c:2
@ TRANSPOSE_CLOCK_FLIP
Definition: transpose.h:34
@ TRANSPOSE_CCLOCK
Definition: transpose.h:33
@ TRANSPOSE_CCLOCK_FLIP
Definition: transpose.h:31
@ TRANSPOSE_CLOCK
Definition: transpose.h:32
@ TRANSPOSE_PT_TYPE_NONE
Definition: transpose.h:25
@ TRANSPOSE_PT_TYPE_PORTRAIT
Definition: transpose.h:27
@ TRANSPOSE_PT_TYPE_LANDSCAPE
Definition: transpose.h:26
static const AVFilterPad transpose_opencl_inputs[]
static int transpose_opencl_init(AVFilterContext *avctx)
static int transpose_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
static const AVFilterPad transpose_opencl_outputs[]
static int transpose_opencl_config_output(AVFilterLink *outlink)
#define FLAGS
static const AVOption transpose_opencl_options[]
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
AVFilter ff_vf_transpose_opencl
#define OFFSET(x)
AVFILTER_DEFINE_CLASS(transpose_opencl)
static av_cold void transpose_opencl_uninit(AVFilterContext *avctx)
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:39
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:104
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:99