diff --git a/libavfilter/vf_cuda_grid.c b/libavfilter/vf_cuda_grid.c index b95b606..aaa2a31 100644 --- a/libavfilter/vf_cuda_grid.c +++ b/libavfilter/vf_cuda_grid.c @@ -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 } };