From 8cc5a5acfb7a810a4dabfac921533c9f9dc6cecb Mon Sep 17 00:00:00 2001 From: Evgeny Demchenko Date: Wed, 3 Jun 2026 05:34:07 +0100 Subject: [PATCH] =?UTF-8?q?Phase=203b:=20PNG=20icon=20overlays=20=D1=87?= =?UTF-8?q?=D0=B5=D1=80=D0=B5=D0=B7=20libpng=20+=20alpha=5Fblit=5Frgba=5Fn?= =?UTF-8?q?v12?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 | 3 + examples/grid_record.c | 51 +++++++- include/cuframes_composer/cugrid.h | 23 ++++ include/cuframes_composer/overlay.h | 24 ++++ src/CMakeLists.txt | 1 + src/cugrid/cugrid.cu | 89 ++++++++++++++ src/overlay.c | 178 +++++++++++++++++++++++++++- 7 files changed, 366 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5159193..6526707 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/examples/grid_record.c b/examples/grid_record.c index 65d41b7..6ff3df8 100644 --- a/examples/grid_record.c +++ b/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) { diff --git a/include/cuframes_composer/cugrid.h b/include/cuframes_composer/cugrid.h index 25eb44c..26d66a5 100644 --- a/include/cuframes_composer/cugrid.h +++ b/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'а должны быть одним diff --git a/include/cuframes_composer/overlay.h b/include/cuframes_composer/overlay.h index b34c414..faaf94c 100644 --- a/include/cuframes_composer/overlay.h +++ b/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. В рамках одного diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 548b848..5d6638a 100644 --- a/src/CMakeLists.txt +++ b/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. diff --git a/src/cugrid/cugrid.cu b/src/cugrid/cugrid.cu index 30a64a9..d85e2a0 100644 --- a/src/cugrid/cugrid.cu +++ b/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<<>>( + (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<<>>( + (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, diff --git a/src/overlay.c b/src/overlay.c index fe24794..d94eba6 100644 --- a/src/overlay.c +++ b/src/overlay.c @@ -17,14 +17,24 @@ #include "../include/cuframes_composer/overlay.h" #include "../include/cuframes_composer/cugrid.h" +#include +#include +#include #include #include +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; }