From 777ad490882029d0afb1cd0b9c3b6dd4c28daabd Mon Sep 17 00:00:00 2001 From: Dmitry Sergeev Date: Sun, 13 Sep 2026 16:36:42 +0300 Subject: [PATCH] Remove excessive --- cmd/pattern-gen/main.go | 159 ---------------------- cmd/wgpu-sample/main.go | 282 ---------------------------------------- kernels/v210_pack.cl | 72 ---------- scripts/gen_font.py | 53 -------- 4 files changed, 566 deletions(-) delete mode 100644 cmd/pattern-gen/main.go delete mode 100644 cmd/wgpu-sample/main.go delete mode 100644 kernels/v210_pack.cl delete mode 100644 scripts/gen_font.py diff --git a/cmd/pattern-gen/main.go b/cmd/pattern-gen/main.go deleted file mode 100644 index 212351d..0000000 --- a/cmd/pattern-gen/main.go +++ /dev/null @@ -1,159 +0,0 @@ -package main - -import ( - "encoding/binary" - "log" - flowdef "mxl-pattern-generator/internal/flow-def" - "mxl-pattern-generator/internal/generator" - "os" - "os/signal" - "syscall" - - "github.com/qvest-digital/go-mxl/mxl" -) - -func main() { - var width, height uint = 1920, 1080 - var fps_num, fps_den uint = 25, 1 - flowUUID := "5fbec3b1-1b0f-417d-9059-8b94a47197ed" - log.Printf("Go Pattern Gen") - log.Printf("Video: %dx%dp%d/%d", width, height, fps_num, fps_den) - log.Printf("UUID: %s", flowUUID) - domain := "/dev/shm/mxl" - - inst, err := mxl.NewInstance(domain, "") - if err != nil { - log.Fatalf("MXL Init Failed: %v", err) - } - defer inst.Close() - - flowDef, err := flowdef.NewFlowDefJSON( - flowdef.TYPE_VIDEO, - flowUUID, - width, - height, - fps_num, - fps_den, - ) - if err != nil { - log.Fatalf("Could not create Flow Definition: %v", err) - } - writer, isCreated, err := inst.NewWriter(flowDef) - if err != nil { - log.Fatalf("Failed to create MXL writer: %v", err) - } - if !isCreated { - log.Printf("reusing existing flow: %s, domain: %s", flowUUID, domain) - } - defer writer.Close() - - var gen *generator.OpenCLGenerator - gen, err = generator.NewOpenCLGenerator(width, height, "/home/itten/codeproj/go-mxl-pattern-gen/kernels/v210_pack.cl") - gen, err = generator.NewOpenCLGenerator(width, height, "/home/itten/codeproj/go-mxl-pattern-gen/kernels/v210_pack.cl") - if err != nil { - log.Fatalf("OpenCL init failed: %v", err) - } - defer gen.Close() - gen.SetText(generator.TextConfig{ - Str: "MXL PATTERN GEN", - X: 0, Y: 0, Scale: 8, - FG: [3]uint32{940, 512, 512}, - BG: [3]uint32{64, 512, 512}, - Boxed: true, - }) - - flowCfg := writer.Config() - rate := flowCfg.Common.GrainRate - idx := mxl.CurrentIndex(rate) - log.Printf("writing flow grainRate=%d/%d starting at idx=%d", rate.Num, rate.Den, idx) - - stop := make(chan os.Signal, 1) - signal.Notify(stop, os.Interrupt, syscall.SIGTERM) - - var written int64 - for { - select { - case <-stop: - log.Printf("stopping after %d grains", written) - return - default: - } - - gwa, err := writer.OpenGrain(idx) - if err != nil { - log.Fatalf("OpenGrain(%d): %v", idx, err) - } - // fillGrainRamp(gwa.Payload, idx) - // smpteBars(gwa.Payload) - gen.GenerateFrame(gwa.Payload, int(idx)) - // if err != nil { - // log.Fatalf("Frame Generator: %v", err) - // } - if err := gwa.Commit(gwa.TotalSlices, 0); err != nil { - log.Fatalf("Commit(%d): %v", idx, err) - } - - written++ - idx++ - if written%100 == 0 { - log.Printf("grains written=%d, index=%d", written, idx) - } - // Pace ourselves to roughly the grain rate - mxl.SleepNs(mxl.NsUntilIndex(idx, rate)) - } -} - -// fillGrainRamp writes byte((i+start)&0xFF) across buf. A matching reader -// uses the same pattern to verify grain continuit -func fillGrainRamp(buf []byte, start uint64) { - for i := range buf { - buf[i] = byte((uint64(i) + start) & 0xFF) - } -} - -func smpteBars(dest []byte) { - const width = 1920 - - // SMPTE 75% bars, Rec.709, 10-bit Y'CbCr: {Y, Cb, Cr} - bars := [7][3]uint32{ - {721, 512, 512}, // 75% white - {674, 176, 543}, // 75% yellow - {581, 589, 176}, // 75% cyan - {534, 253, 207}, // 75% green - {251, 771, 817}, // 75% magenta - {204, 435, 848}, // 75% red - {111, 848, 481}, // 75% blue - } - - barIdx := func(x int) int { - idx := x * len(bars) / width - if idx >= len(bars) { - return len(bars) - 1 - } - return idx - } - - // v210: each 16-byte block carries 6 pixels (Y0..Y5, chroma co-sited at 0/2/4) - for block := 0; block*16+16 <= len(dest); block++ { - baseX := (block * 6) % width - - var y [6]uint32 - for i := range y { - y[i] = bars[barIdx(baseX+i)][0] - } - cb0, cr0 := bars[barIdx(baseX)][1], bars[barIdx(baseX)][2] - cb2, cr2 := bars[barIdx(baseX+2)][1], bars[barIdx(baseX+2)][2] - cb4, cr4 := bars[barIdx(baseX+4)][1], bars[barIdx(baseX+4)][2] - - w0 := (cb0 & 0x3FF) | ((y[0] & 0x3FF) << 10) | ((cr0 & 0x3FF) << 20) - w1 := (y[1] & 0x3FF) | ((cb2 & 0x3FF) << 10) | ((y[2] & 0x3FF) << 20) - w2 := (cr2 & 0x3FF) | ((y[3] & 0x3FF) << 10) | ((cb4 & 0x3FF) << 20) - w3 := (y[4] & 0x3FF) | ((cr4 & 0x3FF) << 10) | ((y[5] & 0x3FF) << 20) - - byteIdx := block * 16 - binary.LittleEndian.PutUint32(dest[byteIdx+0:], w0) - binary.LittleEndian.PutUint32(dest[byteIdx+4:], w1) - binary.LittleEndian.PutUint32(dest[byteIdx+8:], w2) - binary.LittleEndian.PutUint32(dest[byteIdx+12:], w3) - } -} diff --git a/cmd/wgpu-sample/main.go b/cmd/wgpu-sample/main.go deleted file mode 100644 index 785a8c8..0000000 --- a/cmd/wgpu-sample/main.go +++ /dev/null @@ -1,282 +0,0 @@ -// Command wgpu-sample renders 75% SMPTE color bars (Rec.709, 10-bit v210) -// with a wgpu compute shader, verifies the output against the reference -// table, and reports per-frame timings. -// -// Run from the repo root: go run ./cmd/wgpu-sample -package main - -import ( - "context" - "encoding/binary" - "fmt" - "log" - "os" - "time" - - "github.com/gogpu/gputypes" - "github.com/gogpu/wgpu" - - // Register all available GPU backends (Vulkan, DX12, GLES, Metal, etc.) - _ "github.com/gogpu/wgpu/hal/allbackends" -) - -const ( - width, height = 1920, 1080 - wgSize = 64 - blocks = width * height / 6 - frameSize = uint64(blocks * 16) - benchFrames = 100 -) - -// 75% SMPTE bars, Rec.709, 10-bit {Y, Cb, Cr} -var bars = [7][3]uint32{ - {721, 512, 512}, {674, 176, 543}, {581, 589, 176}, - {534, 253, 207}, {251, 771, 817}, {204, 435, 848}, {111, 848, 481}, -} - -func main() { - if err := run(); err != nil { - log.Fatalf("FATAL: %v", err) - } -} - -func run() error { - instance, err := wgpu.CreateInstance(nil) - if err != nil { - return fmt.Errorf("CreateInstance: %w", err) - } - defer instance.Release() - - adapter, err := instance.RequestAdapter(nil) - if err != nil { - return fmt.Errorf("RequestAdapter: %w", err) - } - defer adapter.Release() - - info := adapter.Info() - fmt.Printf("adapter: %s (%s, %s, driver %s)\n", info.Name, info.Vendor, info.DeviceType, info.Driver) - - device, err := adapter.RequestDevice(nil) - if err != nil { - return fmt.Errorf("RequestDevice: %w", err) - } - defer device.Release() - queue := device.Queue() - - wgsl, err := os.ReadFile("kernels/v210_bars.wgsl") - if err != nil { - return fmt.Errorf("read kernel: %w", err) - } - - out, err := device.CreateBuffer(&wgpu.BufferDescriptor{ - Label: "v210-out", Size: frameSize, - Usage: wgpu.BufferUsageStorage | wgpu.BufferUsageCopySrc, - }) - if err != nil { - return err - } - defer out.Release() - - staging, err := device.CreateBuffer(&wgpu.BufferDescriptor{ - Label: "v210-staging", Size: frameSize, - Usage: wgpu.BufferUsageCopyDst | wgpu.BufferUsageMapRead, - }) - if err != nil { - return err - } - defer staging.Release() - - params := make([]byte, 16) - binary.LittleEndian.PutUint32(params[0:], width) - binary.LittleEndian.PutUint32(params[4:], height) - uniform, err := device.CreateBuffer(&wgpu.BufferDescriptor{ - Label: "params", Size: uint64(len(params)), - Usage: wgpu.BufferUsageUniform | wgpu.BufferUsageCopyDst, - }) - if err != nil { - return err - } - defer uniform.Release() - if err := queue.WriteBuffer(uniform, 0, params); err != nil { - return fmt.Errorf("write params: %w", err) - } - - shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{ - Label: "bars-shader", WGSL: string(wgsl), - }) - if err != nil { - return fmt.Errorf("shader: %w", err) - } - defer shader.Release() - - bgl, err := device.CreateBindGroupLayout(&wgpu.BindGroupLayoutDescriptor{ - Label: "bars-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}}, - }, - }) - if err != nil { - return err - } - defer bgl.Release() - - bg, err := device.CreateBindGroup(&wgpu.BindGroupDescriptor{ - Label: "bars-bg", Layout: bgl, - Entries: []wgpu.BindGroupEntry{ - {Binding: 0, Buffer: out, Size: frameSize}, - {Binding: 1, Buffer: uniform, Size: uint64(len(params))}, - }, - }) - if err != nil { - return err - } - defer bg.Release() - - pl, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{ - Label: "bars-pl", BindGroupLayouts: []*wgpu.BindGroupLayout{bgl}, - }) - if err != nil { - return err - } - defer pl.Release() - - pipeline, err := device.CreateComputePipeline(&wgpu.ComputePipelineDescriptor{ - Label: "bars-pipeline", Layout: pl, Module: shader, EntryPoint: "main", - }) - if err != nil { - return fmt.Errorf("pipeline: %w", err) - } - defer pipeline.Release() - - render := func(frame uint32) error { - binary.LittleEndian.PutUint32(params[8:], frame) - if err := queue.WriteBuffer(uniform, 0, params); err != nil { - return err - } - encoder, err := device.CreateCommandEncoder(nil) - if err != nil { - return err - } - pass, err := encoder.BeginComputePass(nil) - if err != nil { - return err - } - pass.SetPipeline(pipeline) - pass.SetBindGroup(0, bg, nil) - pass.Dispatch((blocks+wgSize-1)/wgSize, 1, 1) - if err := pass.End(); err != nil { - return err - } - encoder.CopyBufferToBuffer(out, 0, staging, 0, frameSize) - cmd, err := encoder.Finish() - if err != nil { - return err - } - if _, err := queue.Submit(cmd); err != nil { - return err - } - return nil - } - - readback := func() ([]byte, error) { - ctx, cancel := context.WithTimeout(context.Background(), 5*time.Second) - defer cancel() - if err := staging.Map(ctx, wgpu.MapModeRead, 0, frameSize); err != nil { - return nil, fmt.Errorf("map: %w", err) - } - rng, err := staging.MappedRange(0, frameSize) - if err != nil { - _ = staging.Unmap() - return nil, fmt.Errorf("mapped range: %w", err) - } - data := make([]byte, frameSize) - copy(data, rng.Bytes()) - if err := staging.Unmap(); err != nil { - return nil, err - } - return data, nil - } - - // Correctness: render one frame and verify every pixel. - fmt.Printf("frame: %dx%d v210, %d bytes\n", width, height, frameSize) - if err := render(0); err != nil { - return fmt.Errorf("render: %w", err) - } - data, err := readback() - if err != nil { - return fmt.Errorf("readback: %w", err) - } - if err := verify(data); err != nil { - return err - } - fmt.Println("verify: PASS (all 1920x1080 pixels match the 75% Rec.709 table)") - - // Rough timing: full dispatch + copy + map/unmap per frame. - for i := 0; i < 10; i++ { - if err := render(uint32(i)); err != nil { - return err - } - if _, err := readback(); err != nil { - return err - } - } - start := time.Now() - for i := 0; i < benchFrames; i++ { - if err := render(uint32(i)); err != nil { - return err - } - if _, err := readback(); err != nil { - return err - } - } - perFrame := time.Since(start) / benchFrames - fmt.Printf("timing: %v/frame (dispatch+copy+readback), %.1f fps equivalent\n", - perFrame, float64(time.Second)/float64(perFrame)) - - return nil -} - -// verify decodes the v210 frame and checks Y at every pixel and the -// co-sited chroma (even columns) against the reference table. -func verify(data []byte) error { - barOf := func(x int) int { - if b := x * 7 / width; b < 7 { - return b - } - return 6 - } - for y := 0; y < height; y++ { - for x := 0; x < width; x++ { - p := y*width + x - off := (p / 6) * 16 - w0 := binary.LittleEndian.Uint32(data[off:]) - w1 := binary.LittleEndian.Uint32(data[off+4:]) - w2 := binary.LittleEndian.Uint32(data[off+8:]) - w3 := binary.LittleEndian.Uint32(data[off+12:]) - var yv, cb, cr uint32 - switch p % 6 { - case 0: - yv, cb, cr = (w0>>10)&0x3FF, w0&0x3FF, (w0>>20)&0x3FF - case 1: - yv = w1 & 0x3FF - case 2: - yv, cb, cr = (w1>>20)&0x3FF, (w1>>10)&0x3FF, w2&0x3FF - case 3: - yv = (w2 >> 10) & 0x3FF - case 4: - yv, cb, cr = w3&0x3FF, (w2>>20)&0x3FF, (w3>>10)&0x3FF - case 5: - yv = (w3 >> 20) & 0x3FF - } - b := bars[barOf(x)] - if yv != b[0] { - return fmt.Errorf("pixel (%d,%d): Y=%d want %d", x, y, yv, b[0]) - } - if x%2 == 0 && (cb != b[1] || cr != b[2]) { - return fmt.Errorf("pixel (%d,%d): Cb=%d Cr=%d want %d/%d", x, y, cb, cr, b[1], b[2]) - } - } - } - return nil -} diff --git a/kernels/v210_pack.cl b/kernels/v210_pack.cl deleted file mode 100644 index 938f2ef..0000000 --- a/kernels/v210_pack.cl +++ /dev/null @@ -1,72 +0,0 @@ -// 75% SMPTE color bars, Rec.709, 10-bit Y'CbCr, packed as v210 (4:2:2), -// with an 8x16 bitmap-font text overlay (burn-in). -// Each work-item packs 6 pixels (6 Y + 3 Cb + 3 Cr) into one 16-byte block. - -#define GW 8 -#define GH 16 - -// Text overlay lookup. Atlas: 96 glyphs (ASCII 32..127) x GH rows, -// one byte per row, bit n = pixel at column n (LSB = leftmost). -// Returns: 1 = glyph pixel (foreground), 0 = inside text box (background), -// -1 = outside the text box (pattern shows through). -int text_pixel(__global const uchar* atlas, __global const uchar* text, int text_len, - int px, int py, int txt_x, int txt_y, int scale) { - if (text_len <= 0) return -1; - int tx = px - txt_x; - int ty = py - txt_y; - if (tx < 0 || ty < 0) return -1; - if (tx >= text_len * GW * scale || ty >= GH * scale) return -1; - int c = text[tx / (GW * scale)]; - if (c < 32 || c > 127) return 0; - uchar row = atlas[(c - 32) * GH + ty / scale]; - return (row >> ((tx / scale) % GW)) & 1; -} - -__kernel void generate_v210_pattern( - __global uint* output, int width, int height, int frame_index, - __global const uchar* atlas, __global const uchar* text, int text_len, - int txt_x, int txt_y, int scale, int has_bg, - uint fg_y, uint fg_cb, uint fg_cr, - uint bg_y, uint bg_cb, uint bg_cr) { - int g_id = get_global_id(0); - int total_blocks = (width * height) / 6; - if (g_id >= total_blocks) return; - - int x = (g_id * 6) % width; - int py = (g_id * 6) / width; - - // Bar order: 75% white, yellow, cyan, green, magenta, red, blue - ushort Y_table[7] = {721, 674, 581, 534, 251, 204, 111}; - ushort U_table[7] = {512, 176, 589, 253, 771, 435, 848}; - ushort V_table[7] = {512, 543, 176, 207, 817, 848, 481}; - - #define BAR(px) min(((px) * 7) / width, 6) - - uint y[6], cb[6], cr[6]; - for (int i = 0; i < 6; i++) { - int px = x + i; - int b = BAR(px); - uint yv = Y_table[b], cbv = U_table[b], crv = V_table[b]; - int st = text_pixel(atlas, text, text_len, px, py, txt_x, txt_y, scale); - if (st > 0) { - yv = fg_y; cbv = fg_cb; crv = fg_cr; - } else if (st == 0 && has_bg) { - yv = bg_y; cbv = bg_cb; crv = bg_cr; - } - y[i] = yv; cb[i] = cbv; cr[i] = crv; - } - - // 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 - uint word0 = (cb[0] & 0x3FF) | ((y[0] & 0x3FF) << 10) | ((cr[0] & 0x3FF) << 20); - uint word1 = (y[1] & 0x3FF) | ((cb[2] & 0x3FF) << 10) | ((y[2] & 0x3FF) << 20); - uint word2 = (cr[2] & 0x3FF) | ((y[3] & 0x3FF) << 10) | ((cb[4] & 0x3FF) << 20); - uint word3 = (y[4] & 0x3FF) | ((cr[4] & 0x3FF) << 10) | ((y[5] & 0x3FF) << 20); - - int out_idx = g_id * 4; - output[out_idx + 0] = word0; - output[out_idx + 1] = word1; - output[out_idx + 2] = word2; - output[out_idx + 3] = word3; -} diff --git a/scripts/gen_font.py b/scripts/gen_font.py deleted file mode 100644 index bcacd7c..0000000 --- a/scripts/gen_font.py +++ /dev/null @@ -1,53 +0,0 @@ -#!/usr/bin/env python3 -"""Bake an 8x16 1-bpp bitmap font (ASCII 32..127) from DejaVu Sans Mono 13px. - -Output: Go source file with the atlas. One byte per glyph row, bit n = pixel -at column n (LSB = leftmost), matching the kernel's (row >> col) & 1 lookup. -""" -from PIL import Image, ImageDraw, ImageFont - -FONT = "/usr/share/fonts/truetype/dejavu/DejaVuSansMono.ttf" -SIZE = 13 -BASELINE = 12 # baseline row inside the 16px cell -THRESH = 110 - -font = ImageFont.truetype(FONT, SIZE) -atlas = [] -for c in range(32, 128): - img = Image.new("L", (8, 16), 0) - d = ImageDraw.Draw(img) - d.text((0, BASELINE), chr(c), font=font, fill=255, anchor="ls") - px = img.load() - for y in range(16): - b = 0 - for x in range(8): - if px[x, y] >= THRESH: - b |= 1 << x - atlas.append(b) - -def art(c): - print(f"--- '{chr(c)}' ---") - for y in range(16): - row = atlas[(c - 32) * 16 + y] - print("".join("#" if (row >> x) & 1 else "." for x in range(8))) - -for c in "AgM%m0(": - art(ord(c)) - -lines = [] -for i in range(0, len(atlas), 16): - chunk = ", ".join(f"0x{b:02X}" for b in atlas[i:i+16]) - lines.append(f"\t{chunk},") -body = "\n".join(lines) - -src = f'''package generator - -// 8x16 1-bpp bitmap font baked from DejaVu Sans Mono {SIZE}px (public-domain -// style console glyphs). ASCII 32..127, one byte per row, LSB = leftmost pixel. -var fontAtlas8x16 = [96 * 16]byte{{ -{body} -}} -''' -with open("internal/generator/font.go", "w") as f: - f.write(src) -print("wrote internal/generator/font.go,", len(atlas), "bytes of glyph data")