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