gx/ffmpeg-patched
1
0
Fork 0
Code Issues Pull Requests 1 Actions Packages Projects Releases Wiki Activity

vf_cuda_grid: + placeholder_icon / placeholder_timeout_ms — Level 2 resilience

Browse Source 
Закрывает архитектурный недостаток: камера / cuframes-pub умирает →
pipeline filter показывает "застывшую картинку" этого input (framesync
EXT_INFINITY mode), TV видит зависшее изображение и не понимает что не так.

Now: per-input wall-clock detection of frame staleness. Если PTS не меняется
> placeholder_timeout_ms (default 2000), filter:
  1. Black-fill cell rect (Alpha_Fill_Y/UV kernels)
  2. Blit placeholder atlas (default "offline" icon, loaded из icon_dir)
     centered внутри cell

Pipeline продолжает работать. Остальные cells normal video. User видит
explicit "НЕТ СИГНАЛА" placeholder вместо stuck frame.

Options:
  placeholder_timeout_ms (default 2000, 0=disable)
  placeholder_icon (default "offline" — resolved через icon_dir)

Per-input state:
  last_frame_pts[MAX_CELLS] — last seen PTS per pad
  last_change_wall_us[MAX_CELLS] — wall-clock когда PTS changed

Тест: docker stop cuframes-pub-back_yard → через ~2s в back_yard cell
появляется "НЕТ СИГНАЛА" placeholder, остальные 3 cells продолжают
normal video с 25 fps.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
gx
2026-05-25 11:24:12 +01:00
parent d7fd75657d
commit 4795b7a7f1
1 changed files with 100 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
libavfilter/vf_cuda_grid.c
+100 -1
View File
@@ -39,6 +39,7 @@
#include "libavutil/hwcontext.h"
#include "libavutil/hwcontext_cuda_internal.h"
#include "libavutil/log.h"
#include "libavutil/time.h"
#include "libavutil/mem.h"
#include "libavutil/opt.h"
#include "libavutil/pixdesc.h"
@@ -242,6 +243,12 @@ typedef struct CudaGridContext {
char *icon_dir;
int max_cells; /* Phase 7: число input pads creates в init() */
int aspect_mode; /* 0=stretch, 1=fit (letterbox/pillarbox) */
int placeholder_timeout_ms; /* Stale input → blit placeholder. 0=disable */
char *placeholder_icon; /* icon name (resolved via icon_dir) для placeholder */
/* Per-input stale detection (wall-clock based) */
int64_t last_frame_pts[MAX_CELLS];
int64_t last_change_wall_us[MAX_CELLS];
/* Resolved layout (после init) */
const LayoutTemplate *layout;
@@ -1296,6 +1303,7 @@ static int cuda_grid_compose(FFFrameSync *fs)
if (ret < 0)
goto fail;
int64_t now_us = av_gettime();
for (i = 0; i < nb; i++) {
int pad = active_cell_map[i];
AVFrame *src;
@@ -1303,12 +1311,99 @@ static int cuda_grid_compose(FFFrameSync *fs)
int cy = active_px[i].y;
int cw = active_px[i].w;
int ch = active_px[i].h;
int is_stale = 0;
int64_t age_us;
int dx, dy, dw, dh;
if (pad < 0 || pad >= s->max_cells) pad = i; /* fallback identity */
src = in[pad];
/* Stale input detection — wall-clock based.
* Framesync mode EXT_INFINITY возвращает last frame даже после EOF,
* так src->pts не advance при mёртвом publisher. Detect "frozen" по
* wall-clock возрасту с last PTS change → blit placeholder вместо
* stale frame. Закрывает архитектурный недостаток где одна камера
* dying показывает "застывшую картинку" — теперь "Нет сигнала". */
if (s->placeholder_timeout_ms > 0) {
if (src->pts != s->last_frame_pts[pad]) {
s->last_frame_pts[pad] = src->pts;
s->last_change_wall_us[pad] = now_us;
}
age_us = now_us - s->last_change_wall_us[pad];
if (s->last_change_wall_us[pad] != 0 &&
age_us > (int64_t)s->placeholder_timeout_ms * 1000) {
is_stale = 1;
}
}
if (is_stale) {
/* Black-fill cell + blit placeholder icon (если loaded).
* Если placeholder PNG отсутствует — просто черный cell. */
int fill_y_val = 16, fill_u = 128, fill_v = 128, alpha = 255;
CUdeviceptr d_y_fill = (CUdeviceptr)out->data[0];
CUdeviceptr d_uv_fill = (CUdeviceptr)out->data[1];
int dp_y_fill = out->linesize[0];
int dp_uv_fill = out->linesize[1];
int half_cx = cx / 2, half_cy = cy / 2;
int half_cw = cw / 2, half_ch = ch / 2;
IconAtlas *pa = NULL;
{
void *args[] = { &d_y_fill, &dp_y_fill,
&cx, &cy, &cw, &ch, &fill_y_val, &alpha };
CHECK_CU(s->hwctx->internal->cuda_dl->cuLaunchKernel(
s->cu_func_alpha_fill_y,
DIV_UP(cw, BLOCKX), DIV_UP(ch, BLOCKY), 1,
BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
}
{
void *args[] = { &d_uv_fill, &dp_uv_fill,
&half_cx, &half_cy, &half_cw, &half_ch,
&fill_u, &fill_v, &alpha };
CHECK_CU(s->hwctx->internal->cuda_dl->cuLaunchKernel(
s->cu_func_alpha_fill_uv,
DIV_UP(half_cw, BLOCKX), DIV_UP(half_ch, BLOCKY), 1,
BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
}
/* Blit placeholder centered (если loaded). Используем существующий
* icon atlas machinery — placeholder = icon с именем s->placeholder_icon. */
if (s->placeholder_icon && s->placeholder_icon[0]) {
if (ensure_icon_atlas(ctx, s->placeholder_icon, &pa) == 0 && pa) {
int px = cx + (cw - pa->w) / 2;
int py = cy + (ch - pa->h) / 2;
px &= ~1; py &= ~1;
if (px >= cx && py >= cy &&
px + pa->w <= cx + cw && py + pa->h <= cy + ch) {
int aw = pa->w, ah = pa->h, ealpha = 255;
size_t ap = pa->device_pitch;
CUdeviceptr atlas_dev = pa->device_ptr;
int dp_y_b = out->linesize[0];
int dp_uv_b = out->linesize[1];
CUdeviceptr dst_y_b = (CUdeviceptr)out->data[0];
CUdeviceptr dst_uv_b = (CUdeviceptr)out->data[1];
{
void *args[] = { &dst_y_b, &dp_y_b, &px, &py,
&atlas_dev, &ap, &aw, &ah, &ealpha };
CHECK_CU(s->hwctx->internal->cuda_dl->cuLaunchKernel(
s->cu_func_alpha_blit_rgba_y,
DIV_UP(aw, BLOCKX), DIV_UP(ah, BLOCKY), 1,
BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
}
{
void *args[] = { &dst_uv_b, &dp_uv_b, &px, &py,
&atlas_dev, &ap, &aw, &ah, &ealpha };
CHECK_CU(s->hwctx->internal->cuda_dl->cuLaunchKernel(
s->cu_func_alpha_blit_rgba_uv,
DIV_UP(aw / 2, BLOCKX), DIV_UP(ah / 2, BLOCKY), 1,
BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
}
}
}
}
continue; /* skip normal scale path */
}
/* aspect_mode=1 (fit) — compute inscribed rect preserving aspect.
* Remainder of cell стане letterbox/pillarbox (black). */
int dx = cx, dy = cy, dw = cw, dh = ch; /* default — stretch к cell */
dx = cx; dy = cy; dw = cw; dh = ch; /* default — stretch к cell */
if (s->aspect_mode == 1 && src->width > 0 && src->height > 0) {
float src_ar = (float)src->width / (float)src->height;
float cell_ar = (float)cw / (float)ch;
@@ -1848,6 +1943,10 @@ static const AVOption cuda_grid_options[] = {
OFFSET(max_cells), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, MAX_CELLS, FLAGS },
{ "aspect_mode", "0=stretch (default, fill cell, может исказить), 1=fit (letterbox/pillarbox, preserve aspect, black borders)",
OFFSET(aspect_mode), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 1, FLAGS },
{ "placeholder_timeout_ms", "если input pad не дал new frame N ms — blit placeholder (default 2000, 0=disable)",
OFFSET(placeholder_timeout_ms), AV_OPT_TYPE_INT, { .i64 = 2000 }, 0, 60000, FLAGS },
{ "placeholder_icon", "icon name (через icon_dir) для placeholder при stale input. Default 'offline'",
OFFSET(placeholder_icon), AV_OPT_TYPE_STRING, { .str = "offline" }, 0, 0, FLAGS },
{ NULL }
};