diff --git a/cmd/mxl-pattern/main.go b/cmd/mxl-pattern/main.go index 0f9d8cb..e2785d1 100644 --- a/cmd/mxl-pattern/main.go +++ b/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 { - // v210 stores 6 pixels in each 16-byte block. - return fmt.Errorf("video width must be greater than zero and divisible by 6") + if args.videoWidth == 0 || args.videoWidth%2 != 0 { + return fmt.Errorf("video width must be greater than zero and even for 4:2:2 video") } if args.videoHeight == 0 { return fmt.Errorf("video height must be greater than zero") diff --git a/internal/flowdef/flowdef.go b/internal/flowdef/flowdef.go index c916250..0e0b108 100644 --- a/internal/flowdef/flowdef.go +++ b/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 { - return fmt.Errorf("frame_width must be greater than zero and divisible by 6, got %d", v.FrameWidth) + if v.FrameWidth == 0 || v.FrameWidth%2 != 0 { + 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") diff --git a/internal/flowdef/flowdef_test.go b/internal/flowdef/flowdef_test.go index d30bc54..dbf8fe7 100644 --- a/internal/flowdef/flowdef_test.go +++ b/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") { - t.Fatalf("error = %v, want width divisibility error", err) + if err == nil || !strings.Contains(err.Error(), "even") { + t.Fatalf("error = %v, want even-width error", err) } } diff --git a/internal/generator/cpu.go b/internal/generator/cpu.go index a85222b..0ea0580 100644 --- a/internal/generator/cpu.go +++ b/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 { - return nil, fmt.Errorf("cpu: width must be divisible by 6, got %d", width) + if width%2 != 0 { + 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 { diff --git a/internal/generator/cpu_moving_square.go b/internal/generator/cpu_moving_square.go index 84c39b1..2ab4c24 100644 --- a/internal/generator/cpu_moving_square.go +++ b/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) } } diff --git a/internal/generator/cpu_patterns.go b/internal/generator/cpu_patterns.go index 25a6ccd..7bba31e 100644 --- a/internal/generator/cpu_patterns.go +++ b/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) } } diff --git a/internal/generator/cpu_test.go b/internal/generator/cpu_test.go index 3f27962..f89723e 100644 --- a/internal/generator/cpu_test.go +++ b/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) + } + } + } +} diff --git a/internal/generator/text.go b/internal/generator/text.go index a9e2bf9..cadec6b 100644 --- a/internal/generator/text.go +++ 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 - if need := (lastPixel/6 + 1) * 16; len(dest) < need { + frameStride := V210LineSize(o.frameW) + 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 } diff --git a/internal/generator/text_test.go b/internal/generator/text_test.go index 45c06f9..6db73a1 100644 --- a/internal/generator/text_test.go +++ b/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]) diff --git a/internal/generator/v210.go b/internal/generator/v210.go new file mode 100644 index 0000000..059896d --- /dev/null +++ b/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 +} diff --git a/internal/generator/v210_test.go b/internal/generator/v210_test.go index c7fc830..09affc3 100644 --- a/internal/generator/v210_test.go +++ b/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 := (pixel / 6) * 16 + offset := y*V210LineSize(width) + x/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: diff --git a/internal/generator/wgpu.go b/internal/generator/wgpu.go index dc0a3eb..9793506 100644 --- a/internal/generator/wgpu.go +++ b/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 - adapter *wgpu.Adapter - device *wgpu.Device - queue *wgpu.Queue - shader *wgpu.ShaderModule - bgl *wgpu.BindGroupLayout - bg *wgpu.BindGroup - pl *wgpu.PipelineLayout - pipeline *wgpu.ComputePipeline - out *wgpu.Buffer - host *wgpu.Buffer - uniform *wgpu.Buffer - params []byte - width int - height int - blocks int - frameSize uint64 + instance *wgpu.Instance + adapter *wgpu.Adapter + device *wgpu.Device + queue *wgpu.Queue + shader *wgpu.ShaderModule + bgl *wgpu.BindGroupLayout + bg *wgpu.BindGroup + pl *wgpu.PipelineLayout + pipeline *wgpu.ComputePipeline + out *wgpu.Buffer + host *wgpu.Buffer + uniform *wgpu.Buffer + params []byte + width int + height int + blocks int + 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), - height: int(height), - blocks: int(width*height) / 6, - params: make([]byte, 16), + width: int(width), + height: int(height), + blocks: ((int(width) + 5) / 6) * int(height), + 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() } diff --git a/internal/generator/wgpu_test.go b/internal/generator/wgpu_test.go index cc7f27a..d85b9ee 100644 --- a/internal/generator/wgpu_test.go +++ b/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) + } + } + } +} diff --git a/internal/video/runner.go b/internal/video/runner.go index e5b77fa..aa3f5b5 100644 --- a/internal/video/runner.go +++ 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 { diff --git a/kernels/dynamic/ebu100.wgsl b/kernels/dynamic/ebu100.wgsl index 06d8160..3b668f7 100644 --- a/kernels/dynamic/ebu100.wgsl +++ b/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) { 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 py = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + let py = block / blocks_per_row; // Bar order: 100% white, yellow, cyan, green, magenta, red, blue, black let y_tab = array(940u, 877u, 754u, 691u, 313u, 250u, 127u, 64u); diff --git a/kernels/dynamic/ebu75.wgsl b/kernels/dynamic/ebu75.wgsl index d4526bd..65efebf 100644 --- a/kernels/dynamic/ebu75.wgsl +++ b/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) { 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 py = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + let py = block / blocks_per_row; // Bar order: 75% white, yellow, cyan, green, magenta, red, blue, black let y_tab = array(721u, 674u, 581u, 534u, 251u, 204u, 111u, 64u); diff --git a/kernels/dynamic/smpteBars.wgsl b/kernels/dynamic/smpteBars.wgsl index d7c47e2..323b2a4 100644 --- a/kernels/dynamic/smpteBars.wgsl +++ b/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) { 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 y_px = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + let y_px = block / blocks_per_row; // 75% Color Table: white, yellow, cyan, green, magenta, red, blue let y_tab = array(721u, 674u, 581u, 534u, 251u, 204u, 111u); diff --git a/kernels/dynamic/yBars.wgsl b/kernels/dynamic/yBars.wgsl index 0e6b9fd..e9a7dd1 100644 --- a/kernels/dynamic/yBars.wgsl +++ b/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) { 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(64u, 137u, 210u, 283u, 356u, 429u, 502u, 575u, 648u, 721u, 794u, 867u, 940u); - let x = (block * 6u) % params.width; - let py = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + 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. diff --git a/kernels/dynamic/yRamp.wgsl b/kernels/dynamic/yRamp.wgsl index 85a9a93..22ee795 100644 --- a/kernels/dynamic/yRamp.wgsl +++ b/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) { 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 py = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + let py = block / blocks_per_row; // Bar order: 75% white, yellow, cyan, green, magenta, red, blue let y_tab = array(721u, 674u, 581u, 534u, 251u, 204u, 111u); diff --git a/kernels/static/ebu100.wgsl b/kernels/static/ebu100.wgsl index e7dd9d7..ebff8ac 100644 --- a/kernels/static/ebu100.wgsl +++ b/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) { 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(940u, 877u, 754u, 691u, 313u, 250u, 127u, 64u); diff --git a/kernels/static/ebu75.wgsl b/kernels/static/ebu75.wgsl index 7102826..077227b 100644 --- a/kernels/static/ebu75.wgsl +++ b/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) { 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(721u, 674u, 581u, 534u, 251u, 204u, 111u, 64u); diff --git a/kernels/static/smpteBars.wgsl b/kernels/static/smpteBars.wgsl index d9742c1..a353903 100644 --- a/kernels/static/smpteBars.wgsl +++ b/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) { 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 y_px = (block * 6u) / params.width; + let x = (block % blocks_per_row) * 6u; + let y_px = block / blocks_per_row; // 75% Color Table: white, yellow, cyan, green, magenta, red, blue let y_tab = array(721u, 674u, 581u, 534u, 251u, 204u, 111u); diff --git a/kernels/static/yBars.wgsl b/kernels/static/yBars.wgsl index 89b610d..5830a44 100644 --- a/kernels/static/yBars.wgsl +++ b/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) { 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) { let y_tab = array(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; var cb: array; var cr: array; diff --git a/kernels/static/yRamp.wgsl b/kernels/static/yRamp.wgsl index 7ce195a..b413b5b 100644 --- a/kernels/static/yRamp.wgsl +++ b/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) { 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) { let cb_tab = array(512u, 176u, 589u, 253u, 771u, 435u, 848u); let cr_tab = array(512u, 543u, 176u, 207u, 817u, 848u, 481u); - let x = (block * 6u) % params.width; + let x = (block % blocks_per_row) * 6u; var y: array; var cb: array; var cr: array;