WGPU is coming

2026-09-07 17:18:03 +03:00
commit f48195a8bf
19 changed files with 2230 additions and 0 deletions
@@ -0,0 +1,50 @@
+// 75% SMPTE color bars, Rec.709, 10-bit Y'CbCr, packed as v210 (4:2:2).
+// One work-item per 16-byte block = 6 pixels (6 Y + 3 Cb + 3 Cr).
+
+struct Params {
+    width: u32,
+    height: u32,
+    frame: u32,
+    _pad0: u32,
+};
+
+@group(0) @binding(0) var<storage, read_write> out: array<u32>;
+@group(0) @binding(1) var<uniform> params: Params;
+
+fn bar_index(px: u32) -> u32 {
+    return min((px * 7u) / params.width, 6u);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let block = gid.x;
+    let total = (params.width * params.height) / 6u;
+    if (block >= total) {
+        return;
+    }
+
+    let x = (block * 6u) % params.width;
+
+    // Bar order: 75% white, yellow, cyan, green, magenta, red, blue
+    let y_tab  = array<u32, 7>(721u, 674u, 581u, 534u, 251u, 204u, 111u);
+    let cb_tab = array<u32, 7>(512u, 176u, 589u, 253u, 771u, 435u, 848u);
+    let cr_tab = array<u32, 7>(512u, 543u, 176u, 207u, 817u, 848u, 481u);
+
+    var y: array<u32, 6>;
+    var cb: array<u32, 6>;
+    var cr: array<u32, 6>;
+    for (var i = 0u; i < 6u; i++) {
+        let b = bar_index(x + i);
+        y[i] = y_tab[b];
+        cb[i] = cb_tab[b];
+        cr[i] = cr_tab[b];
+    }
+
+    // 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
+    out[block * 4u + 0u] = (cb[0] & 0x3FFu) | ((y[0] & 0x3FFu) << 10u) | ((cr[0] & 0x3FFu) << 20u);
+    out[block * 4u + 1u] = (y[1] & 0x3FFu) | ((cb[2] & 0x3FFu) << 10u) | ((y[2] & 0x3FFu) << 20u);
+    out[block * 4u + 2u] = (cr[2] & 0x3FFu) | ((y[3] & 0x3FFu) << 10u) | ((cb[4] & 0x3FFu) << 20u);
+    out[block * 4u + 3u] = (y[4] & 0x3FFu) | ((cr[4] & 0x3FFu) << 10u) | ((y[5] & 0x3FFu) << 20u);
+}
@@ -0,0 +1,77 @@
+// 75% SMPTE color bars + moving square, Rec.709, 10-bit Y'CbCr, packed as v210 (4:2:2).
+// One work-item per 16-byte block = 6 pixels (6 Y + 3 Cb + 3 Cr).
+
+struct Params {
+    width: u32,
+    height: u32,
+    frame: u32, // animation tick (0,1,2,... per generated frame), NOT the raw grain index
+    _pad0: u32,
+};
+
+@group(0) @binding(0) var<storage, read_write> out: array<u32>;
+@group(0) @binding(1) var<uniform> params: Params;
+
+fn bar_index(px: u32) -> u32 {
+    return min((px * 7u) / params.width, 6u);
+}
+
+@compute @workgroup_size(64)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let block = gid.x;
+    let total = (params.width * params.height) / 6u;
+    if (block >= total) {
+        return;
+    }
+
+    let x = (block * 6u) % params.width;
+    let py = (block * 6u) / params.width;
+
+    // Bar order: 75% white, yellow, cyan, green, magenta, red, blue
+    let y_tab  = array<u32, 7>(721u, 674u, 581u, 534u, 251u, 204u, 111u);
+    let cb_tab = array<u32, 7>(512u, 176u, 589u, 253u, 771u, 435u, 848u);
+    let cr_tab = array<u32, 7>(512u, 543u, 176u, 207u, 817u, 848u, 481u);
+
+    // Moving square: horizontal oscillation around screen center.
+    // frame is a small tick; converting the huge raw grain index here
+    // would destroy f32 precision and freeze the motion.
+    let centerX = f32(params.width) / 2.0;
+    let centerY = f32(params.height) / 2.0;
+    let fps = 25.0;
+    let t = f32(params.frame) / fps;
+    const squareSize = 150u;
+    let half = f32(squareSize) / 2.0;
+    let offsetPixels = sin(t * 0.5) * (centerX - half);
+
+    let sq_x_min = centerX - half + offsetPixels;
+    let sq_x_max = centerX + half + offsetPixels;
+    let sq_y_min = centerY - half;
+    let sq_y_max = centerY + half;
+    let py_f = f32(py);
+
+    var y: array<u32, 6>;
+    var cb: array<u32, 6>;
+    var cr: array<u32, 6>;
+    for (var i = 0u; i < 6u; i++) {
+        let px = f32(x + i);
+        let b = bar_index(x + i);
+        if (px >= sq_x_min && px < sq_x_max && py_f >= sq_y_min && py_f < sq_y_max) {
+            // inverted bar: luma mirrored across studio range (64..940),
+            // chroma mirrored across neutral 512
+            y[i]  = 1004u - y_tab[b];
+            cb[i] = 1024u - cb_tab[b];
+            cr[i] = 1024u - cr_tab[b];
+        } else {
+            y[i] = y_tab[b];
+            cb[i] = cb_tab[b];
+            cr[i] = cr_tab[b];
+        }
+    }
+
+    // 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
+    out[block * 4u + 0u] = (cb[0] & 0x3FFu) | ((y[0] & 0x3FFu) << 10u) | ((cr[0] & 0x3FFu) << 20u);
+    out[block * 4u + 1u] = (y[1] & 0x3FFu) | ((cb[2] & 0x3FFu) << 10u) | ((y[2] & 0x3FFu) << 20u);
+    out[block * 4u + 2u] = (cr[2] & 0x3FFu) | ((y[3] & 0x3FFu) << 10u) | ((cb[4] & 0x3FFu) << 20u);
+    out[block * 4u + 3u] = (y[4] & 0x3FFu) | ((cr[4] & 0x3FFu) << 10u) | ((y[5] & 0x3FFu) << 20u);
+}
@@ -0,0 +1,72 @@
+// 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;
+}