diff --git a/assets/fonts/JetBrainsMonoNLNerdFontMono-Regular.ttf b/assets/fonts/JetBrainsMonoNLNerdFontMono-Regular.ttf new file mode 100644 index 0000000..ec7c67b Binary files /dev/null and b/assets/fonts/JetBrainsMonoNLNerdFontMono-Regular.ttf differ diff --git a/assets/fonts/OFL.txt b/assets/fonts/OFL.txt new file mode 100644 index 0000000..8bee414 --- /dev/null +++ b/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. diff --git a/assets/fonts/README.md b/assets/fonts/README.md new file mode 100644 index 0000000..d467cca --- /dev/null +++ b/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: + +The font is distributed under the SIL Open Font License 1.1. The complete +license and copyright notice are included in `OFL.txt`. diff --git a/assets/mxl.png b/assets/mxl.png new file mode 100644 index 0000000..ed6a040 Binary files /dev/null and b/assets/mxl.png differ diff --git a/cmd/wgpu-gen/main.go b/cmd/wgpu-gen/main.go index 2c2743f..9d617cb 100644 --- a/cmd/wgpu-gen/main.go +++ b/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 { - log.Printf("grains written=%d, index=%d", written, idx) + if grainsWritten%100 == 0 { + log.Printf("grains written=%d, index=%d", grainsWritten, idx) } // Pace ourselves to roughly the grain rate mxl.SleepNs(mxl.NsUntilIndex(idx, rate)) diff --git a/go.mod b/go.mod index 63f8bdd..d136af6 100644 --- a/go.mod +++ b/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 ) diff --git a/go.sum b/go.sum index 690359c..310f0eb 100644 --- a/go.sum +++ b/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.3/go.mod h1:erhaIjD9psJKjHqAhn5MWdJETcACR1NGoqLhkTAHkFM= +github.com/gogpu/wgpu v0.34.5 h1:fgfyyhHjhD1mSkScWbeeKXZ2MLDxzC9bemKWPTneyUs= +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/sys v0.47.0/go.mod h1:4GL1E5IUh+htKOUEOaiffhrAeqysfVGipDYzABqnCmw= +golang.org/x/image v0.46.0 h1:b1+oYj0Jbp6K5MDT4i4/eZpYlk3V8SJhhDKh6LBHAyQ= +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= diff --git a/internal/generator/logo.go b/internal/generator/logo.go new file mode 100644 index 0000000..ef11ee2 --- /dev/null +++ b/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) +} diff --git a/internal/generator/logo_test.go b/internal/generator/logo_test.go new file mode 100644 index 0000000..83c1995 --- /dev/null +++ b/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 +} diff --git a/internal/generator/text.go b/internal/generator/text.go new file mode 100644 index 0000000..04f4de4 --- /dev/null +++ b/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 +} diff --git a/internal/generator/text_test.go b/internal/generator/text_test.go new file mode 100644 index 0000000..5a12ba3 --- /dev/null +++ b/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]) + } + } + } +} diff --git a/internal/generator/wgpu.go b/internal/generator/wgpu.go index d1f0ea7..c957a01 100644 --- a/internal/generator/wgpu.go +++ b/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", - 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}}, - }, + Label: "v210-bgl", + Entries: bglEntries, }); 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{ - {Binding: 0, Buffer: g.out, Size: g.frameSize}, - {Binding: 1, Buffer: g.uniform, Size: uint64(len(g.params))}, - }, + Entries: bgEntries, }); 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 diff --git a/internal/generator/wgpu_dvd_logo_test.go b/internal/generator/wgpu_dvd_logo_test.go new file mode 100644 index 0000000..56840cb --- /dev/null +++ b/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 +} diff --git a/kernels/v210_dvd_logo.wgsl b/kernels/v210_dvd_logo.wgsl new file mode 100644 index 0000000..5258c4f --- /dev/null +++ b/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 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); +} diff --git a/kernels/v210_y_ramp.wgsl b/kernels/v210_y_ramp.wgsl new file mode 100644 index 0000000..85a9a93 --- /dev/null +++ b/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 out: array; +@group(0) @binding(1) var 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) { + 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(721u, 674u, 581u, 534u, 251u, 204u, 111u); + let cb_tab = array(512u, 176u, 589u, 253u, 771u, 435u, 848u); + let cr_tab = array(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; + var cb: array; + var cr: array; + 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); +}