Phase 3a: BORDER overlays через 4 fill_nv12 операции

Live-validated на rtsp://192.168.88.23:554/cfc-grid: цветные рамки
вокруг каждой ячейки реализованы через 4 cfc_cugrid_fill_nv12
(top/bottom/left/right) — отдельного kernel'я не понадобилось,
переиспользуем existing fill_nv12 для region α-blend.

Содержимое:

- include/cuframes_composer/overlay.h — opaque cfc_overlay_t + типы
  (BORDER реализован, PNG/TEXT skeleton'ы для Phase 3b/3c).
- src/overlay.c — реализация BORDER: clamp rect в границы кадра,
  выравнивание координат на чётные (4:2:0 chroma требование),
  thickness clamp если 2*thickness > w/h.
- composer.c — список overlays (max 64), z-order = порядок добавления,
  draw поверх grid'а перед stream sync.
- examples/grid_record — флаг --border N (толщина пикселей) добавляет
  серо-голубой border (Y=180,U=120,V=110 alpha=220) для каждой
  ячейки автоматически.

Phase 3b (PNG icons через stb_image + cugrid alpha_blit_rgba) и Phase
3c (text через FreeType + font atlas) — отдельные commit'ы.

examples/grid_record.c

@@ -20,6 +20,7 @@
 
 #include "../include/cuframes_composer/composer.h"
 #include "../include/cuframes_composer/nvenc.h"
+#include "../include/cuframes_composer/overlay.h"
 
 #include <cuda.h>
 
@@ -89,6 +90,7 @@ int main(int argc, char **argv)
     const char *out_path = NULL;
     int fps = 25, bitrate = 10000, max_seconds = 0;
     int out_w = 3840, out_h = 2160;
+    int border_thickness = 0;          /* 0 = без border'ов */
     cfc_composer_cell_t cells[MAX_CELLS] = { 0 };
     static char cell_keys[MAX_CELLS][64];
     int num_cells = 0;
@@ -101,10 +103,11 @@ int main(int argc, char **argv)
         {"width",   required_argument, 0, 'W'},
         {"height",  required_argument, 0, 'H'},
         {"seconds", required_argument, 0, 's'},
+        {"border",  required_argument, 0, 'r'},  /* толщина border'ов */
         {0, 0, 0, 0},
     };
     int c;
-    while ((c = getopt_long(argc, argv, "o:c:f:b:W:H:s:", opts, NULL)) != -1) {
+    while ((c = getopt_long(argc, argv, "o:c:f:b:W:H:s:r:", opts, NULL)) != -1) {
         switch (c) {
         case 'o': out_path = optarg; break;
         case 'c':
@@ -123,6 +126,7 @@ int main(int argc, char **argv)
         case 'W': out_w = atoi(optarg); break;
         case 'H': out_h = atoi(optarg); break;
         case 's': max_seconds = atoi(optarg); break;
+        case 'r': border_thickness = atoi(optarg); break;
         default: return 1;
         }
     }
@@ -163,6 +167,27 @@ int main(int argc, char **argv)
     fprintf(stderr, "[grid_record] composer %dx%d, %d ячеек\n",
             out_w, out_h, num_cells);
 
+    /* Border'ы вокруг каждой ячейки если --border задан.
+     * Цвет: серо-голубой (BT.709 limited): Y=180, U=120, V=110. */
+    if (border_thickness > 0) {
+        for (int i = 0; i < num_cells; i++) {
+            cfc_overlay_border_config_t bc = {
+                .x = cells[i].x, .y = cells[i].y,
+                .w = cells[i].w, .h = cells[i].h,
+                .thickness = border_thickness,
+                .color_y = 180, .color_u = 120, .color_v = 110,
+                .alpha = 220,
+                .visible = 1,
+            };
+            cfc_overlay_t *ov = NULL;
+            if (cfc_overlay_create_border(&bc, &ov) == 0) {
+                cfc_composer_add_overlay(comp, ov);
+            }
+        }
+        fprintf(stderr, "[grid_record] добавлены border'ы (толщина %d) для %d ячеек\n",
+                border_thickness, num_cells);
+    }
+
     /* Encoder. */
     cfc_encoder_config_t ecfg = {
         .cuda_ctx = ctx,

include/cuframes_composer/composer.h

@@ -25,6 +25,7 @@
 #define CUFRAMES_COMPOSER_COMPOSER_H
 
 #include "source.h"
+#include "overlay.h"
 
 #include <cuda.h>
 #include <stdint.h>
@@ -78,6 +79,11 @@ int cfc_composer_compose(
     int *out_height
 );
 
+/* Зарегистрировать overlay (border/png/text) для отрисовки поверх grid'а.
+ * composer takes ownership — на cfc_composer_destroy всё освобождается.
+ * Порядок добавления = z-order (последний рисуется поверх). Лимит — 64. */
+int cfc_composer_add_overlay(cfc_composer_t *comp, cfc_overlay_t *ov);
+
 /* Получить layout статистику по источникам — для debug / health-репортов. */
 typedef struct cfc_composer_health {
     int total;                        /* всего источников */

include/cuframes_composer/overlay.h

@@ -0,0 +1,85 @@
+/* cuframes-composer — overlay'и (border, png-иконка, текст) поверх
+ * композиции.
+ *
+ * Архитектура:
+ *   cfc_overlay_t — opaque объект одного overlay'я (border, png, text).
+ *   composer ведёт список overlays + порядок (z-order); на каждом compose'е
+ *   проходит по списку и рисует их поверх NV12 grid'а.
+ *
+ * Типы overlay'ев:
+ *   BORDER — цветная рамка вокруг rect'а (Phase 3a). Реализован через 4
+ *            fill_nv12 операции (top/bottom/left/right). Полезно для:
+ *              - визуальной отбивки ячеек grid'а
+ *              - подсветки активных/неактивных источников (красная при DEAD)
+ *   PNG    — RGBA-изображение (Phase 3b). Decode-раз через stb_image,
+ *            upload в VRAM, на каждом кадре alpha-blit через
+ *            cfc_kern_alpha_blit_rgba_*. Для статичных иконок (NO SIGNAL,
+ *            offline, recording dot).
+ *   TEXT   — динамический рендер строки (Phase 3c). Требует FreeType +
+ *            font atlas. Сделать позже когда FreeType добавлен как зависимость.
+ *
+ * Жизненный цикл:
+ *   cfc_overlay_create_*() — выделяет ресурсы (для PNG/TEXT — VRAM-атлас).
+ *   cfc_overlay_update_*() — обновляет параметры (положение, видимость, текст).
+ *   cfc_overlay_draw()      — вызывается composer'ом на каждом кадре.
+ *   cfc_overlay_destroy()   — освобождает VRAM.
+ *
+ * Phase 3a лимит: только BORDER через fill_nv12.
+ *
+ * Лицензия: LGPL-2.1+
+ */
+
+#ifndef CUFRAMES_COMPOSER_OVERLAY_H
+#define CUFRAMES_COMPOSER_OVERLAY_H
+
+#include <cuda.h>
+#include <stdint.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef enum cfc_overlay_type {
+    CFC_OVERLAY_BORDER = 0,
+    CFC_OVERLAY_PNG    = 1,   /* Phase 3b */
+    CFC_OVERLAY_TEXT   = 2,   /* Phase 3c */
+} cfc_overlay_type_t;
+
+typedef struct cfc_overlay cfc_overlay_t;
+
+/* Параметры BORDER overlay'я. */
+typedef struct cfc_overlay_border_config {
+    int x, y, w, h;                   /* прямоугольник в full-res пикселях */
+    int thickness;                    /* толщина рамки в пикселях */
+    int color_y, color_u, color_v;    /* цвет в BT.709 (Y:16-235, UV:16-240) */
+    int alpha;                        /* 0..255; 0 = invisible */
+    int visible;                      /* 0/1 — выводить ли вообще */
+} cfc_overlay_border_config_t;
+
+/* Создать BORDER overlay. */
+int cfc_overlay_create_border(const cfc_overlay_border_config_t *cfg,
+                                cfc_overlay_t **out);
+
+/* Обновить параметры BORDER overlay'я (можно переключить visible,
+ * сменить цвет, изменить позицию). Thread-safe? Нет — caller должен сам
+ * заботиться о том, чтобы update не пересекался с draw. В рамках одного
+ * compose-thread'а это естественно. */
+int cfc_overlay_update_border(cfc_overlay_t *ov,
+                                const cfc_overlay_border_config_t *cfg);
+
+/* Нарисовать overlay поверх NV12 кадра. Вызывается из cfc_composer_compose.
+ * Возвращает 0 при успехе. */
+int cfc_overlay_draw(cfc_overlay_t *ov,
+                      CUstream stream,
+                      CUdeviceptr dst_y, int pitch_y,
+                      CUdeviceptr dst_uv, int pitch_uv,
+                      int frame_w, int frame_h);
+
+/* Уничтожить overlay (освободить VRAM если был). */
+int cfc_overlay_destroy(cfc_overlay_t *ov);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* CUFRAMES_COMPOSER_OVERLAY_H */

src/CMakeLists.txt

@@ -15,6 +15,7 @@ set(COMPOSER_SOURCES_C
     nvenc_loader.c
     nvenc.c
     composer.c
+    overlay.c
 )
 set(COMPOSER_SOURCES_CU
     cugrid/cugrid.cu

src/composer.c

@@ -25,6 +25,7 @@
 
 #include "../include/cuframes_composer/composer.h"
 #include "../include/cuframes_composer/cugrid.h"
+#include "../include/cuframes_composer/overlay.h"
 
 #include <cuda_runtime.h>
 #include <stdio.h>
@@ -33,6 +34,7 @@
 #include <time.h>
 
 #define CFC_COMPOSER_MAX_CELLS 64
+#define CFC_COMPOSER_MAX_OVERLAYS 64
 
 struct cfc_composer {
     cfc_composer_config_t cfg;
@@ -55,6 +57,10 @@ struct cfc_composer {
 
     /* CUDA stream для compose (Phase 2 — default stream = 0). */
     cudaStream_t stream;
+
+    /* Overlays — в порядке добавления (= z-order). composer take ownership. */
+    cfc_overlay_t *overlays[CFC_COMPOSER_MAX_OVERLAYS];
+    int num_overlays;
 };
 
 /* ── Helpers ──────────────────────────────────────────────────────────── */
@@ -223,6 +229,20 @@ int cfc_composer_compose(cfc_composer_t *comp,
         }
     }
 
+    /* Overlays поверх grid'а — в порядке добавления. */
+    CUdeviceptr uv = comp->output_ptr +
+                     (size_t)comp->output_pitch_y * comp->cfg.height;
+    for (int i = 0; i < comp->num_overlays; i++) {
+        if (cfc_overlay_draw(comp->overlays[i],
+                              (CUstream)comp->stream,
+                              comp->output_ptr, comp->output_pitch_y,
+                              uv, comp->output_pitch_uv,
+                              comp->cfg.width, comp->cfg.height) != 0) {
+            fprintf(stderr, "[cfc/composer] overlay %d draw failed\n", i);
+            /* Не fatal. */
+        }
+    }
+
     cudaError_t e = cudaStreamSynchronize(comp->stream);
     if (e != cudaSuccess) {
         fprintf(stderr, "[cfc/composer] stream sync failed: %s\n",
@@ -237,6 +257,18 @@ int cfc_composer_compose(cfc_composer_t *comp,
     return 0;
 }
 
+int cfc_composer_add_overlay(cfc_composer_t *comp, cfc_overlay_t *ov)
+{
+    if (!comp || !ov) return -1;
+    if (comp->num_overlays >= CFC_COMPOSER_MAX_OVERLAYS) {
+        fprintf(stderr, "[cfc/composer] overlay limit %d reached\n",
+                CFC_COMPOSER_MAX_OVERLAYS);
+        return -1;
+    }
+    comp->overlays[comp->num_overlays++] = ov;
+    return 0;
+}
+
 int cfc_composer_get_health(cfc_composer_t *comp, cfc_composer_health_t *out)
 {
     if (!comp || !out) return -1;
@@ -264,6 +296,9 @@ int cfc_composer_destroy(cfc_composer_t *comp)
     for (int i = 0; i < comp->num_cells; i++) {
         if (comp->sources[i]) cfc_source_destroy(comp->sources[i]);
     }
+    for (int i = 0; i < comp->num_overlays; i++) {
+        cfc_overlay_destroy(comp->overlays[i]);
+    }
     if (comp->output_ptr) cuMemFree(comp->output_ptr);
     free(comp);
     return 0;

src/overlay.c

@@ -0,0 +1,132 @@
+/* Реализация cfc_overlay_t — Phase 3a: только BORDER overlays.
+ *
+ * BORDER рисуется как 4 fill_nv12 операции:
+ *   top:    rect (x,            y,            w, thickness)
+ *   bottom: rect (x,            y + h - thick, w, thickness)
+ *   left:   rect (x,            y + thick,    thickness, h - 2*thick)
+ *   right:  rect (x + w - thick, y + thick,   thickness, h - 2*thick)
+ *
+ * Это работает на уже скомпонованном NV12 кадре, не требует отдельного kernel'я
+ * (используем cfc_cugrid_fill_nv12).
+ *
+ * Phase 3b/3c (PNG/TEXT) — потом, требует доп. ресурсов в VRAM.
+ *
+ * Лицензия: LGPL-2.1+
+ */
+
+#include "../include/cuframes_composer/overlay.h"
+#include "../include/cuframes_composer/cugrid.h"
+
+#include <stdlib.h>
+#include <string.h>
+
+struct cfc_overlay {
+    cfc_overlay_type_t type;
+    union {
+        cfc_overlay_border_config_t border;
+        /* PNG / TEXT — Phase 3b/3c, добавится позже. */
+    } u;
+};
+
+/* ── BORDER ───────────────────────────────────────────────────────────── */
+
+int cfc_overlay_create_border(const cfc_overlay_border_config_t *cfg,
+                                cfc_overlay_t **out)
+{
+    if (!cfg || !out) return -1;
+    if (cfg->w <= 0 || cfg->h <= 0) return -1;
+    if (cfg->thickness <= 0) return -1;
+
+    cfc_overlay_t *ov = calloc(1, sizeof(*ov));
+    if (!ov) return -1;
+    ov->type = CFC_OVERLAY_BORDER;
+    ov->u.border = *cfg;
+    if (ov->u.border.alpha == 0) ov->u.border.alpha = 255;
+    *out = ov;
+    return 0;
+}
+
+int cfc_overlay_update_border(cfc_overlay_t *ov,
+                                const cfc_overlay_border_config_t *cfg)
+{
+    if (!ov || !cfg) return -1;
+    if (ov->type != CFC_OVERLAY_BORDER) return -1;
+    ov->u.border = *cfg;
+    return 0;
+}
+
+static int draw_border(cfc_overlay_t *ov,
+                         CUstream stream,
+                         CUdeviceptr dst_y, int pitch_y,
+                         CUdeviceptr dst_uv, int pitch_uv,
+                         int frame_w, int frame_h)
+{
+    const cfc_overlay_border_config_t *b = &ov->u.border;
+    if (!b->visible) return 0;
+    if (b->alpha <= 0) return 0;
+
+    int t = b->thickness;
+    int x = b->x, y = b->y, w = b->w, h = b->h;
+
+    /* Clamp rect в границы кадра. */
+    if (x < 0) { w += x; x = 0; }
+    if (y < 0) { h += y; y = 0; }
+    if (x + w > frame_w) w = frame_w - x;
+    if (y + h > frame_h) h = frame_h - y;
+    if (w <= 0 || h <= 0) return 0;
+    if (t * 2 > w) t = w / 2;
+    if (t * 2 > h) t = h / 2;
+    /* Все координаты выравниваем на чётные (требование 4:2:0 chroma). */
+    x &= ~1; y &= ~1; w &= ~1; h &= ~1; t &= ~1;
+    if (t == 0) t = 2;
+
+    /* Top */
+    if (cfc_cugrid_fill_nv12(stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                              x, y, w, t,
+                              b->color_y, b->color_u, b->color_v, b->alpha) != 0)
+        return -1;
+    /* Bottom */
+    if (cfc_cugrid_fill_nv12(stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                              x, y + h - t, w, t,
+                              b->color_y, b->color_u, b->color_v, b->alpha) != 0)
+        return -1;
+    /* Left */
+    if (cfc_cugrid_fill_nv12(stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                              x, y + t, t, h - 2 * t,
+                              b->color_y, b->color_u, b->color_v, b->alpha) != 0)
+        return -1;
+    /* Right */
+    if (cfc_cugrid_fill_nv12(stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                              x + w - t, y + t, t, h - 2 * t,
+                              b->color_y, b->color_u, b->color_v, b->alpha) != 0)
+        return -1;
+    return 0;
+}
+
+/* ── Public dispatch ─────────────────────────────────────────────────── */
+
+int cfc_overlay_draw(cfc_overlay_t *ov,
+                      CUstream stream,
+                      CUdeviceptr dst_y, int pitch_y,
+                      CUdeviceptr dst_uv, int pitch_uv,
+                      int frame_w, int frame_h)
+{
+    if (!ov) return -1;
+    switch (ov->type) {
+    case CFC_OVERLAY_BORDER:
+        return draw_border(ov, stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                            frame_w, frame_h);
+    case CFC_OVERLAY_PNG:
+    case CFC_OVERLAY_TEXT:
+        /* Phase 3b/3c — пока no-op. */
+        return 0;
+    }
+    return -1;
+}
+
+int cfc_overlay_destroy(cfc_overlay_t *ov)
+{
+    if (!ov) return 0;
+    free(ov);
+    return 0;
+}