diff --git a/cmd/mxl-pattern/main.go b/cmd/mxl-pattern/main.go index 726bcce..24db40e 100644 --- a/cmd/mxl-pattern/main.go +++ b/cmd/mxl-pattern/main.go @@ -210,12 +210,6 @@ var patterns = map[string]pattern{ kernelPath: "kernels/dynamic/yRamp.wgsl", motion: false, }, - "dvd": { - name: "dvd", - description: "DVD logo pattern, but with MXL logo", - kernelPath: "kernels/dynamic/dvdLogo.wgsl", - motion: true, - }, } func listPatterns(f *os.File) { @@ -319,15 +313,7 @@ func main() { // TODO: if init failed -> CPU generator videoPattern := patterns[args.pattern] - var genOpts []generator.WGPUOption - if videoPattern.name == "dvd" { - plane, err := generator.LoadLogoPNG("assets/mxl.png") - if err != nil { - log.Fatalf("logo load failed: %v", err) - } - genOpts = append(genOpts, generator.WithLogo(plane)) - } - gen, err := generator.NewWGPUGenerator(vi.width, vi.height, videoPattern.kernelPath, genOpts...) + gen, err := generator.NewWGPUGenerator(vi.width, vi.height, videoPattern.kernelPath) if err != nil { log.Fatalf("wgpu init failed: %v", err) } diff --git a/internal/generator/logo.go b/internal/generator/logo.go deleted file mode 100644 index ef11ee2..0000000 --- a/internal/generator/logo.go +++ /dev/null @@ -1,76 +0,0 @@ -// Logo plane construction for texture-sampling kernels (v210_dvd_logo.wgsl). -package generator - -import ( - "fmt" - "image" - "math" - "os" - - // PNG decoder registration for LoadLogoPNG. - _ "image/png" -) - -// BuildLogoPlane converts an image into the packed logo plane consumed by -// v210_dvd_logo.wgsl as read-only storage binding 2: -// -// logo[0] = width | height<<16 -// logo[1+dy*w+dx] = opaque<<30 | Cr<<20 | Cb<<10 | Y -// -// Rec.709 full-range RGB -> 10-bit studio swing, the same transform behind -// the color-bar tables in the kernels (verified against all 21 bar values). -// Semi-transparent pixels (alpha >= 50%) count as opaque; the rest show the -// kernel background through. -func BuildLogoPlane(img image.Image) ([]uint32, error) { - b := img.Bounds() - w, h := b.Dx(), b.Dy() - if w <= 0 || h <= 0 || w > 0xFFFF || h > 0xFFFF { - return nil, fmt.Errorf("logo: bad dimensions %dx%d", w, h) - } - plane := make([]uint32, 1+w*h) - plane[0] = uint32(w) | uint32(h)<<16 - for yy := 0; yy < h; yy++ { - for xx := 0; xx < w; xx++ { - r, g, bl, a := img.At(b.Min.X+xx, b.Min.Y+yy).RGBA() - if a < 1<<15 { - continue // transparent: shader falls back to background - } - // At().RGBA() is premultiplied; unpremultiply to full-range RGB. - rf := float64(r) / float64(a) - gf := float64(g) / float64(a) - bf := float64(bl) / float64(a) - luma := 0.2126*rf + 0.7152*gf + 0.0722*bf - yc := clamp10(64 + math.Round(876*luma)) - cb := clamp10(512 + math.Round(448*(bf-luma)/(1-0.0722))) - cr := clamp10(512 + math.Round(448*(rf-luma)/(1-0.2126))) - plane[1+yy*w+xx] = 1<<30 | cr<<20 | cb<<10 | yc - } - } - return plane, nil -} - -// clamp10 keeps a 10-bit code inside the legal 4..1019 range. -func clamp10(v float64) uint32 { - switch { - case v < 4: - return 4 - case v > 1019: - return 1019 - default: - return uint32(v) - } -} - -// LoadLogoPNG reads a PNG file and builds the logo plane. -func LoadLogoPNG(path string) ([]uint32, error) { - f, err := os.Open(path) - if err != nil { - return nil, fmt.Errorf("logo: %w", err) - } - defer f.Close() - img, _, err := image.Decode(f) - if err != nil { - return nil, fmt.Errorf("logo: decode %s: %w", path, err) - } - return BuildLogoPlane(img) -} diff --git a/internal/generator/logo_test.go b/internal/generator/logo_test.go deleted file mode 100644 index 83c1995..0000000 --- a/internal/generator/logo_test.go +++ /dev/null @@ -1,64 +0,0 @@ -package generator - -import ( - "image" - "image/color" - "testing" -) - -func TestBuildLogoPlane(t *testing.T) { - // 8x4 logo: white left half, near-75% red right half (8-bit 191 is - // 0.749, a hair under the exact-0.75 bars value, hence Y=203 not 204), - // one transparent pixel at (2,1). - img := image.NewRGBA(image.Rect(0, 0, 8, 4)) - for yy := 0; yy < 4; yy++ { - for xx := 0; xx < 8; xx++ { - if xx < 4 { - img.Set(xx, yy, color.RGBA{255, 255, 255, 255}) - } else { - img.Set(xx, yy, color.RGBA{191, 0, 0, 255}) - } - } - } - img.Set(2, 1, color.RGBA{0, 0, 0, 0}) - - plane, err := BuildLogoPlane(img) - if err != nil { - t.Fatalf("BuildLogoPlane: %v", err) - } - if len(plane) != 1+8*4 { - t.Fatalf("plane len=%d, want 33", len(plane)) - } - if hdr := plane[0]; hdr != 8|4<<16 { - t.Fatalf("header=%#x, want w=8 h=4", hdr) - } - cases := []struct { - dx, dy int - opaque bool - y, cb, cr uint32 - }{ - {1, 1, true, 940, 512, 512}, // white - {5, 2, true, 203, 435, 848}, // near-75% red - {2, 1, false, 0, 0, 0}, // transparent - } - for _, c := range cases { - s := plane[1+c.dy*8+c.dx] - if got := (s >> 30) & 1; got != boolToBit(c.opaque) { - t.Fatalf("pixel (%d,%d): opaque=%d, want %v", c.dx, c.dy, got, c.opaque) - } - if !c.opaque { - continue - } - if s&0x3FF != c.y || (s>>10)&0x3FF != c.cb || (s>>20)&0x3FF != c.cr { - t.Fatalf("pixel (%d,%d): got %d/%d/%d, want %d/%d/%d", - c.dx, c.dy, s&0x3FF, (s>>10)&0x3FF, (s>>20)&0x3FF, c.y, c.cb, c.cr) - } - } -} - -func boolToBit(b bool) uint32 { - if b { - return 1 - } - return 0 -} diff --git a/internal/generator/wgpu.go b/internal/generator/wgpu.go index c957a01..df8b02f 100644 --- a/internal/generator/wgpu.go +++ b/internal/generator/wgpu.go @@ -35,8 +35,6 @@ type WGPUGenerator struct { out *wgpu.Buffer host *wgpu.Buffer uniform *wgpu.Buffer - logo *wgpu.Buffer - logoPlane []uint32 params []byte width int height int @@ -49,14 +47,6 @@ var _ FrameGenerator = (*WGPUGenerator)(nil) // WGPUOption customizes NewWGPUGenerator. type WGPUOption func(*WGPUGenerator) -// WithLogo attaches a packed logo plane (see BuildLogoPlane) as read-only -// storage binding 2 for kernels that sample it (v210_dvd_logo.wgsl). -func WithLogo(plane []uint32) WGPUOption { - return func(g *WGPUGenerator) { - g.logoPlane = plane - } -} - func NewWGPUGenerator(width, height uint, kernelPath string, opts ...WGPUOption) (*WGPUGenerator, error) { g := &WGPUGenerator{ width: int(width), @@ -121,48 +111,22 @@ func NewWGPUGenerator(width, height uint, kernelPath string, opts ...WGPUOption) g.Close() return nil, fmt.Errorf("wgpu: write params: %w", err) } - bglEntries := []gputypes.BindGroupLayoutEntry{ - {Binding: 0, Visibility: wgpu.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}}, - {Binding: 1, Visibility: wgpu.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeUniform}}, - } - bgEntries := []wgpu.BindGroupEntry{ - {Binding: 0, Buffer: g.out, Size: g.frameSize}, - {Binding: 1, Buffer: g.uniform, Size: uint64(len(g.params))}, - } - if g.logoPlane != nil { - logoBytes := make([]byte, 4*len(g.logoPlane)) - for i, v := range g.logoPlane { - binary.LittleEndian.PutUint32(logoBytes[4*i:], v) - } - if g.logo, err = g.device.CreateBuffer(&wgpu.BufferDescriptor{ - Label: "v210-logo", Size: uint64(len(logoBytes)), - Usage: wgpu.BufferUsageStorage | wgpu.BufferUsageCopyDst, - }); err != nil { - g.Close() - return nil, fmt.Errorf("wgpu: logo buffer: %w", err) - } - if err := g.queue.WriteBuffer(g.logo, 0, logoBytes); err != nil { - g.Close() - return nil, fmt.Errorf("wgpu: write logo: %w", err) - } - bglEntries = append(bglEntries, gputypes.BindGroupLayoutEntry{ - Binding: 2, Visibility: wgpu.ShaderStageCompute, - Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeReadOnlyStorage}, - }) - bgEntries = append(bgEntries, wgpu.BindGroupEntry{ - Binding: 2, Buffer: g.logo, Size: uint64(len(logoBytes)), - }) - } if g.bgl, err = g.device.CreateBindGroupLayout(&wgpu.BindGroupLayoutDescriptor{ - Label: "v210-bgl", - Entries: bglEntries, + Label: "v210-bgl", + Entries: []gputypes.BindGroupLayoutEntry{ + {Binding: 0, Visibility: wgpu.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}}, + {Binding: 1, Visibility: wgpu.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeUniform}}, + }, }); err != nil { g.Close() return nil, fmt.Errorf("wgpu: bind group layout: %w", err) } if g.bg, err = g.device.CreateBindGroup(&wgpu.BindGroupDescriptor{ Label: "v210-bg", Layout: g.bgl, - Entries: bgEntries, + Entries: []wgpu.BindGroupEntry{ + {Binding: 0, Buffer: g.out, Size: g.frameSize}, + {Binding: 1, Buffer: g.uniform, Size: uint64(len(g.params))}, + }, }); err != nil { g.Close() return nil, fmt.Errorf("wgpu: bind group: %w", err) @@ -254,10 +218,6 @@ func (g *WGPUGenerator) Close() error { g.uniform.Release() g.uniform = nil } - if g.logo != nil { - g.logo.Release() - g.logo = nil - } if g.host != nil { g.host.Release() g.host = nil diff --git a/internal/generator/wgpu_dvd_logo_test.go b/internal/generator/wgpu_dvd_logo_test.go deleted file mode 100644 index 56840cb..0000000 --- a/internal/generator/wgpu_dvd_logo_test.go +++ /dev/null @@ -1,169 +0,0 @@ -package generator - -import ( - "encoding/binary" - "image" - "image/color" - "math" - "path/filepath" - "testing" -) - -const logoW, logoH = 8, 4 - -func TestWGPUDVDLogo(t *testing.T) { - const width, height = 1920, 1080 - - // Synthetic 8x4 logo: white left half, near-75% red right half, one - // transparent pixel at (2,1) to exercise the alpha fallback. - img := image.NewRGBA(image.Rect(0, 0, logoW, logoH)) - for yy := 0; yy < logoH; yy++ { - for xx := 0; xx < logoW; xx++ { - if xx < 4 { - img.Set(xx, yy, color.RGBA{255, 255, 255, 255}) - } else { - img.Set(xx, yy, color.RGBA{191, 0, 0, 255}) - } - } - } - img.Set(2, 1, color.RGBA{0, 0, 0, 0}) - plane, err := BuildLogoPlane(img) - if err != nil { - t.Fatalf("plane: %v", err) - } - - g, err := NewWGPUGenerator(width, height, - filepath.Join("..", "..", "kernels", "v210_dvd_logo.wgsl"), WithLogo(plane)) - if err != nil { - t.Fatalf("init: %v", err) - } - defer g.Close() - - buf := make([]byte, width*height*8/3) - sample := func(x, y int) (yc, cb, cr uint32) { - p := y*width + x - off := (p / 6) * 16 - w0 := binary.LittleEndian.Uint32(buf[off:]) - w1 := binary.LittleEndian.Uint32(buf[off+4:]) - w2 := binary.LittleEndian.Uint32(buf[off+8:]) - w3 := binary.LittleEndian.Uint32(buf[off+12:]) - switch p % 6 { - case 0: - return (w0 >> 10) & 0x3FF, w0 & 0x3FF, (w0 >> 20) & 0x3FF - case 1: - return w1 & 0x3FF, w0 & 0x3FF, (w0 >> 20) & 0x3FF - case 2: - return (w1 >> 20) & 0x3FF, (w1 >> 10) & 0x3FF, w2 & 0x3FF - case 3: - return (w2 >> 10) & 0x3FF, (w1 >> 10) & 0x3FF, w2 & 0x3FF - case 4: - return w3 & 0x3FF, (w2 >> 20) & 0x3FF, (w3 >> 10) & 0x3FF - default: - return (w3 >> 20) & 0x3FF, (w2 >> 20) & 0x3FF, (w3 >> 10) & 0x3FF - } - } - - // Float32 mirrors of the shader position math and bilinear tap blend - // (kernels/v210_dvd_logo.wgsl), so expectations track the shader exactly. - tri01 := func(x float32) float32 { - f := x - float32(math.Floor(float64(x))) - if f > 0.5 { - return 2 - 2*f - } - return 2 * f - } - type logoPos struct { - bx, by int - fx, fy float32 - } - logoOrigin := func(frame uint32) logoPos { - tick := float32(frame) / 25.0 - lx := tri01(tick*0.061) * (float32(width) - logoW) - ly := tri01(tick*0.088) * (float32(height) - logoH) - return logoPos{int(lx), int(ly), lx - float32(int(lx)), ly - float32(int(ly))} - } - tap := func(dx, dy int) [3]float32 { - if dx < 0 || dy < 0 || dx >= logoW || dy >= logoH { - return [3]float32{} - } - s := plane[1+dy*logoW+dx] - if s>>30 == 0 { - return [3]float32{} - } - return [3]float32{float32(s&0x3FF) - 64, float32((s>>10)&0x3FF) - 512, float32((s>>20)&0x3FF) - 512} - } - mix3 := func(a, b [3]float32, k float32) [3]float32 { - return [3]float32{a[0] + k*(b[0]-a[0]), a[1] + k*(b[1]-a[1]), a[2] + k*(b[2]-a[2])} - } - expect := func(p logoPos, px, py int) (yc, cb, cr uint32) { - m, n := px-p.bx, py-p.by - top := mix3(tap(m, n-1), tap(m-1, n-1), p.fx) - bot := mix3(tap(m, n), tap(m-1, n), p.fx) - d := mix3(bot, top, p.fy) - return uint32(64 + d[0] + 0.5), uint32(512 + d[1] + 0.5), uint32(512 + d[2] + 0.5) - } - // tol allows +-1 code for blended pixels (FMA contraction on GPU may - // differ from Go's separately-rounded float32 ops by 1 ulp). - check := func(p logoPos, px, py int, tol int, desc string) { - t.Helper() - yc, cb, cr := sample(px, py) - wy, wcb, wcr := expect(p, px, py) - dy, dcb, dcr := int(yc)-int(wy), int(cb)-int(wcb), int(cr)-int(wcr) - if abs(dy) > tol || abs(dcb) > tol || abs(dcr) > tol { - t.Fatalf("%s: (%d,%d) got %d/%d/%d, want %d/%d/%d (diff %d/%d/%d, tol %d)", - desc, px, py, yc, cb, cr, wy, wcb, wcr, dy, dcb, dcr, tol) - } - } - - // tick 0: fx=fy=0 -> pure nearest lookup, exact values known upfront - if err := g.GenerateFrame(buf, 0); err != nil { - t.Fatalf("tick 0: %v", err) - } - p0 := logoOrigin(0) - if p0.bx != 0 || p0.by != 0 || p0.fx != 0 || p0.fy != 0 { - t.Fatalf("tick 0 origin %+v, want (0,0,0,0)", p0) - } - check(p0, 1, 1, 0, "tick 0 white pixel") - if y, cb, cr := sample(1, 1); y != 940 || cb != 512 || cr != 512 { - t.Fatalf("tick 0 white pixel: got %d/%d/%d, want 940/512/512", y, cb, cr) - } - check(p0, 2, 1, 0, "tick 0 transparent pixel") - if y, _, _ := sample(2, 1); y != 64 { - t.Fatalf("tick 0 transparent pixel: Y=%d, want background 64", y) - } - check(p0, 5, 2, 0, "tick 0 red pixel") - if y, cb, cr := sample(5, 2); y != 203 || cb != 435 || cr != 848 { - t.Fatalf("tick 0 red pixel: got %d/%d/%d, want 203/435/848", y, cb, cr) - } - check(p0, 10, 1, 0, "tick 0 right of logo") - check(p0, 1, 5, 0, "tick 0 below logo") - - // tick 100: origin (933,757), fx~0.056 fy~0.504 -> bilinear blending. - // (934,758) lands on solid white (all 4 taps white) and must reproduce - // the texel exactly; (935,759) straddles the transparent pixel (2,1) - // and must match the mirrored blend within +-1. - if err := g.GenerateFrame(buf, 100); err != nil { - t.Fatalf("tick 100: %v", err) - } - p100 := logoOrigin(100) - if p100.bx != 933 || p100.by != 757 { - t.Fatalf("tick 100 origin (%d,%d), want (933,757)", p100.bx, p100.by) - } - check(p100, 934, 758, 0, "tick 100 solid pixel exact") - if y, _, _ := sample(934, 758); y != 940 { - t.Fatalf("tick 100 solid pixel: Y=%d, want exact texel 940", y) - } - check(p100, 935, 759, 1, "tick 100 blended over transparent pixel") - check(p100, 960, 540, 0, "tick 100 frame center") - if y, _, _ := sample(960, 540); y != 64 { - t.Fatalf("tick 100 frame center: Y=%d, want background 64", y) - } - check(p100, 2, 1, 0, "tick 100 old logo position") -} - -func abs(v int) int { - if v < 0 { - return -v - } - return v -} diff --git a/internal/generator/wgpu_move_test.go b/internal/generator/wgpu_move_test.go index 67f715e..6f1962f 100644 --- a/internal/generator/wgpu_move_test.go +++ b/internal/generator/wgpu_move_test.go @@ -8,7 +8,7 @@ import ( func TestWGPUMoveSquare(t *testing.T) { const width, height = 1920, 1080 - g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "v210_bars_move.wgsl")) + g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "dynamic", "smpteBars.wgsl")) if err != nil { t.Fatalf("init: %v", err) } diff --git a/internal/generator/wgpu_test.go b/internal/generator/wgpu_test.go index 10db231..f857ae3 100644 --- a/internal/generator/wgpu_test.go +++ b/internal/generator/wgpu_test.go @@ -8,7 +8,7 @@ import ( func TestWGPUGenerator(t *testing.T) { const width, height = 1920, 1080 - g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "v210_bars.wgsl")) + g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "static", "smpteBars.wgsl")) if err != nil { t.Fatalf("init: %v", err) } diff --git a/kernels/dynamic/dvdLogo.wgsl b/kernels/dynamic/dvdLogo.wgsl deleted file mode 100644 index 5258c4f..0000000 --- a/kernels/dynamic/dvdLogo.wgsl +++ /dev/null @@ -1,101 +0,0 @@ -// DVD-style bouncing logo (image texture, e.g. assets/mxl.png) on studio black, -// 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). -// Stateless bounce: logo position is a pure function of the frame tick -// (two triangle waves), so no velocity/position state is carried between frames. -// -// Sub-pixel motion: the logo plane is bilinearly resampled at the fractional -// position, so edges glide continuously instead of stepping whole pixels. -// The blend is built from mix() chains, which reproduce texel values exactly -// in solid regions (mix of equal values is exact in f32). -// -// Logo plane (binding 2, read-only), one u32 per logo pixel: -// logo[0] = logo_w | logo_h << 16 -// logo[1 + dy*logo_w + dx] = opaque<<30 | Cr<<20 | Cb<<10 | Y - -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 out: array; -@group(0) @binding(1) var params: Params; -@group(0) @binding(2) var logo: array; - -// 0 -> 1 -> 0 triangle wave; input is elapsed bounce cycles. -fn tri01(x: f32) -> f32 { - let f = fract(x); - return select(2.0 * f, 2.0 - 2.0 * f, f > 0.5); -} - -// Deviation of a logo texel from the background (64/512/512); zero when the -// texel is transparent or out of the plane. u32 wraparound in the callers -// makes every lower/upper bound case fail the range test, so out-of-bounds -// needs no special handling. -fn tap(lw: u32, lh: u32, dx: u32, dy: u32) -> vec3 { - if (dx < lw && dy < lh) { - let s = logo[1u + dy * lw + dx]; - if ((s >> 30u) & 1u == 1u) { - return vec3(f32(s & 0x3FFu) - 64.0, - f32((s >> 10u) & 0x3FFu) - 512.0, - f32((s >> 20u) & 0x3FFu) - 512.0); - } - } - return vec3(0.0, 0.0, 0.0); -} - -@compute @workgroup_size(64) -fn main(@builtin(global_invocation_id) gid: vec3) { - 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; - - // Logo geometry comes from the plane header. - let hdr = logo[0]; - let lw = hdr & 0xFFFFu; - let lh = hdr >> 16u; - - // Slow glide: ~8.6 px/frame horizontal, ~6.5 px/frame vertical at 25 fps. - // Near-irrational speed ratio (~sqrt(2)) covers the whole frame and - // rarely repeats the path, like the DVD logo. - let t = f32(params.frame) / 25.0; - let lx = tri01(t * 0.061) * max(f32(params.width) - f32(lw), 0.0); - let ly = tri01(t * 0.088) * max(f32(params.height) - f32(lh), 0.0); - let bx = u32(lx); // integer part of the position - let by = u32(ly); - let fx = lx - f32(bx); // fractional part, [0,1) - let fy = ly - f32(by); - - var y: array; - var cb: array; - var cr: array; - for (var i = 0u; i < 6u; i++) { - // Screen pixel (x+i, py) covers logo-plane footprint [m-f, m+1-f) x - // [n-f, n+1-f): texel m-1 gets weight fx, texel m gets 1-fx (same - // vertically). Chroma is blended the same way and packed at even - // columns, matching v210 4:2:2 co-siting. - let m = (x + i) - bx; - let n = py - by; - let top = mix(tap(lw, lh, m, n - 1u), tap(lw, lh, m - 1u, n - 1u), fx); - let bot = mix(tap(lw, lh, m, n), tap(lw, lh, m - 1u, n), fx); - let d = mix(bot, top, fy); - y[i] = u32(64.0 + d.x + 0.5); - cb[i] = u32(512.0 + d.y + 0.5); - cr[i] = u32(512.0 + d.z + 0.5); - } - - // 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); -}