diff --git a/examples/grid_record.c b/examples/grid_record.c index 50a9f55..65d41b7 100644 --- a/examples/grid_record.c +++ b/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 @@ -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, diff --git a/include/cuframes_composer/composer.h b/include/cuframes_composer/composer.h index 12d350a..086d32d 100644 --- a/include/cuframes_composer/composer.h +++ b/include/cuframes_composer/composer.h @@ -25,6 +25,7 @@ #define CUFRAMES_COMPOSER_COMPOSER_H #include "source.h" +#include "overlay.h" #include #include @@ -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; /* всего источников */ diff --git a/include/cuframes_composer/overlay.h b/include/cuframes_composer/overlay.h new file mode 100644 index 0000000..b34c414 --- /dev/null +++ b/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 +#include + +#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 */ diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 17867e1..548b848 100644 --- a/src/CMakeLists.txt +++ b/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 diff --git a/src/composer.c b/src/composer.c index 3b999d0..1ae96ee 100644 --- a/src/composer.c +++ b/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 #include @@ -33,6 +34,7 @@ #include #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; diff --git a/src/overlay.c b/src/overlay.c new file mode 100644 index 0000000..fe24794 --- /dev/null +++ b/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 +#include + +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; +}