V210 issue fix + correct packing

cmd/mxl-pattern/main.go

@@ -146,9 +146,8 @@ func validateVideoArgs(args *appArgs) error {
 	if args.videoFlowDefFile != "" {
 		return nil
 	}
-	if args.videoWidth == 0 || args.videoWidth%6 != 0 {
+	if args.videoWidth == 0 || args.videoWidth%2 != 0 {
-		// v210 stores 6 pixels in each 16-byte block.
+		return fmt.Errorf("video width must be greater than zero and even for 4:2:2 video")
-		return fmt.Errorf("video width must be greater than zero and divisible by 6")
 	}
 	if args.videoHeight == 0 {
 		return fmt.Errorf("video height must be greater than zero")

internal/flowdef/flowdef.go

@@ -132,8 +132,8 @@ func (v Video) Validate() error {
 	if v.ColorSpace != ColorSpaceBT709 {
 		return fmt.Errorf("colorspace must be %q, got %q", ColorSpaceBT709, v.ColorSpace)
 	}
-	if v.FrameWidth == 0 || v.FrameWidth%6 != 0 {
+	if v.FrameWidth == 0 || v.FrameWidth%2 != 0 {
-		return fmt.Errorf("frame_width must be greater than zero and divisible by 6, got %d", v.FrameWidth)
+		return fmt.Errorf("frame_width must be greater than zero and even for 4:2:2 video, got %d", v.FrameWidth)
 	}
 	if v.FrameHeight == 0 {
 		return fmt.Errorf("frame_height must be greater than zero")

internal/flowdef/flowdef_test.go

@@ -28,10 +28,10 @@ func TestNewV210Video(t *testing.T) {
 	}
 }
 
-func TestNewV210VideoRejectsInvalidWidth(t *testing.T) {
+func TestNewV210VideoRejectsOddWidth(t *testing.T) {
 	_, err := NewV210Video(testVideoID, 1919, 1080, Rational{Numerator: 25, Denominator: 1})
-	if err == nil || !strings.Contains(err.Error(), "divisible by 6") {
+	if err == nil || !strings.Contains(err.Error(), "even") {
-		t.Fatalf("error = %v, want width divisibility error", err)
+		t.Fatalf("error = %v, want even-width error", err)
 	}
 }
 

internal/generator/cpu.go

@@ -29,8 +29,8 @@ func NewCPUGenerator(
 	if width == 0 || height == 0 {
 		return nil, fmt.Errorf("cpu: width and height must be greater than zero, got %dx%d", width, height)
 	}
-	if width%6 != 0 {
+	if width%2 != 0 {
-		return nil, fmt.Errorf("cpu: width must be divisible by 6, got %d", width)
+		return nil, fmt.Errorf("cpu: width must be even for 4:2:2 video, got %d", width)
 	}
 	if baseRenderer == nil {
 		return nil, fmt.Errorf("cpu: base renderer is nil")
@@ -39,7 +39,7 @@ func NewCPUGenerator(
 	g := &CPUGenerator{
 		width:  int(width),
 		height: int(height),
-		base:   make([]byte, int(width/6*height)*16),
+		base:   make([]byte, V210FrameSize(int(width), int(height))),
 		patch:  patch,
 	}
 	if err := baseRenderer(g.base, g.width, g.height, 0); err != nil {

internal/generator/cpu_moving_square.go

@@ -45,6 +45,7 @@ func patchMovingSquare(
 	lastBlockX := min(width, (lastPixelX+5)/6*6)
 	firstY := max(0, int(math.Floor(bounds.minY)))
 	lastY := min(height, int(math.Ceil(bounds.maxY)))
+	stride := V210LineSize(width)
 
 	for y := firstY; y < lastY; y++ {
 		for x := firstBlockX; x < lastBlockX; x += 6 {
@@ -57,7 +58,7 @@ func patchMovingSquare(
 				}
 				pixels[i] = color
 			}
-			offset := (y*width + x) / 6 * 16
+			offset := y*stride + x/6*16
 			packV210Block(dst[offset:offset+16], pixels)
 		}
 	}

internal/generator/cpu_patterns.go

@@ -127,13 +127,14 @@ func movingSquarePatch(baseColor baseColorFunc) FrameRenderer {
 }
 
 func renderBasePattern(dst []byte, width, height int, baseColor baseColorFunc) error {
+	stride := V210LineSize(width)
 	for y := 0; y < height; y++ {
 		for x := 0; x < width; x += 6 {
 			var pixels [6]YCbCr10
 			for i := range pixels {
 				pixels[i] = baseColor(x+i, y, width, height)
 			}
-			offset := (y*width + x) / 6 * 16
+			offset := y*stride + x/6*16
 			packV210Block(dst[offset:offset+16], pixels)
 		}
 	}

internal/generator/cpu_test.go

@@ -26,7 +26,7 @@ func TestNewCPUGeneratorValidation(t *testing.T) {
 	}{
 		{name: "zero width", height: 1, renderer: renderer, wantErrSub: "greater than zero"},
 		{name: "zero height", width: 6, renderer: renderer, wantErrSub: "greater than zero"},
-		{name: "unaligned width", width: 7, height: 1, renderer: renderer, wantErrSub: "divisible by 6"},
+		{name: "odd width", width: 7, height: 1, renderer: renderer, wantErrSub: "even"},
 		{name: "nil renderer", width: 6, height: 1, wantErrSub: "renderer is nil"},
 	}
 	for _, tc := range tests {
@@ -54,7 +54,7 @@ func TestCPUGeneratorCopiesBaseAndRestoresBeforePatch(t *testing.T) {
 		t.Fatalf("NewCPUGenerator: %v", err)
 	}
 
-	frame := make([]byte, 32)
+	frame := make([]byte, V210FrameSize(12, 1))
 	if err := g.GenerateFrame(frame, 1); err != nil {
 		t.Fatalf("GenerateFrame(1): %v", err)
 	}
@@ -113,10 +113,35 @@ func TestCPUGeneratorErrors(t *testing.T) {
 	if err != nil {
 		t.Fatalf("NewCPUGenerator: %v", err)
 	}
-	if err := g.GenerateFrame(make([]byte, 15), 0); err == nil || !strings.Contains(err.Error(), "too small") {
+	if err := g.GenerateFrame(make([]byte, V210FrameSize(6, 1)-1), 0); err == nil || !strings.Contains(err.Error(), "too small") {
 		t.Fatalf("small destination error = %v", err)
 	}
-	if err := g.GenerateFrame(make([]byte, 16), 4); !errors.Is(err, wantErr) {
+	if err := g.GenerateFrame(make([]byte, V210FrameSize(6, 1)), 4); !errors.Is(err, wantErr) {
 		t.Fatalf("patch error = %v, want wrapped %v", err, wantErr)
 	}
 }
+
+func TestCPUGeneratorUsesPaddedV210Rows(t *testing.T) {
+	const width, height = 100, 2
+	g, err := NewCPUPatternGenerator(width, height, "gray-ramp")
+	if err != nil {
+		t.Fatalf("NewCPUPatternGenerator: %v", err)
+	}
+	frame := make([]byte, V210FrameSize(width, height))
+	if err := g.GenerateFrame(frame, 0); err != nil {
+		t.Fatalf("GenerateFrame: %v", err)
+	}
+
+	for y := 0; y < height; y++ {
+		gotY, gotCb, gotCr := sampleV210(frame, width, 0, y)
+		if gotY != 64 || gotCb != 512 || gotCr != 512 {
+			t.Fatalf("row %d first pixel = %d/%d/%d, want 64/512/512", y, gotY, gotCb, gotCr)
+		}
+		padding := frame[y*V210LineSize(width)+V210ActiveLineSize(width) : (y+1)*V210LineSize(width)]
+		for i, b := range padding {
+			if b != 0 {
+				t.Fatalf("row %d padding byte %d = %#x, want 0", y, i, b)
+			}
+		}
+	}
+}

internal/generator/text.go

@@ -189,15 +189,16 @@ func (o *TextOverlay) pack() {
 // 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
+	frameStride := V210LineSize(o.frameW)
-	if need := (lastPixel/6 + 1) * 16; len(dest) < need {
+	tileStride := (o.w / 6) * 16
+	need := (o.y+o.h-1)*frameStride + o.x/6*16 + tileStride
+	if 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
+		frameOffset := (o.y+row)*frameStride + o.x/6*16
 		src := o.blocks[row*tileStride : (row+1)*tileStride]
-		copy(dest[frameBlock*16:frameBlock*16+tileStride], src)
+		copy(dest[frameOffset:frameOffset+tileStride], src)
 	}
 	return nil
 }

internal/generator/text_test.go

@@ -163,7 +163,7 @@ func TestTextOverlayApplyV210(t *testing.T) {
 		t.Fatalf("NewTextOverlay: %v", err)
 	}
 
-	frame := make([]byte, frameW*frameH*8/3)
+	frame := make([]byte, V210FrameSize(frameW, frameH))
 	for i := range frame {
 		frame[i] = 0x5A // marker: untouched regions must survive
 	}
@@ -192,7 +192,7 @@ func TestTextOverlayApplyV210(t *testing.T) {
 		{o.x + o.w, o.y + o.h}, // corner
 		{100, 1000},            // far away
 	} {
-		off := ((p[1]*frameW + p[0]) / 6) * 16
+		off := p[1]*V210LineSize(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/v210.go

@@ -0,0 +1,20 @@
+package generator
+
+const v210RowAlignment = 128
+
+// V210ActiveLineSize returns the number of bytes containing pixel data in one
+// v210 row, including the final partial six-pixel block when needed.
+func V210ActiveLineSize(width int) int {
+	return (width + 5) / 6 * 16
+}
+
+// V210LineSize returns the MXL v210 row stride. MXL stores every row at a
+// 128-byte boundary, equivalent to rounding the width up to 48 pixels.
+func V210LineSize(width int) int {
+	return ((width + 47) / 48) * v210RowAlignment
+}
+
+// V210FrameSize returns the complete MXL payload size for a v210 frame.
+func V210FrameSize(width, height int) int {
+	return V210LineSize(width) * height
+}

internal/generator/v210_test.go

@@ -1,16 +1,43 @@
 package generator
 
-import "encoding/binary"
+import (
+	"encoding/binary"
+	"testing"
+)
+
+func TestV210Sizes(t *testing.T) {
+	tests := []struct {
+		width      int
+		activeLine int
+		line       int
+	}{
+		{width: 1920, activeLine: 5120, line: 5120},
+		{width: 1280, activeLine: 3424, line: 3456},
+		{width: 100, activeLine: 272, line: 384},
+		{width: 54, activeLine: 144, line: 256},
+		{width: 48, activeLine: 128, line: 128},
+	}
+	for _, tc := range tests {
+		if got := V210ActiveLineSize(tc.width); got != tc.activeLine {
+			t.Errorf("V210ActiveLineSize(%d) = %d, want %d", tc.width, got, tc.activeLine)
+		}
+		if got := V210LineSize(tc.width); got != tc.line {
+			t.Errorf("V210LineSize(%d) = %d, want %d", tc.width, got, tc.line)
+		}
+		if got := V210FrameSize(tc.width, 2); got != tc.line*2 {
+			t.Errorf("V210FrameSize(%d, 2) = %d, want %d", tc.width, got, tc.line*2)
+		}
+	}
+}
 
 func sampleV210(buf []byte, width, x, y int) (yc, cb, cr uint32) {
-	pixel := y*width + x
+	offset := y*V210LineSize(width) + x/6*16
-	offset := (pixel / 6) * 16
 	w0 := binary.LittleEndian.Uint32(buf[offset:])
 	w1 := binary.LittleEndian.Uint32(buf[offset+4:])
 	w2 := binary.LittleEndian.Uint32(buf[offset+8:])
 	w3 := binary.LittleEndian.Uint32(buf[offset+12:])
 
-	switch pixel % 6 {
+	switch x % 6 {
 	case 0:
 		return (w0 >> 10) & 0x3FF, w0 & 0x3FF, (w0 >> 20) & 0x3FF
 	case 1:

internal/generator/wgpu.go

@@ -22,23 +22,26 @@ const wgpuWorkgroupSize = 64
 // buffer, the GPU DMAs it into a persistent host-visible buffer, and the
 // mapped contents are copied straight into the destination grain.
 type WGPUGenerator struct {
-	instance  *wgpu.Instance
+	instance         *wgpu.Instance
-	adapter   *wgpu.Adapter
+	adapter          *wgpu.Adapter
-	device    *wgpu.Device
+	device           *wgpu.Device
-	queue     *wgpu.Queue
+	queue            *wgpu.Queue
-	shader    *wgpu.ShaderModule
+	shader           *wgpu.ShaderModule
-	bgl       *wgpu.BindGroupLayout
+	bgl              *wgpu.BindGroupLayout
-	bg        *wgpu.BindGroup
+	bg               *wgpu.BindGroup
-	pl        *wgpu.PipelineLayout
+	pl               *wgpu.PipelineLayout
-	pipeline  *wgpu.ComputePipeline
+	pipeline         *wgpu.ComputePipeline
-	out       *wgpu.Buffer
+	out              *wgpu.Buffer
-	host      *wgpu.Buffer
+	host             *wgpu.Buffer
-	uniform   *wgpu.Buffer
+	uniform          *wgpu.Buffer
-	params    []byte
+	params           []byte
-	width     int
+	width            int
-	height    int
+	height           int
-	blocks    int
+	blocks           int
-	frameSize uint64
+	activeLineSize   int
+	lineSize         int
+	compactFrameSize uint64
+	frameSize        uint64
 }
 
 var _ FrameGenerator = (*WGPUGenerator)(nil)
@@ -48,15 +51,18 @@ type WGPUOption func(*WGPUGenerator)
 
 func NewWGPUGenerator(width, height uint, wgsl string, opts ...WGPUOption) (*WGPUGenerator, error) {
 	g := &WGPUGenerator{
-		width:  int(width),
+		width:          int(width),
-		height: int(height),
+		height:         int(height),
-		blocks: int(width*height) / 6,
+		blocks:         ((int(width) + 5) / 6) * int(height),
-		params: make([]byte, 16),
+		activeLineSize: V210ActiveLineSize(int(width)),
+		lineSize:       V210LineSize(int(width)),
+		params:         make([]byte, 16),
 	}
 	for _, opt := range opts {
 		opt(g)
 	}
-	g.frameSize = uint64(g.blocks) * 16
+	g.compactFrameSize = uint64(g.activeLineSize * g.height)
+	g.frameSize = uint64(g.lineSize * g.height)
 	binary.LittleEndian.PutUint32(g.params[0:], uint32(width))
 	binary.LittleEndian.PutUint32(g.params[4:], uint32(height))
 
@@ -81,14 +87,14 @@ func NewWGPUGenerator(width, height uint, wgsl string, opts ...WGPUOption) (*WGP
 		return nil, fmt.Errorf("wgpu: shader: %w", err)
 	}
 	if g.out, err = g.device.CreateBuffer(&wgpu.BufferDescriptor{
-		Label: "v210-out", Size: g.frameSize,
+		Label: "v210-out", Size: g.compactFrameSize,
 		Usage: wgpu.BufferUsageStorage | wgpu.BufferUsageCopySrc,
 	}); err != nil {
 		g.Close()
 		return nil, fmt.Errorf("wgpu: out buffer: %w", err)
 	}
 	if g.host, err = g.device.CreateBuffer(&wgpu.BufferDescriptor{
-		Label: "v210-host", Size: g.frameSize,
+		Label: "v210-host", Size: g.compactFrameSize,
 		Usage: wgpu.BufferUsageCopyDst | wgpu.BufferUsageMapRead,
 	}); err != nil {
 		g.Close()
@@ -118,7 +124,7 @@ func NewWGPUGenerator(width, height uint, wgsl string, opts ...WGPUOption) (*WGP
 	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: 0, Buffer: g.out, Size: g.compactFrameSize},
 			{Binding: 1, Buffer: g.uniform, Size: uint64(len(g.params))},
 		},
 	}); err != nil {
@@ -164,7 +170,7 @@ func (g *WGPUGenerator) GenerateFrame(dest []byte, frameIndex int) error {
 	if err := pass.End(); err != nil {
 		return fmt.Errorf("wgpu: end pass: %w", err)
 	}
-	encoder.CopyBufferToBuffer(g.out, 0, g.host, 0, g.frameSize)
+	encoder.CopyBufferToBuffer(g.out, 0, g.host, 0, g.compactFrameSize)
 	cmd, err := encoder.Finish()
 	if err != nil {
 		return fmt.Errorf("wgpu: finish: %w", err)
@@ -175,15 +181,25 @@ func (g *WGPUGenerator) GenerateFrame(dest []byte, frameIndex int) error {
 
 	ctx, cancel := context.WithTimeout(context.Background(), 5*time.Second)
 	defer cancel()
-	if err := g.host.Map(ctx, wgpu.MapModeRead, 0, g.frameSize); err != nil {
+	if err := g.host.Map(ctx, wgpu.MapModeRead, 0, g.compactFrameSize); err != nil {
 		return fmt.Errorf("wgpu: map: %w", err)
 	}
-	rng, err := g.host.MappedRange(0, g.frameSize)
+	rng, err := g.host.MappedRange(0, g.compactFrameSize)
 	if err != nil {
 		_ = g.host.Unmap()
 		return fmt.Errorf("wgpu: mapped range: %w", err)
 	}
-	copy(dest, rng.Bytes())
+	mapped := rng.Bytes()
+	if g.activeLineSize == g.lineSize {
+		copy(dest[:g.frameSize], mapped)
+	} else {
+		for y := 0; y < g.height; y++ {
+			src := mapped[y*g.activeLineSize : (y+1)*g.activeLineSize]
+			dst := dest[y*g.lineSize : (y+1)*g.lineSize]
+			copy(dst, src)
+			clear(dst[g.activeLineSize:])
+		}
+	}
 	return g.host.Unmap()
 }
 

internal/generator/wgpu_test.go

@@ -54,3 +54,37 @@ func TestWGPUSMPTEPattern(t *testing.T) {
 		})
 	}
 }
+
+func TestWGPUGeneratorUsesPaddedV210Rows(t *testing.T) {
+	const width, height = 100, 2
+	shader, err := kernels.Read("static/ebu75.wgsl")
+	if err != nil {
+		t.Fatalf("read shader: %v", err)
+	}
+	g, err := NewWGPUGenerator(width, height, string(shader))
+	if err != nil {
+		t.Fatalf("init: %v", err)
+	}
+	defer g.Close()
+
+	buf := make([]byte, V210FrameSize(width, height))
+	for i := range buf {
+		buf[i] = 0xff
+	}
+	if err := g.GenerateFrame(buf, 0); err != nil {
+		t.Fatalf("GenerateFrame: %v", err)
+	}
+
+	for y := 0; y < height; y++ {
+		gotY, gotCb, gotCr := sampleV210(buf, width, 0, y)
+		if gotY != 721 || gotCb != 512 || gotCr != 512 {
+			t.Fatalf("row %d first pixel = %d/%d/%d, want 721/512/512", y, gotY, gotCb, gotCr)
+		}
+		padding := buf[y*V210LineSize(width)+V210ActiveLineSize(width) : (y+1)*V210LineSize(width)]
+		for i, b := range padding {
+			if b != 0 {
+				t.Fatalf("row %d padding byte %d = %#x, want 0", y, i, b)
+			}
+		}
+	}
+}

internal/video/runner.go

@@ -37,7 +37,7 @@ func Run(ctx context.Context, inst *mxl.Instance, cfg Config) (runErr error) {
 
 	var staticFrame []byte
 	if !pattern.dynamic {
-		frameSize := int(cfg.Width()*cfg.Height()) * 8 / 3
+		frameSize := generator.V210FrameSize(int(cfg.Width()), int(cfg.Height()))
 		staticFrame = make([]byte, frameSize)
 
 		if err := gen.GenerateFrame(staticFrame, 0); err != nil {

kernels/dynamic/ebu100.wgsl

@@ -18,13 +18,14 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let py = (block * 6u) / params.width;
+    let py = block / blocks_per_row;
 
     // Bar order: 100% white, yellow, cyan, green, magenta, red, blue, black
     let y_tab  = array<u32, 8>(940u, 877u, 754u, 691u, 313u, 250u, 127u, 64u);

kernels/dynamic/ebu75.wgsl

@@ -18,13 +18,14 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let py = (block * 6u) / params.width;
+    let py = block / blocks_per_row;
 
     // Bar order: 75% white, yellow, cyan, green, magenta, red, blue, black
     let y_tab  = array<u32, 8>(721u, 674u, 581u, 534u, 251u, 204u, 111u, 64u);

kernels/dynamic/smpteBars.wgsl

@@ -18,13 +18,14 @@ fn bar_index(px: u32, bars_75_width: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let y_px = (block * 6u) / params.width;
+    let y_px = block / blocks_per_row;
 
     // 75% Color Table: white, yellow, cyan, green, magenta, red, blue
     let y_tab  = array<u32, 7>(721u, 674u, 581u, 534u, 251u, 204u, 111u);

kernels/dynamic/yBars.wgsl

@@ -19,15 +19,16 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
     let y_tab  = array<u32, 13>(64u, 137u, 210u, 283u, 356u, 429u, 502u, 575u, 648u, 721u, 794u, 867u, 940u);
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let py = (block * 6u) / params.width;
+    let py = block / blocks_per_row;
     // 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.

kernels/dynamic/yRamp.wgsl

@@ -18,13 +18,14 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let py = (block * 6u) / params.width;
+    let py = block / blocks_per_row;
 
     // Bar order: 75% white, yellow, cyan, green, magenta, red, blue
     let y_tab  = array<u32, 7>(721u, 674u, 581u, 534u, 251u, 204u, 111u);

kernels/static/ebu100.wgsl

@@ -18,12 +18,13 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
 
     // Bar order: 100% white, yellow, cyan, green, magenta, red, blue, black
     let y_tab  = array<u32, 8>(940u, 877u, 754u, 691u, 313u, 250u, 127u, 64u);

kernels/static/ebu75.wgsl

@@ -18,12 +18,13 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
 
     // Bar order: 75% white, yellow, cyan, green, magenta, red, blue, black
     let y_tab  = array<u32, 8>(721u, 674u, 581u, 534u, 251u, 204u, 111u, 64u);

kernels/static/smpteBars.wgsl

@@ -18,13 +18,14 @@ fn bar_index(px: u32, bars_75_width: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
-    let y_px = (block * 6u) / params.width;
+    let y_px = block / blocks_per_row;
 
     // 75% Color Table: white, yellow, cyan, green, magenta, red, blue
     let y_tab  = array<u32, 7>(721u, 674u, 581u, 534u, 251u, 204u, 111u);

kernels/static/yBars.wgsl

@@ -19,7 +19,8 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
@@ -27,7 +28,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let y_tab  = array<u32, 13>(64u, 137u, 210u, 283u, 356u, 429u, 502u, 575u, 648u, 721u, 794u, 867u, 940u);
 
     // 6 = pixels per v210 block (NOT the bar count)
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
     var y: array<u32, 6>;
     var cb: array<u32, 6>;
     var cr: array<u32, 6>;

kernels/static/yRamp.wgsl

@@ -18,7 +18,8 @@ fn bar_index(px: u32) -> u32 {
 @compute @workgroup_size(64)
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     let block = gid.x;
-    let total = (params.width * params.height) / 6u;
+    let blocks_per_row = (params.width + 5u) / 6u;
+    let total = blocks_per_row * params.height;
     if (block >= total) {
         return;
     }
@@ -28,7 +29,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
     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);
 
-    let x = (block * 6u) % params.width;
+    let x = (block % blocks_per_row) * 6u;
     var y: array<u32, 6>;
     var cb: array<u32, 6>;
     var cr: array<u32, 6>;