Phase 3b: PNG icon overlays через libpng + alpha_blit_rgba_nv12

Live-validated на rtsp://192.168.88.23:554/cfc-grid с иконками
temp_outside.png и offline.png из cuda_grid_icons volume.

Содержимое:

- cugrid.h/cugrid.cu — cfc_cugrid_blit_rgba_nv12 (Y+UV α-blend) +
  два новых kernel'я blit_rgba_y/uv (BT.709 limited-range RGB→YUV
  conversion, 4:2:0 chroma 2x2 averaging).
- overlay.h — cfc_overlay_create_png + cfc_overlay_png_size +
  cfc_overlay_update_png.
- overlay.c — libpng decode (paletted/gray/16-bit → 8-bit RGBA),
  cuMemAlloc atlas в VRAM, cuMemcpyHtoD один раз, draw_png через
  cfc_cugrid_blit_rgba_nv12.
- CMakeLists.txt — find_package(PNG REQUIRED) + PNG::PNG в link.
- examples/grid_record — флаг --icon path,x,y[,alpha] (несколько раз
  можно). Атлас грузится один раз при старте.

PNG из cuda_grid_icons volume переиспользуются (offline, temp_outside,
grafana_gpu, и пр.). PNG decode'ятся одинаково — paletted, RGBA, gray.

Phase 3c (text rendering через FreeType + font atlas) — отдельный
commit. Атлас text'а тоже окажется в VRAM как RGBA через тот же
cfc_cugrid_blit_rgba_nv12 — kernels уже готовы.

CMakeLists.txt

@@ -35,6 +35,9 @@ find_package(Threads REQUIRED)
 # dl — для dlopen libnvidia-encode.so в runtime
 find_library(LIBDL_LIBRARY dl REQUIRED)
 
+# PNG — для декода RGBA-иконок overlay'ев (Phase 3b)
+find_package(PNG REQUIRED)
+
 # ── Сторонние библиотеки (subomodules в third_party/) ───────────────────
 
 # cuframes — статически линкуем libcuframes. cuframes_static — это static lib

examples/grid_record.c

@@ -95,6 +95,11 @@ int main(int argc, char **argv)
     static char cell_keys[MAX_CELLS][64];
     int num_cells = 0;
 
+    /* --icon path,x,y[,alpha] */
+    typedef struct { const char *path; int x, y, alpha; } icon_spec_t;
+    icon_spec_t icons[MAX_CELLS] = { 0 };
+    int num_icons = 0;
+
     static struct option opts[] = {
         {"out",     required_argument, 0, 'o'},
         {"cell",    required_argument, 0, 'c'},
@@ -104,10 +109,11 @@ int main(int argc, char **argv)
         {"height",  required_argument, 0, 'H'},
         {"seconds", required_argument, 0, 's'},
         {"border",  required_argument, 0, 'r'},  /* толщина border'ов */
+        {"icon",    required_argument, 0, 'i'},  /* path,x,y[,alpha] */
         {0, 0, 0, 0},
     };
     int c;
-    while ((c = getopt_long(argc, argv, "o:c:f:b:W:H:s:r:", opts, NULL)) != -1) {
+    while ((c = getopt_long(argc, argv, "o:c:f:b:W:H:s:r:i:", opts, NULL)) != -1) {
         switch (c) {
         case 'o': out_path = optarg; break;
         case 'c':
@@ -127,6 +133,28 @@ int main(int argc, char **argv)
         case 'H': out_h = atoi(optarg); break;
         case 's': max_seconds = atoi(optarg); break;
         case 'r': border_thickness = atoi(optarg); break;
+        case 'i': {
+            if (num_icons >= MAX_CELLS) {
+                fprintf(stderr, "max %d icons\n", MAX_CELLS);
+                return 1;
+            }
+            /* Парсим path,x,y[,alpha] аналогично --cell.
+             * Спецификация хранится статично — указатель уйдёт в overlay. */
+            static char icon_paths[MAX_CELLS][256];
+            char buf[300]; strncpy(buf, optarg, sizeof(buf) - 1);
+            buf[sizeof(buf) - 1] = '\0';
+            char *tok = strtok(buf, ","); if (!tok) { fprintf(stderr, "bad --icon\n"); return 1; }
+            strncpy(icon_paths[num_icons], tok, 255);
+            icons[num_icons].path = icon_paths[num_icons];
+            tok = strtok(NULL, ","); if (!tok) { fprintf(stderr, "bad --icon\n"); return 1; }
+            icons[num_icons].x = atoi(tok);
+            tok = strtok(NULL, ","); if (!tok) { fprintf(stderr, "bad --icon\n"); return 1; }
+            icons[num_icons].y = atoi(tok);
+            tok = strtok(NULL, ",");
+            icons[num_icons].alpha = tok ? atoi(tok) : 255;
+            num_icons++;
+            break;
+        }
         default: return 1;
         }
     }
@@ -167,6 +195,27 @@ int main(int argc, char **argv)
     fprintf(stderr, "[grid_record] composer %dx%d, %d ячеек\n",
             out_w, out_h, num_cells);
 
+    /* PNG иконки. */
+    for (int i = 0; i < num_icons; i++) {
+        cfc_overlay_png_config_t pc = {
+            .path = icons[i].path,
+            .x = icons[i].x,
+            .y = icons[i].y,
+            .extra_alpha = icons[i].alpha,
+            .visible = 1,
+        };
+        cfc_overlay_t *ov = NULL;
+        if (cfc_overlay_create_png(&pc, &ov) != 0) {
+            fprintf(stderr, "[grid_record] PNG '%s' load failed\n", icons[i].path);
+            continue;
+        }
+        int iw = 0, ih = 0;
+        cfc_overlay_png_size(ov, &iw, &ih);
+        cfc_composer_add_overlay(comp, ov);
+        fprintf(stderr, "[grid_record] icon '%s' %dx%d @ (%d,%d) alpha=%d\n",
+                icons[i].path, iw, ih, icons[i].x, icons[i].y, icons[i].alpha);
+    }
+
     /* Border'ы вокруг каждой ячейки если --border задан.
      * Цвет: серо-голубой (BT.709 limited): Y=180, U=120, V=110. */
     if (border_thickness > 0) {

include/cuframes_composer/cugrid.h

@@ -51,6 +51,29 @@ int cfc_cugrid_fill_nv12(
     int color_y, int color_u, int color_v, int alpha
 );
 
+/* Blit RGBA atlas → NV12 frame с α-blending.
+ *
+ * BT.709 limited-range conversion (R/G/B → Y/U/V). UV pixel = среднее 2x2
+ * пикселей RGBA (4:2:0 chroma subsampling).
+ *
+ *   dst_y/uv, pitch_y/uv  — NV12 destination.
+ *   dst_x, dst_y_off      — позиция top-left на dst (full-res, чётные).
+ *   atlas                 — CUdeviceptr RGBA байт (4 байта/пиксель,
+ *                            R,G,B,A interleaved).
+ *   atlas_w, atlas_h      — размер atlas в пикселях.
+ *   atlas_pitch           — pitch atlas в байтах.
+ *   extra_alpha           — 0..255, общий множитель прозрачности (для
+ *                            анимаций fade-in/fade-out).
+ */
+int cfc_cugrid_blit_rgba_nv12(
+    CUstream stream,
+    CUdeviceptr dst_y, int dst_pitch_y,
+    CUdeviceptr dst_uv, int dst_pitch_uv,
+    int dst_x, int dst_y_off,
+    CUdeviceptr atlas, int atlas_w, int atlas_h, int atlas_pitch,
+    int extra_alpha
+);
+
 /* Resize NV12 src → rect (dst_x, dst_y, dst_w, dst_h) на NV12 dst.
  *
  * Bilinear interpolation на Y и UV. Y и UV plane src'а должны быть одним

include/cuframes_composer/overlay.h

@@ -60,6 +60,30 @@ typedef struct cfc_overlay_border_config {
 int cfc_overlay_create_border(const cfc_overlay_border_config_t *cfg,
                                 cfc_overlay_t **out);
 
+/* Параметры PNG overlay'я. */
+typedef struct cfc_overlay_png_config {
+    const char *path;                 /* путь к PNG-файлу (декод один раз) */
+    int x, y;                         /* позиция top-left на output буфере (чётные) */
+    int extra_alpha;                  /* 0..255 общий множитель прозрачности */
+    int visible;                      /* 0/1 — выводить ли */
+} cfc_overlay_png_config_t;
+
+/* Создать PNG overlay. Декодирует файл через libpng, аллоцирует RGBA-атлас
+ * в VRAM (cuMemAlloc), копирует туда декодированные пиксели один раз.
+ *
+ * Размер изображения берётся из PNG header'а — caller узнаёт width/height
+ * через cfc_overlay_png_size после create. */
+int cfc_overlay_create_png(const cfc_overlay_png_config_t *cfg,
+                             cfc_overlay_t **out);
+
+/* Получить реальный размер декодированного PNG. После create — стабильны
+ * до destroy. */
+int cfc_overlay_png_size(cfc_overlay_t *ov, int *width, int *height);
+
+/* Обновить параметры PNG overlay'я (без re-decode). */
+int cfc_overlay_update_png(cfc_overlay_t *ov,
+                             const cfc_overlay_png_config_t *cfg);
+
 /* Обновить параметры BORDER overlay'я (можно переключить visible,
  * сменить цвет, изменить позицию). Thread-safe? Нет — caller должен сам
  * заботиться о том, чтобы update не пересекался с draw. В рамках одного

src/CMakeLists.txt

@@ -51,6 +51,7 @@ foreach(target cuframes_composer cuframes_composer_static)
             CUDA::cuda_driver
             Threads::Threads
             ${LIBDL_LIBRARY}           # для dlopen libnvidia-encode.so
+            PNG::PNG                   # для PNG overlay'ев (Phase 3b)
             rt
     )
     # CUDA properties.

src/cugrid/cugrid.cu

@@ -59,6 +59,59 @@ __global__ static void cfc_kern_fill_uv(unsigned char *dst, int dst_pitch,
     p[1] = (unsigned char)((fill_v * alpha + cv * (255 - alpha)) / 255);
 }
 
+__global__ static void cfc_kern_blit_rgba_y(unsigned char *dst, int dst_pitch,
+                                              int dx, int dy,
+                                              const unsigned char *atlas, int atlas_pitch,
+                                              int w, int h, int extra_alpha)
+{
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
+    if (x >= w || y >= h) return;
+    const unsigned char *sp = atlas + y * atlas_pitch + x * 4;
+    int r = sp[0], g = sp[1], b = sp[2], a = sp[3];
+    a = a * extra_alpha / 255;
+    if (a == 0) return;
+    int Y = (int)(0.183f * r + 0.614f * g + 0.062f * b) + 16;
+    Y = Y < 0 ? 0 : (Y > 255 ? 255 : Y);
+    unsigned char *p = dst + (dy + y) * dst_pitch + (dx + x);
+    int cur = *p;
+    *p = (unsigned char)((Y * a + cur * (255 - a)) / 255);
+}
+
+__global__ static void cfc_kern_blit_rgba_uv(unsigned char *dst, int dst_pitch,
+                                               int dx, int dy,
+                                               const unsigned char *atlas, int atlas_pitch,
+                                               int w, int h, int extra_alpha)
+{
+    int x = blockIdx.x * blockDim.x + threadIdx.x;
+    int y = blockIdx.y * blockDim.y + threadIdx.y;
+    int hw = w / 2, hh = h / 2;
+    if (x >= hw || y >= hh) return;
+
+    int sx = x * 2, sy = y * 2;
+    const unsigned char *row0 = atlas + sy * atlas_pitch;
+    const unsigned char *row1 = atlas + (sy + 1) * atlas_pitch;
+
+    int r = (row0[sx*4+0] + row0[(sx+1)*4+0] + row1[sx*4+0] + row1[(sx+1)*4+0]) >> 2;
+    int g = (row0[sx*4+1] + row0[(sx+1)*4+1] + row1[sx*4+1] + row1[(sx+1)*4+1]) >> 2;
+    int b = (row0[sx*4+2] + row0[(sx+1)*4+2] + row1[sx*4+2] + row1[(sx+1)*4+2]) >> 2;
+    int a = (row0[sx*4+3] + row0[(sx+1)*4+3] + row1[sx*4+3] + row1[(sx+1)*4+3]) >> 2;
+    a = a * extra_alpha / 255;
+    if (a == 0) return;
+
+    int U = (int)(-0.101f * r - 0.339f * g + 0.439f * b) + 128;
+    int V = (int)( 0.439f * r - 0.399f * g - 0.040f * b) + 128;
+    U = U < 0 ? 0 : (U > 255 ? 255 : U);
+    V = V < 0 ? 0 : (V > 255 ? 255 : V);
+
+    int du = dx / 2 + x;
+    int dv = dy / 2 + y;
+    unsigned char *p = dst + dv * dst_pitch + du * 2;
+    int cu = p[0], cv = p[1];
+    p[0] = (unsigned char)((U * a + cu * (255 - a)) / 255);
+    p[1] = (unsigned char)((V * a + cv * (255 - a)) / 255);
+}
+
 __global__ static void cfc_kern_resize_y(
     const unsigned char *src, int src_w, int src_h, int src_pitch,
     unsigned char *dst, int dst_x, int dst_y, int dst_w, int dst_h, int dst_pitch)
@@ -157,6 +210,42 @@ int cfc_cugrid_fill_nv12(
     return 0;
 }
 
+int cfc_cugrid_blit_rgba_nv12(
+    CUstream stream,
+    CUdeviceptr dst_y, int dst_pitch_y,
+    CUdeviceptr dst_uv, int dst_pitch_uv,
+    int dst_x, int dst_y_off,
+    CUdeviceptr atlas, int atlas_w, int atlas_h, int atlas_pitch,
+    int extra_alpha)
+{
+    if (atlas_w <= 0 || atlas_h <= 0) return 0;
+    if (extra_alpha <= 0) return 0;
+    if (extra_alpha > 255) extra_alpha = 255;
+
+    dim3 block_y_(CFC_BLOCK_X, CFC_BLOCK_Y);
+    dim3 grid_y_((atlas_w + CFC_BLOCK_X - 1) / CFC_BLOCK_X,
+                 (atlas_h + CFC_BLOCK_Y - 1) / CFC_BLOCK_Y);
+    cfc_kern_blit_rgba_y<<<grid_y_, block_y_, 0, (cudaStream_t)stream>>>(
+        (unsigned char *)dst_y, dst_pitch_y,
+        dst_x, dst_y_off,
+        (const unsigned char *)atlas, atlas_pitch,
+        atlas_w, atlas_h, extra_alpha);
+    CHECK_CUDA(cudaGetLastError());
+
+    int uv_w = atlas_w / 2, uv_h = atlas_h / 2;
+    dim3 block_uv_(CFC_BLOCK_X, CFC_BLOCK_Y);
+    dim3 grid_uv_((uv_w + CFC_BLOCK_X - 1) / CFC_BLOCK_X,
+                  (uv_h + CFC_BLOCK_Y - 1) / CFC_BLOCK_Y);
+    cfc_kern_blit_rgba_uv<<<grid_uv_, block_uv_, 0, (cudaStream_t)stream>>>(
+        (unsigned char *)dst_uv, dst_pitch_uv,
+        dst_x, dst_y_off,
+        (const unsigned char *)atlas, atlas_pitch,
+        atlas_w, atlas_h, extra_alpha);
+    CHECK_CUDA(cudaGetLastError());
+
+    return 0;
+}
+
 int cfc_cugrid_resize_nv12(
     CUstream stream,
     CUdeviceptr src_y, int src_w, int src_h, int src_pitch_y,

src/overlay.c

@@ -17,14 +17,24 @@
 #include "../include/cuframes_composer/overlay.h"
 #include "../include/cuframes_composer/cugrid.h"
 
+#include <cuda.h>
+#include <png.h>
+#include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
 
+typedef struct png_data {
+    cfc_overlay_png_config_t cfg;
+    int width, height;
+    int pitch;                          /* 4 * width, выровнено на 16 */
+    CUdeviceptr atlas;                  /* RGBA buffer в VRAM */
+} png_data_t;
+
 struct cfc_overlay {
     cfc_overlay_type_t type;
     union {
         cfc_overlay_border_config_t border;
-        /* PNG / TEXT — Phase 3b/3c, добавится позже. */
+        png_data_t png;
     } u;
 };
 
@@ -103,6 +113,165 @@ static int draw_border(cfc_overlay_t *ov,
     return 0;
 }
 
+/* ── PNG ──────────────────────────────────────────────────────────────── */
+
+/* Decode PNG в RGBA на CPU (через libpng). Возвращает malloc'нутый buffer.
+ * Caller освобождает через free(). */
+static unsigned char *decode_png(const char *path, int *out_w, int *out_h)
+{
+    FILE *fp = fopen(path, "rb");
+    if (!fp) {
+        fprintf(stderr, "[cfc/overlay/png] fopen(%s) failed\n", path);
+        return NULL;
+    }
+
+    unsigned char header[8];
+    if (fread(header, 1, 8, fp) != 8 || png_sig_cmp(header, 0, 8)) {
+        fprintf(stderr, "[cfc/overlay/png] не PNG: %s\n", path);
+        fclose(fp);
+        return NULL;
+    }
+
+    png_structp png = png_create_read_struct(PNG_LIBPNG_VER_STRING, NULL, NULL, NULL);
+    if (!png) { fclose(fp); return NULL; }
+    png_infop info = png_create_info_struct(png);
+    if (!info) { png_destroy_read_struct(&png, NULL, NULL); fclose(fp); return NULL; }
+
+    if (setjmp(png_jmpbuf(png))) {
+        png_destroy_read_struct(&png, &info, NULL);
+        fclose(fp);
+        return NULL;
+    }
+
+    png_init_io(png, fp);
+    png_set_sig_bytes(png, 8);
+    png_read_info(png, info);
+
+    int w = png_get_image_width(png, info);
+    int h = png_get_image_height(png, info);
+    png_byte color = png_get_color_type(png, info);
+    png_byte depth = png_get_bit_depth(png, info);
+
+    /* Нормализуем в 8-bit RGBA. */
+    if (depth == 16) png_set_strip_16(png);
+    if (color == PNG_COLOR_TYPE_PALETTE) png_set_palette_to_rgb(png);
+    if (color == PNG_COLOR_TYPE_GRAY && depth < 8) png_set_expand_gray_1_2_4_to_8(png);
+    if (png_get_valid(png, info, PNG_INFO_tRNS)) png_set_tRNS_to_alpha(png);
+    if (color == PNG_COLOR_TYPE_RGB || color == PNG_COLOR_TYPE_GRAY ||
+        color == PNG_COLOR_TYPE_PALETTE)
+        png_set_filler(png, 0xFF, PNG_FILLER_AFTER);
+    if (color == PNG_COLOR_TYPE_GRAY || color == PNG_COLOR_TYPE_GRAY_ALPHA)
+        png_set_gray_to_rgb(png);
+    png_read_update_info(png, info);
+
+    int rowbytes = png_get_rowbytes(png, info);
+    if (rowbytes != w * 4) {
+        fprintf(stderr, "[cfc/overlay/png] unexpected rowbytes %d (want %d)\n",
+                rowbytes, w * 4);
+        png_destroy_read_struct(&png, &info, NULL);
+        fclose(fp);
+        return NULL;
+    }
+
+    unsigned char *buf = malloc((size_t)rowbytes * h);
+    if (!buf) {
+        png_destroy_read_struct(&png, &info, NULL);
+        fclose(fp);
+        return NULL;
+    }
+    png_bytep *rows = malloc(sizeof(png_bytep) * h);
+    for (int i = 0; i < h; i++) rows[i] = buf + (size_t)i * rowbytes;
+    png_read_image(png, rows);
+    free(rows);
+
+    png_destroy_read_struct(&png, &info, NULL);
+    fclose(fp);
+
+    *out_w = w;
+    *out_h = h;
+    return buf;
+}
+
+int cfc_overlay_create_png(const cfc_overlay_png_config_t *cfg,
+                             cfc_overlay_t **out)
+{
+    if (!cfg || !cfg->path || !out) return -1;
+
+    int w = 0, h = 0;
+    unsigned char *rgba = decode_png(cfg->path, &w, &h);
+    if (!rgba) return -1;
+
+    cfc_overlay_t *ov = calloc(1, sizeof(*ov));
+    if (!ov) { free(rgba); return -1; }
+    ov->type = CFC_OVERLAY_PNG;
+    ov->u.png.cfg = *cfg;
+    if (ov->u.png.cfg.extra_alpha == 0) ov->u.png.cfg.extra_alpha = 255;
+    ov->u.png.width = w;
+    ov->u.png.height = h;
+    ov->u.png.pitch = w * 4;
+
+    /* Atlas в VRAM. */
+    size_t size = (size_t)ov->u.png.pitch * h;
+    CUresult cr = cuMemAlloc(&ov->u.png.atlas, size);
+    if (cr != CUDA_SUCCESS) {
+        fprintf(stderr, "[cfc/overlay/png] cuMemAlloc(%zu) failed\n", size);
+        free(rgba); free(ov); return -1;
+    }
+    cr = cuMemcpyHtoD(ov->u.png.atlas, rgba, size);
+    free(rgba);
+    if (cr != CUDA_SUCCESS) {
+        fprintf(stderr, "[cfc/overlay/png] cuMemcpyHtoD failed\n");
+        cuMemFree(ov->u.png.atlas);
+        free(ov);
+        return -1;
+    }
+
+    *out = ov;
+    return 0;
+}
+
+int cfc_overlay_png_size(cfc_overlay_t *ov, int *width, int *height)
+{
+    if (!ov || ov->type != CFC_OVERLAY_PNG) return -1;
+    if (width) *width = ov->u.png.width;
+    if (height) *height = ov->u.png.height;
+    return 0;
+}
+
+int cfc_overlay_update_png(cfc_overlay_t *ov,
+                             const cfc_overlay_png_config_t *cfg)
+{
+    if (!ov || !cfg) return -1;
+    if (ov->type != CFC_OVERLAY_PNG) return -1;
+    /* Path игнорируем — re-decode не делаем. Меняем только runtime поля. */
+    ov->u.png.cfg.x = cfg->x;
+    ov->u.png.cfg.y = cfg->y;
+    ov->u.png.cfg.extra_alpha = cfg->extra_alpha;
+    ov->u.png.cfg.visible = cfg->visible;
+    return 0;
+}
+
+static int draw_png(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 png_data_t *p = &ov->u.png;
+    if (!p->cfg.visible) return 0;
+    if (p->cfg.extra_alpha <= 0) return 0;
+    if (p->cfg.x >= frame_w || p->cfg.y >= frame_h) return 0;
+
+    /* Координаты выравниваем на чётные. */
+    int x = p->cfg.x & ~1;
+    int y = p->cfg.y & ~1;
+    return cfc_cugrid_blit_rgba_nv12(
+        stream,
+        dst_y, pitch_y, dst_uv, pitch_uv,
+        x, y,
+        p->atlas, p->width, p->height, p->pitch,
+        p->cfg.extra_alpha);
+}
+
 /* ── Public dispatch ─────────────────────────────────────────────────── */
 
 int cfc_overlay_draw(cfc_overlay_t *ov,
@@ -117,8 +286,10 @@ int cfc_overlay_draw(cfc_overlay_t *ov,
         return draw_border(ov, stream, dst_y, pitch_y, dst_uv, pitch_uv,
                             frame_w, frame_h);
     case CFC_OVERLAY_PNG:
+        return draw_png(ov, stream, dst_y, pitch_y, dst_uv, pitch_uv,
+                         frame_w, frame_h);
     case CFC_OVERLAY_TEXT:
-        /* Phase 3b/3c — пока no-op. */
+        /* Phase 3c — пока no-op. */
         return 0;
     }
     return -1;
@@ -127,6 +298,9 @@ int cfc_overlay_draw(cfc_overlay_t *ov,
 int cfc_overlay_destroy(cfc_overlay_t *ov)
 {
     if (!ov) return 0;
+    if (ov->type == CFC_OVERLAY_PNG && ov->u.png.atlas) {
+        cuMemFree(ov->u.png.atlas);
+    }
     free(ov);
     return 0;
 }