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 "internal.h"
26 #include "video.h"
27 
28 #define IN_A 0
29 #define IN_B 1
30 #define IN_NB 2
31 
32 typedef struct XFadeParameters {
33  float progress;
34 } XFadeParameters;
35 
36 typedef struct XFadeVulkanContext {
37  FFVulkanContext vkctx;
38 
39  int transition;
40  int64_t duration;
41  int64_t offset;
42 
43  int initialized;
44  FFVulkanPipeline pl;
45  FFVkExecPool e;
46  FFVkQueueFamilyCtx qf;
47  FFVkSPIRVShader shd;
48  VkSampler sampler;
49 
50  // PTS when the fade should start (in IN_A timebase)
51  int64_t start_pts;
52 
53  // PTS offset between IN_A and IN_B
54  int64_t inputs_offset_pts;
55 
56  // Duration of the transition
57  int64_t duration_pts;
58 
59  // Current PTS of the first input (IN_A)
60  int64_t pts;
61 
62  // If frames are currently just passed through
63  // unmodified, like before and after the actual
64  // transition.
65  int passthrough;
66 
67  int status[IN_NB];
68 } XFadeVulkanContext;
69 
70 enum XFadeTransitions {
71  FADE,
72  WIPELEFT,
73  WIPERIGHT,
74  WIPEUP,
75  WIPEDOWN,
76  SLIDEDOWN,
77  SLIDEUP,
78  SLIDELEFT,
79  SLIDERIGHT,
80  CIRCLEOPEN,
81  CIRCLECLOSE,
82  DISSOLVE,
83  PIXELIZE,
84  WIPETL,
85  WIPETR,
86  WIPEBL,
87  WIPEBR,
88  NB_TRANSITIONS,
89 };
90 
91 static const char transition_fade[] = {
92  C(0, void transition(int idx, ivec2 pos, float progress) )
93  C(0, { )
94  C(1, vec4 a = texture(a_images[idx], pos); )
95  C(1, vec4 b = texture(b_images[idx], pos); )
96  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
97  C(0, } )
98 };
99 
100 static const char transition_wipeleft[] = {
101  C(0, void transition(int idx, ivec2 pos, float progress) )
102  C(0, { )
103  C(1, ivec2 size = imageSize(output_images[idx]); )
104  C(1, int s = int(size.x * (1.0 - progress)); )
105  C(1, vec4 a = texture(a_images[idx], pos); )
106  C(1, vec4 b = texture(b_images[idx], pos); )
107  C(1, imageStore(output_images[idx], pos, pos.x > s ? b : a); )
108  C(0, } )
109 };
110 
111 static const char transition_wiperight[] = {
112  C(0, void transition(int idx, ivec2 pos, float progress) )
113  C(0, { )
114  C(1, ivec2 size = imageSize(output_images[idx]); )
115  C(1, int s = int(size.x * progress); )
116  C(1, vec4 a = texture(a_images[idx], pos); )
117  C(1, vec4 b = texture(b_images[idx], pos); )
118  C(1, imageStore(output_images[idx], pos, pos.x > s ? a : b); )
119  C(0, } )
120 };
121 
122 static const char transition_wipeup[] = {
123  C(0, void transition(int idx, ivec2 pos, float progress) )
124  C(0, { )
125  C(1, ivec2 size = imageSize(output_images[idx]); )
126  C(1, int s = int(size.y * (1.0 - progress)); )
127  C(1, vec4 a = texture(a_images[idx], pos); )
128  C(1, vec4 b = texture(b_images[idx], pos); )
129  C(1, imageStore(output_images[idx], pos, pos.y > s ? b : a); )
130  C(0, } )
131 };
132 
133 static const char transition_wipedown[] = {
134  C(0, void transition(int idx, ivec2 pos, float progress) )
135  C(0, { )
136  C(1, ivec2 size = imageSize(output_images[idx]); )
137  C(1, int s = int(size.y * progress); )
138  C(1, vec4 a = texture(a_images[idx], pos); )
139  C(1, vec4 b = texture(b_images[idx], pos); )
140  C(1, imageStore(output_images[idx], pos, pos.y > s ? a : b); )
141  C(0, } )
142 };
143 
144 #define SHADER_SLIDE_COMMON \
145  C(0, void slide(int idx, ivec2 pos, float progress, ivec2 direction) ) \
146  C(0, { ) \
147  C(1, ivec2 size = imageSize(output_images[idx]); ) \
148  C(1, ivec2 pi = ivec2(progress * size); ) \
149  C(1, ivec2 p = pos + pi * direction; ) \
150  C(1, ivec2 f = p % size; ) \
151  C(1, f = f + size * ivec2(f.x < 0, f.y < 0); ) \
152  C(1, vec4 a = texture(a_images[idx], f); ) \
153  C(1, vec4 b = texture(b_images[idx], f); ) \
154  C(1, vec4 r = (p.y >= 0 && p.x >= 0 && size.y > p.y && size.x > p.x) ? a : b; ) \
155  C(1, imageStore(output_images[idx], pos, r); ) \
156  C(0, } )
157 
158 static const char transition_slidedown[] = {
159  SHADER_SLIDE_COMMON
160  C(0, void transition(int idx, ivec2 pos, float progress) )
161  C(0, { )
162  C(1, slide(idx, pos, progress, ivec2(0, -1)); )
163  C(0, } )
164 };
165 
166 static const char transition_slideup[] = {
167  SHADER_SLIDE_COMMON
168  C(0, void transition(int idx, ivec2 pos, float progress) )
169  C(0, { )
170  C(1, slide(idx, pos, progress, ivec2(0, +1)); )
171  C(0, } )
172 };
173 
174 static const char transition_slideleft[] = {
175  SHADER_SLIDE_COMMON
176  C(0, void transition(int idx, ivec2 pos, float progress) )
177  C(0, { )
178  C(1, slide(idx, pos, progress, ivec2(+1, 0)); )
179  C(0, } )
180 };
181 
182 static const char transition_slideright[] = {
183  SHADER_SLIDE_COMMON
184  C(0, void transition(int idx, ivec2 pos, float progress) )
185  C(0, { )
186  C(1, slide(idx, pos, progress, ivec2(-1, 0)); )
187  C(0, } )
188 };
189 
190 #define SHADER_CIRCLE_COMMON \
191  C(0, void circle(int idx, ivec2 pos, float progress, bool open) ) \
192  C(0, { ) \
193  C(1, const ivec2 half_size = imageSize(output_images[idx]) / 2; ) \
194  C(1, const float z = dot(half_size, half_size); ) \
195  C(1, float p = ((open ? (1.0 - progress) : progress) - 0.5) * 3.0; ) \
196  C(1, ivec2 dsize = pos - half_size; ) \
197  C(1, float sm = dot(dsize, dsize) / z + p; ) \
198  C(1, vec4 a = texture(a_images[idx], pos); ) \
199  C(1, vec4 b = texture(b_images[idx], pos); ) \
200  C(1, imageStore(output_images[idx], pos, \
201  mix(open ? b : a, open ? a : b, \
202  smoothstep(0.f, 1.f, sm))); ) \
203  C(0, } )
204 
205 static const char transition_circleopen[] = {
206  SHADER_CIRCLE_COMMON
207  C(0, void transition(int idx, ivec2 pos, float progress) )
208  C(0, { )
209  C(1, circle(idx, pos, progress, true); )
210  C(0, } )
211 };
212 
213 static const char transition_circleclose[] = {
214  SHADER_CIRCLE_COMMON
215  C(0, void transition(int idx, ivec2 pos, float progress) )
216  C(0, { )
217  C(1, circle(idx, pos, progress, false); )
218  C(0, } )
219 };
220 
221 #define SHADER_FRAND_FUNC \
222  C(0, float frand(vec2 v) ) \
223  C(0, { ) \
224  C(1, return fract(sin(dot(v, vec2(12.9898, 78.233))) * 43758.545); ) \
225  C(0, } )
226 
227 static const char transition_dissolve[] = {
228  SHADER_FRAND_FUNC
229  C(0, void transition(int idx, ivec2 pos, float progress) )
230  C(0, { )
231  C(1, float sm = frand(pos) * 2.0 + (1.0 - progress) * 2.0 - 1.5; )
232  C(1, vec4 a = texture(a_images[idx], pos); )
233  C(1, vec4 b = texture(b_images[idx], pos); )
234  C(1, imageStore(output_images[idx], pos, sm >= 0.5 ? a : b); )
235  C(0, } )
236 };
237 
238 static const char transition_pixelize[] = {
239  C(0, void transition(int idx, ivec2 pos, float progress) )
240  C(0, { )
241  C(1, ivec2 size = imageSize(output_images[idx]); )
242  C(1, float d = min(progress, 1.0 - progress); )
243  C(1, float dist = ceil(d * 50.0) / 50.0; )
244  C(1, float sq = 2.0 * dist * min(size.x, size.y) / 20.0; )
245  C(1, float sx = dist > 0.0 ? min((floor(pos.x / sq) + 0.5) * sq, size.x - 1) : pos.x; )
246  C(1, float sy = dist > 0.0 ? min((floor(pos.y / sq) + 0.5) * sq, size.y - 1) : pos.y; )
247  C(1, vec4 a = texture(a_images[idx], vec2(sx, sy)); )
248  C(1, vec4 b = texture(b_images[idx], vec2(sx, sy)); )
249  C(1, imageStore(output_images[idx], pos, mix(a, b, progress)); )
250  C(0, } )
251 };
252 
253 static const char transition_wipetl[] = {
254  C(0, void transition(int idx, ivec2 pos, float progress) )
255  C(0, { )
256  C(1, ivec2 size = imageSize(output_images[idx]); )
257  C(1, float zw = size.x * (1.0 - progress); )
258  C(1, float zh = size.y * (1.0 - progress); )
259  C(1, vec4 a = texture(a_images[idx], pos); )
260  C(1, vec4 b = texture(b_images[idx], pos); )
261  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x <= zw) ? a : b); )
262  C(0, } )
263 };
264 
265 static const char transition_wipetr[] = {
266  C(0, void transition(int idx, ivec2 pos, float progress) )
267  C(0, { )
268  C(1, ivec2 size = imageSize(output_images[idx]); )
269  C(1, float zw = size.x * (progress); )
270  C(1, float zh = size.y * (1.0 - progress); )
271  C(1, vec4 a = texture(a_images[idx], pos); )
272  C(1, vec4 b = texture(b_images[idx], pos); )
273  C(1, imageStore(output_images[idx], pos, (pos.y <= zh && pos.x > zw) ? a : b); )
274  C(0, } )
275 };
276 
277 static const char transition_wipebl[] = {
278  C(0, void transition(int idx, ivec2 pos, float progress) )
279  C(0, { )
280  C(1, ivec2 size = imageSize(output_images[idx]); )
281  C(1, float zw = size.x * (1.0 - progress); )
282  C(1, float zh = size.y * (progress); )
283  C(1, vec4 a = texture(a_images[idx], pos); )
284  C(1, vec4 b = texture(b_images[idx], pos); )
285  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x <= zw) ? a : b); )
286  C(0, } )
287 };
288 
289 static const char transition_wipebr[] = {
290  C(0, void transition(int idx, ivec2 pos, float progress) )
291  C(0, { )
292  C(1, ivec2 size = imageSize(output_images[idx]); )
293  C(1, float zw = size.x * (progress); )
294  C(1, float zh = size.y * (progress); )
295  C(1, vec4 a = texture(a_images[idx], pos); )
296  C(1, vec4 b = texture(b_images[idx], pos); )
297  C(1, imageStore(output_images[idx], pos, (pos.y > zh && pos.x > zw) ? a : b); )
298  C(0, } )
299 };
300 
301 static const char* transitions_map[NB_TRANSITIONS] = {
302  [FADE] = transition_fade,
303  [WIPELEFT] = transition_wipeleft,
304  [WIPERIGHT] = transition_wiperight,
305  [WIPEUP] = transition_wipeup,
306  [WIPEDOWN] = transition_wipedown,
307  [SLIDEDOWN] = transition_slidedown,
308  [SLIDEUP] = transition_slideup,
309  [SLIDELEFT] = transition_slideleft,
310  [SLIDERIGHT] = transition_slideright,
311  [CIRCLEOPEN] = transition_circleopen,
312  [CIRCLECLOSE] = transition_circleclose,
313  [DISSOLVE] = transition_dissolve,
314  [PIXELIZE] = transition_pixelize,
315  [WIPETL] = transition_wipetl,
316  [WIPETR] = transition_wipetr,
317  [WIPEBL] = transition_wipebl,
318  [WIPEBR] = transition_wipebr,
319 };
320 
321 static av_cold int init_vulkan(AVFilterContext *avctx)
322 {
323  int err = 0;
324  uint8_t *spv_data;
325  size_t spv_len;
326  void *spv_opaque = NULL;
327  XFadeVulkanContext *s = avctx->priv;
328  FFVulkanContext *vkctx = &s->vkctx;
329  const int planes = av_pix_fmt_count_planes(s->vkctx.output_format);
330  FFVkSPIRVShader *shd = &s->shd;
331  FFVkSPIRVCompiler *spv;
332  FFVulkanDescriptorSetBinding *desc;
333 
334  spv = ff_vk_spirv_init();
335  if (!spv) {
336  av_log(avctx, AV_LOG_ERROR, "Unable to initialize SPIR-V compiler!\n");
337  return AVERROR_EXTERNAL;
338  }
339 
340  ff_vk_qf_init(vkctx, &s->qf, VK_QUEUE_COMPUTE_BIT);
341  RET(ff_vk_exec_pool_init(vkctx, &s->qf, &s->e, s->qf.nb_queues*4, 0, 0, 0, NULL));
342  RET(ff_vk_init_sampler(vkctx, &s->sampler, 1, VK_FILTER_NEAREST));
343  RET(ff_vk_shader_init(&s->pl, &s->shd, "xfade_compute",
344  VK_SHADER_STAGE_COMPUTE_BIT, 0));
345 
346  ff_vk_shader_set_compute_sizes(&s->shd, 32, 32, 1);
347 
348  desc = (FFVulkanDescriptorSetBinding []) {
349  {
350  .name = "a_images",
351  .type = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER,
352  .dimensions = 2,
353  .elems = planes,
354  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
355  .samplers = DUP_SAMPLER(s->sampler),
356  },
357  {
358  .name = "b_images",
359  .type = VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER,
360  .dimensions = 2,
361  .elems = planes,
362  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
363  .samplers = DUP_SAMPLER(s->sampler),
364  },
365  {
366  .name = "output_images",
367  .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
368  .mem_layout = ff_vk_shader_rep_fmt(s->vkctx.output_format),
369  .mem_quali = "writeonly",
370  .dimensions = 2,
371  .elems = planes,
372  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
373  },
374  };
375 
376  RET(ff_vk_pipeline_descriptor_set_add(vkctx, &s->pl, shd, desc, 3, 0, 0));
377 
378  GLSLC(0, layout(push_constant, std430) uniform pushConstants { );
379  GLSLC(1, float progress; );
380  GLSLC(0, }; );
381 
382  ff_vk_add_push_constant(&s->pl, 0, sizeof(XFadeParameters),
383  VK_SHADER_STAGE_COMPUTE_BIT);
384 
385  // Add the right transition type function to the shader
386  GLSLD(transitions_map[s->transition]);
387 
388  GLSLC(0, void main() );
389  GLSLC(0, { );
390  GLSLC(1, ivec2 pos = ivec2(gl_GlobalInvocationID.xy); );
391  GLSLF(1, int planes = %i; ,planes);
392  GLSLC(1, for (int i = 0; i < planes; i++) { );
393  GLSLC(2, transition(i, pos, progress); );
394  GLSLC(1, } );
395  GLSLC(0, } );
396 
397  RET(spv->compile_shader(spv, avctx, shd, &spv_data, &spv_len, "main",
398  &spv_opaque));
399  RET(ff_vk_shader_create(vkctx, shd, spv_data, spv_len, "main"));
400 
401  RET(ff_vk_init_compute_pipeline(vkctx, &s->pl, shd));
402  RET(ff_vk_exec_pipeline_register(vkctx, &s->e, &s->pl));
403 
404  s->initialized = 1;
405 
406 fail:
407  if (spv_opaque)
408  spv->free_shader(spv, &spv_opaque);
409  if (spv)
410  spv->uninit(&spv);
411 
412  return err;
413 }
414 
415 static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
416 {
417  int err;
418  AVFilterLink *outlink = avctx->outputs[0];
419  XFadeVulkanContext *s = avctx->priv;
420  float progress;
421 
422  AVFrame *output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
423  if (!output) {
424  err = AVERROR(ENOMEM);
425  goto fail;
426  }
427 
428  if (!s->initialized) {
429  AVHWFramesContext *a_fc = (AVHWFramesContext*)frame_a->hw_frames_ctx->data;
430  AVHWFramesContext *b_fc = (AVHWFramesContext*)frame_b->hw_frames_ctx->data;
431  if (a_fc->sw_format != b_fc->sw_format) {
432  av_log(avctx, AV_LOG_ERROR,
433  "Currently the sw format of the first input needs to match the second!\n");
434  return AVERROR(EINVAL);
435  }
436  RET(init_vulkan(avctx));
437  }
438 
439  RET(av_frame_copy_props(output, frame_a));
440  output->pts = s->pts;
441 
442  progress = av_clipf((float)(s->pts - s->start_pts) / s->duration_pts,
443  0.f, 1.f);
444 
445  RET(ff_vk_filter_process_Nin(&s->vkctx, &s->e, &s->pl, output,
446  (AVFrame *[]){ frame_a, frame_b }, 2, s->sampler,
447  &(XFadeParameters){ progress }, sizeof(XFadeParameters)));
448 
449  return ff_filter_frame(outlink, output);
450 
451 fail:
452  av_frame_free(&output);
453  return err;
454 }
455 
456 static int config_props_output(AVFilterLink *outlink)
457 {
458  int err;
459  AVFilterContext *avctx = outlink->src;
460  XFadeVulkanContext *s = avctx->priv;
461  AVFilterLink *inlink_a = avctx->inputs[IN_A];
462  AVFilterLink *inlink_b = avctx->inputs[IN_B];
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  outlink->frame_rate = inlink_a->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_pipeline_free(vkctx, &s->pl);
635  ff_vk_shader_free(vkctx, &s->shd);
636 
637  if (s->sampler)
638  vk->DestroySampler(vkctx->hwctx->act_dev, s->sampler,
639  vkctx->hwctx->alloc);
640 
641  ff_vk_uninit(&s->vkctx);
642 
643  s->initialized = 0;
644 }
645 
646 static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
647 {
648  XFadeVulkanContext *s = inlink->dst->priv;
649 
650  return s->passthrough ?
651  ff_null_get_video_buffer (inlink, w, h) :
652  ff_default_get_video_buffer(inlink, w, h);
653 }
654 
655 #define OFFSET(x) offsetof(XFadeVulkanContext, x)
656 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
657 
658 static const AVOption xfade_vulkan_options[] = {
659  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=FADE}, 0, NB_TRANSITIONS-1, FLAGS, "transition" },
660  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, "transition" },
661  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, "transition" },
662  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, "transition" },
663  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, "transition" },
664  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, "transition" },
665  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, "transition" },
666  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, "transition" },
667  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, "transition" },
668  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, "transition" },
669  { "circleopen", "circleopen transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLEOPEN}, 0, 0, FLAGS, "transition" },
670  { "circleclose", "circleclose transition", 0, AV_OPT_TYPE_CONST, {.i64=CIRCLECLOSE}, 0, 0, FLAGS, "transition" },
671  { "dissolve", "dissolve transition", 0, AV_OPT_TYPE_CONST, {.i64=DISSOLVE}, 0, 0, FLAGS, "transition" },
672  { "pixelize", "pixelize transition", 0, AV_OPT_TYPE_CONST, {.i64=PIXELIZE}, 0, 0, FLAGS, "transition" },
673  { "wipetl", "wipe top left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETL}, 0, 0, FLAGS, "transition" },
674  { "wipetr", "wipe top right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPETR}, 0, 0, FLAGS, "transition" },
675  { "wipebl", "wipe bottom left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBL}, 0, 0, FLAGS, "transition" },
676  { "wipebr", "wipe bottom right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEBR}, 0, 0, FLAGS, "transition" },
677  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
678  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
679  { NULL }
680 };
681 
682 AVFILTER_DEFINE_CLASS(xfade_vulkan);
683 
684 static const AVFilterPad xfade_vulkan_inputs[] = {
685  {
686  .name = "main",
687  .type = AVMEDIA_TYPE_VIDEO,
688  .get_buffer.video = &get_video_buffer,
689  .config_props = &ff_vk_filter_config_input,
690  },
691  {
692  .name = "xfade",
693  .type = AVMEDIA_TYPE_VIDEO,
694  .get_buffer.video = &get_video_buffer,
695  .config_props = &ff_vk_filter_config_input,
696  },
697 };
698 
699 static const AVFilterPad xfade_vulkan_outputs[] = {
700  {
701  .name = "default",
702  .type = AVMEDIA_TYPE_VIDEO,
703  .config_props = &config_props_output,
704  },
705 };
706 
707 const AVFilter ff_vf_xfade_vulkan = {
708  .name = "xfade_vulkan",
709  .description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
710  .priv_size = sizeof(XFadeVulkanContext),
711  .init = &ff_vk_filter_init,
712  .uninit = &uninit,
713  .activate = &activate,
714  FILTER_INPUTS(xfade_vulkan_inputs),
715  FILTER_OUTPUTS(xfade_vulkan_outputs),
716  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VULKAN),
717  .priv_class = &xfade_vulkan_class,
718  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
719  .flags = AVFILTER_FLAG_HWDEVICE,
720 };
IN_A
#define IN_A
Definition: vf_xfade_vulkan.c:28
WIPETR
@ WIPETR
Definition: vf_xfade_vulkan.c:85
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:108
IN_B
#define IN_B
Definition: vf_xfade_vulkan.c:29
WIPELEFT
@ WIPELEFT
Definition: vf_xfade_vulkan.c:72
ff_vk_pipeline_free
void ff_vk_pipeline_free(FFVulkanContext *s, FFVulkanPipeline *pl)
Definition: vulkan.c:1868
mix
static int mix(int c0, int c1)
Definition: 4xm.c:717
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:227
xfade_frame
static int xfade_frame(AVFilterContext *avctx, AVFrame *frame_a, AVFrame *frame_b)
Definition: vf_xfade_vulkan.c:415
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: internal.h:364
transition_wipeup
static const char transition_wipeup[]
Definition: vf_xfade_vulkan.c:122
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:978
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:225
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:166
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:100
XFadeVulkanContext::start_pts
int64_t start_pts
Definition: vf_xfade_vulkan.c:51
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:340
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:452
ff_vk_filter_init
int ff_vk_filter_init(AVFilterContext *avctx)
General lavfi IO functions.
Definition: vulkan_filter.c:219
w
uint8_t w
Definition: llviddspenc.c:38
frand
static float frand(int x, int y)
Definition: vf_deband.c:115
XFadeVulkanContext::status
int status[IN_NB]
Definition: vf_xfade_vulkan.c:67
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:385
transitions_map
static const char * transitions_map[NB_TRANSITIONS]
Definition: vf_xfade_vulkan.c:301
XFadeVulkanContext::transition
int transition
Definition: vf_xfade_vulkan.c:39
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:1413
AVOption
AVOption.
Definition: opt.h:251
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:646
CIRCLEOPEN
@ CIRCLEOPEN
Definition: vf_xfade_vulkan.c:80
AV_OPT_TYPE_DURATION
@ AV_OPT_TYPE_DURATION
Definition: opt.h:239
transition_wiperight
static const char transition_wiperight[]
Definition: vf_xfade_vulkan.c:111
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:45
transition_wipebl
static const char transition_wipebl[]
Definition: vf_xfade_vulkan.c:277
XFadeVulkanContext::qf
FFVkQueueFamilyCtx qf
Definition: vf_xfade_vulkan.c:46
XFadeVulkanContext::initialized
int initialized
Definition: vf_xfade_vulkan.c:43
ff_vk_uninit
void ff_vk_uninit(FFVulkanContext *s)
Frees main context.
Definition: vulkan.c:1897
transition_wipebr
static const char transition_wipebr[]
Definition: vf_xfade_vulkan.c:289
FFVkSPIRVCompiler::uninit
void(* uninit)(struct FFVkSPIRVCompiler **ctx)
Definition: vulkan_spirv.h:33
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:170
ff_vk_pipeline_descriptor_set_add
int ff_vk_pipeline_descriptor_set_add(FFVulkanContext *s, FFVulkanPipeline *pl, FFVkSPIRVShader *shd, FFVulkanDescriptorSetBinding *desc, int nb, int read_only, int print_to_shader_only)
Add descriptor to a pipeline.
Definition: vulkan.c:1463
ff_vk_shader_set_compute_sizes
void ff_vk_shader_set_compute_sizes(FFVkSPIRVShader *shd, int x, int y, int z)
Definition: vulkan.c:1371
video.h
AV_PIX_FMT_VULKAN
@ AV_PIX_FMT_VULKAN
Vulkan hardware images.
Definition: pixfmt.h:376
ff_vf_xfade_vulkan
const AVFilter ff_vf_xfade_vulkan
Definition: vf_xfade_vulkan.c:707
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:1383
ff_default_get_video_buffer
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:103
WIPEBL
@ WIPEBL
Definition: vf_xfade_vulkan.c:86
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3004
SLIDELEFT
@ SLIDELEFT
Definition: vf_xfade_vulkan.c:78
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:212
transition_slideleft
static const char transition_slideleft[]
Definition: vf_xfade_vulkan.c:174
AVVulkanDeviceContext::alloc
const VkAllocationCallbacks * alloc
Custom memory allocator, else NULL.
Definition: hwcontext_vulkan.h:48
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:1143
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:412
fail
#define fail()
Definition: checkasm.h:138
SLIDEDOWN
@ SLIDEDOWN
Definition: vf_xfade_vulkan.c:76
vulkan_filter.h
SLIDEUP
@ SLIDEUP
Definition: vf_xfade_vulkan.c:77
XFadeVulkanContext::pts
int64_t pts
Definition: vf_xfade_vulkan.c:60
AVRational::num
int num
Numerator.
Definition: rational.h:59
XFadeVulkanContext::passthrough
int passthrough
Definition: vf_xfade_vulkan.c:65
SHADER_FRAND_FUNC
#define SHADER_FRAND_FUNC
Definition: vf_xfade_vulkan.c:221
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:47
transition_circleclose
static const char transition_circleclose[]
Definition: vf_xfade_vulkan.c:213
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:404
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
duration
int64_t duration
Definition: movenc.c:64
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:189
ff_inlink_request_frame
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
Definition: avfilter.c:1506
s
#define s(width, name)
Definition: cbs_vp9.c:198
floor
static __device__ float floor(float a)
Definition: cuda_runtime.h:173
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts_bsf.c:365
filters.h
transition_wipedown
static const char transition_wipedown[]
Definition: vf_xfade_vulkan.c:133
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:256
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:192
frame
static AVFrame * frame
Definition: demux_decode.c:54
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:1422
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:222
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:736
GLSLD
#define GLSLD(D)
Definition: vulkan.h:60
transition_pixelize
static const char transition_pixelize[]
Definition: vf_xfade_vulkan.c:238
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:405
activate
static int activate(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:535
transition_circleopen
static const char transition_circleopen[]
Definition: vf_xfade_vulkan.c:205
ff_vk_filter_config_output
int ff_vk_filter_config_output(AVFilterLink *outlink)
Definition: vulkan_filter.c:196
transition_wipetl
static const char transition_wipetl[]
Definition: vf_xfade_vulkan.c:253
transition_fade
static const char transition_fade[]
Definition: vf_xfade_vulkan.c:91
xfade_vulkan_outputs
static const AVFilterPad xfade_vulkan_outputs[]
Definition: vf_xfade_vulkan.c:699
ff_vk_init_compute_pipeline
int ff_vk_init_compute_pipeline(FFVulkanContext *s, FFVulkanPipeline *pl, FFVkSPIRVShader *shd)
Definition: vulkan.c:1809
xfade_vulkan_options
static const AVOption xfade_vulkan_options[]
Definition: vf_xfade_vulkan.c:658
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:296
av_clipf
av_clipf
Definition: af_crystalizer.c:121
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:1337
FFVulkanPipeline
Definition: vulkan.h:132
PIXELIZE
@ PIXELIZE
Definition: vf_xfade_vulkan.c:83
DISSOLVE
@ DISSOLVE
Definition: vf_xfade_vulkan.c:82
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:1345
main
int main(int argc, char **argv)
Definition: avio_http_serve_files.c:99
XFadeVulkanContext::pl
FFVulkanPipeline pl
Definition: vf_xfade_vulkan.c:44
FLAGS
#define FLAGS
Definition: vf_xfade_vulkan.c:656
ff_inlink_set_status
void ff_inlink_set_status(AVFilterLink *link, int status)
Set the status on an input link.
Definition: avfilter.c:1514
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:1357
FFVulkanDescriptorSetBinding
Definition: vulkan.h:84
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:106
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:138
SHADER_CIRCLE_COMMON
#define SHADER_CIRCLE_COMMON
Definition: vf_xfade_vulkan.c:190
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:111
XFadeVulkanContext::sampler
VkSampler sampler
Definition: vf_xfade_vulkan.c:48
XFadeVulkanContext::duration
int64_t duration
Definition: vf_xfade_vulkan.c:40
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:41
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:54
ff_null_get_video_buffer
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:43
internal.h
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
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:182
xfade_vulkan_inputs
static const AVFilterPad xfade_vulkan_inputs[]
Definition: vf_xfade_vulkan.c:684
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:255
XFadeVulkanContext::shd
FFVkSPIRVShader shd
Definition: vf_xfade_vulkan.c:47
NB_TRANSITIONS
@ NB_TRANSITIONS
Definition: vf_xfade_vulkan.c:88
DUP_SAMPLER
#define DUP_SAMPLER(x)
Definition: vulkan.h:74
planes
static const struct @363 planes[]
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:1207
vulkan_spirv.h
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:53
CIRCLECLOSE
@ CIRCLECLOSE
Definition: vf_xfade_vulkan.c:81
GLSLF
#define GLSLF(N, S,...)
Definition: vulkan.h:55
XFadeParameters
Definition: vf_xfade_vulkan.c:32
FFVkSPIRVCompiler::free_shader
void(* free_shader)(struct FFVkSPIRVCompiler *ctx, void **opaque)
Definition: vulkan_spirv.h:32
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
ret
ret
Definition: filter_design.txt:187
WIPERIGHT
@ WIPERIGHT
Definition: vf_xfade_vulkan.c:73
FADE
@ FADE
Definition: vf_xfade_vulkan.c:71
FFVulkanContext::vkfn
FFVulkanFunctions vkfn
Definition: vulkan.h:232
SHADER_SLIDE_COMMON
#define SHADER_SLIDE_COMMON
Definition: vf_xfade_vulkan.c:144
FFVkExecPool
Definition: vulkan.h:211
pos
unsigned int pos
Definition: spdifenc.c:413
XFadeVulkanContext::duration_pts
int64_t duration_pts
Definition: vf_xfade_vulkan.c:57
XFadeVulkanContext::vkctx
FFVulkanContext vkctx
Definition: vf_xfade_vulkan.c:37
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:752
status
ov_status_e status
Definition: dnn_backend_openvino.c:119
random_seed.h
transition_slideright
static const char transition_slideright[]
Definition: vf_xfade_vulkan.c:182
FFVkSPIRVShader
Definition: vulkan.h:76
AVRational::den
int den
Denominator.
Definition: rational.h:60
WIPEBR
@ WIPEBR
Definition: vf_xfade_vulkan.c:87
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
config_props_output
static int config_props_output(AVFilterLink *outlink)
Definition: vf_xfade_vulkan.c:456
init_vulkan
static av_cold int init_vulkan(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:321
transition_wipetr
static const char transition_wipetr[]
Definition: vf_xfade_vulkan.c:265
XFadeVulkanContext
Definition: vf_xfade_vulkan.c:36
AVFilterContext
An instance of a filter.
Definition: avfilter.h:397
desc
const char * desc
Definition: libsvtav1.c:83
GLSLC
#define GLSLC(N, S)
Definition: vulkan.h:45
ff_vk_filter_config_input
int ff_vk_filter_config_input(AVFilterLink *inlink)
Definition: vulkan_filter.c:165
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
FFVulkanContext::hwctx
AVVulkanDeviceContext * hwctx
Definition: vulkan.h:254
WIPEDOWN
@ WIPEDOWN
Definition: vf_xfade_vulkan.c:75
AVVulkanDeviceContext::act_dev
VkDevice act_dev
Active device.
Definition: hwcontext_vulkan.h:70
XFadeTransitions
XFadeTransitions
Definition: vf_xfade.c:30
WIPETL
@ WIPETL
Definition: vf_xfade_vulkan.c:84
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:193
ff_vk_init_sampler
int ff_vk_init_sampler(FFVulkanContext *s, VkSampler *sampler, int unnorm_coords, VkFilter filt)
Create a sampler.
Definition: vulkan.c:1163
SLIDERIGHT
@ SLIDERIGHT
Definition: vf_xfade_vulkan.c:79
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:1577
d
d
Definition: ffmpeg_filter.c:368
XFadeParameters::progress
float progress
Definition: vf_xfade_vulkan.c:33
transition_wipeleft
static const char transition_wipeleft[]
Definition: vf_xfade_vulkan.c:100
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
IN_NB
#define IN_NB
Definition: vf_xfade_vulkan.c:30
h
h
Definition: vp9dsp_template.c:2038
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:74
transition_slidedown
static const char transition_slidedown[]
Definition: vf_xfade_vulkan.c:158
ff_vk_shader_free
void ff_vk_shader_free(FFVulkanContext *s, FFVkSPIRVShader *shd)
Definition: vulkan.c:1404
uninit
static av_cold void uninit(AVFilterContext *avctx)
Definition: vf_xfade_vulkan.c:627
OFFSET
#define OFFSET(x)
Definition: vf_xfade_vulkan.c:655
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
RET
#define RET(x)
Definition: vulkan.h:68
FFVulkanFunctions
Definition: vulkan_functions.h:226
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:216
min
float min
Definition: vorbis_enc_data.h:429
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:409
Generated on Thu Apr 18 2024 22:42:31 for FFmpeg by   doxygen 1.8.17