FFmpeg
vf_xfade_vulkan.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 
19 #include "libavutil/avassert.h"
20 #include "libavutil/random_seed.h"
21 #include "libavutil/opt.h"
22 #include "libavutil/vulkan_spirv.h"
23 #include "vulkan_filter.h"
24 #include "filters.h"
25 #include "video.h"
26 
27 #define IN_A 0
28 #define IN_B 1
29 #define IN_NB 2
30 
31 typedef struct XFadeParameters {
32  float progress;
33 } XFadeParameters;
34 
35 typedef struct XFadeVulkanContext {
36  FFVulkanContext vkctx;
37 
38  int transition;
39  int64_t duration;
40  int64_t offset;
41 
42  int initialized;
43  FFVkExecPool e;
44  FFVkQueueFamilyCtx qf;
45  FFVulkanShader shd;
46  VkSampler sampler;
47 
48  // PTS when the fade should start (in IN_A timebase)
49  int64_t start_pts;
50 
51  // PTS offset between IN_A and IN_B
52  int64_t inputs_offset_pts;
53 
54  // Duration of the transition
55  int64_t duration_pts;
56 
57  // Current PTS of the first input (IN_A)
58  int64_t pts;
59 
60  // If frames are currently just passed through
61  // unmodified, like before and after the actual
62  // transition.
63  int passthrough;
64 
65  int status[IN_NB];
66 } XFadeVulkanContext;
67 
68 enum XFadeTransitions {
69  FADE,
70  WIPELEFT,
71  WIPERIGHT,
72  WIPEUP,
73  WIPEDOWN,
74  SLIDEDOWN,
75  SLIDEUP,
76  SLIDELEFT,
77  SLIDERIGHT,
78  CIRCLEOPEN,
79  CIRCLECLOSE,
80  DISSOLVE,
81  PIXELIZE,
82  WIPETL,
83  WIPETR,
84  WIPEBL,
85  WIPEBR,
86  NB_TRANSITIONS,
87 };
88 
89 static const char transition_fade[] = {
90  C(0, void transition(int idx, ivec2 pos, float progress) )
91  C(0, { )
92  C(1, vec4 a = texture(a_images[idx], pos); )
93  C(1, vec4 b = texture(b_images[idx], pos); )
94  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
95  C(0, } )
96 };
97 
98 static const char transition_wipeleft[] = {
99  C(0, void transition(int idx, ivec2 pos, float progress) )
100  C(0, { )
101  C(1, ivec2 size = imageSize(output_images[idx]); )
102  C(1, int s = int(size.x * (1.0 - progress)); )
103  C(1, vec4 a = texture(a_images[idx], pos); )
104  C(1, vec4 b = texture(b_images[idx], pos); )
105  C(1, imageStore(output_images[idx], pos, pos.x > s ? b : a); )
106  C(0, } )
107 };
108 
109 static const char transition_wiperight[] = {
110  C(0, void transition(int idx, ivec2 pos, float progress) )
111  C(0, { )
112  C(1, ivec2 size = imageSize(output_images[idx]); )
113  C(1, int s = int(size.x * progress); )
114  C(1, vec4 a = texture(a_images[idx], pos); )
115  C(1, vec4 b = texture(b_images[idx], pos); )
116  C(1, imageStore(output_images[idx], pos, pos.x > s ? a : b); )
117  C(0, } )
118 };
119 
120 static const char transition_wipeup[] = {
121  C(0, void transition(int idx, ivec2 pos, float progress) )
122  C(0, { )
123  C(1, ivec2 size = imageSize(output_images[idx]); )
124  C(1, int s = int(size.y * (1.0 - progress)); )
125  C(1, vec4 a = texture(a_images[idx], pos); )
126  C(1, vec4 b = texture(b_images[idx], pos); )
127  C(1, imageStore(output_images[idx], pos, pos.y > s ? b : a); )
128  C(0, } )
129 };
130 
131 static const char transition_wipedown[] = {
132  C(0, void transition(int idx, ivec2 pos, float progress) )
133  C(0, { )
134  C(1, ivec2 size = imageSize(output_images[idx]); )
135  C(1, int s = int(size.y * progress); )
136  C(1, vec4 a = texture(a_images[idx], pos); )
137  C(1, vec4 b = texture(b_images[idx], pos); )
138  C(1, imageStore(output_images[idx], pos, pos.y > s ? a : b); )
139  C(0, } )
140 };
141 
142 #define SHADER_SLIDE_COMMON \
143  C(0, void slide(int idx, ivec2 pos, float progress, ivec2 direction) ) \
144  C(0, { ) \
145  C(1, ivec2 size = imageSize(output_images[idx]); ) \
146  C(1, ivec2 pi = ivec2(progress * size); ) \
147  C(1, ivec2 p = pos + pi * direction; ) \
148  C(1, ivec2 f = p % size; ) \
149  C(1, f = f + size * ivec2(f.x < 0, f.y < 0); ) \
150  C(1, vec4 a = texture(a_images[idx], f); ) \
151  C(1, vec4 b = texture(b_images[idx], f); ) \
152  C(1, vec4 r = (p.y >= 0 && p.x >= 0 && size.y > p.y && size.x > p.x) ? a : b; ) \
153  C(1, imageStore(output_images[idx], pos, r); ) \
154  C(0, } )
155 
156 static const char transition_slidedown[] = {
157  SHADER_SLIDE_COMMON
158  C(0, void transition(int idx, ivec2 pos, float progress) )
159  C(0, { )
160  C(1, slide(idx, pos, progress, ivec2(0, -1)); )
161  C(0, } )
162 };
163 
164 static const char transition_slideup[] = {
165  SHADER_SLIDE_COMMON
166  C(0, void transition(int idx, ivec2 pos, float progress) )
167  C(0, { )
168  C(1, slide(idx, pos, progress, ivec2(0, +1)); )
169  C(0, } )
170 };
171 
172 static const char transition_slideleft[] = {
173  SHADER_SLIDE_COMMON
174  C(0, void transition(int idx, ivec2 pos, float progress) )
175  C(0, { )
176  C(1, slide(idx, pos, progress, ivec2(+1, 0)); )
177  C(0, } )
178 };
179 
180 static const char transition_slideright[] = {
181  SHADER_SLIDE_COMMON
182  C(0, void transition(int idx, ivec2 pos, float progress) )
183  C(0, { )
184  C(1, slide(idx, pos, progress, ivec2(-1, 0)); )
185  C(0, } )
186 };
187 
188 #define SHADER_CIRCLE_COMMON \
189  C(0, void circle(int idx, ivec2 pos, float progress, bool open) ) \
190  C(0, { ) \
191  C(1, const ivec2 half_size = imageSize(output_images[idx]) / 2; ) \
192  C(1, const float z = dot(half_size, half_size); ) \
193  C(1, float p = ((open ? (1.0 - progress) : progress) - 0.5) * 3.0; ) \
194  C(1, ivec2 dsize = pos - half_size; ) \
195  C(1, float sm = dot(dsize, dsize) / z + p; ) \
196  C(1, vec4 a = texture(a_images[idx], pos); ) \
197  C(1, vec4 b = texture(b_images[idx], pos); ) \
198  C(1, imageStore(output_images[idx], pos, \
199  mix(open ? b : a, open ? a : b, \
200  smoothstep(0.f, 1.f, sm))); ) \
201  C(0, } )
202 
203 static const char transition_circleopen[] = {
204  SHADER_CIRCLE_COMMON
205  C(0, void transition(int idx, ivec2 pos, float progress) )
206  C(0, { )
207  C(1, circle(idx, pos, progress, true); )
208  C(0, } )
209 };
210 
211 static const char transition_circleclose[] = {
212  SHADER_CIRCLE_COMMON
213  C(0, void transition(int idx, ivec2 pos, float progress) )
214  C(0, { )
215  C(1, circle(idx, pos, progress, false); )
216  C(0, } )
217 };
218 
219 #define SHADER_FRAND_FUNC \
220  C(0, float frand(vec2 v) ) \
221  C(0, { ) \
222  C(1, return fract(sin(dot(v, vec2(12.9898, 78.233))) * 43758.545); ) \
223  C(0, } )
224 
225 static const char transition_dissolve[] = {
226  SHADER_FRAND_FUNC
227  C(0, void transition(int idx, ivec2 pos, float progress) )
228  C(0, { )
229  C(1, float sm = frand(pos) * 2.0 + (1.0 - progress) * 2.0 - 1.5; )
230  C(1, vec4 a = texture(a_images[idx], pos); )
231  C(1, vec4 b = texture(b_images[idx], pos); )
232  C(1, imageStore(output_images[idx], pos, sm >= 0.5 ? a : b); )
233  C(0, } )
234 };
235 
236 static const char transition_pixelize[] = {
237  C(0, void transition(int idx, ivec2 pos, float progress) )
238  C(0, { )
239  C(1, ivec2 size = imageSize(output_images[idx]); )
240  C(1, float d = min(progress, 1.0 - progress); )
241  C(1, float dist = ceil(d * 50.0) / 50.0; )
242  C(1, float sq = 2.0 * dist * min(size.x, size.y) / 20.0; )
243  C(1, float sx = dist > 0.0 ? min((floor(pos.x / sq) + 0.5) * sq, size.x - 1) : pos.x; )
244  C(1, float sy = dist > 0.0 ? min((floor(pos.y / sq) + 0.5) * sq, size.y - 1) : pos.y; )
245  C(1, vec4 a = texture(a_images[idx], vec2(sx, sy)); )
246  C(1, vec4 b = texture(b_images[idx], vec2(sx, sy)); )
247  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
248  C(0, } )
249 };
250 
251 static const char transition_wipetl[] = {
252  C(0, void transition(int idx, ivec2 pos, float progress) )
253  C(0, { )
254  C(1, ivec2 size = imageSize(output_images[idx]); )
255  C(1, float zw = size.x * (1.0 - progress); )
256  C(1, float zh = size.y * (1.0 - progress); )
257  C(1, vec4 a = texture(a_images[idx], pos); )
258  C(1, vec4 b = texture(b_images[idx], pos); )
259  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x <= zw) ? a : b); )
260  C(0, } )
261 };
262 
263 static const char transition_wipetr[] = {
264  C(0, void transition(int idx, ivec2 pos, float progress) )
265  C(0, { )
266  C(1, ivec2 size = imageSize(output_images[idx]); )
267  C(1, float zw = size.x * (progress); )
268  C(1, float zh = size.y * (1.0 - progress); )
269  C(1, vec4 a = texture(a_images[idx], pos); )
270  C(1, vec4 b = texture(b_images[idx], pos); )
271  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x > zw) ? a : b); )
272  C(0, } )
273 };
274 
275 static const char transition_wipebl[] = {
276  C(0, void transition(int idx, ivec2 pos, float progress) )
277  C(0, { )
278  C(1, ivec2 size = imageSize(output_images[idx]); )
279  C(1, float zw = size.x * (1.0 - progress); )
280  C(1, float zh = size.y * (progress); )
281  C(1, vec4 a = texture(a_images[idx], pos); )
282  C(1, vec4 b = texture(b_images[idx], pos); )
283  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x <= zw) ? a : b); )
284  C(0, } )
285 };
286 
287 static const char transition_wipebr[] = {
288  C(0, void transition(int idx, ivec2 pos, float progress) )
289  C(0, { )
290  C(1, ivec2 size = imageSize(output_images[idx]); )
291  C(1, float zw = size.x * (progress); )
292  C(1, float zh = size.y * (progress); )
293  C(1, vec4 a = texture(a_images[idx], pos); )
294  C(1, vec4 b = texture(b_images[idx], pos); )
295  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x > zw) ? a : b); )
296  C(0, } )
297 };
298 
299 static const char* transitions_map[NB_TRANSITIONS] = {
300  [FADE] = transition_fade,
301  [WIPELEFT] = transition_wipeleft,
302  [WIPERIGHT] = transition_wiperight,
303  [WIPEUP] = transition_wipeup,
304  [WIPEDOWN] = transition_wipedown,
305  [SLIDEDOWN] = transition_slidedown,
306  [SLIDEUP] = transition_slideup,
307  [SLIDELEFT] = transition_slideleft,
308  [SLIDERIGHT] = transition_slideright,
309  [CIRCLEOPEN] = transition_circleopen,
310  [CIRCLECLOSE] = transition_circleclose,
311  [DISSOLVE] = transition_dissolve,
312  [PIXELIZE] = transition_pixelize,
313  [WIPETL] = transition_wipetl,
314  [WIPETR] = transition_wipetr,
315  [WIPEBL] = transition_wipebl,
316  [WIPEBR] = transition_wipebr,
317 };
318 
319 static av_cold int init_vulkan(AVFilterContext *avctx)
320 {
321  int err = 0;
322  uint8_t *spv_data;
323  size_t spv_len;
324  void *spv_opaque = NULL;
325  XFadeVulkanContext *s = avctx->priv;
326  FFVulkanContext *vkctx = &s->vkctx;
327  const int planes = av_pix_fmt_count_planes(s->vkctx.output_format);
328  FFVulkanShader *shd = &s->shd;
329  FFVkSPIRVCompiler *spv;
330  FFVulkanDescriptorSetBinding *desc;
331 
332  spv = ff_vk_spirv_init();
333  if (!spv) {
334  av_log(avctx, AV_LOG_ERROR, "Unable to initialize SPIR-V compiler!\n");
335  return AVERROR_EXTERNAL;
336  }
337 
338  ff_vk_qf_init(vkctx, &s->qf, VK_QUEUE_COMPUTE_BIT);
339  RET(ff_vk_exec_pool_init(vkctx, &s->qf, &s->e, s->qf.nb_queues*4, 0, 0, 0, NULL));
340  RET(ff_vk_init_sampler(vkctx, &s->sampler, 1, VK_FILTER_NEAREST));
341  RET(ff_vk_shader_init(vkctx, &s->shd, "xfade",
342  VK_SHADER_STAGE_COMPUTE_BIT,
343  NULL, 0,
344  32, 32, 1,
345  0));
346 
347  desc = (FFVulkanDescriptorSetBinding []) {
348  {
349  .name = "a_images",
350  .type = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER,
351  .dimensions = 2,
352  .elems = planes,
353  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
354  .samplers = DUP_SAMPLER(s->sampler),
355  },
356  {
357  .name = "b_images",
358  .type = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER,
359  .dimensions = 2,
360  .elems = planes,
361  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
362  .samplers = DUP_SAMPLER(s->sampler),
363  },
364  {
365  .name = "output_images",
366  .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
367  .mem_layout = ff_vk_shader_rep_fmt(s->vkctx.output_format, FF_VK_REP_FLOAT),
368  .mem_quali = "writeonly",
369  .dimensions = 2,
370  .elems = planes,
371  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
372  },
373  };
374 
375  RET(ff_vk_shader_add_descriptor_set(vkctx, &s->shd, desc, 3, 0, 0));
376 
377  GLSLC(0, layout(push_constant, std430) uniform pushConstants { );
378  GLSLC(1, float progress; );
379  GLSLC(0, }; );
380 
381  ff_vk_shader_add_push_const(&s->shd, 0, sizeof(XFadeParameters),
382  VK_SHADER_STAGE_COMPUTE_BIT);
383 
384  // Add the right transition type function to the shader
385  GLSLD(transitions_map[s->transition]);
386 
387  GLSLC(0, void main() );
388  GLSLC(0, { );
389  GLSLC(1, ivec2 pos = ivec2(gl_GlobalInvocationID.xy); );
390  GLSLF(1, int planes = %i; ,planes);
391  GLSLC(1, for (int i = 0; i < planes; i++) { );
392  GLSLC(2, transition(i, pos, progress); );
393  GLSLC(1, } );
394  GLSLC(0, } );
395 
396  RET(spv->compile_shader(vkctx, spv, shd, &spv_data, &spv_len, "main",
397  &spv_opaque));
398  RET(ff_vk_shader_link(vkctx, shd, spv_data, spv_len, "main"));
399 
400  RET(ff_vk_shader_register_exec(vkctx, &s->e, &s->shd));
401 
402  s->initialized = 1;
403 
404 fail:
405  if (spv_opaque)
406  spv->free_shader(spv, &spv_opaque);
407  if (spv)
408  spv->uninit(&spv);
409 
410  return err;
411 }
412 
413 static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
414 {
415  int err;
416  AVFilterLink *outlink = avctx->outputs[0];
417  XFadeVulkanContext *s = avctx->priv;
418  float progress;
419 
420  AVFrame *output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
421  if (!output) {
422  err = AVERROR(ENOMEM);
423  goto fail;
424  }
425 
426  if (!s->initialized) {
427  AVHWFramesContext *a_fc = (AVHWFramesContext*)frame_a->hw_frames_ctx->data;
428  AVHWFramesContext *b_fc = (AVHWFramesContext*)frame_b->hw_frames_ctx->data;
429  if (a_fc->sw_format != b_fc->sw_format) {
430  av_log(avctx, AV_LOG_ERROR,
431  "Currently the sw format of the first input needs to match the second!\n");
432  return AVERROR(EINVAL);
433  }
434  RET(init_vulkan(avctx));
435  }
436 
437  RET(av_frame_copy_props(output, frame_a));
438  output->pts = s->pts;
439 
440  progress = av_clipf((float)(s->pts - s->start_pts) / s->duration_pts,
441  0.f, 1.f);
442 
443  RET(ff_vk_filter_process_Nin(&s->vkctx, &s->e, &s->shd, output,
444  (AVFrame *[]){ frame_a, frame_b }, 2, s->sampler,
445  &(XFadeParameters){ progress }, sizeof(XFadeParameters)));
446 
447  return ff_filter_frame(outlink, output);
448 
449 fail:
450  av_frame_free(&output);
451  return err;
452 }
453 
454 static int config_props_output(AVFilterLink *outlink)
455 {
456  int err;
457  AVFilterContext *avctx = outlink->src;
458  XFadeVulkanContext *s = avctx->priv;
459  AVFilterLink *inlink_a = avctx->inputs[IN_A];
460  AVFilterLink *inlink_b = avctx->inputs[IN_B];
461  FilterLink *il = ff_filter_link(inlink_a);
462  FilterLink *ol = ff_filter_link(outlink);
463 
464  if (inlink_a->w != inlink_b->w || inlink_a->h != inlink_b->h) {
465  av_log(avctx, AV_LOG_ERROR, "First input link %s parameters "
466  "(size %dx%d) do not match the corresponding "
467  "second input link %s parameters (size %dx%d)\n",
468  avctx->input_pads[IN_A].name, inlink_a->w, inlink_a->h,
469  avctx->input_pads[IN_B].name, inlink_b->w, inlink_b->h);
470  return AVERROR(EINVAL);
471  }
472 
473  if (inlink_a->time_base.num != inlink_b->time_base.num ||
474  inlink_a->time_base.den != inlink_b->time_base.den) {
475  av_log(avctx, AV_LOG_ERROR, "First input link %s timebase "
476  "(%d/%d) does not match the corresponding "
477  "second input link %s timebase (%d/%d)\n",
478  avctx->input_pads[IN_A].name, inlink_a->time_base.num, inlink_a->time_base.den,
479  avctx->input_pads[IN_B].name, inlink_b->time_base.num, inlink_b->time_base.den);
480  return AVERROR(EINVAL);
481  }
482 
483  s->start_pts = s->inputs_offset_pts = AV_NOPTS_VALUE;
484 
485  outlink->time_base = inlink_a->time_base;
486  ol->frame_rate = il->frame_rate;
487  outlink->sample_aspect_ratio = inlink_a->sample_aspect_ratio;
488 
489  if (s->duration)
490  s->duration_pts = av_rescale_q(s->duration, AV_TIME_BASE_Q, inlink_a->time_base);
491  RET(ff_vk_filter_config_output(outlink));
492 
493 fail:
494  return err;
495 }
496 
497 static int forward_frame(XFadeVulkanContext *s,
498  AVFilterLink *inlink, AVFilterLink *outlink)
499 {
500  int64_t status_pts;
501  int ret = 0, status;
502  AVFrame *frame = NULL;
503 
504  ret = ff_inlink_consume_frame(inlink, &frame);
505  if (ret < 0)
506  return ret;
507 
508  if (ret > 0) {
509  // If we do not have an offset yet, it's because we
510  // never got a first input. Just offset to 0
511  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
512  s->inputs_offset_pts = -frame->pts;
513 
514  // We got a frame, nothing to do other than adjusting the timestamp
515  frame->pts += s->inputs_offset_pts;
516  return ff_filter_frame(outlink, frame);
517  }
518 
519  // Forward status with our timestamp
520  if (ff_inlink_acknowledge_status(inlink, &status, &status_pts)) {
521  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
522  s->inputs_offset_pts = -status_pts;
523 
524  ff_outlink_set_status(outlink, status, status_pts + s->inputs_offset_pts);
525  return 0;
526  }
527 
528  // No frame available, request one if needed
529  if (ff_outlink_frame_wanted(outlink))
530  ff_inlink_request_frame(inlink);
531 
532  return 0;
533 }
534 
535 static int activate(AVFilterContext *avctx)
536 {
537  XFadeVulkanContext *s = avctx->priv;
538  AVFilterLink *in_a = avctx->inputs[IN_A];
539  AVFilterLink *in_b = avctx->inputs[IN_B];
540  AVFilterLink *outlink = avctx->outputs[0];
541  int64_t status_pts;
542 
543  FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, avctx);
544 
545  // Check if we already transitioned or IN_A ended prematurely,
546  // in which case just forward the frames from IN_B with adjusted
547  // timestamps until EOF.
548  if (s->status[IN_A] && !s->status[IN_B])
549  return forward_frame(s, in_b, outlink);
550 
551  // We did not finish transitioning yet and the first stream
552  // did not end either, so check if there are more frames to consume.
553  if (ff_inlink_check_available_frame(in_a)) {
554  AVFrame *peeked_frame = ff_inlink_peek_frame(in_a, 0);
555  s->pts = peeked_frame->pts;
556 
557  if (s->start_pts == AV_NOPTS_VALUE)
558  s->start_pts =
559  s->pts + av_rescale_q(s->offset, AV_TIME_BASE_Q, in_a->time_base);
560 
561  // Check if we are not yet transitioning, in which case
562  // just request and forward the input frame.
563  if (s->start_pts > s->pts) {
564  AVFrame *frame_a = NULL;
565  s->passthrough = 1;
566  ff_inlink_consume_frame(in_a, &frame_a);
567  return ff_filter_frame(outlink, frame_a);
568  }
569  s->passthrough = 0;
570 
571  // We are transitioning, so we need a frame from IN_B
572  if (ff_inlink_check_available_frame(in_b)) {
573  int ret;
574  AVFrame *frame_a = NULL, *frame_b = NULL;
575  ff_inlink_consume_frame(avctx->inputs[IN_A], &frame_a);
576  ff_inlink_consume_frame(avctx->inputs[IN_B], &frame_b);
577 
578  // Calculate PTS offset to first input
579  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
580  s->inputs_offset_pts = s->pts - frame_b->pts;
581 
582  // Check if we finished transitioning, in which case we
583  // report back EOF to IN_A as it is no longer needed.
584  if (s->pts - s->start_pts > s->duration_pts) {
585  s->status[IN_A] = AVERROR_EOF;
586  ff_inlink_set_status(in_a, AVERROR_EOF);
587  s->passthrough = 1;
588  }
589  ret = xfade_frame(avctx, frame_a, frame_b);
590  av_frame_free(&frame_a);
591  av_frame_free(&frame_b);
592  return ret;
593  }
594 
595  // We did not get a frame from IN_B, check its status.
596  if (ff_inlink_acknowledge_status(in_b, &s->status[IN_B], &status_pts)) {
597  // We should transition, but IN_B is EOF so just report EOF output now.
598  ff_outlink_set_status(outlink, s->status[IN_B], s->pts);
599  return 0;
600  }
601 
602  // We did not get a frame for IN_B but no EOF either, so just request more.
603  if (ff_outlink_frame_wanted(outlink)) {
604  ff_inlink_request_frame(in_b);
605  return 0;
606  }
607  }
608 
609  // We did not get a frame from IN_A, check its status.
610  if (ff_inlink_acknowledge_status(in_a, &s->status[IN_A], &status_pts)) {
611  // No more frames from IN_A, do not report EOF though, we will just
612  // forward the IN_B frames in the next activate calls.
613  s->passthrough = 1;
614  ff_filter_set_ready(avctx, 100);
615  return 0;
616  }
617 
618  // We have no frames yet from IN_A and no EOF, so request some.
619  if (ff_outlink_frame_wanted(outlink)) {
620  ff_inlink_request_frame(in_a);
621  return 0;
622  }
623 
624  return FFERROR_NOT_READY;
625 }
626 
627 static av_cold void uninit(AVFilterContext *avctx)
628 {
629  XFadeVulkanContext *s = avctx->priv;
630  FFVulkanContext *vkctx = &s->vkctx;
631  FFVulkanFunctions *vk = &vkctx->vkfn;
632 
633  ff_vk_exec_pool_free(vkctx, &s->e);
634  ff_vk_shader_free(vkctx, &s->shd);
635 
636  if (s->sampler)
637  vk->DestroySampler(vkctx->hwctx->act_dev, s->sampler,
638  vkctx->hwctx->alloc);
639 
640  ff_vk_uninit(&s->vkctx);
641 
642  s->initialized = 0;
643 }
644 
645 static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
646 {
647  XFadeVulkanContext *s = inlink->dst->priv;
648 
649  return s->passthrough ?
650  ff_null_get_video_buffer (inlink, w, h) :
651  ff_default_get_video_buffer(inlink, w, h);
652 }
653 
654 #define OFFSET(x) offsetof(XFadeVulkanContext, x)
655 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
656 
657 static const AVOption xfade_vulkan_options[] = {
658  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=FADE}, 0, NB_TRANSITIONS-1, FLAGS, .unit = "transition" },
659  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, .unit = "transition" },
660  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, .unit = "transition" },
661  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, .unit = "transition" },
662  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, .unit = "transition" },
663  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, .unit = "transition" },
664  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, .unit = "transition" },
665  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, .unit = "transition" },
666  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, .unit = "transition" },
667  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, .unit = "transition" },
668  { "circleopen", "circleopen transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLEOPEN}, 0, 0, FLAGS, .unit = "transition" },
669  { "circleclose", "circleclose transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLECLOSE}, 0, 0, FLAGS, .unit = "transition" },
670  { "dissolve", "dissolve transition", 0, AV_OPT_TYPE_CONST, {.i64=DISSOLVE}, 0, 0, FLAGS, .unit = "transition" },
671  { "pixelize", "pixelize transition", 0, AV_OPT_TYPE_CONST, {.i64=PIXELIZE}, 0, 0, FLAGS, .unit = "transition" },
672  { "wipetl", "wipe top left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETL}, 0, 0, FLAGS, .unit = "transition" },
673  { "wipetr", "wipe top right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETR}, 0, 0, FLAGS, .unit = "transition" },
674  { "wipebl", "wipe bottom left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBL}, 0, 0, FLAGS, .unit = "transition" },
675  { "wipebr", "wipe bottom right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBR}, 0, 0, FLAGS, .unit = "transition" },
676  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
677  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
678  { NULL }
679 };
680 
681 AVFILTER_DEFINE_CLASS(xfade_vulkan);
682 
683 static const AVFilterPad xfade_vulkan_inputs[] = {
684  {
685  .name = "main",
686  .type = AVMEDIA_TYPE_VIDEO,
687  .get_buffer.video = &get_video_buffer,
688  .config_props = &ff_vk_filter_config_input,
689  },
690  {
691  .name = "xfade",
692  .type = AVMEDIA_TYPE_VIDEO,
693  .get_buffer.video = &get_video_buffer,
694  .config_props = &ff_vk_filter_config_input,
695  },
696 };
697 
698 static const AVFilterPad xfade_vulkan_outputs[] = {
699  {
700  .name = "default",
701  .type = AVMEDIA_TYPE_VIDEO,
702  .config_props = &config_props_output,
703  },
704 };
705 
706 const AVFilter ff_vf_xfade_vulkan = {
707  .name = "xfade_vulkan",
708  .description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
709  .priv_size = sizeof(XFadeVulkanContext),
710  .init = &ff_vk_filter_init,
711  .uninit = &uninit,
712  .activate = &activate,
713  FILTER_INPUTS(xfade_vulkan_inputs),
714  FILTER_OUTPUTS(xfade_vulkan_outputs),
715  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VULKAN),
716  .priv_class = &xfade_vulkan_class,
717  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
718  .flags = AVFILTER_FLAG_HWDEVICE,
719 };
IN_A
#define IN_A
Definition: vf_xfade_vulkan.c:27
WIPETR
@ WIPETR
Definition: vf_xfade_vulkan.c:83
ff_get_video_buffer
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:116
IN_B
#define IN_B
Definition: vf_xfade_vulkan.c:28
WIPELEFT
@ WIPELEFT
Definition: vf_xfade_vulkan.c:70
mix
static int mix(int c0, int c1)
Definition: 4xm.c:716
AVERROR
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
opt.h
transition_dissolve
static const char transition_dissolve[]
Definition: vf_xfade_vulkan.c:225
ff_vk_shader_free
void ff_vk_shader_free(FFVulkanContext *s, FFVulkanShader *shd)
Free a shader.
Definition: vulkan.c:2529
ff_vk_shader_init
int ff_vk_shader_init(FFVulkanContext *s, FFVulkanShader *shd, const char *name, VkPipelineStageFlags stage, const char *extensions[], int nb_extensions, int lg_x, int lg_y, int lg_z, uint32_t required_subgroup_size)
Initialize a shader object, with a specific set of extensions, type+bind, local group size,...
Definition: vulkan.c:1680
xfade_frame
static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
Definition: vf_xfade_vulkan.c:413
transition_wipeup
static const char transition_wipeup[]
Definition: vf_xfade_vulkan.c:120
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1062
AVERROR_EOF
#define AVERROR_EOF
End of file.
Definition: error.h:57
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
RET
#define RET(x)
Definition: vulkan.h:67
FFERROR_NOT_READY
return FFERROR_NOT_READY
Definition: filter_design.txt:204
AV_TIME_BASE_Q
#define AV_TIME_BASE_Q
Internal time base represented as fractional value.
Definition: avutil.h:264
ff_vk_qf_init
int ff_vk_qf_init(FFVulkanContext *s, FFVkQueueFamilyCtx *qf, VkQueueFlagBits dev_family)
Chooses a QF and loads it into a context.
Definition: vulkan.c:228
int64_t
long long int64_t
Definition: coverity.c:34
output
filter_frame For filters that do not use the this method is called when a frame is pushed to the filter s input It can be called at any time except in a reentrant way If the input frame is enough to produce output
Definition: filter_design.txt:225
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
transition_slideup
static const char transition_slideup[]
Definition: vf_xfade_vulkan.c:164
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:162
XFadeVulkanContext::start_pts
int64_t start_pts
Definition: vf_xfade_vulkan.c:49
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:501
ff_vk_filter_init
int ff_vk_filter_init(AVFilterContext *avctx)
General lavfi IO functions.
Definition: vulkan_filter.c:233
w
uint8_t w
Definition: llviddspenc.c:38
frand
static float frand(int x, int y)
Definition: vf_deband.c:119
XFadeVulkanContext::status
int status[IN_NB]
Definition: vf_xfade_vulkan.c:65
transitions_map
static const char * transitions_map[NB_TRANSITIONS]
Definition: vf_xfade_vulkan.c:299
XFadeVulkanContext::transition
int transition
Definition: vf_xfade_vulkan.c:38
AVOption
AVOption.
Definition: opt.h:429
b
#define b
Definition: input.c:41
get_video_buffer
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
Definition: vf_xfade_vulkan.c:645
CIRCLEOPEN
@ CIRCLEOPEN
Definition: vf_xfade_vulkan.c:78
AV_OPT_TYPE_DURATION
@ AV_OPT_TYPE_DURATION
Underlying C type is int64_t.
Definition: opt.h:319
transition_wiperight
static const char transition_wiperight[]
Definition: vf_xfade_vulkan.c:109
forward_frame
static int forward_frame(XFadeVulkanContext *s, AVFilterLink *inlink, AVFilterLink *outlink)
Definition: vf_xfade_vulkan.c:497
XFadeVulkanContext::e
FFVkExecPool e
Definition: vf_xfade_vulkan.c:43
transition_wipebl
static const char transition_wipebl[]
Definition: vf_xfade_vulkan.c:275
XFadeVulkanContext::qf
FFVkQueueFamilyCtx qf
Definition: vf_xfade_vulkan.c:44
XFadeVulkanContext::initialized
int initialized
Definition: vf_xfade_vulkan.c:42
ff_vk_uninit
void ff_vk_uninit(FFVulkanContext *s)
Frees main context.
Definition: vulkan.c:2568
transition_wipebr
static const char transition_wipebr[]
Definition: vf_xfade_vulkan.c:287
FFVkSPIRVCompiler::uninit
void(* uninit)(struct FFVkSPIRVCompiler **ctx)
Definition: vulkan_spirv.h:32
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:205
video.h
AV_PIX_FMT_VULKAN
@ AV_PIX_FMT_VULKAN
Vulkan hardware images.
Definition: pixfmt.h:379
ff_vf_xfade_vulkan
const AVFilter ff_vf_xfade_vulkan
Definition: vf_xfade_vulkan.c:706
ff_inlink_consume_frame
int ff_inlink_consume_frame(AVFilterLink *link, AVFrame **rframe)
Take a frame from the link's FIFO and update the link's stats.
Definition: avfilter.c:1491
ff_default_get_video_buffer
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:111
WIPEBL
@ WIPEBL
Definition: vf_xfade_vulkan.c:84
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3210
SLIDELEFT
@ SLIDELEFT
Definition: vf_xfade_vulkan.c:76
FF_FILTER_FORWARD_STATUS_BACK_ALL
#define FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, filter)
Forward the status on an output link to all input links.
Definition: filters.h:447
transition_slideleft
static const char transition_slideleft[]
Definition: vf_xfade_vulkan.c:172
AVVulkanDeviceContext::alloc
const VkAllocationCallbacks * alloc
Custom memory allocator, else NULL.
Definition: hwcontext_vulkan.h:63
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:472
fail
#define fail()
Definition: checkasm.h:188
SLIDEDOWN
@ SLIDEDOWN
Definition: vf_xfade_vulkan.c:74
vulkan_filter.h
ff_vk_shader_register_exec
int ff_vk_shader_register_exec(FFVulkanContext *s, FFVkExecPool *pool, FFVulkanShader *shd)
Register a shader with an exec pool.
Definition: vulkan.c:2169
SLIDEUP
@ SLIDEUP
Definition: vf_xfade_vulkan.c:75
ff_vk_shader_add_descriptor_set
int ff_vk_shader_add_descriptor_set(FFVulkanContext *s, FFVulkanShader *shd, FFVulkanDescriptorSetBinding *desc, int nb, int singular, int print_to_shader_only)
Add descriptor to a shader.
Definition: vulkan.c:2044
XFadeVulkanContext::pts
int64_t pts
Definition: vf_xfade_vulkan.c:58
AVRational::num
int num
Numerator.
Definition: rational.h:59
XFadeVulkanContext::passthrough
int passthrough
Definition: vf_xfade_vulkan.c:63
SHADER_FRAND_FUNC
#define SHADER_FRAND_FUNC
Definition: vf_xfade_vulkan.c:219
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:38
transition_circleclose
static const char transition_circleclose[]
Definition: vf_xfade_vulkan.c:211
C
s EdgeDetect Foobar g libavfilter vf_edgedetect c libavfilter vf_foobar c edit libavfilter and add an entry for foobar following the pattern of the other filters edit libavfilter allfilters and add an entry for foobar following the pattern of the other filters configure make j< whatever > ffmpeg ffmpeg i you should get a foobar png with Lena edge detected That s your new playground is ready Some little details about what s going which in turn will define variables for the build system and the C
Definition: writing_filters.txt:58
AVFilterContext::input_pads
AVFilterPad * input_pads
array of input pads
Definition: avfilter.h:464
avassert.h
ceil
static __device__ float ceil(float a)
Definition: cuda_runtime.h:176
GLSLC
#define GLSLC(N, S)
Definition: vulkan.h:44
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:209
av_cold
#define av_cold
Definition: attributes.h:90
main
int main
Definition: dovi_rpuenc.c:37
duration
int64_t duration
Definition: movenc.c:65
ff_outlink_set_status
static void ff_outlink_set_status(AVFilterLink *link, int status, int64_t pts)
Set the status field of a link from the source filter.
Definition: filters.h:424
ff_inlink_request_frame
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
Definition: avfilter.c:1594
s
#define s(width, name)
Definition: cbs_vp9.c:198
floor
static __device__ float floor(float a)
Definition: cuda_runtime.h:173
filters.h
FF_VK_REP_FLOAT
@ FF_VK_REP_FLOAT
Definition: vulkan.h:367
transition_wipedown
static const char transition_wipedown[]
Definition: vf_xfade_vulkan.c:131
av_rescale_q
int64_t av_rescale_q(int64_t a, AVRational bq, AVRational cq)
Rescale a 64-bit integer by 2 rational numbers.
Definition: mathematics.c:142
GLSLD
#define GLSLD(D)
Definition: vulkan.h:59
ff_vk_exec_pool_free
void ff_vk_exec_pool_free(FFVulkanContext *s, FFVkExecPool *pool)
Definition: vulkan.c:238
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
ff_inlink_peek_frame
AVFrame * ff_inlink_peek_frame(AVFilterLink *link, size_t idx)
Access a frame in the link fifo without consuming it.
Definition: avfilter.c:1532
planes
static const struct @465 planes[]
ff_vk_shader_rep_fmt
const char * ff_vk_shader_rep_fmt(enum AVPixelFormat pix_fmt, enum FFVkShaderRepFormat rep_fmt)
Definition: vulkan.c:1291
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:210
av_frame_copy_props
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:713
DUP_SAMPLER
#define DUP_SAMPLER(x)
Definition: vulkan.h:73
transition_pixelize
static const char transition_pixelize[]
Definition: vf_xfade_vulkan.c:236
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:465
activate
static int activate(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:535
transition_circleopen
static const char transition_circleopen[]
Definition: vf_xfade_vulkan.c:203
ff_vk_filter_config_output
int ff_vk_filter_config_output(AVFilterLink *outlink)
Definition: vulkan_filter.c:209
transition_wipetl
static const char transition_wipetl[]
Definition: vf_xfade_vulkan.c:251
transition_fade
static const char transition_fade[]
Definition: vf_xfade_vulkan.c:89
xfade_vulkan_outputs
static const AVFilterPad xfade_vulkan_outputs[]
Definition: vf_xfade_vulkan.c:698
xfade_vulkan_options
static const AVOption xfade_vulkan_options[]
Definition: vf_xfade_vulkan.c:657
ff_vk_exec_pool_init
int ff_vk_exec_pool_init(FFVulkanContext *s, FFVkQueueFamilyCtx *qf, FFVkExecPool *pool, int nb_contexts, int nb_queries, VkQueryType query_type, int query_64bit, const void *query_create_pnext)
Allocates/frees an execution pool.
Definition: vulkan.c:300
av_clipf
av_clipf
Definition: af_crystalizer.c:122
FFVulkanContext
Definition: vulkan.h:263
ff_inlink_acknowledge_status
int ff_inlink_acknowledge_status(AVFilterLink *link, int *rstatus, int64_t *rpts)
Test and acknowledge the change of status on the link.
Definition: avfilter.c:1438
XFadeVulkanContext::shd
FFVulkanShader shd
Definition: vf_xfade_vulkan.c:45
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
PIXELIZE
@ PIXELIZE
Definition: vf_xfade_vulkan.c:81
DISSOLVE
@ DISSOLVE
Definition: vf_xfade_vulkan.c:80
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: filters.h:206
ff_vk_filter_process_Nin
int ff_vk_filter_process_Nin(FFVulkanContext *vkctx, FFVkExecPool *e, FFVulkanShader *shd, AVFrame *out, AVFrame *in[], int nb_in, VkSampler sampler, void *push_src, size_t push_size)
Up to 16 inputs, one output.
Definition: vulkan_filter.c:401
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:368
FLAGS
#define FLAGS
Definition: vf_xfade_vulkan.c:655
ff_inlink_set_status
void ff_inlink_set_status(AVFilterLink *link, int status)
Set the status on an input link.
Definition: avfilter.c:1603
ff_inlink_check_available_frame
int ff_inlink_check_available_frame(AVFilterLink *link)
Test if a frame is available on the link.
Definition: avfilter.c:1460
FFVulkanDescriptorSetBinding
Definition: vulkan.h:75
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:94
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:173
SHADER_CIRCLE_COMMON
#define SHADER_CIRCLE_COMMON
Definition: vf_xfade_vulkan.c:188
size
int size
Definition: twinvq_data.h:10344
AV_NOPTS_VALUE
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:248
FFVkQueueFamilyCtx
Definition: vulkan.h:102
FFVulkanShader
Definition: vulkan.h:179
XFadeVulkanContext::sampler
VkSampler sampler
Definition: vf_xfade_vulkan.c:46
FFVkSPIRVCompiler::compile_shader
int(* compile_shader)(FFVulkanContext *s, struct FFVkSPIRVCompiler *ctx, FFVulkanShader *shd, uint8_t **data, size_t *size, const char *entrypoint, void **opaque)
Definition: vulkan_spirv.h:28
XFadeVulkanContext::duration
int64_t duration
Definition: vf_xfade_vulkan.c:39
a
The reader does not expect b to be semantically here and if the code is changed by maybe adding a a division or other the signedness will almost certainly be mistaken To avoid this confusion a new type was SUINT is the C unsigned type but it holds a signed int to use the same example SUINT a
Definition: undefined.txt:41
XFadeVulkanContext::offset
int64_t offset
Definition: vf_xfade_vulkan.c:40
AVERROR_EXTERNAL
#define AVERROR_EXTERNAL
Generic error in an external library.
Definition: error.h:59
offset
it s the only field you need to keep assuming you have a context There is some magic you don t need to care about around this just let it vf offset
Definition: writing_filters.txt:86
XFadeVulkanContext::inputs_offset_pts
int64_t inputs_offset_pts
Definition: vf_xfade_vulkan.c:52
ff_null_get_video_buffer
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:44
FFVkSPIRVCompiler
Definition: vulkan_spirv.h:26
layout
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel layout
Definition: filter_design.txt:18
xfade_vulkan_inputs
static const AVFilterPad xfade_vulkan_inputs[]
Definition: vf_xfade_vulkan.c:683
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
NB_TRANSITIONS
@ NB_TRANSITIONS
Definition: vf_xfade_vulkan.c:86
ff_vk_shader_link
int ff_vk_shader_link(FFVulkanContext *s, FFVulkanShader *shd, uint8_t *spirv, size_t spirv_len, const char *entrypoint)
Link a shader into an executable.
Definition: vulkan.c:1969
vulkan_spirv.h
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
CIRCLECLOSE
@ CIRCLECLOSE
Definition: vf_xfade_vulkan.c:79
XFadeParameters
Definition: vf_xfade_vulkan.c:31
FFVkSPIRVCompiler::free_shader
void(* free_shader)(struct FFVkSPIRVCompiler *ctx, void **opaque)
Definition: vulkan_spirv.h:31
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
ret
ret
Definition: filter_design.txt:187
WIPERIGHT
@ WIPERIGHT
Definition: vf_xfade_vulkan.c:71
FADE
@ FADE
Definition: vf_xfade_vulkan.c:69
frame
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several the filter must be ready for frames arriving randomly on any input any filter with several inputs will most likely require some kind of queuing mechanism It is perfectly acceptable to have a limited queue and to drop frames when the inputs are too unbalanced request_frame For filters that do not use the this method is called when a frame is wanted on an output For a it should directly call filter_frame on the corresponding output For a if there are queued frames already one of these frames should be pushed If the filter should request a frame on one of its repeatedly until at least one frame has been pushed Return or at least make progress towards producing a frame
Definition: filter_design.txt:264
FFVulkanContext::vkfn
FFVulkanFunctions vkfn
Definition: vulkan.h:267
SHADER_SLIDE_COMMON
#define SHADER_SLIDE_COMMON
Definition: vf_xfade_vulkan.c:142
FFVkExecPool
Definition: vulkan.h:241
pos
unsigned int pos
Definition: spdifenc.c:414
XFadeVulkanContext::duration_pts
int64_t duration_pts
Definition: vf_xfade_vulkan.c:55
ff_vk_shader_add_push_const
int ff_vk_shader_add_push_const(FFVulkanShader *shd, int offset, int size, VkShaderStageFlagBits stage)
Add/update push constants for execution.
Definition: vulkan.c:1223
XFadeVulkanContext::vkctx
FFVulkanContext vkctx
Definition: vf_xfade_vulkan.c:36
AVFrame::hw_frames_ctx
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
Definition: frame.h:740
status
ov_status_e status
Definition: dnn_backend_openvino.c:100
random_seed.h
transition_slideright
static const char transition_slideright[]
Definition: vf_xfade_vulkan.c:180
GLSLF
#define GLSLF(N, S,...)
Definition: vulkan.h:54
AVRational::den
int den
Denominator.
Definition: rational.h:60
WIPEBR
@ WIPEBR
Definition: vf_xfade_vulkan.c:85
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
config_props_output
static int config_props_output(AVFilterLink *outlink)
Definition: vf_xfade_vulkan.c:454
init_vulkan
static av_cold int init_vulkan(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:319
transition_wipetr
static const char transition_wipetr[]
Definition: vf_xfade_vulkan.c:263
XFadeVulkanContext
Definition: vf_xfade_vulkan.c:35
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
desc
const char * desc
Definition: libsvtav1.c:79
ff_vk_filter_config_input
int ff_vk_filter_config_input(AVFilterLink *inlink)
Definition: vulkan_filter.c:176
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
FFVulkanContext::hwctx
AVVulkanDeviceContext * hwctx
Definition: vulkan.h:291
WIPEDOWN
@ WIPEDOWN
Definition: vf_xfade_vulkan.c:73
AVVulkanDeviceContext::act_dev
VkDevice act_dev
Active device.
Definition: hwcontext_vulkan.h:84
XFadeTransitions
XFadeTransitions
Definition: vf_xfade.c:29
WIPETL
@ WIPETL
Definition: vf_xfade_vulkan.c:82
ff_vk_init_sampler
int ff_vk_init_sampler(FFVulkanContext *s, VkSampler *sampler, int unnorm_coords, VkFilter filt)
Create a sampler.
Definition: vulkan.c:1244
SLIDERIGHT
@ SLIDERIGHT
Definition: vf_xfade_vulkan.c:77
XFadeParameters::progress
float progress
Definition: vf_xfade_vulkan.c:32
transition_wipeleft
static const char transition_wipeleft[]
Definition: vf_xfade_vulkan.c:98
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
IN_NB
#define IN_NB
Definition: vf_xfade_vulkan.c:29
h
h
Definition: vp9dsp_template.c:2070
ff_outlink_frame_wanted
the definition of that something depends on the semantic of the filter The callback must examine the status of the filter s links and proceed accordingly The status of output links is stored in the status_in and status_out fields and tested by the ff_outlink_frame_wanted() function. If this function returns true
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(xfade_vulkan)
WIPEUP
@ WIPEUP
Definition: vf_xfade_vulkan.c:72
transition_slidedown
static const char transition_slidedown[]
Definition: vf_xfade_vulkan.c:156
uninit
static av_cold void uninit(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:627
OFFSET
#define OFFSET(x)
Definition: vf_xfade_vulkan.c:654
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
FFVulkanFunctions
Definition: vulkan_functions.h:263
ff_filter_set_ready
void ff_filter_set_ready(AVFilterContext *filter, unsigned priority)
Mark a filter ready and schedule it for activation.
Definition: avfilter.c:239
min
float min
Definition: vorbis_enc_data.h:429
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:469
Generated on Wed Nov 6 2024 19:22:46 for FFmpeg by   doxygen 1.8.17