FFmpeg
Loading...
Searching...
No Matches
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
28extern const unsigned char ff_scdet_comp_spv_data[];
29extern const unsigned int ff_scdet_comp_spv_len;
30
48
49typedef 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, FF_VK_DEFAULT_EXEC_CONTEXTS, 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);
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
107fail:
108 return err;
109}
110
111static double evaluate(AVFilterContext *ctx, const SceneDetectBuf *buf)
112{
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;
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 FFVkBuffer *buf_vk = NULL;
148
149 SceneDetectBuf *sad;
150 double score = 0.0;
151 char str[64];
152
153 if (!s->initialized)
155
156 av_frame_free(&s->prev);
157 s->prev = s->cur;
158 s->cur = av_frame_clone(in);
159 if (!s->prev)
160 goto done;
161
162 RET(ff_vk_get_pooled_buffer(vkctx, &s->det_buf_pool, &buf_vk,
163 VK_BUFFER_USAGE_TRANSFER_DST_BIT |
164 VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
165 NULL,
166 sizeof(SceneDetectBuf),
167 VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT |
168 VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT |
169 VK_MEMORY_PROPERTY_HOST_COHERENT_BIT));
170 sad = (SceneDetectBuf *) buf_vk->mapped_mem;
171
172 exec = ff_vk_exec_get(vkctx, &s->e);
173 err = ff_vk_exec_start(vkctx, exec);
174 if (err < 0) {
175 av_frame_free(&in);
176 av_refstruct_unref(&buf_vk);
177 return err;
178 }
179
180 RET(ff_vk_exec_add_dep_frame(vkctx, exec, s->prev,
181 VK_PIPELINE_STAGE_2_NONE,
182 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
183 RET(ff_vk_create_imageviews(vkctx, exec, prev_views, s->prev, FF_VK_REP_UINT));
184
185 ff_vk_shader_update_img_array(vkctx, exec, &s->shd, s->prev, prev_views, 0, 0,
186 VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
187
188 ff_vk_frame_barrier(vkctx, exec, s->prev, img_bar, &nb_img_bar,
189 VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
190 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
191 VK_ACCESS_SHADER_READ_BIT,
192 VK_IMAGE_LAYOUT_GENERAL,
193 VK_QUEUE_FAMILY_IGNORED);
194
195 RET(ff_vk_exec_add_dep_frame(vkctx, exec, s->cur,
196 VK_PIPELINE_STAGE_2_NONE,
197 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
198 RET(ff_vk_create_imageviews(vkctx, exec, cur_views, s->cur, FF_VK_REP_UINT));
199
200 ff_vk_shader_update_img_array(vkctx, exec, &s->shd, s->cur, cur_views, 0, 1,
201 VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
202
203 ff_vk_frame_barrier(vkctx, exec, s->cur, img_bar, &nb_img_bar,
204 VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT,
205 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
206 VK_ACCESS_SHADER_READ_BIT,
207 VK_IMAGE_LAYOUT_GENERAL,
208 VK_QUEUE_FAMILY_IGNORED);
209
210 /* zero buffer */
211 vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
212 .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
213 .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
214 .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
215 .srcStageMask = VK_PIPELINE_STAGE_2_NONE,
216 .dstStageMask = VK_PIPELINE_STAGE_2_TRANSFER_BIT,
217 .dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
218 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
219 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
220 .buffer = buf_vk->buf,
221 .size = buf_vk->size,
222 .offset = 0,
223 },
224 .bufferMemoryBarrierCount = 1,
225 });
226
227 vk->CmdFillBuffer(exec->buf, buf_vk->buf, 0, buf_vk->size, 0x0);
228
229 vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
230 .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
231 .pImageMemoryBarriers = img_bar,
232 .imageMemoryBarrierCount = nb_img_bar,
233 .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
234 .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
235 .srcStageMask = VK_PIPELINE_STAGE_2_TRANSFER_BIT,
236 .dstStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
237 .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
238 .dstAccessMask = VK_ACCESS_2_SHADER_STORAGE_READ_BIT |
239 VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT,
240 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
241 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
242 .buffer = buf_vk->buf,
243 .size = buf_vk->size,
244 .offset = 0,
245 },
246 .bufferMemoryBarrierCount = 1,
247 });
248
249 RET(ff_vk_shader_update_desc_buffer(&s->vkctx, exec, &s->shd, 0, 2, 0,
250 buf_vk, 0, buf_vk->size,
251 VK_FORMAT_UNDEFINED));
252
253 ff_vk_exec_bind_shader(vkctx, exec, &s->shd);
254
255 vk->CmdDispatch(exec->buf,
256 FFALIGN(in->width, s->shd.lg_size[0]) / s->shd.lg_size[0],
257 FFALIGN(in->height, s->shd.lg_size[1]) / s->shd.lg_size[1],
258 s->shd.lg_size[2]);
259
260 vk->CmdPipelineBarrier2(exec->buf, &(VkDependencyInfo) {
261 .sType = VK_STRUCTURE_TYPE_DEPENDENCY_INFO,
262 .pBufferMemoryBarriers = &(VkBufferMemoryBarrier2) {
263 .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER_2,
264 .srcStageMask = VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT,
265 .dstStageMask = VK_PIPELINE_STAGE_2_HOST_BIT,
266 .srcAccessMask = VK_ACCESS_2_SHADER_STORAGE_READ_BIT |
267 VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT,
268 .dstAccessMask = VK_ACCESS_HOST_READ_BIT,
269 .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
270 .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED,
271 .buffer = buf_vk->buf,
272 .size = buf_vk->size,
273 .offset = 0,
274 },
275 .bufferMemoryBarrierCount = 1,
276 });
277
278 err = ff_vk_exec_submit(vkctx, exec);
279 if (err < 0) {
280 av_frame_free(&in);
281 av_refstruct_unref(&buf_vk);
282 return err;
283 }
284 ff_vk_exec_wait(vkctx, exec);
285 score = evaluate(ctx, sad);
286
287done:
288 snprintf(str, sizeof(str), "%0.3f", s->prev_mafd);
289 av_dict_set(&in->metadata, "lavfi.scd.mafd", str, 0);
290 snprintf(str, sizeof(str), "%0.3f", score);
291 av_dict_set(&in->metadata, "lavfi.scd.score", str, 0);
292
293 if (score >= s->threshold) {
294 const char *pts = av_ts2timestr(in->pts, &link->time_base);
295 av_dict_set(&in->metadata, "lavfi.scd.time", pts, 0);
296 av_log(ctx, AV_LOG_INFO, "lavfi.scd.score: %.3f, lavfi.scd.time: %s\n",
297 score, pts);
298 }
299
300 av_refstruct_unref(&buf_vk);
301 if (!s->sc_pass || score >= s->threshold)
302 return ff_filter_frame(outlink, in);
303 else {
304 av_frame_free(&in);
305 return 0;
306 }
307
308fail:
309 if (exec)
310 ff_vk_exec_discard(&s->vkctx, exec);
311 av_frame_free(&in);
312 av_refstruct_unref(&buf_vk);
313 return err;
314}
315
317{
319 FFVulkanContext *vkctx = &s->vkctx;
320
321 av_frame_free(&s->prev);
322 av_frame_free(&s->cur);
323
324 ff_vk_exec_pool_free(vkctx, &s->e);
325 ff_vk_shader_free(vkctx, &s->shd);
326
327 av_refstruct_pool_uninit(&s->det_buf_pool);
328
329 ff_vk_uninit(&s->vkctx);
330
331 s->initialized = 0;
332}
333
334#define OFFSET(x) offsetof(SceneDetectVulkanContext, x)
335#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
337 { "threshold", "set scene change detect threshold", OFFSET(threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 10.}, 0, 100., FLAGS },
338 { "t", "set scene change detect threshold", OFFSET(threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 10.}, 0, 100., FLAGS },
339 { "sc_pass", "Set the flag to pass scene change frames", OFFSET(sc_pass), AV_OPT_TYPE_BOOL, {.i64 = 0 }, 0, 1, FLAGS },
340 { "s", "Set the flag to pass scene change frames", OFFSET(sc_pass), AV_OPT_TYPE_BOOL, {.i64 = 0 }, 0, 1, FLAGS },
341 { NULL }
342};
343
345
347 {
348 .name = "default",
349 .type = AVMEDIA_TYPE_VIDEO,
350 .filter_frame = &scdet_vulkan_filter_frame,
351 .config_props = &ff_vk_filter_config_input,
352 },
353};
354
356 {
357 .name = "default",
358 .type = AVMEDIA_TYPE_VIDEO,
359 .config_props = &ff_vk_filter_config_output,
360 },
361};
362
364 .p.name = "scdet_vulkan",
365 .p.description = NULL_IF_CONFIG_SMALL("Detect video scene change"),
366 .p.priv_class = &scdet_vulkan_class,
367 .p.flags = AVFILTER_FLAG_HWDEVICE,
368 .priv_size = sizeof(SceneDetectVulkanContext),
374 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
375};
const FFFilter ff_vf_scdet_vulkan
static AVFormatContext * ctx
simple assert() macros that are a bit more flexible than ISO C assert().
#define av_assert2(cond)
assert() equivalent, that does lie in speed critical code.
Definition avassert.h:68
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition avfilter.c:1068
#define i(width, name, range_min, range_max)
Definition cbs_h264.c:63
#define s(width, name)
Definition cbs_vp9.c:198
#define FLAGS
Definition cmdutils.c:598
#define av_clipf
Definition common.h:145
#define NULL
Definition coverity.c:32
static __device__ float fabs(float a)
int(* init)(AVBSFContext *ctx)
Definition dts2pts.c:608
#define AV_NUM_DATA_POINTERS
Definition frame.h:473
#define fail
Definition test.h:479
@ AV_OPT_TYPE_DOUBLE
Underlying C type is double.
Definition opt.h:266
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
Definition opt.h:326
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition avfilter.h:187
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
#define AVERROR(e)
Definition error.h:45
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition frame.c:64
AVFrame * av_frame_clone(const AVFrame *src)
Create a new frame that references the same data as src.
Definition frame.c:483
#define AV_LOG_INFO
Standard information.
Definition log.h:221
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition log.h:210
@ AVMEDIA_TYPE_VIDEO
Definition avutil.h:200
static av_cold void uninit(AVBitStreamFilterContext *ctx)
#define FILTER_INPUTS(array)
Definition filters.h:264
#define FILTER_OUTPUTS(array)
Definition filters.h:265
#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
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition filters.h:254
#define AVFILTER_DEFINE_CLASS(fname)
Definition filters.h:478
#define av_cold
Definition attributes.h:117
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition internal.h:88
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:2719
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:2240
void ff_vk_shader_add_descriptor_set(FFVulkanContext *s, FFVulkanShader *shd, const FFVulkanDescriptorSetBinding *desc, int nb, int singular)
Add descriptor to a shader.
Definition vulkan.c:2550
void ff_vk_exec_pool_free(FFVulkanContext *s, FFVkExecPool *pool)
Definition vulkan.c:331
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:395
void ff_vk_exec_wait(FFVulkanContext *s, FFVkExecContext *e)
Definition vulkan.c:645
void ff_vk_uninit(FFVulkanContext *s)
Frees main context.
Definition vulkan.c:2815
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:2197
int ff_vk_exec_start(FFVulkanContext *s, FFVkExecContext *e)
Start/submit/wait an execution.
Definition vulkan.c:660
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:2128
void ff_vk_shader_free(FFVulkanContext *s, FFVulkanShader *shd)
Free a shader.
Definition vulkan.c:2791
int ff_vk_shader_register_exec(FFVulkanContext *s, FFVkExecPool *pool, FFVulkanShader *shd)
Register a shader with an exec pool.
Definition vulkan.c:2584
FFVkExecContext * ff_vk_exec_get(FFVulkanContext *s, FFVkExecPool *pool)
Retrieve an execution pool.
Definition vulkan.c:622
int ff_vk_exec_submit(FFVulkanContext *s, FFVkExecContext *e)
Definition vulkan.c:969
void ff_vk_exec_bind_shader(FFVulkanContext *s, FFVkExecContext *e, const FFVulkanShader *shd)
Bind a shader.
Definition vulkan.c:2768
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:2732
AVVulkanDeviceQueueFamily * ff_vk_qf_find(FFVulkanContext *s, VkQueueFlagBits dev_family, VkVideoCodecOperationFlagBitsKHR vid_ops)
Chooses an appropriate QF.
Definition vulkan.c:316
int ff_vk_get_pooled_buffer(FFVulkanContext *ctx, AVRefStructPool **buf_pool, FFVkBuffer **buf, VkBufferUsageFlags usage, void *create_pNext, size_t size, VkMemoryPropertyFlagBits mem_props)
Initialize a pool and create AVBufferRefs containing FFVkBuffer.
Definition vulkan.c:1426
void ff_vk_exec_discard(FFVulkanContext *s, FFVkExecContext *e)
Definition vulkan.c:763
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:2444
int ff_vk_exec_add_dep_frame(FFVulkanContext *s, FFVkExecContext *e, AVFrame *f, VkPipelineStageFlagBits2 wait_stage, VkPipelineStageFlagBits2 signal_stage)
Definition vulkan.c:867
const char * desc
Definition libsvtav1.c:83
#define FFMIN(a, b)
Definition macros.h:49
#define FFALIGN(x, a)
Definition macros.h:78
AVOptions.
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition pixdesc.c:3500
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition pixdesc.c:3460
#define AV_PIX_FMT_FLAG_RGB
The pixel format contains RGB-like data (as opposed to YUV/grayscale).
Definition pixdesc.h:136
#define AV_PIX_FMT_FLAG_PLANAR
At least one pixel component is not in the first data plane.
Definition pixdesc.h:132
@ AV_PIX_FMT_VULKAN
Vulkan hardware images.
Definition pixfmt.h:379
void av_refstruct_unref(void *objp)
Decrement the reference count of the underlying object and automatically free the object if there are...
Definition refstruct.c:120
static void av_refstruct_pool_uninit(AVRefStructPool **poolp)
Mark the pool as being available for freeing.
Definition refstruct.h:292
#define snprintf
Definition snprintf.h:34
An instance of a filter.
Definition avfilter.h:273
void * priv
private data for use by the filter
Definition avfilter.h:288
A filter pad used for either input or output.
Definition filters.h:40
This structure describes decoded (raw) audio or video data.
Definition frame.h:472
AVOption.
Definition opt.h:428
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition pixdesc.h:69
uint64_t flags
Combination of AV_PIX_FMT_FLAG_... flags.
Definition pixdesc.h:94
AVRefStructPool is an API for a thread-safe pool of objects managed via the RefStruct API.
Definition refstruct.c:183
uint8_t * mapped_mem
Definition vulkan.h:103
VkCommandBuffer buf
Definition vulkan.h:142
FFVulkanFunctions vkfn
Definition vulkan.h:298
uint32_t frame_sad[SLICES]
AVVulkanDeviceQueueFamily * qf
AVRefStructPool * det_buf_pool
#define av_log(a,...)
timestamp utils, mostly useful for debugging/logging purposes
#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
static int64_t pts
#define SLICES
static av_always_inline int diff(const struct color_info *a, const struct color_info *b, const int trans_thresh)
static const AVFilterPad scdet_vulkan_inputs[]
static void scdet_vulkan_uninit(AVFilterContext *avctx)
static av_cold int init_filter(AVFilterContext *ctx)
static int scdet_vulkan_filter_frame(AVFilterLink *link, AVFrame *in)
#define SLICES
const unsigned char ff_scdet_comp_spv_data[]
static double evaluate(AVFilterContext *ctx, const SceneDetectBuf *buf)
const unsigned int ff_scdet_comp_spv_len
static const AVOption scdet_vulkan_options[]
static const AVFilterPad scdet_vulkan_outputs[]
#define OFFSET(x)
@ FF_VK_REP_UINT
Definition vulkan.h:445
#define FF_VK_DEFAULT_EXEC_CONTEXTS
Definition vulkan.h:121
#define RET(x)
Definition vulkan.h:37
#define SPEC_LIST_ADD(name, idx, val_bits, val)
Definition vulkan.h:55
#define SPEC_LIST_CREATE(name, max_length, max_size)
Definition vulkan.h:45
int ff_vk_filter_config_input(AVFilterLink *inlink)
int ff_vk_filter_config_output(AVFilterLink *outlink)
int ff_vk_filter_init(AVFilterContext *avctx)
General lavfi IO functions.