79 (
int []) { 32, 32, 1 }, 0);
83 .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
84 .stages = VK_SHADER_STAGE_COMPUTE_BIT,
88 .type = VK_DESCRIPTOR_TYPE_STORAGE_IMAGE,
89 .stages = VK_SHADER_STAGE_COMPUTE_BIT,
93 .type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER,
94 .stages = VK_SHADER_STAGE_COMPUTE_BIT,
124 count =
s->nb_planes * inlink->
w * inlink->
h;
125 mafd = (
double) sad * 100.0 / count / (1ULL <<
desc->comp[0].depth);
141 VkImageMemoryBarrier2 img_bar[8];
163 VK_BUFFER_USAGE_TRANSFER_DST_BIT |
164 VK_BUFFER_USAGE_STORAGE_BUFFER_BIT,
167 VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT |
168 VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT |
169 VK_MEMORY_PROPERTY_HOST_COHERENT_BIT));
181 VK_PIPELINE_STAGE_2_NONE,
182 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
186 VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
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);
196 VK_PIPELINE_STAGE_2_NONE,
197 VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT));
201 VK_IMAGE_LAYOUT_GENERAL, VK_NULL_HANDLE);
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);
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,
224 .bufferMemoryBarrierCount = 1,
227 vk->CmdFillBuffer(exec->buf, buf_vk->buf, 0, buf_vk->size, 0x0);
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,
246 .bufferMemoryBarrierCount = 1,
250 buf_vk, 0, buf_vk->size,
251 VK_FORMAT_UNDEFINED));
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],
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,
275 .bufferMemoryBarrierCount = 1,
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);
293 if (score >=
s->threshold) {
301 if (!
s->sc_pass || score >=
s->threshold)
334#define OFFSET(x) offsetof(SceneDetectVulkanContext, x)
335#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
364 .p.name =
"scdet_vulkan",
366 .p.priv_class = &scdet_vulkan_class,
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.
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
#define i(width, name, range_min, range_max)
static __device__ float fabs(float a)
int(* init)(AVBSFContext *ctx)
#define AV_NUM_DATA_POINTERS
@ AV_OPT_TYPE_DOUBLE
Underlying C type is double.
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
int av_dict_set(AVDictionary **pm, const char *key, const char *value, int flags)
Set the given entry in *pm, overwriting an existing entry.
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
AVFrame * av_frame_clone(const AVFrame *src)
Create a new frame that references the same data as src.
#define AV_LOG_INFO
Standard information.
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
static av_cold void uninit(AVBitStreamFilterContext *ctx)
#define FILTER_INPUTS(array)
#define FILTER_OUTPUTS(array)
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
#define AVFILTER_DEFINE_CLASS(fname)
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
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.
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.
void ff_vk_shader_add_descriptor_set(FFVulkanContext *s, FFVulkanShader *shd, const FFVulkanDescriptorSetBinding *desc, int nb, int singular)
Add descriptor to a shader.
void ff_vk_exec_pool_free(FFVulkanContext *s, FFVkExecPool *pool)
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.
void ff_vk_exec_wait(FFVulkanContext *s, FFVkExecContext *e)
void ff_vk_uninit(FFVulkanContext *s)
Frees main context.
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)
int ff_vk_exec_start(FFVulkanContext *s, FFVkExecContext *e)
Start/submit/wait an execution.
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.
void ff_vk_shader_free(FFVulkanContext *s, FFVulkanShader *shd)
Free a shader.
int ff_vk_shader_register_exec(FFVulkanContext *s, FFVkExecPool *pool, FFVulkanShader *shd)
Register a shader with an exec pool.
FFVkExecContext * ff_vk_exec_get(FFVulkanContext *s, FFVkExecPool *pool)
Retrieve an execution pool.
int ff_vk_exec_submit(FFVulkanContext *s, FFVkExecContext *e)
void ff_vk_exec_bind_shader(FFVulkanContext *s, FFVkExecContext *e, const FFVulkanShader *shd)
Bind a shader.
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.
AVVulkanDeviceQueueFamily * ff_vk_qf_find(FFVulkanContext *s, VkQueueFlagBits dev_family, VkVideoCodecOperationFlagBitsKHR vid_ops)
Chooses an appropriate QF.
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.
void ff_vk_exec_discard(FFVulkanContext *s, FFVkExecContext *e)
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.
int ff_vk_exec_add_dep_frame(FFVulkanContext *s, FFVkExecContext *e, AVFrame *f, VkPipelineStageFlagBits2 wait_stage, VkPipelineStageFlagBits2 signal_stage)
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
#define AV_PIX_FMT_FLAG_RGB
The pixel format contains RGB-like data (as opposed to YUV/grayscale).
#define AV_PIX_FMT_FLAG_PLANAR
At least one pixel component is not in the first data plane.
@ AV_PIX_FMT_VULKAN
Vulkan hardware images.
void av_refstruct_unref(void *objp)
Decrement the reference count of the underlying object and automatically free the object if there are...
static void av_refstruct_pool_uninit(AVRefStructPool **poolp)
Mark the pool as being available for freeing.
void * priv
private data for use by the filter
A link between two filters.
int w
agreed upon image width
int h
agreed upon image height
AVFilterContext * dst
dest filter
A filter pad used for either input or output.
This structure describes decoded (raw) audio or video data.
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
uint64_t flags
Combination of AV_PIX_FMT_FLAG_... flags.
AVRefStructPool is an API for a thread-safe pool of objects managed via the RefStruct API.
uint32_t frame_sad[SLICES]
AVVulkanDeviceQueueFamily * qf
AVRefStructPool * det_buf_pool
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...
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)
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 FF_VK_DEFAULT_EXEC_CONTEXTS
#define SPEC_LIST_ADD(name, idx, val_bits, val)
#define SPEC_LIST_CREATE(name, max_length, max_size)
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.