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 "vulkan_filter.h"
23 #include "vulkan_spirv.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  FFVulkanPipeline pl;
44  FFVkExecPool e;
45  FFVkQueueFamilyCtx qf;
46  FFVkSPIRVShader shd;
47  VkSampler sampler;
48 
49  // PTS when the fade should start (in IN_A timebase)
50  int64_t start_pts;
51 
52  // PTS offset between IN_A and IN_B
53  int64_t inputs_offset_pts;
54 
55  // Duration of the transition
56  int64_t duration_pts;
57 
58  // Current PTS of the first input (IN_A)
59  int64_t pts;
60 
61  // If frames are currently just passed through
62  // unmodified, like before and after the actual
63  // transition.
64  int passthrough;
65 
66  int status[IN_NB];
67 } XFadeVulkanContext;
68 
69 enum XFadeTransitions {
70  FADE,
71  WIPELEFT,
72  WIPERIGHT,
73  WIPEUP,
74  WIPEDOWN,
75  SLIDEDOWN,
76  SLIDEUP,
77  SLIDELEFT,
78  SLIDERIGHT,
79  CIRCLEOPEN,
80  CIRCLECLOSE,
81  DISSOLVE,
82  PIXELIZE,
83  WIPETL,
84  WIPETR,
85  WIPEBL,
86  WIPEBR,
87  NB_TRANSITIONS,
88 };
89 
90 static const char transition_fade[] = {
91  C(0, void transition(int idx, ivec2 pos, float progress) )
92  C(0, { )
93  C(1, vec4 a = texture(a_images[idx], pos); )
94  C(1, vec4 b = texture(b_images[idx], pos); )
95  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
96  C(0, } )
97 };
98 
99 static const char transition_wipeleft[] = {
100  C(0, void transition(int idx, ivec2 pos, float progress) )
101  C(0, { )
102  C(1, ivec2 size = imageSize(output_images[idx]); )
103  C(1, int s = int(size.x * (1.0 - progress)); )
104  C(1, vec4 a = texture(a_images[idx], pos); )
105  C(1, vec4 b = texture(b_images[idx], pos); )
106  C(1, imageStore(output_images[idx], pos, pos.x > s ? b : a); )
107  C(0, } )
108 };
109 
110 static const char transition_wiperight[] = {
111  C(0, void transition(int idx, ivec2 pos, float progress) )
112  C(0, { )
113  C(1, ivec2 size = imageSize(output_images[idx]); )
114  C(1, int s = int(size.x * progress); )
115  C(1, vec4 a = texture(a_images[idx], pos); )
116  C(1, vec4 b = texture(b_images[idx], pos); )
117  C(1, imageStore(output_images[idx], pos, pos.x > s ? a : b); )
118  C(0, } )
119 };
120 
121 static const char transition_wipeup[] = {
122  C(0, void transition(int idx, ivec2 pos, float progress) )
123  C(0, { )
124  C(1, ivec2 size = imageSize(output_images[idx]); )
125  C(1, int s = int(size.y * (1.0 - progress)); )
126  C(1, vec4 a = texture(a_images[idx], pos); )
127  C(1, vec4 b = texture(b_images[idx], pos); )
128  C(1, imageStore(output_images[idx], pos, pos.y > s ? b : a); )
129  C(0, } )
130 };
131 
132 static const char transition_wipedown[] = {
133  C(0, void transition(int idx, ivec2 pos, float progress) )
134  C(0, { )
135  C(1, ivec2 size = imageSize(output_images[idx]); )
136  C(1, int s = int(size.y * progress); )
137  C(1, vec4 a = texture(a_images[idx], pos); )
138  C(1, vec4 b = texture(b_images[idx], pos); )
139  C(1, imageStore(output_images[idx], pos, pos.y > s ? a : b); )
140  C(0, } )
141 };
142 
143 #define SHADER_SLIDE_COMMON \
144  C(0, void slide(int idx, ivec2 pos, float progress, ivec2 direction) ) \
145  C(0, { ) \
146  C(1, ivec2 size = imageSize(output_images[idx]); ) \
147  C(1, ivec2 pi = ivec2(progress * size); ) \
148  C(1, ivec2 p = pos + pi * direction; ) \
149  C(1, ivec2 f = p % size; ) \
150  C(1, f = f + size * ivec2(f.x < 0, f.y < 0); ) \
151  C(1, vec4 a = texture(a_images[idx], f); ) \
152  C(1, vec4 b = texture(b_images[idx], f); ) \
153  C(1, vec4 r = (p.y >= 0 && p.x >= 0 && size.y > p.y && size.x > p.x) ? a : b; ) \
154  C(1, imageStore(output_images[idx], pos, r); ) \
155  C(0, } )
156 
157 static const char transition_slidedown[] = {
158  SHADER_SLIDE_COMMON
159  C(0, void transition(int idx, ivec2 pos, float progress) )
160  C(0, { )
161  C(1, slide(idx, pos, progress, ivec2(0, -1)); )
162  C(0, } )
163 };
164 
165 static const char transition_slideup[] = {
166  SHADER_SLIDE_COMMON
167  C(0, void transition(int idx, ivec2 pos, float progress) )
168  C(0, { )
169  C(1, slide(idx, pos, progress, ivec2(0, +1)); )
170  C(0, } )
171 };
172 
173 static const char transition_slideleft[] = {
174  SHADER_SLIDE_COMMON
175  C(0, void transition(int idx, ivec2 pos, float progress) )
176  C(0, { )
177  C(1, slide(idx, pos, progress, ivec2(+1, 0)); )
178  C(0, } )
179 };
180 
181 static const char transition_slideright[] = {
182  SHADER_SLIDE_COMMON
183  C(0, void transition(int idx, ivec2 pos, float progress) )
184  C(0, { )
185  C(1, slide(idx, pos, progress, ivec2(-1, 0)); )
186  C(0, } )
187 };
188 
189 #define SHADER_CIRCLE_COMMON \
190  C(0, void circle(int idx, ivec2 pos, float progress, bool open) ) \
191  C(0, { ) \
192  C(1, const ivec2 half_size = imageSize(output_images[idx]) / 2; ) \
193  C(1, const float z = dot(half_size, half_size); ) \
194  C(1, float p = ((open ? (1.0 - progress) : progress) - 0.5) * 3.0; ) \
195  C(1, ivec2 dsize = pos - half_size; ) \
196  C(1, float sm = dot(dsize, dsize) / z + p; ) \
197  C(1, vec4 a = texture(a_images[idx], pos); ) \
198  C(1, vec4 b = texture(b_images[idx], pos); ) \
199  C(1, imageStore(output_images[idx], pos, \
200  mix(open ? b : a, open ? a : b, \
201  smoothstep(0.f, 1.f, sm))); ) \
202  C(0, } )
203 
204 static const char transition_circleopen[] = {
205  SHADER_CIRCLE_COMMON
206  C(0, void transition(int idx, ivec2 pos, float progress) )
207  C(0, { )
208  C(1, circle(idx, pos, progress, true); )
209  C(0, } )
210 };
211 
212 static const char transition_circleclose[] = {
213  SHADER_CIRCLE_COMMON
214  C(0, void transition(int idx, ivec2 pos, float progress) )
215  C(0, { )
216  C(1, circle(idx, pos, progress, false); )
217  C(0, } )
218 };
219 
220 #define SHADER_FRAND_FUNC \
221  C(0, float frand(vec2 v) ) \
222  C(0, { ) \
223  C(1, return fract(sin(dot(v, vec2(12.9898, 78.233))) * 43758.545); ) \
224  C(0, } )
225 
226 static const char transition_dissolve[] = {
227  SHADER_FRAND_FUNC
228  C(0, void transition(int idx, ivec2 pos, float progress) )
229  C(0, { )
230  C(1, float sm = frand(pos) * 2.0 + (1.0 - progress) * 2.0 - 1.5; )
231  C(1, vec4 a = texture(a_images[idx], pos); )
232  C(1, vec4 b = texture(b_images[idx], pos); )
233  C(1, imageStore(output_images[idx], pos, sm >= 0.5 ? a : b); )
234  C(0, } )
235 };
236 
237 static const char transition_pixelize[] = {
238  C(0, void transition(int idx, ivec2 pos, float progress) )
239  C(0, { )
240  C(1, ivec2 size = imageSize(output_images[idx]); )
241  C(1, float d = min(progress, 1.0 - progress); )
242  C(1, float dist = ceil(d * 50.0) / 50.0; )
243  C(1, float sq = 2.0 * dist * min(size.x, size.y) / 20.0; )
244  C(1, float sx = dist > 0.0 ? min((floor(pos.x / sq) + 0.5) * sq, size.x - 1) : pos.x; )
245  C(1, float sy = dist > 0.0 ? min((floor(pos.y / sq) + 0.5) * sq, size.y - 1) : pos.y; )
246  C(1, vec4 a = texture(a_images[idx], vec2(sx, sy)); )
247  C(1, vec4 b = texture(b_images[idx], vec2(sx, sy)); )
248  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
249  C(0, } )
250 };
251 
252 static const char transition_wipetl[] = {
253  C(0, void transition(int idx, ivec2 pos, float progress) )
254  C(0, { )
255  C(1, ivec2 size = imageSize(output_images[idx]); )
256  C(1, float zw = size.x * (1.0 - progress); )
257  C(1, float zh = size.y * (1.0 - progress); )
258  C(1, vec4 a = texture(a_images[idx], pos); )
259  C(1, vec4 b = texture(b_images[idx], pos); )
260  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x <= zw) ? a : b); )
261  C(0, } )
262 };
263 
264 static const char transition_wipetr[] = {
265  C(0, void transition(int idx, ivec2 pos, float progress) )
266  C(0, { )
267  C(1, ivec2 size = imageSize(output_images[idx]); )
268  C(1, float zw = size.x * (progress); )
269  C(1, float zh = size.y * (1.0 - progress); )
270  C(1, vec4 a = texture(a_images[idx], pos); )
271  C(1, vec4 b = texture(b_images[idx], pos); )
272  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x > zw) ? a : b); )
273  C(0, } )
274 };
275 
276 static const char transition_wipebl[] = {
277  C(0, void transition(int idx, ivec2 pos, float progress) )
278  C(0, { )
279  C(1, ivec2 size = imageSize(output_images[idx]); )
280  C(1, float zw = size.x * (1.0 - progress); )
281  C(1, float zh = size.y * (progress); )
282  C(1, vec4 a = texture(a_images[idx], pos); )
283  C(1, vec4 b = texture(b_images[idx], pos); )
284  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x <= zw) ? a : b); )
285  C(0, } )
286 };
287 
288 static const char transition_wipebr[] = {
289  C(0, void transition(int idx, ivec2 pos, float progress) )
290  C(0, { )
291  C(1, ivec2 size = imageSize(output_images[idx]); )
292  C(1, float zw = size.x * (progress); )
293  C(1, float zh = size.y * (progress); )
294  C(1, vec4 a = texture(a_images[idx], pos); )
295  C(1, vec4 b = texture(b_images[idx], pos); )
296  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x > zw) ? a : b); )
297  C(0, } )
298 };
299 
300 static const char* transitions_map[NB_TRANSITIONS] = {
301  [FADE] = transition_fade,
302  [WIPELEFT] = transition_wipeleft,
303  [WIPERIGHT] = transition_wiperight,
304  [WIPEUP] = transition_wipeup,
305  [WIPEDOWN] = transition_wipedown,
306  [SLIDEDOWN] = transition_slidedown,
307  [SLIDEUP] = transition_slideup,
308  [SLIDELEFT] = transition_slideleft,
309  [SLIDERIGHT] = transition_slideright,
310  [CIRCLEOPEN] = transition_circleopen,
311  [CIRCLECLOSE] = transition_circleclose,
312  [DISSOLVE] = transition_dissolve,
313  [PIXELIZE] = transition_pixelize,
314  [WIPETL] = transition_wipetl,
315  [WIPETR] = transition_wipetr,
316  [WIPEBL] = transition_wipebl,
317  [WIPEBR] = transition_wipebr,
318 };
319 
320 static av_cold int init_vulkan(AVFilterContext *avctx)
321 {
322  int err = 0;
323  uint8_t *spv_data;
324  size_t spv_len;
325  void *spv_opaque = NULL;
326  XFadeVulkanContext *s = avctx->priv;
327  FFVulkanContext *vkctx = &s->vkctx;
328  const int planes = av_pix_fmt_count_planes(s->vkctx.output_format);
329  FFVkSPIRVShader *shd = &s->shd;
330  FFVkSPIRVCompiler *spv;
331  FFVulkanDescriptorSetBinding *desc;
332 
333  spv = ff_vk_spirv_init();
334  if (!spv) {
335  av_log(avctx, AV_LOG_ERROR, "Unable to initialize SPIR-V compiler!\n");
336  return AVERROR_EXTERNAL;
337  }
338 
339  ff_vk_qf_init(vkctx, &s->qf, VK_QUEUE_COMPUTE_BIT);
340  RET(ff_vk_exec_pool_init(vkctx, &s->qf, &s->e, s->qf.nb_queues*4, 0, 0, 0, NULL));
341  RET(ff_vk_init_sampler(vkctx, &s->sampler, 1, VK_FILTER_NEAREST));
342  RET(ff_vk_shader_init(&s->pl, &s->shd, "xfade_compute",
343  VK_SHADER_STAGE_COMPUTE_BIT, 0));
344 
345  ff_vk_shader_set_compute_sizes(&s->shd, 32, 32, 1);
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),
368  .mem_quali = "writeonly",
369  .dimensions = 2,
370  .elems = planes,
371  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
372  },
373  };
374 
375  RET(ff_vk_pipeline_descriptor_set_add(vkctx, &s->pl, 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_add_push_constant(&s->pl, 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(spv, avctx, shd, &spv_data, &spv_len, "main",
397  &spv_opaque));
398  RET(ff_vk_shader_create(vkctx, shd, spv_data, spv_len, "main"));
399 
400  RET(ff_vk_init_compute_pipeline(vkctx, &s->pl, shd));
401  RET(ff_vk_exec_pipeline_register(vkctx, &s->e, &s->pl));
402 
403  s->initialized = 1;
404 
405 fail:
406  if (spv_opaque)
407  spv->free_shader(spv, &spv_opaque);
408  if (spv)
409  spv->uninit(&spv);
410 
411  return err;
412 }
413 
414 static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
415 {
416  int err;
417  AVFilterLink *outlink = avctx->outputs[0];
418  XFadeVulkanContext *s = avctx->priv;
419  float progress;
420 
421  AVFrame *output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
422  if (!output) {
423  err = AVERROR(ENOMEM);
424  goto fail;
425  }
426 
427  if (!s->initialized) {
428  AVHWFramesContext *a_fc = (AVHWFramesContext*)frame_a->hw_frames_ctx->data;
429  AVHWFramesContext *b_fc = (AVHWFramesContext*)frame_b->hw_frames_ctx->data;
430  if (a_fc->sw_format != b_fc->sw_format) {
431  av_log(avctx, AV_LOG_ERROR,
432  "Currently the sw format of the first input needs to match the second!\n");
433  return AVERROR(EINVAL);
434  }
435  RET(init_vulkan(avctx));
436  }
437 
438  RET(av_frame_copy_props(output, frame_a));
439  output->pts = s->pts;
440 
441  progress = av_clipf((float)(s->pts - s->start_pts) / s->duration_pts,
442  0.f, 1.f);
443 
444  RET(ff_vk_filter_process_Nin(&s->vkctx, &s->e, &s->pl, output,
445  (AVFrame *[]){ frame_a, frame_b }, 2, s->sampler,
446  &(XFadeParameters){ progress }, sizeof(XFadeParameters)));
447 
448  return ff_filter_frame(outlink, output);
449 
450 fail:
451  av_frame_free(&output);
452  return err;
453 }
454 
455 static int config_props_output(AVFilterLink *outlink)
456 {
457  int err;
458  AVFilterContext *avctx = outlink->src;
459  XFadeVulkanContext *s = avctx->priv;
460  AVFilterLink *inlink_a = avctx->inputs[IN_A];
461  AVFilterLink *inlink_b = avctx->inputs[IN_B];
462  FilterLink *il = ff_filter_link(inlink_a);
463  FilterLink *ol = ff_filter_link(outlink);
464 
465  if (inlink_a->w != inlink_b->w || inlink_a->h != inlink_b->h) {
466  av_log(avctx, AV_LOG_ERROR, "First input link %s parameters "
467  "(size %dx%d) do not match the corresponding "
468  "second input link %s parameters (size %dx%d)\n",
469  avctx->input_pads[IN_A].name, inlink_a->w, inlink_a->h,
470  avctx->input_pads[IN_B].name, inlink_b->w, inlink_b->h);
471  return AVERROR(EINVAL);
472  }
473 
474  if (inlink_a->time_base.num != inlink_b->time_base.num ||
475  inlink_a->time_base.den != inlink_b->time_base.den) {
476  av_log(avctx, AV_LOG_ERROR, "First input link %s timebase "
477  "(%d/%d) does not match the corresponding "
478  "second input link %s timebase (%d/%d)\n",
479  avctx->input_pads[IN_A].name, inlink_a->time_base.num, inlink_a->time_base.den,
480  avctx->input_pads[IN_B].name, inlink_b->time_base.num, inlink_b->time_base.den);
481  return AVERROR(EINVAL);
482  }
483 
484  s->start_pts = s->inputs_offset_pts = AV_NOPTS_VALUE;
485 
486  outlink->time_base = inlink_a->time_base;
487  ol->frame_rate = il->frame_rate;
488  outlink->sample_aspect_ratio = inlink_a->sample_aspect_ratio;
489 
490  if (s->duration)
491  s->duration_pts = av_rescale_q(s->duration, AV_TIME_BASE_Q, inlink_a->time_base);
492  RET(ff_vk_filter_config_output(outlink));
493 
494 fail:
495  return err;
496 }
497 
498 static int forward_frame(XFadeVulkanContext *s,
499  AVFilterLink *inlink, AVFilterLink *outlink)
500 {
501  int64_t status_pts;
502  int ret = 0, status;
503  AVFrame *frame = NULL;
504 
505  ret = ff_inlink_consume_frame(inlink, &frame);
506  if (ret < 0)
507  return ret;
508 
509  if (ret > 0) {
510  // If we do not have an offset yet, it's because we
511  // never got a first input. Just offset to 0
512  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
513  s->inputs_offset_pts = -frame->pts;
514 
515  // We got a frame, nothing to do other than adjusting the timestamp
516  frame->pts += s->inputs_offset_pts;
517  return ff_filter_frame(outlink, frame);
518  }
519 
520  // Forward status with our timestamp
521  if (ff_inlink_acknowledge_status(inlink, &status, &status_pts)) {
522  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
523  s->inputs_offset_pts = -status_pts;
524 
525  ff_outlink_set_status(outlink, status, status_pts + s->inputs_offset_pts);
526  return 0;
527  }
528 
529  // No frame available, request one if needed
530  if (ff_outlink_frame_wanted(outlink))
531  ff_inlink_request_frame(inlink);
532 
533  return 0;
534 }
535 
536 static int activate(AVFilterContext *avctx)
537 {
538  XFadeVulkanContext *s = avctx->priv;
539  AVFilterLink *in_a = avctx->inputs[IN_A];
540  AVFilterLink *in_b = avctx->inputs[IN_B];
541  AVFilterLink *outlink = avctx->outputs[0];
542  int64_t status_pts;
543 
544  FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, avctx);
545 
546  // Check if we already transitioned or IN_A ended prematurely,
547  // in which case just forward the frames from IN_B with adjusted
548  // timestamps until EOF.
549  if (s->status[IN_A] && !s->status[IN_B])
550  return forward_frame(s, in_b, outlink);
551 
552  // We did not finish transitioning yet and the first stream
553  // did not end either, so check if there are more frames to consume.
554  if (ff_inlink_check_available_frame(in_a)) {
555  AVFrame *peeked_frame = ff_inlink_peek_frame(in_a, 0);
556  s->pts = peeked_frame->pts;
557 
558  if (s->start_pts == AV_NOPTS_VALUE)
559  s->start_pts =
560  s->pts + av_rescale_q(s->offset, AV_TIME_BASE_Q, in_a->time_base);
561 
562  // Check if we are not yet transitioning, in which case
563  // just request and forward the input frame.
564  if (s->start_pts > s->pts) {
565  AVFrame *frame_a = NULL;
566  s->passthrough = 1;
567  ff_inlink_consume_frame(in_a, &frame_a);
568  return ff_filter_frame(outlink, frame_a);
569  }
570  s->passthrough = 0;
571 
572  // We are transitioning, so we need a frame from IN_B
573  if (ff_inlink_check_available_frame(in_b)) {
574  int ret;
575  AVFrame *frame_a = NULL, *frame_b = NULL;
576  ff_inlink_consume_frame(avctx->inputs[IN_A], &frame_a);
577  ff_inlink_consume_frame(avctx->inputs[IN_B], &frame_b);
578 
579  // Calculate PTS offset to first input
580  if (s->inputs_offset_pts == AV_NOPTS_VALUE)
581  s->inputs_offset_pts = s->pts - frame_b->pts;
582 
583  // Check if we finished transitioning, in which case we
584  // report back EOF to IN_A as it is no longer needed.
585  if (s->pts - s->start_pts > s->duration_pts) {
586  s->status[IN_A] = AVERROR_EOF;
587  ff_inlink_set_status(in_a, AVERROR_EOF);
588  s->passthrough = 1;
589  }
590  ret = xfade_frame(avctx, frame_a, frame_b);
591  av_frame_free(&frame_a);
592  av_frame_free(&frame_b);
593  return ret;
594  }
595 
596  // We did not get a frame from IN_B, check its status.
597  if (ff_inlink_acknowledge_status(in_b, &s->status[IN_B], &status_pts)) {
598  // We should transition, but IN_B is EOF so just report EOF output now.
599  ff_outlink_set_status(outlink, s->status[IN_B], s->pts);
600  return 0;
601  }
602 
603  // We did not get a frame for IN_B but no EOF either, so just request more.
604  if (ff_outlink_frame_wanted(outlink)) {
605  ff_inlink_request_frame(in_b);
606  return 0;
607  }
608  }
609 
610  // We did not get a frame from IN_A, check its status.
611  if (ff_inlink_acknowledge_status(in_a, &s->status[IN_A], &status_pts)) {
612  // No more frames from IN_A, do not report EOF though, we will just
613  // forward the IN_B frames in the next activate calls.
614  s->passthrough = 1;
615  ff_filter_set_ready(avctx, 100);
616  return 0;
617  }
618 
619  // We have no frames yet from IN_A and no EOF, so request some.
620  if (ff_outlink_frame_wanted(outlink)) {
621  ff_inlink_request_frame(in_a);
622  return 0;
623  }
624 
625  return FFERROR_NOT_READY;
626 }
627 
628 static av_cold void uninit(AVFilterContext *avctx)
629 {
630  XFadeVulkanContext *s = avctx->priv;
631  FFVulkanContext *vkctx = &s->vkctx;
632  FFVulkanFunctions *vk = &vkctx->vkfn;
633 
634  ff_vk_exec_pool_free(vkctx, &s->e);
635  ff_vk_pipeline_free(vkctx, &s->pl);
636  ff_vk_shader_free(vkctx, &s->shd);
637 
638  if (s->sampler)
639  vk->DestroySampler(vkctx->hwctx->act_dev, s->sampler,
640  vkctx->hwctx->alloc);
641 
642  ff_vk_uninit(&s->vkctx);
643 
644  s->initialized = 0;
645 }
646 
647 static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
648 {
649  XFadeVulkanContext *s = inlink->dst->priv;
650 
651  return s->passthrough ?
652  ff_null_get_video_buffer (inlink, w, h) :
653  ff_default_get_video_buffer(inlink, w, h);
654 }
655 
656 #define OFFSET(x) offsetof(XFadeVulkanContext, x)
657 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
658 
659 static const AVOption xfade_vulkan_options[] = {
660  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=FADE}, 0, NB_TRANSITIONS-1, FLAGS, .unit = "transition" },
661  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, .unit = "transition" },
662  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, .unit = "transition" },
663  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, .unit = "transition" },
664  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, .unit = "transition" },
665  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, .unit = "transition" },
666  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, .unit = "transition" },
667  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, .unit = "transition" },
668  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, .unit = "transition" },
669  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, .unit = "transition" },
670  { "circleopen", "circleopen transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLEOPEN}, 0, 0, FLAGS, .unit = "transition" },
671  { "circleclose", "circleclose transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLECLOSE}, 0, 0, FLAGS, .unit = "transition" },
672  { "dissolve", "dissolve transition", 0, AV_OPT_TYPE_CONST, {.i64=DISSOLVE}, 0, 0, FLAGS, .unit = "transition" },
673  { "pixelize", "pixelize transition", 0, AV_OPT_TYPE_CONST, {.i64=PIXELIZE}, 0, 0, FLAGS, .unit = "transition" },
674  { "wipetl", "wipe top left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETL}, 0, 0, FLAGS, .unit = "transition" },
675  { "wipetr", "wipe top right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETR}, 0, 0, FLAGS, .unit = "transition" },
676  { "wipebl", "wipe bottom left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBL}, 0, 0, FLAGS, .unit = "transition" },
677  { "wipebr", "wipe bottom right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBR}, 0, 0, FLAGS, .unit = "transition" },
678  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
679  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
680  { NULL }
681 };
682 
683 AVFILTER_DEFINE_CLASS(xfade_vulkan);
684 
685 static const AVFilterPad xfade_vulkan_inputs[] = {
686  {
687  .name = "main",
688  .type = AVMEDIA_TYPE_VIDEO,
689  .get_buffer.video = &get_video_buffer,
690  .config_props = &ff_vk_filter_config_input,
691  },
692  {
693  .name = "xfade",
694  .type = AVMEDIA_TYPE_VIDEO,
695  .get_buffer.video = &get_video_buffer,
696  .config_props = &ff_vk_filter_config_input,
697  },
698 };
699 
700 static const AVFilterPad xfade_vulkan_outputs[] = {
701  {
702  .name = "default",
703  .type = AVMEDIA_TYPE_VIDEO,
704  .config_props = &config_props_output,
705  },
706 };
707 
708 const AVFilter ff_vf_xfade_vulkan = {
709  .name = "xfade_vulkan",
710  .description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
711  .priv_size = sizeof(XFadeVulkanContext),
712  .init = &ff_vk_filter_init,
713  .uninit = &uninit,
714  .activate = &activate,
715  FILTER_INPUTS(xfade_vulkan_inputs),
716  FILTER_OUTPUTS(xfade_vulkan_outputs),
717  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VULKAN),
718  .priv_class = &xfade_vulkan_class,
719  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
720  .flags = AVFILTER_FLAG_HWDEVICE,
721 };
IN_A
#define IN_A
Definition: vf_xfade_vulkan.c:27
WIPETR
@ WIPETR
Definition: vf_xfade_vulkan.c:84
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:71
ff_vk_pipeline_free
void ff_vk_pipeline_free(FFVulkanContext *s, FFVulkanPipeline *pl)
Definition: vulkan.c:1809
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:226
xfade_frame
static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
Definition: vf_xfade_vulkan.c:414
transition_wipeup
static const char transition_wipeup[]
Definition: vf_xfade_vulkan.c:121
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1023
AVERROR_EOF
#define AVERROR_EOF
End of file.
Definition: error.h:57
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
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:165
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:160
XFadeVulkanContext::start_pts
int64_t start_pts
Definition: vf_xfade_vulkan.c:50
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:374
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:486
ff_vk_filter_init
int ff_vk_filter_init(AVFilterContext *avctx)
General lavfi IO functions.
Definition: vulkan_filter.c:243
w
uint8_t w
Definition: llviddspenc.c:38
frand
static float frand(int x, int y)
Definition: vf_deband.c:116
XFadeVulkanContext::status
int status[IN_NB]
Definition: vf_xfade_vulkan.c:66
ff_vk_filter_process_Nin
int ff_vk_filter_process_Nin(FFVulkanContext *vkctx, FFVkExecPool *e, FFVulkanPipeline *pl, 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:409
transitions_map
static const char * transitions_map[NB_TRANSITIONS]
Definition: vf_xfade_vulkan.c:300
XFadeVulkanContext::transition
int transition
Definition: vf_xfade_vulkan.c:38
ff_vk_shader_create
int ff_vk_shader_create(FFVulkanContext *s, FFVkSPIRVShader *shd, uint8_t *spirv, size_t spirv_size, const char *entrypoint)
Definition: vulkan.c:1378
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:647
CIRCLEOPEN
@ CIRCLEOPEN
Definition: vf_xfade_vulkan.c:79
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:110
forward_frame
static int forward_frame(XFadeVulkanContext *s, AVFilterLink *inlink, AVFilterLink *outlink)
Definition: vf_xfade_vulkan.c:498
XFadeVulkanContext::e
FFVkExecPool e
Definition: vf_xfade_vulkan.c:44
transition_wipebl
static const char transition_wipebl[]
Definition: vf_xfade_vulkan.c:276
XFadeVulkanContext::qf
FFVkQueueFamilyCtx qf
Definition: vf_xfade_vulkan.c:45
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:1838
transition_wipebr
static const char transition_wipebr[]
Definition: vf_xfade_vulkan.c:288
FFVkSPIRVCompiler::uninit
void(* uninit)(struct FFVkSPIRVCompiler **ctx)
Definition: vulkan_spirv.h:33
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:205
ff_vk_shader_set_compute_sizes
void ff_vk_shader_set_compute_sizes(FFVkSPIRVShader *shd, int x, int y, int z)
Definition: vulkan.c:1336
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:708
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:1451
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:85
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3005
SLIDELEFT
@ SLIDELEFT
Definition: vf_xfade_vulkan.c:77
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:173
AVVulkanDeviceContext::alloc
const VkAllocationCallbacks * alloc
Custom memory allocator, else NULL.
Definition: hwcontext_vulkan.h:63
ff_vk_add_push_constant
int ff_vk_add_push_constant(FFVulkanPipeline *pl, int offset, int size, VkShaderStageFlagBits stage)
Add/update push constants for execution.
Definition: vulkan.c:1106
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:75
vulkan_filter.h
SLIDEUP
@ SLIDEUP
Definition: vf_xfade_vulkan.c:76
XFadeVulkanContext::pts
int64_t pts
Definition: vf_xfade_vulkan.c:59
AVRational::num
int num
Numerator.
Definition: rational.h:59
XFadeVulkanContext::passthrough
int passthrough
Definition: vf_xfade_vulkan.c:64
SHADER_FRAND_FUNC
#define SHADER_FRAND_FUNC
Definition: vf_xfade_vulkan.c:220
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:212
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
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
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:1578
s
#define s(width, name)
Definition: cbs_vp9.c:198
floor
static __device__ float floor(float a)
Definition: cuda_runtime.h:173
filters.h
transition_wipedown
static const char transition_wipedown[]
Definition: vf_xfade_vulkan.c:132
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
FFVkSPIRVCompiler::compile_shader
int(* compile_shader)(struct FFVkSPIRVCompiler *ctx, void *avctx, struct FFVkSPIRVShader *shd, uint8_t **data, size_t *size, const char *entrypoint, void **opaque)
Definition: vulkan_spirv.h:29
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:1492
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:711
GLSLD
#define GLSLD(D)
Definition: vulkan.h:59
transition_pixelize
static const char transition_pixelize[]
Definition: vf_xfade_vulkan.c:237
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:536
transition_circleopen
static const char transition_circleopen[]
Definition: vf_xfade_vulkan.c:204
ff_vk_filter_config_output
int ff_vk_filter_config_output(AVFilterLink *outlink)
Definition: vulkan_filter.c:219
transition_wipetl
static const char transition_wipetl[]
Definition: vf_xfade_vulkan.c:252
transition_fade
static const char transition_fade[]
Definition: vf_xfade_vulkan.c:90
xfade_vulkan_outputs
static const AVFilterPad xfade_vulkan_outputs[]
Definition: vf_xfade_vulkan.c:700
ff_vk_init_compute_pipeline
int ff_vk_init_compute_pipeline(FFVulkanContext *s, FFVulkanPipeline *pl, FFVkSPIRVShader *shd)
Definition: vulkan.c:1750
xfade_vulkan_options
static const AVOption xfade_vulkan_options[]
Definition: vf_xfade_vulkan.c:659
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:278
av_clipf
av_clipf
Definition: af_crystalizer.c:122
FFVulkanContext
Definition: vulkan.h:229
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:1398
FFVulkanPipeline
Definition: vulkan.h:132
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
PIXELIZE
@ PIXELIZE
Definition: vf_xfade_vulkan.c:82
DISSOLVE
@ DISSOLVE
Definition: vf_xfade_vulkan.c:81
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_shader_init
int ff_vk_shader_init(FFVulkanPipeline *pl, FFVkSPIRVShader *shd, const char *name, VkShaderStageFlags stage, uint32_t required_subgroup_size)
Shader management.
Definition: vulkan.c:1310
XFadeVulkanContext::pl
FFVulkanPipeline pl
Definition: vf_xfade_vulkan.c:43
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:366
FLAGS
#define FLAGS
Definition: vf_xfade_vulkan.c:657
ff_inlink_set_status
void ff_inlink_set_status(AVFilterLink *link, int status)
Set the status on an input link.
Definition: avfilter.c:1587
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:1420
FFVulkanDescriptorSetBinding
Definition: vulkan.h:83
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:189
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:110
XFadeVulkanContext::sampler
VkSampler sampler
Definition: vf_xfade_vulkan.c:47
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:53
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:27
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:685
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
XFadeVulkanContext::shd
FFVkSPIRVShader shd
Definition: vf_xfade_vulkan.c:46
NB_TRANSITIONS
@ NB_TRANSITIONS
Definition: vf_xfade_vulkan.c:87
DUP_SAMPLER
#define DUP_SAMPLER(x)
Definition: vulkan.h:73
ff_vk_shader_rep_fmt
const char * ff_vk_shader_rep_fmt(enum AVPixelFormat pixfmt)
Returns the format to use for images in shaders.
Definition: vulkan.c:1170
vulkan_spirv.h
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
CIRCLECLOSE
@ CIRCLECLOSE
Definition: vf_xfade_vulkan.c:80
GLSLF
#define GLSLF(N, S,...)
Definition: vulkan.h:54
XFadeParameters
Definition: vf_xfade_vulkan.c:31
FFVkSPIRVCompiler::free_shader
void(* free_shader)(struct FFVkSPIRVCompiler *ctx, void **opaque)
Definition: vulkan_spirv.h:32
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
ff_vk_pipeline_descriptor_set_add
int ff_vk_pipeline_descriptor_set_add(FFVulkanContext *s, FFVulkanPipeline *pl, FFVkSPIRVShader *shd, FFVulkanDescriptorSetBinding *desc, int nb, int singular, int print_to_shader_only)
Add descriptor to a pipeline.
Definition: vulkan.c:1428
ret
ret
Definition: filter_design.txt:187
WIPERIGHT
@ WIPERIGHT
Definition: vf_xfade_vulkan.c:72
FADE
@ FADE
Definition: vf_xfade_vulkan.c:70
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:233
SHADER_SLIDE_COMMON
#define SHADER_SLIDE_COMMON
Definition: vf_xfade_vulkan.c:143
FFVkExecPool
Definition: vulkan.h:211
pos
unsigned int pos
Definition: spdifenc.c:414
XFadeVulkanContext::duration_pts
int64_t duration_pts
Definition: vf_xfade_vulkan.c:56
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:725
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:181
FFVkSPIRVShader
Definition: vulkan.h:75
AVRational::den
int den
Denominator.
Definition: rational.h:60
WIPEBR
@ WIPEBR
Definition: vf_xfade_vulkan.c:86
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:455
init_vulkan
static av_cold int init_vulkan(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:320
transition_wipetr
static const char transition_wipetr[]
Definition: vf_xfade_vulkan.c:264
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
GLSLC
#define GLSLC(N, S)
Definition: vulkan.h:44
ff_vk_filter_config_input
int ff_vk_filter_config_input(AVFilterLink *inlink)
Definition: vulkan_filter.c:186
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
FFVulkanContext::hwctx
AVVulkanDeviceContext * hwctx
Definition: vulkan.h:257
WIPEDOWN
@ WIPEDOWN
Definition: vf_xfade_vulkan.c:74
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:83
ff_vk_init_sampler
int ff_vk_init_sampler(FFVulkanContext *s, VkSampler *sampler, int unnorm_coords, VkFilter filt)
Create a sampler.
Definition: vulkan.c:1126
SLIDERIGHT
@ SLIDERIGHT
Definition: vf_xfade_vulkan.c:78
ff_vk_exec_pipeline_register
int ff_vk_exec_pipeline_register(FFVulkanContext *s, FFVkExecPool *pool, FFVulkanPipeline *pl)
Register a pipeline with an exec pool.
Definition: vulkan.c:1542
XFadeParameters::progress
float progress
Definition: vf_xfade_vulkan.c:32
transition_wipeleft
static const char transition_wipeleft[]
Definition: vf_xfade_vulkan.c:99
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:73
transition_slidedown
static const char transition_slidedown[]
Definition: vf_xfade_vulkan.c:157
ff_vk_shader_free
void ff_vk_shader_free(FFVulkanContext *s, FFVkSPIRVShader *shd)
Definition: vulkan.c:1369
uninit
static av_cold void uninit(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:628
OFFSET
#define OFFSET(x)
Definition: vf_xfade_vulkan.c:656
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
RET
#define RET(x)
Definition: vulkan.h:67
FFVulkanFunctions
Definition: vulkan_functions.h:249
planes
static const struct @450 planes[]
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:237
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 Sat Sep 14 2024 19:22:37 for FFmpeg by   doxygen 1.8.17