Text + MXL logo pattern

assets/fonts/OFL.txt

@@ -0,0 +1,93 @@
+Copyright 2020 The JetBrains Mono Project Authors (https://github.com/JetBrains/JetBrainsMono)
+
+This Font Software is licensed under the SIL Open Font License, Version 1.1.
+This license is copied below, and is also available with a FAQ at:
+https://scripts.sil.org/OFL
+
+
+-----------------------------------------------------------
+SIL OPEN FONT LICENSE Version 1.1 - 26 February 2007
+-----------------------------------------------------------
+
+PREAMBLE
+The goals of the Open Font License (OFL) are to stimulate worldwide
+development of collaborative font projects, to support the font creation
+efforts of academic and linguistic communities, and to provide a free and
+open framework in which fonts may be shared and improved in partnership
+with others.
+
+The OFL allows the licensed fonts to be used, studied, modified and
+redistributed freely as long as they are not sold by themselves. The
+fonts, including any derivative works, can be bundled, embedded, 
+redistributed and/or sold with any software provided that any reserved
+names are not used by derivative works. The fonts and derivatives,
+however, cannot be released under any other type of license. The
+requirement for fonts to remain under this license does not apply
+to any document created using the fonts or their derivatives.
+
+DEFINITIONS
+"Font Software" refers to the set of files released by the Copyright
+Holder(s) under this license and clearly marked as such. This may
+include source files, build scripts and documentation.
+
+"Reserved Font Name" refers to any names specified as such after the
+copyright statement(s).
+
+"Original Version" refers to the collection of Font Software components as
+distributed by the Copyright Holder(s).
+
+"Modified Version" refers to any derivative made by adding to, deleting,
+or substituting -- in part or in whole -- any of the components of the
+Original Version, by changing formats or by porting the Font Software to a
+new environment.
+
+"Author" refers to any designer, engineer, programmer, technical
+writer or other person who contributed to the Font Software.
+
+PERMISSION & CONDITIONS
+Permission is hereby granted, free of charge, to any person obtaining
+a copy of the Font Software, to use, study, copy, merge, embed, modify,
+redistribute, and sell modified and unmodified copies of the Font
+Software, subject to the following conditions:
+
+1) Neither the Font Software nor any of its individual components,
+in Original or Modified Versions, may be sold by itself.
+
+2) Original or Modified Versions of the Font Software may be bundled,
+redistributed and/or sold with any software, provided that each copy
+contains the above copyright notice and this license. These can be
+included either as stand-alone text files, human-readable headers or
+in the appropriate machine-readable metadata fields within text or
+binary files as long as those fields can be easily viewed by the user.
+
+3) No Modified Version of the Font Software may use the Reserved Font
+Name(s) unless explicit written permission is granted by the corresponding
+Copyright Holder. This restriction only applies to the primary font name as
+presented to the users.
+
+4) The name(s) of the Copyright Holder(s) or the Author(s) of the Font
+Software shall not be used to promote, endorse or advertise any
+Modified Version, except to acknowledge the contribution(s) of the
+Copyright Holder(s) and the Author(s) or with their explicit written
+permission.
+
+5) The Font Software, modified or unmodified, in part or in whole,
+must be distributed entirely under this license, and must not be
+distributed under any other license. The requirement for fonts to
+remain under this license does not apply to any document created
+using the Font Software.
+
+TERMINATION
+This license becomes null and void if any of the above conditions are
+not met.
+
+DISCLAIMER
+THE FONT SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO ANY WARRANTIES OF
+MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT
+OF COPYRIGHT, PATENT, TRADEMARK, OR OTHER RIGHT. IN NO EVENT SHALL THE
+COPYRIGHT HOLDER BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
+INCLUDING ANY GENERAL, SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL
+DAMAGES, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+FROM, OUT OF THE USE OR INABILITY TO USE THE FONT SOFTWARE OR FROM
+OTHER DEALINGS IN THE FONT SOFTWARE.

assets/fonts/README.md

@@ -0,0 +1,9 @@
+# JetBrainsMonoNL Nerd Font Mono
+
+`JetBrainsMonoNLNerdFontMono-Regular.ttf` is the regular, no-ligatures,
+monospaced variant of JetBrains Mono patched by the Nerd Fonts project.
+
+Source: <https://github.com/ryanoasis/nerd-fonts>
+
+The font is distributed under the SIL Open Font License 1.1. The complete
+license and copyright notice are included in `OFL.txt`.

cmd/wgpu-gen/main.go

@@ -182,7 +182,7 @@ func flagSetAddFlags(fs *pflag.FlagSet, args *appArgs) {
 	// Video pattern flags
 	fs.Uint8VarP(&args.pattern, "pattern", "p", 0, "Video pattern type [TODO]")
 	fs.BoolVar(&args.listPatterns, "list-patterns", false, "List video available video patterns and exit [TODO]")
-	fs.StringVarP(&args.textOverlay, "text", "t", "", "Text overlay above video pattern [TODO]")
+	fs.StringVarP(&args.textOverlay, "text", "t", "", "Text overlay above video pattern")
 	fs.UintVar(&args.videoWidth, "width", 1920, "Video pattern width")
 	fs.UintVar(&args.videoHeight, "height", 1080, "Video pattern height")
 	fs.StringVar(&args.videoFPS, "fps", "25", "Video pattern FPS")
@@ -261,12 +261,37 @@ func main() {
 	// TODO: if init failed -> CPU generator
 	// But before i should create same patterns for both generators
 	// and text overlays
-	gen, err := generator.NewWGPUGenerator(vi.width, vi.height, "kernels/v210_bars_move.wgsl")
+	kernel := "kernels/v210_bars_move.wgsl"
+	// kernel = "kernels/v210_dvd_logo.wgsl"
+	var genOpts []generator.WGPUOption
+	if kernel == "kernels/v210_dvd_logo.wgsl" {
+		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, kernel, genOpts...)
 	if err != nil {
 		log.Fatalf("wgpu init failed: %v", err)
 	}
 	defer gen.Close()
 
+	// Static text overlay: rasterized + pre-packed once, stamped on each
+	// frame after the GPU render (microseconds per frame).
+	var overlay *generator.TextOverlay
+	if args.textOverlay != "" {
+		face, err := generator.LoadFace("assets/fonts/JetBrainsMonoNLNerdFontMono-Regular.ttf", 48)
+		if err != nil {
+			log.Fatalf("text overlay init failed: %v", err)
+		}
+		defer face.Close()
+		overlay, err = generator.NewTextOverlay(args.textOverlay, int(vi.width), int(vi.height), face)
+		if err != nil {
+			log.Fatalf("text overlay init failed: %v", err)
+		}
+	}
+
 	inst, err := mxl.NewInstance(mxlDomain, "")
 	if err != nil {
 		log.Fatalf("MXL Init Failed: %v", err)
@@ -291,14 +316,14 @@ func main() {
 	signal.Notify(stop, os.Interrupt, syscall.SIGTERM)
 
 	// core loop
-	var written int64
+	var grainsWritten int64
 	// animation clock: small counter, not the huge grain index.
 	// Reason: current wgpu shaders limitations
 	var tick uint32
 	for {
 		select {
 		case <-stop:
-			log.Printf("stopping after %d grains", written)
+			log.Printf("stopping after %d grains", grainsWritten)
 			return
 		default:
 		}
@@ -310,15 +335,20 @@ func main() {
 		if err := gen.GenerateFrame(gwa.Payload, int(tick)); err != nil {
 			log.Fatalf("GenerateFrame(%d): %v", idx, err)
 		}
+		if overlay != nil {
+			if err := overlay.ApplyV210(gwa.Payload); err != nil {
+				log.Fatalf("text overlay: %v", err)
+			}
+		}
 		if err := gwa.Commit(gwa.TotalSlices, 0); err != nil {
 			log.Fatalf("Commit(%d): %v", idx, err)
 		}
 
-		written++
+		grainsWritten++
 		idx++
 		tick++
-		if written%100 == 0 {
+		if grainsWritten%100 == 0 {
-			log.Printf("grains written=%d, index=%d", written, idx)
+			log.Printf("grains written=%d, index=%d", grainsWritten, idx)
 		}
 		// Pace ourselves to roughly the grain rate
 		mxl.SleepNs(mxl.NsUntilIndex(idx, rate))

go.mod

@@ -4,10 +4,11 @@ go 1.26.0
 
 require (
 	github.com/gogpu/gputypes v0.8.0
-	github.com/gogpu/wgpu v0.34.3
+	github.com/gogpu/wgpu v0.34.5
 	github.com/google/uuid v1.6.0
 	github.com/qvest-digital/go-mxl v1.1.0-rc.4
 	github.com/spf13/pflag v1.0.10
+	golang.org/x/image v0.46.0
 )
 
 require (
@@ -15,5 +16,6 @@ require (
 	github.com/go-webgpu/webgpu v0.5.5 // indirect
 	github.com/gogpu/gpucontext v0.31.3 // indirect
 	github.com/gogpu/naga v0.19.0 // indirect
-	golang.org/x/sys v0.47.0 // indirect
+	golang.org/x/sys v0.48.0 // indirect
+	golang.org/x/text v0.42.0 // indirect
 )

go.sum

@@ -8,13 +8,17 @@ github.com/gogpu/gputypes v0.8.0 h1:Iw4g9lqydTcAFKKQWfgbUpdY6WxObGtjf9SuRkzYMJw=
 github.com/gogpu/gputypes v0.8.0/go.mod h1:cnXrDMwTpWTvJLW1Vreop3PcT6a2YP/i3s91rPaOavw=
 github.com/gogpu/naga v0.19.0 h1:+Pu7kahdMhvRWtpEbTW5YFQZPoklccMO4vryEk6w9zU=
 github.com/gogpu/naga v0.19.0/go.mod h1:15sQaHKkbqXcwTN+hHYGLsA0WBBnkmYzne/eF5p5WEg=
-github.com/gogpu/wgpu v0.34.3 h1:oV7K5ueBqxZuvp1w9KpArzj15+eOrm/l9ZnM3jcKbXA=
+github.com/gogpu/wgpu v0.34.5 h1:fgfyyhHjhD1mSkScWbeeKXZ2MLDxzC9bemKWPTneyUs=
-github.com/gogpu/wgpu v0.34.3/go.mod h1:erhaIjD9psJKjHqAhn5MWdJETcACR1NGoqLhkTAHkFM=
+github.com/gogpu/wgpu v0.34.5/go.mod h1:erhaIjD9psJKjHqAhn5MWdJETcACR1NGoqLhkTAHkFM=
 github.com/google/uuid v1.6.0 h1:NIvaJDMOsjHA8n1jAhLSgzrAzy1Hgr+hNrb57e+94F0=
 github.com/google/uuid v1.6.0/go.mod h1:TIyPZe4MgqvfeYDBFedMoGGpEw/LqOeaOT+nhxU+yHo=
 github.com/qvest-digital/go-mxl v1.1.0-rc.4 h1:qgOGSIv53HYQLTB06m39c6WE2xqMceObViZy5A3rZlA=
 github.com/qvest-digital/go-mxl v1.1.0-rc.4/go.mod h1:cYzyT+S/AONytsKr/DozzFQP2OrNrg/JsQOSjvwn+Rk=
 github.com/spf13/pflag v1.0.10 h1:4EBh2KAYBwaONj6b2Ye1GiHfwjqyROoF4RwYO+vPwFk=
 github.com/spf13/pflag v1.0.10/go.mod h1:McXfInJRrz4CZXVZOBLb0bTZqETkiAhM9Iw0y3An2Bg=
-golang.org/x/sys v0.47.0 h1:o7XGOvZQCADBQQ4Y7VNq2dRWQR7JmOUW8Kxx4ZsNgWs=
+golang.org/x/image v0.46.0 h1:b1+oYj0Jbp6K5MDT4i4/eZpYlk3V8SJhhDKh6LBHAyQ=
-golang.org/x/sys v0.47.0/go.mod h1:4GL1E5IUh+htKOUEOaiffhrAeqysfVGipDYzABqnCmw=
+golang.org/x/image v0.46.0/go.mod h1:3B3W05VGVQyuXucLINLjXKrqISASfi4Xj+iCVkLMwew=
+golang.org/x/sys v0.48.0 h1:bbX/i/6MgT9BVLM9RT1thmxL04yeTAhbEz4SyadbXoo=
+golang.org/x/sys v0.48.0/go.mod h1:hNLxWAXmnKAxqDtdwIYC4bM9oQPEecfsnNMuSxOs3og=
+golang.org/x/text v0.42.0 h1:JbOZXgfeCPU9gacVtYliJqOhD+zhrEqK4LfdpmlUZqI=
+golang.org/x/text v0.42.0/go.mod h1:ojzP1Z+2QtioaF8DTtO8K5q7JWVVYwZKenzujK0Zd0E=

internal/generator/logo.go

@@ -0,0 +1,76 @@
+// 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)
+}

internal/generator/logo_test.go

@@ -0,0 +1,64 @@
+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
+}

internal/generator/text.go

@@ -0,0 +1,151 @@
+// Static text overlay for v210 frames: text is rasterized with real font
+// anti-aliasing onto a black backing box, pre-packed into whole v210 blocks
+// once, and stamped per frame with a block-aligned copy (microseconds).
+package generator
+
+import (
+	"encoding/binary"
+	"fmt"
+	"image"
+	"os"
+
+	"golang.org/x/image/font"
+	"golang.org/x/image/font/opentype"
+	"golang.org/x/image/math/fixed"
+)
+
+type TextOverlay struct {
+	frameW int
+	x, y   int    // box top-left in frame pixels
+	w, h   int    // box size in pixels (w is a multiple of 6)
+	cov    []byte // 8-bit anti-alias coverage, row-major w*h
+	blocks []byte // pre-packed v210 tile: h rows x (w/6) blocks x 16 bytes
+}
+
+// LoadFace parses a TTF/OTF file and builds a render-ready face at the given
+// pixel size (DPI 72, full hinting for crisp video text). Call once and
+// reuse for every overlay.
+func LoadFace(path string, size float64) (font.Face, error) {
+	data, err := os.ReadFile(path)
+	if err != nil {
+		return nil, fmt.Errorf("text: %w", err)
+	}
+	f, err := opentype.Parse(data)
+	if err != nil {
+		return nil, fmt.Errorf("text: parse %s: %w", path, err)
+	}
+	face, err := opentype.NewFace(f, &opentype.FaceOptions{
+		Size:    size,
+		DPI:     72,
+		Hinting: font.HintingFull,
+	})
+	if err != nil {
+		return nil, fmt.Errorf("text: face %s: %w", path, err)
+	}
+	return face, nil
+}
+
+// NewTextOverlay rasterizes text (white, anti-aliased) onto a black backing
+// box, centered horizontally. Box width and x are multiples of 6 so the box
+// covers whole v210 blocks and the per-frame copy is word-aligned.
+func NewTextOverlay(text string, frameW, frameH int, face font.Face) (*TextOverlay, error) {
+	if text == "" {
+		return nil, fmt.Errorf("text: empty string")
+	}
+	if frameW <= 0 || frameW%6 != 0 {
+		return nil, fmt.Errorf("text: frame width %d not a multiple of 6", frameW)
+	}
+	pix := func(v fixed.Int26_6) int { return int(v+63) / 64 } // ceil to px
+	m := face.Metrics()
+	textH := pix(m.Ascent+m.Descent) + 2 // slack for AA/hinting overshoot
+
+	// Box width: sum of per-rune advances (uniform for monospace fonts).
+	spaceAdv, ok := face.GlyphAdvance(' ')
+	if !ok {
+		spaceAdv = m.Height
+	}
+	var total fixed.Int26_6
+	for _, r := range text {
+		a, ok := face.GlyphAdvance(r)
+		if !ok {
+			a = spaceAdv
+		}
+		total += a
+	}
+	textW := pix(total)
+
+	padX := pix(m.Ascent) / 3
+	padY := padX / 2
+	w := 2*padX + textW
+	if r := w % 6; r != 0 {
+		w += 6 - r
+	}
+	if w > frameW {
+		return nil, fmt.Errorf("text: box %dpx wider than frame %dpx", w, frameW)
+	}
+	h := 2*padY + textH
+	const topMargin = 48
+	if topMargin+h > frameH {
+		return nil, fmt.Errorf("text: box %dpx does not fit frame height %dpx", topMargin+h, frameH)
+	}
+
+	o := &TextOverlay{
+		frameW: frameW,
+		x:      ((frameW - w) / 2 / 6) * 6,
+		y:      topMargin,
+		w:      w,
+		h:      h,
+		cov:    make([]byte, w*h),
+	}
+
+	// Rasterize: drawing white onto *image.Alpha writes pure glyph coverage.
+	dst := image.NewAlpha(image.Rect(0, 0, textW, textH))
+	(&font.Drawer{
+		Dst: dst, Src: image.White, Face: face,
+		Dot: fixed.P(0, pix(m.Ascent)),
+	}).DrawString(text)
+	for row := 0; row < textH; row++ {
+		copy(o.cov[(padY+row)*w+padX:], dst.Pix[row*dst.Stride:row*dst.Stride+textW])
+	}
+
+	o.pack()
+	return o, nil
+}
+
+// pack pre-renders the box into v210 blocks: white text (Y=940) blended by
+// coverage over a black box (Y=64), neutral chroma (512). Same word layout
+// as the kernels, chroma co-sited with luma samples 0/2/4.
+func (o *TextOverlay) pack() {
+	const neutral = uint32(512)
+	o.blocks = make([]byte, o.h*(o.w/6)*16)
+	for row := 0; row < o.h; row++ {
+		for blk := 0; blk < o.w/6; blk++ {
+			var y [6]uint32
+			for i := 0; i < 6; i++ {
+				c := int(o.cov[row*o.w+blk*6+i])
+				y[i] = 64 + uint32((c*876+127)/255)
+			}
+			off := (row*(o.w/6) + blk) * 16
+			binary.LittleEndian.PutUint32(o.blocks[off+0:], neutral|(y[0]<<10)|(neutral<<20))
+			binary.LittleEndian.PutUint32(o.blocks[off+4:], y[1]|(neutral<<10)|(y[2]<<20))
+			binary.LittleEndian.PutUint32(o.blocks[off+8:], neutral|(y[3]<<10)|(neutral<<20))
+			binary.LittleEndian.PutUint32(o.blocks[off+12:], y[4]|(neutral<<10)|(y[5]<<20))
+		}
+	}
+}
+
+// ApplyV210 stamps the pre-packed text tile over a packed v210 frame.
+// The tile occupies whole 16-byte blocks, so this is a row-wise copy.
+func (o *TextOverlay) ApplyV210(dest []byte) error {
+	lastPixel := (o.y+o.h-1)*o.frameW + o.x + o.w - 1
+	if need := (lastPixel/6 + 1) * 16; len(dest) < need {
+		return fmt.Errorf("text: dest %d bytes too small, need %d", len(dest), need)
+	}
+	tileStride := (o.w / 6) * 16
+	for row := 0; row < o.h; row++ {
+		frameBlock := ((o.y+row)*o.frameW + o.x) / 6
+		src := o.blocks[row*tileStride : (row+1)*tileStride]
+		copy(dest[frameBlock*16:frameBlock*16+tileStride], src)
+	}
+	return nil
+}

internal/generator/text_test.go

@@ -0,0 +1,156 @@
+package generator
+
+import (
+	"encoding/binary"
+	"path/filepath"
+	"testing"
+
+	"golang.org/x/image/font"
+)
+
+func testFace(t *testing.T, size float64) font.Face {
+	t.Helper()
+	face, err := LoadFace(filepath.Join("..", "..", "assets", "fonts",
+		"JetBrainsMonoNLNerdFontMono-Regular.ttf"), size)
+	if err != nil {
+		t.Fatalf("LoadFace: %v", err)
+	}
+	t.Cleanup(func() { face.Close() })
+	return face
+}
+
+func TestNewTextOverlay(t *testing.T) {
+	const frameW, frameH = 1920, 1080
+
+	o, err := NewTextOverlay("AB", frameW, frameH, testFace(t, 48))
+	if err != nil {
+		t.Fatalf("NewTextOverlay: %v", err)
+	}
+	if o.w%6 != 0 {
+		t.Fatalf("w=%d not a multiple of 6", o.w)
+	}
+	if o.x%6 != 0 {
+		t.Fatalf("x=%d not a multiple of 6", o.x)
+	}
+	if center := o.x + o.w/2; center < frameW/2-3 || center > frameW/2+3 {
+		t.Fatalf("box center %d not near frame center %d", center, frameW/2)
+	}
+	if o.y != 48 {
+		t.Fatalf("y=%d, want 48", o.y)
+	}
+	if o.blocks == nil || len(o.blocks) != o.h*(o.w/6)*16 {
+		t.Fatalf("pre-packed tile has wrong size")
+	}
+	if o.cov[2*o.w+2] != 0 {
+		t.Fatalf("padding pixel (2,2) covered: %d", o.cov[2*o.w+2])
+	}
+	nonzero, full := 0, 0
+	for _, c := range o.cov {
+		if c > 0 {
+			nonzero++
+		}
+		if c == 255 {
+			full++
+		}
+	}
+	if nonzero < 100 {
+		t.Fatalf("too few covered pixels: %d", nonzero)
+	}
+	if full == 0 {
+		t.Fatalf("no fully covered (255) pixels")
+	}
+
+	// Text without descenders ("AB") must leave the bottom quarter empty;
+	// text with a descender ("Ag") must have coverage there.
+	for _, tc := range []struct {
+		text    string
+		wantAny bool
+	}{
+		{"AB", false},
+		{"Ag", true},
+	} {
+		o, err := NewTextOverlay(tc.text, frameW, frameH, testFace(t, 48))
+		if err != nil {
+			t.Fatalf("NewTextOverlay(%q): %v", tc.text, err)
+		}
+		any := false
+		for row := o.h * 3 / 4; row < o.h; row++ {
+			for col := 0; col < o.w; col++ {
+				if o.cov[row*o.w+col] > 0 {
+					any = true
+				}
+			}
+		}
+		if any != tc.wantAny {
+			t.Fatalf("%q: coverage in bottom quarter = %v, want %v", tc.text, any, tc.wantAny)
+		}
+	}
+}
+
+func TestTextOverlayApplyV210(t *testing.T) {
+	const frameW, frameH = 1920, 1080
+	o, err := NewTextOverlay("MXL", frameW, frameH, testFace(t, 48))
+	if err != nil {
+		t.Fatalf("NewTextOverlay: %v", err)
+	}
+
+	frame := make([]byte, frameW*frameH*8/3)
+	for i := range frame {
+		frame[i] = 0x5A // marker: untouched regions must survive
+	}
+	if err := o.ApplyV210(frame); err != nil {
+		t.Fatalf("ApplyV210: %v", err)
+	}
+
+	sample := func(x, y int) (yc, cb, cr uint32) {
+		p := y*frameW + x
+		off := (p / 6) * 16
+		w0 := binary.LittleEndian.Uint32(frame[off:])
+		w1 := binary.LittleEndian.Uint32(frame[off+4:])
+		w2 := binary.LittleEndian.Uint32(frame[off+8:])
+		w3 := binary.LittleEndian.Uint32(frame[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
+		}
+	}
+
+	// Every box pixel: Y from coverage, neutral chroma at even columns.
+	for row := 0; row < o.h; row++ {
+		for col := 0; col < o.w; col++ {
+			wantY := uint32(64 + (int(o.cov[row*o.w+col])*876+127)/255)
+			yc, cb, cr := sample(o.x+col, o.y+row)
+			if yc != wantY {
+				t.Fatalf("box pixel (%d,%d): Y=%d, want %d", o.x+col, o.y+row, yc, wantY)
+			}
+			if (o.x+col)%2 == 0 && (cb != 512 || cr != 512) {
+				t.Fatalf("box pixel (%d,%d): Cb=%d Cr=%d, want 512/512", o.x+col, o.y+row, cb, cr)
+			}
+		}
+	}
+
+	// Blocks just outside the box (right edge and below) must be untouched.
+	for _, p := range [][2]int{
+		{o.x + o.w, o.y},       // right of box, block-aligned
+		{o.x, o.y + o.h},       // below box
+		{o.x + o.w, o.y + o.h}, // corner
+		{100, 1000},            // far away
+	} {
+		off := ((p[1]*frameW + p[0]) / 6) * 16
+		for i := 0; i < 16; i++ {
+			if frame[off+i] != 0x5A {
+				t.Fatalf("block at (%d,%d) modified outside the box", p[0], p[1])
+			}
+		}
+	}
+}

internal/generator/wgpu.go

@@ -35,6 +35,8 @@ type WGPUGenerator struct {
 	out       *wgpu.Buffer
 	host      *wgpu.Buffer
 	uniform   *wgpu.Buffer
+	logo      *wgpu.Buffer
+	logoPlane []uint32
 	params    []byte
 	width     int
 	height    int
@@ -44,13 +46,27 @@ type WGPUGenerator struct {
 
 var _ FrameGenerator = (*WGPUGenerator)(nil)
 
-func NewWGPUGenerator(width, height uint, kernelPath string) (*WGPUGenerator, error) {
+// 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),
 		height: int(height),
 		blocks: int(width*height) / 6,
 		params: make([]byte, 16),
 	}
+	for _, opt := range opts {
+		opt(g)
+	}
 	g.frameSize = uint64(g.blocks) * 16
 	binary.LittleEndian.PutUint32(g.params[0:], uint32(width))
 	binary.LittleEndian.PutUint32(g.params[4:], uint32(height))
@@ -105,22 +121,48 @@ func NewWGPUGenerator(width, height uint, kernelPath string) (*WGPUGenerator, er
 		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",
+		Label:   "v210-bgl",
-		Entries: []gputypes.BindGroupLayoutEntry{
+		Entries: bglEntries,
-			{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: []wgpu.BindGroupEntry{
+		Entries: bgEntries,
-			{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)
@@ -212,6 +254,10 @@ 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

internal/generator/wgpu_dvd_logo_test.go

@@ -0,0 +1,169 @@
+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
+}

kernels/v210_dvd_logo.wgsl

@@ -0,0 +1,101 @@
+// 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<storage, read_write> out: array<u32>;
+@group(0) @binding(1) var<uniform> params: Params;
+@group(0) @binding(2) var<storage, read> logo: array<u32>;
+
+// 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<f32> {
+    if (dx < lw && dy < lh) {
+        let s = logo[1u + dy * lw + dx];
+        if ((s >> 30u) & 1u == 1u) {
+            return vec3<f32>(f32(s & 0x3FFu) - 64.0,
+                             f32((s >> 10u) & 0x3FFu) - 512.0,
+                             f32((s >> 20u) & 0x3FFu) - 512.0);
+        }
+    }
+    return vec3<f32>(0.0, 0.0, 0.0);
+}
+
+@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;
+
+    // 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<u32, 6>;
+    var cb: array<u32, 6>;
+    var cr: array<u32, 6>;
+    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);
+}

kernels/v210_y_ramp.wgsl

@@ -0,0 +1,75 @@
+// 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);
+        y[i] = 64u + (u32(px) * 876u) / params.width;
+        cb[i] = 512u;
+        cr[i] = 512u;
+        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[i];
+            cb[i] = 1024u - cb[i];
+            cr[i] = 1024u - cr[i];
+        }
+    }
+
+    // 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);
+}