Files
go-mxl-pattern-generator/kernels/v210_pack.cl
T
Dmitry Sergeev f48195a8bf WGPU is coming
2026-09-07 17:18:03 +03:00

73 lines
2.9 KiB
Common Lisp

// 75% SMPTE color bars, Rec.709, 10-bit Y'CbCr, packed as v210 (4:2:2),
// with an 8x16 bitmap-font text overlay (burn-in).
// Each work-item packs 6 pixels (6 Y + 3 Cb + 3 Cr) into one 16-byte block.
#define GW 8
#define GH 16
// Text overlay lookup. Atlas: 96 glyphs (ASCII 32..127) x GH rows,
// one byte per row, bit n = pixel at column n (LSB = leftmost).
// Returns: 1 = glyph pixel (foreground), 0 = inside text box (background),
// -1 = outside the text box (pattern shows through).
int text_pixel(__global const uchar* atlas, __global const uchar* text, int text_len,
int px, int py, int txt_x, int txt_y, int scale) {
if (text_len <= 0) return -1;
int tx = px - txt_x;
int ty = py - txt_y;
if (tx < 0 || ty < 0) return -1;
if (tx >= text_len * GW * scale || ty >= GH * scale) return -1;
int c = text[tx / (GW * scale)];
if (c < 32 || c > 127) return 0;
uchar row = atlas[(c - 32) * GH + ty / scale];
return (row >> ((tx / scale) % GW)) & 1;
}
__kernel void generate_v210_pattern(
__global uint* output, int width, int height, int frame_index,
__global const uchar* atlas, __global const uchar* text, int text_len,
int txt_x, int txt_y, int scale, int has_bg,
uint fg_y, uint fg_cb, uint fg_cr,
uint bg_y, uint bg_cb, uint bg_cr) {
int g_id = get_global_id(0);
int total_blocks = (width * height) / 6;
if (g_id >= total_blocks) return;
int x = (g_id * 6) % width;
int py = (g_id * 6) / width;
// Bar order: 75% white, yellow, cyan, green, magenta, red, blue
ushort Y_table[7] = {721, 674, 581, 534, 251, 204, 111};
ushort U_table[7] = {512, 176, 589, 253, 771, 435, 848};
ushort V_table[7] = {512, 543, 176, 207, 817, 848, 481};
#define BAR(px) min(((px) * 7) / width, 6)
uint y[6], cb[6], cr[6];
for (int i = 0; i < 6; i++) {
int px = x + i;
int b = BAR(px);
uint yv = Y_table[b], cbv = U_table[b], crv = V_table[b];
int st = text_pixel(atlas, text, text_len, px, py, txt_x, txt_y, scale);
if (st > 0) {
yv = fg_y; cbv = fg_cb; crv = fg_cr;
} else if (st == 0 && has_bg) {
yv = bg_y; cbv = bg_cb; crv = bg_cr;
}
y[i] = yv; cb[i] = cbv; cr[i] = crv;
}
// v210 word layout, chroma co-sited with luma samples 0/2/4:
// w0 = Cb0|Y0<<10|Cr0<<20; w1 = Y1|Cb2<<10|Y2<<20;
// w2 = Cr2|Y3<<10|Cb4<<20; w3 = Y4|Cr4<<10|Y5<<20
uint word0 = (cb[0] & 0x3FF) | ((y[0] & 0x3FF) << 10) | ((cr[0] & 0x3FF) << 20);
uint word1 = (y[1] & 0x3FF) | ((cb[2] & 0x3FF) << 10) | ((y[2] & 0x3FF) << 20);
uint word2 = (cr[2] & 0x3FF) | ((y[3] & 0x3FF) << 10) | ((cb[4] & 0x3FF) << 20);
uint word3 = (y[4] & 0x3FF) | ((cr[4] & 0x3FF) << 10) | ((y[5] & 0x3FF) << 20);
int out_idx = g_id * 4;
output[out_idx + 0] = word0;
output[out_idx + 1] = word1;
output[out_idx + 2] = word2;
output[out_idx + 3] = word3;
}