FFmpeg
vf_scdet_vulkan.c
Go to the documentation of this file.
1 /*
2  * Copyright 2025 (c) Niklas Haas
3  *
4  * This file is part of FFmpeg.
5  *
6  * FFmpeg is free software; you can redistribute it and/or
7  * modify it under the terms of the GNU Lesser General Public
8  * License as published by the Free Software Foundation; either
9  * version 2.1 of the License, or (at your option) any later version.
10  *
11  * FFmpeg is distributed in the hope that it will be useful,
12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
14  * Lesser General Public License for more details.
15  *
16  * You should have received a copy of the GNU Lesser General Public
17  * License along with FFmpeg; if not, write to the Free Software
18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19  */
20 
21 #include "libavutil/avassert.h"
22 #include "libavutil/opt.h"
23 #include "libavutil/timestamp.h"
24 #include "vulkan_filter.h"
25 
26 #include "filters.h"
27 
28 extern const unsigned char ff_scdet_comp_spv_data[];
29 extern const unsigned int ff_scdet_comp_spv_len;
30 
31 typedef struct SceneDetectVulkanContext {
33 
39 
40  double threshold;
41  int sc_pass;
42 
43  int nb_planes;
44  double prev_mafd;
48 
49 typedef struct SceneDetectBuf {
50 #define SLICES 16
51  uint32_t frame_sad[SLICES];
53 
55 {
56  int err;
58  FFVulkanContext *vkctx = &s->vkctx;
59 
60  const AVPixFmtDescriptor *pixdesc = av_pix_fmt_desc_get(s->vkctx.input_format);
61  const int lumaonly = !(pixdesc->flags & AV_PIX_FMT_FLAG_RGB) &&
62  (pixdesc->flags & AV_PIX_FMT_FLAG_PLANAR);
63  s->nb_planes = lumaonly ? 1 : av_pix_fmt_count_planes(s->vkctx.input_format);
64 
65  s->qf = ff_vk_qf_find(vkctx, VK_QUEUE_COMPUTE_BIT, 0);
66  if (!s->qf) {
67  av_log(ctx, AV_LOG_ERROR, "Device has no compute queues\n");
68  err = AVERROR(ENOTSUP);
69  goto fail;
70  }
71 
72  RET(ff_vk_exec_pool_init(vkctx, s->qf, &s->e, s->qf->num*4, 0, 0, 0, NULL));
73 
74  SPEC_LIST_CREATE(sl, 2, 2*sizeof(uint32_t))
75  SPEC_LIST_ADD(sl, 0, 32, s->nb_planes);
76  SPEC_LIST_ADD(sl, 1, 32, SLICES);
77 
78  ff_vk_shader_load(&s->shd, VK_SHADER_STAGE_COMPUTE_BIT, sl,
79  (int []) { 32, 32, 1 }, 0);
80 
82  { /* prev_img */
83  .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
84  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
85  .elems = av_pix_fmt_count_planes(s->vkctx.input_format),
86  },
87  { /* cur_img */
88  .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
89  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
90  .elems = av_pix_fmt_count_planes(s->vkctx.input_format),
91  },
92  { /* sad_buffer */
93  .type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
94  .stages = VK_SHADER_STAGE_COMPUTE_BIT,
95  }
96  };
97  ff_vk_shader_add_descriptor_set(vkctx, &s->shd, desc, 3, 0, 0);
98 
99  RET(ff_vk_shader_link(vkctx, &s->shd,
101  ff_scdet_comp_spv_len, "main"));
102 
103  RET(ff_vk_shader_register_exec(vkctx, &s->e, &s->shd));
104 
105  s->initialized = 1;
106 
107 fail:
108  return err;
109 }
110 
111 static double evaluate(AVFilterContext *ctx, const SceneDetectBuf *buf)
112 {
113  SceneDetectVulkanContext *s = ctx->priv;
114  const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(s->vkctx.input_format);
115  const AVFilterLink *inlink = ctx->inputs[0];
116  uint64_t count;
117  double mafd, diff;
118 
119  uint64_t sad = 0;
120  for (int i = 0; i < SLICES; i++)
121  sad += buf->frame_sad[i];
122 
123  av_assert2(s->nb_planes == 1 || !(desc->log2_chroma_w || desc->log2_chroma_h));
124  count = s->nb_planes * inlink->w * inlink->h;
125  mafd = (double) sad * 100.0 / count / (1ULL << desc->comp[0].depth);
126  diff = fabs(mafd - s->prev_mafd);
127  s->prev_mafd = mafd;
128 
129  return av_clipf(FFMIN(mafd, diff), 0.0, 100.0);
130 }
131 
133 {
134  int err;
135  AVFilterContext *ctx = link->dst;
136  SceneDetectVulkanContext *s = ctx->priv;
137  AVFilterLink *outlink = ctx->outputs[0];
138 
139  VkImageView prev_views[AV_NUM_DATA_POINTERS];
140  VkImageView cur_views[AV_NUM_DATA_POINTERS];
141  VkImageMemoryBarrier2 img_bar[8];
142  int nb_img_bar = 0;
143 
144  FFVulkanContext *vkctx = &s->vkctx;
145  FFVulkanFunctions *vk = &vkctx->vkfn;
146  FFVkExecContext *exec = NULL;
147  AVBufferRef *buf = NULL;
148  FFVkBuffer *buf_vk;
149 
150  SceneDetectBuf *sad;
151  double score = 0.0;
152  char str[64];
153 
154  if (!s->initialized)
155  RET(init_filter(ctx));
156 
157  av_frame_free(&s->prev);
158  s->prev = s->cur;
159  s->cur = av_frame_clone(in);
160  if (!s->prev)
161  goto done;
162 
163  RET(ff_vk_get_pooled_buffer(vkctx, &s->det_buf_pool, &buf,
164  VK_BUFFER_USAGE_TRANSFER_DST_BIT |
165  VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
166  NULL,
167  sizeof(SceneDetectBuf),
168  VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT |
169  VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT |
170  VK_MEMORY_PROPERTY_HOST_COHERENT_BIT));
171  buf_vk = (FFVkBuffer *)buf->data;
172  sad = (SceneDetectBuf *) buf_vk->mapped_mem;
173 
174  exec = ff_vk_exec_get(vkctx, &s->e);
175  ff_vk_exec_start(vkctx, exec);
176 
177  RET(ff_vk_exec_add_dep_frame(vkctx, exec, s->prev,
178  VK_PIPELINE_STAGE_2_NONE,
179  VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
180  RET(ff_vk_create_imageviews(vkctx, exec, prev_views, s->prev, FF_VK_REP_UINT));
181 
182  ff_vk_shader_update_img_array(vkctx, exec, &s->shd, s->prev, prev_views, 0, 0,
183  VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
184 
185  ff_vk_frame_barrier(vkctx, exec, s->prev, img_bar, &nb_img_bar,
186  VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
187  VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
188  VK_ACCESS_SHADER_READ_BIT,
189  VK_IMAGE_LAYOUT_GENERAL,
190  VK_QUEUE_FAMILY_IGNORED);
191 
192  RET(ff_vk_exec_add_dep_frame(vkctx, exec, s->cur,
193  VK_PIPELINE_STAGE_2_NONE,
194  VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
195  RET(ff_vk_create_imageviews(vkctx, exec, cur_views, s->cur, FF_VK_REP_UINT));
196 
197  ff_vk_shader_update_img_array(vkctx, exec, &s->shd, s->cur, cur_views, 0, 1,
198  VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
199 
200  ff_vk_frame_barrier(vkctx, exec, s->cur, img_bar, &nb_img_bar,
201  VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
202  VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
203  VK_ACCESS_SHADER_READ_BIT,
204  VK_IMAGE_LAYOUT_GENERAL,
205  VK_QUEUE_FAMILY_IGNORED);
206 
207  /* zero buffer */
208  vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
209  .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
210  .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
211  .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
212  .srcStageMask = VK_PIPELINE_STAGE_2_NONE,
213  .dstStageMask = VK_PIPELINE_STAGE_2_TRANSFER_BIT,
214  .dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
215  .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
216  .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
217  .buffer = buf_vk->buf,
218  .size = buf_vk->size,
219  .offset = 0,
220  },
221  .bufferMemoryBarrierCount = 1,
222  });
223 
224  vk->CmdFillBuffer(exec->buf, buf_vk->buf, 0, buf_vk->size, 0x0);
225 
226  vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
227  .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
228  .pImageMemoryBarriers = img_bar,
229  .imageMemoryBarrierCount = nb_img_bar,
230  .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
231  .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
232  .srcStageMask = VK_PIPELINE_STAGE_2_TRANSFER_BIT,
233  .dstStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
234  .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
235  .dstAccessMask = VK_ACCESS_2_SHADER_STORAGE_READ_BIT |
236  VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT,
237  .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
238  .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
239  .buffer = buf_vk->buf,
240  .size = buf_vk->size,
241  .offset = 0,
242  },
243  .bufferMemoryBarrierCount = 1,
244  });
245 
246  RET(ff_vk_shader_update_desc_buffer(&s->vkctx, exec, &s->shd, 0, 2, 0,
247  buf_vk, 0, buf_vk->size,
248  VK_FORMAT_UNDEFINED));
249 
250  ff_vk_exec_bind_shader(vkctx, exec, &s->shd);
251 
252  vk->CmdDispatch(exec->buf,
253  FFALIGN(in->width, s->shd.lg_size[0]) / s->shd.lg_size[0],
254  FFALIGN(in->height, s->shd.lg_size[1]) / s->shd.lg_size[1],
255  s->shd.lg_size[2]);
256 
257  vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
258  .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
259  .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
260  .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
261  .srcStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
262  .dstStageMask = VK_PIPELINE_STAGE_2_HOST_BIT,
263  .srcAccessMask = VK_ACCESS_2_SHADER_STORAGE_READ_BIT |
264  VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT,
265  .dstAccessMask = VK_ACCESS_HOST_READ_BIT,
266  .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
267  .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
268  .buffer = buf_vk->buf,
269  .size = buf_vk->size,
270  .offset = 0,
271  },
272  .bufferMemoryBarrierCount = 1,
273  });
274 
275  RET(ff_vk_exec_submit(vkctx, exec));
276  ff_vk_exec_wait(vkctx, exec);
277  score = evaluate(ctx, sad);
278 
279 done:
280  snprintf(str, sizeof(str), "%0.3f", s->prev_mafd);
281  av_dict_set(&in->metadata, "lavfi.scd.mafd", str, 0);
282  snprintf(str, sizeof(str), "%0.3f", score);
283  av_dict_set(&in->metadata, "lavfi.scd.score", str, 0);
284 
285  if (score >= s->threshold) {
286  const char *pts = av_ts2timestr(in->pts, &link->time_base);
287  av_dict_set(&in->metadata, "lavfi.scd.time", pts, 0);
288  av_log(ctx, AV_LOG_INFO, "lavfi.scd.score: %.3f, lavfi.scd.time: %s\n",
289  score, pts);
290  }
291 
292  av_buffer_unref(&buf);
293  if (!s->sc_pass || score >= s->threshold)
294  return ff_filter_frame(outlink, in);
295  else {
296  av_frame_free(&in);
297  return 0;
298  }
299 
300 fail:
301  if (exec)
302  ff_vk_exec_discard_deps(&s->vkctx, exec);
303  av_frame_free(&in);
304  av_buffer_unref(&buf);
305  return err;
306 }
307 
309 {
310  SceneDetectVulkanContext *s = avctx->priv;
311  FFVulkanContext *vkctx = &s->vkctx;
312 
313  av_frame_free(&s->prev);
314  av_frame_free(&s->cur);
315 
316  ff_vk_exec_pool_free(vkctx, &s->e);
317  ff_vk_shader_free(vkctx, &s->shd);
318 
319  av_buffer_pool_uninit(&s->det_buf_pool);
320 
321  ff_vk_uninit(&s->vkctx);
322 
323  s->initialized = 0;
324 }
325 
326 #define OFFSET(x) offsetof(SceneDetectVulkanContext, x)
327 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
328 static const AVOption scdet_vulkan_options[] = {
329  { "threshold", "set scene change detect threshold", OFFSET(threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 10.}, 0, 100., FLAGS },
330  { "t", "set scene change detect threshold", OFFSET(threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 10.}, 0, 100., FLAGS },
331  { "sc_pass", "Set the flag to pass scene change frames", OFFSET(sc_pass), AV_OPT_TYPE_BOOL, {.i64 = 0 }, 0, 1, FLAGS },
332  { "s", "Set the flag to pass scene change frames", OFFSET(sc_pass), AV_OPT_TYPE_BOOL, {.i64 = 0 }, 0, 1, FLAGS },
333  { NULL }
334 };
335 
336 AVFILTER_DEFINE_CLASS(scdet_vulkan);
337 
339  {
340  .name = "default",
341  .type = AVMEDIA_TYPE_VIDEO,
342  .filter_frame = &scdet_vulkan_filter_frame,
343  .config_props = &ff_vk_filter_config_input,
344  },
345 };
346 
348  {
349  .name = "default",
350  .type = AVMEDIA_TYPE_VIDEO,
351  .config_props = &ff_vk_filter_config_output,
352  },
353 };
354 
356  .p.name = "scdet_vulkan",
357  .p.description = NULL_IF_CONFIG_SMALL("Detect video scene change"),
358  .p.priv_class = &scdet_vulkan_class,
359  .p.flags = AVFILTER_FLAG_HWDEVICE,
360  .priv_size = sizeof(SceneDetectVulkanContext),
366  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
367 };
SceneDetectVulkanContext::det_buf_pool
AVBufferPool * det_buf_pool
Definition: vf_scdet_vulkan.c:38
scdet_vulkan_uninit
static void scdet_vulkan_uninit(AVFilterContext *avctx)
Definition: vf_scdet_vulkan.c:308
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
ff_vk_shader_free
void ff_vk_shader_free(FFVulkanContext *s, FFVulkanShader *shd)
Free a shader.
Definition: vulkan.c:2810
AVBufferPool
The buffer pool.
Definition: buffer_internal.h:88
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1067
SceneDetectVulkanContext::prev
AVFrame * prev
Definition: vf_scdet_vulkan.c:45
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3456
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
RET
#define RET(x)
Definition: vulkan.h:68
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:357
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(scdet_vulkan)
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
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:208
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:64
SceneDetectVulkanContext::vkctx
FFVulkanContext vkctx
Definition: vf_scdet_vulkan.c:32
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:434
ff_vk_filter_init
int ff_vk_filter_init(AVFilterContext *avctx)
General lavfi IO functions.
Definition: vulkan_filter.c:233
SceneDetectVulkanContext::prev_mafd
double prev_mafd
Definition: vf_scdet_vulkan.c:44
AVOption
AVOption.
Definition: opt.h:429
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:254
filters.h
ff_scdet_comp_spv_len
const unsigned int ff_scdet_comp_spv_len
evaluate
static double evaluate(AVFilterContext *ctx, const SceneDetectBuf *buf)
Definition: vf_scdet_vulkan.c:111
ff_vk_exec_get
FFVkExecContext * ff_vk_exec_get(FFVulkanContext *s, FFVkExecPool *pool)
Retrieve an execution pool.
Definition: vulkan.c:548
ff_vk_uninit
void ff_vk_uninit(FFVulkanContext *s)
Frees main context.
Definition: vulkan.c:2836
SPEC_LIST_ADD
#define SPEC_LIST_ADD(name, idx, val_bits, val)
Definition: vulkan.h:86
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:220
ff_vk_exec_bind_shader
void ff_vk_exec_bind_shader(FFVulkanContext *s, FFVkExecContext *e, const FFVulkanShader *shd)
Bind a shader.
Definition: vulkan.c:2787
AV_PIX_FMT_VULKAN
@ AV_PIX_FMT_VULKAN
Vulkan hardware images.
Definition: pixfmt.h:379
ff_vk_exec_add_dep_frame
int ff_vk_exec_add_dep_frame(FFVulkanContext *s, FFVkExecContext *e, AVFrame *f, VkPipelineStageFlagBits2 wait_stage, VkPipelineStageFlagBits2 signal_stage)
Definition: vulkan.c:780
SceneDetectVulkanContext::initialized
int initialized
Definition: vf_scdet_vulkan.c:34
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3496
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:289
fail
#define fail()
Definition: checkasm.h:224
vulkan_filter.h
ff_vk_shader_update_img_array
void ff_vk_shader_update_img_array(FFVulkanContext *s, FFVkExecContext *e, FFVulkanShader *shd, AVFrame *f, VkImageView *views, int set, int binding, VkImageLayout layout, VkSampler sampler)
Update a descriptor in a buffer with an image array.
Definition: vulkan.c:2738
ff_vk_frame_barrier
void ff_vk_frame_barrier(FFVulkanContext *s, FFVkExecContext *e, AVFrame *pic, VkImageMemoryBarrier2 *bar, int *nb_bar, VkPipelineStageFlags2 src_stage, VkPipelineStageFlags2 dst_stage, VkAccessFlagBits2 new_access, VkImageLayout new_layout, uint32_t new_qf)
Definition: vulkan.c:2050
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:2603
pts
static int64_t pts
Definition: transcode_aac.c:644
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:40
avassert.h
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:210
av_cold
#define av_cold
Definition: attributes.h:111
FFFilter
Definition: filters.h:267
scdet_vulkan_filter_frame
static int scdet_vulkan_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_scdet_vulkan.c:132
s
#define s(width, name)
Definition: cbs_vp9.c:198
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:265
AV_OPT_TYPE_DOUBLE
@ AV_OPT_TYPE_DOUBLE
Underlying C type is double.
Definition: opt.h:267
ff_vk_exec_wait
void ff_vk_exec_wait(FFVulkanContext *s, FFVkExecContext *e)
Definition: vulkan.c:553
ctx
static AVFormatContext * ctx
Definition: movenc.c:49
av_frame_clone
AVFrame * av_frame_clone(const AVFrame *src)
Create a new frame that references the same data as src.
Definition: frame.c:483
SceneDetectVulkanContext::e
FFVkExecPool e
Definition: vf_scdet_vulkan.c:35
ff_vk_exec_pool_free
void ff_vk_exec_pool_free(FFVulkanContext *s, FFVkExecPool *pool)
Definition: vulkan.c:299
link
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 link
Definition: filter_design.txt:23
SceneDetectVulkanContext::sc_pass
int sc_pass
Definition: vf_scdet_vulkan.c:41
fabs
static __device__ float fabs(float a)
Definition: cuda_runtime.h:182
NULL
#define NULL
Definition: coverity.c:32
av_buffer_unref
void av_buffer_unref(AVBufferRef **buf)
Free a given reference and automatically free the buffer if there are no more references to it.
Definition: buffer.c:139
FLAGS
#define FLAGS
Definition: vf_scdet_vulkan.c:327
SceneDetectVulkanContext::nb_planes
int nb_planes
Definition: vf_scdet_vulkan.c:43
av_buffer_pool_uninit
void av_buffer_pool_uninit(AVBufferPool **ppool)
Mark the pool as being available for freeing.
Definition: buffer.c:328
ff_vk_filter_config_output
int ff_vk_filter_config_output(AVFilterLink *outlink)
Definition: vulkan_filter.c:209
ff_vk_shader_link
int ff_vk_shader_link(FFVulkanContext *s, FFVulkanShader *shd, const char *spirv, size_t spirv_len, const char *entrypoint)
Link a shader into an executable.
Definition: vulkan.c:2376
SPEC_LIST_CREATE
#define SPEC_LIST_CREATE(name, max_length, max_size)
Definition: vulkan.h:76
double
double
Definition: af_crystalizer.c:132
av_clipf
av_clipf
Definition: af_crystalizer.c:122
FFVkBuffer::mapped_mem
uint8_t * mapped_mem
Definition: vulkan.h:134
FFVulkanContext
Definition: vulkan.h:312
AVPixFmtDescriptor::flags
uint64_t flags
Combination of AV_PIX_FMT_FLAG_...
Definition: pixdesc.h:94
SceneDetectBuf::frame_sad
uint32_t frame_sad[SLICES]
Definition: vf_scdet_vulkan.c:51
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:550
av_ts2timestr
#define av_ts2timestr(ts, tb)
Convenience macro, the return value should be used only directly in function arguments but never stan...
Definition: timestamp.h:83
FFVulkanDescriptorSetBinding
Definition: vulkan.h:112
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
AV_PIX_FMT_FLAG_RGB
#define AV_PIX_FMT_FLAG_RGB
The pixel format contains RGB-like data (as opposed to YUV/grayscale).
Definition: pixdesc.h:136
i
#define i(width, name, range_min, range_max)
Definition: cbs_h264.c:63
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:188
AV_NUM_DATA_POINTERS
#define AV_NUM_DATA_POINTERS
Definition: frame.h:435
FFVulkanShader
Definition: vulkan.h:225
OFFSET
#define OFFSET(x)
Definition: vf_scdet_vulkan.c:326
AVFrame::time_base
AVRational time_base
Time base for the timestamps in this frame.
Definition: frame.h:551
SceneDetectVulkanContext::qf
AVVulkanDeviceQueueFamily * qf
Definition: vf_scdet_vulkan.c:36
diff
static av_always_inline int diff(const struct color_info *a, const struct color_info *b, const int trans_thresh)
Definition: vf_paletteuse.c:166
FFVkExecContext
Definition: vulkan.h:145
ff_vk_shader_update_desc_buffer
int ff_vk_shader_update_desc_buffer(FFVulkanContext *s, FFVkExecContext *e, FFVulkanShader *shd, int set, int bind, int elem, FFVkBuffer *buf, VkDeviceSize offset, VkDeviceSize len, VkFormat fmt)
Update a descriptor in a buffer with a buffer.
Definition: vulkan.c:2751
AV_LOG_INFO
#define AV_LOG_INFO
Standard information.
Definition: log.h:221
av_assert2
#define av_assert2(cond)
assert() equivalent, that does lie in speed critical code.
Definition: avassert.h:68
uninit
static void uninit(AVBSFContext *ctx)
Definition: pcm_rechunk.c:68
ff_vk_exec_start
int ff_vk_exec_start(FFVulkanContext *s, FFVkExecContext *e)
Start/submit/wait an execution.
Definition: vulkan.c:560
FF_VK_REP_UINT
@ FF_VK_REP_UINT
Definition: vulkan.h:455
ff_scdet_comp_spv_data
const unsigned char ff_scdet_comp_spv_data[]
SceneDetectVulkanContext::threshold
double threshold
Definition: vf_scdet_vulkan.c:40
FFMIN
#define FFMIN(a, b)
Definition: macros.h:49
SceneDetectVulkanContext
Definition: vf_scdet_vulkan.c:31
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:46
SceneDetectVulkanContext::shd
FFVulkanShader shd
Definition: vf_scdet_vulkan.c:37
ff_vk_create_imageviews
int ff_vk_create_imageviews(FFVulkanContext *s, FFVkExecContext *e, VkImageView views[AV_NUM_DATA_POINTERS], AVFrame *f, enum FFVkShaderRepFormat rep_fmt)
Create an imageview and add it as a dependency to an execution.
Definition: vulkan.c:1967
SceneDetectBuf
Definition: vf_scdet_vulkan.c:49
FFVulkanContext::vkfn
FFVulkanFunctions vkfn
Definition: vulkan.h:316
FFVkExecPool
Definition: vulkan.h:290
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:264
ff_vk_qf_find
AVVulkanDeviceQueueFamily * ff_vk_qf_find(FFVulkanContext *s, VkQueueFlagBits dev_family, VkVideoCodecOperationFlagBitsKHR vid_ops)
Chooses an appropriate QF.
Definition: vulkan.c:286
ff_vf_scdet_vulkan
const FFFilter ff_vf_scdet_vulkan
Definition: vf_scdet_vulkan.c:355
FFVkExecContext::buf
VkCommandBuffer buf
Definition: vulkan.h:156
ff_vk_shader_add_descriptor_set
int ff_vk_shader_add_descriptor_set(FFVulkanContext *s, FFVulkanShader *shd, const FFVulkanDescriptorSetBinding *desc, int nb, int singular, int print_to_shader_only)
Add descriptor to a shader.
Definition: vulkan.c:2503
SLICES
#define SLICES
Definition: vf_scdet_vulkan.c:50
AV_PIX_FMT_FLAG_PLANAR
#define AV_PIX_FMT_FLAG_PLANAR
At least one pixel component is not in the first data plane.
Definition: pixdesc.h:132
scdet_vulkan_outputs
static const AVFilterPad scdet_vulkan_outputs[]
Definition: vf_scdet_vulkan.c:347
scdet_vulkan_options
static const AVOption scdet_vulkan_options[]
Definition: vf_scdet_vulkan.c:328
AVFilterContext
An instance of a filter.
Definition: avfilter.h:274
desc
const char * desc
Definition: libsvtav1.c:82
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:200
FFFilter::p
AVFilter p
The public AVFilter.
Definition: filters.h:271
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
init_filter
static av_cold int init_filter(AVFilterContext *ctx)
Definition: vf_scdet_vulkan.c:54
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
ff_vk_exec_discard_deps
void ff_vk_exec_discard_deps(FFVulkanContext *s, FFVkExecContext *e)
Definition: vulkan.c:592
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
Definition: opt.h:327
av_dict_set
int av_dict_set(AVDictionary **pm, const char *key, const char *value, int flags)
Set the given entry in *pm, overwriting an existing entry.
Definition: dict.c:86
SceneDetectVulkanContext::cur
AVFrame * cur
Definition: vf_scdet_vulkan.c:46
FFVkBuffer
Definition: vulkan.h:125
timestamp.h
ff_vk_exec_submit
int ff_vk_exec_submit(FFVulkanContext *s, FFVkExecContext *e)
Definition: vulkan.c:905
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
AVVulkanDeviceQueueFamily
Definition: hwcontext_vulkan.h:33
snprintf
#define snprintf
Definition: snprintf.h:34
scdet_vulkan_inputs
static const AVFilterPad scdet_vulkan_inputs[]
Definition: vf_scdet_vulkan.c:338
FFVulkanFunctions
Definition: vulkan_functions.h:275
ff_vk_shader_load
int ff_vk_shader_load(FFVulkanShader *shd, VkPipelineStageFlags stage, VkSpecializationInfo *spec, uint32_t wg_size[3], uint32_t required_subgroup_size)
Initialize a shader object.
Definition: vulkan.c:2093
ff_vk_get_pooled_buffer
int ff_vk_get_pooled_buffer(FFVulkanContext *ctx, AVBufferPool **buf_pool, AVBufferRef **buf, VkBufferUsageFlags usage, void *create_pNext, size_t size, VkMemoryPropertyFlagBits mem_props)
Initialize a pool and create AVBufferRefs containing FFVkBuffer.
Definition: vulkan.c:1286