cleanup + pattern selector
This commit is contained in:
@@ -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);
|
||||
}
|
||||
Reference in New Issue
Block a user