commit f48195a8bfce612fcacae245a5a6488509530549 Author: Dmitry Sergeev Date: Mon Sep 7 17:18:03 2026 +0300 WGPU is coming diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..0083467 --- /dev/null +++ b/Makefile @@ -0,0 +1,22 @@ +# Binaries that mix libmxl (cgo) and wgpu/goffi must use the internal +# linker: goffi's dlopen stubs (SDYNIMPORT relocations) do not survive +# external linking, which the go tool auto-selects for that call graph. +# Harmless for binaries that don't need it. +LD = -ldflags=-linkmode=internal + +.PHONY: build test vet wgpu-gen clean + +build: + go build $(LD) ./... + +test: + go test $(LD) ./... + +vet: + go vet ./... + +wgpu-gen: + go run $(LD) ./cmd/wgpu-gen + +clean: + go clean diff --git a/cmd/pattern-gen/main.go b/cmd/pattern-gen/main.go new file mode 100644 index 0000000..212351d --- /dev/null +++ b/cmd/pattern-gen/main.go @@ -0,0 +1,159 @@ +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-gen/main.go b/cmd/wgpu-gen/main.go new file mode 100644 index 0000000..c713aee --- /dev/null +++ b/cmd/wgpu-gen/main.go @@ -0,0 +1,99 @@ +// MXL pattern generator feed rendered on the GPU via wgpu (WebGPU/Vulkan). +// Run from the repo root: make wgpu-gen +// (mixing libmxl cgo with wgpu/goffi needs the internal linker) +package main + +import ( + "log" + "os" + "os/signal" + "syscall" + + "github.com/qvest-digital/go-mxl/mxl" + + flowdef "mxl-pattern-generator/internal/flow-def" + "mxl-pattern-generator/internal/generator" +) + +func main() { + var width, height uint = 1920, 1080 + var fpsNum, fpsDen uint = 25, 1 + flowUUID := "8f1d2a4b-6c3e-4f5a-9b2c-1d7e8a3f0b5d" + domain := "/dev/shm/mxl" + + log.Printf("Go Pattern Gen (wgpu)") + log.Printf("Video: %dx%dp%d/%d", width, height, fpsNum, fpsDen) + log.Printf("UUID: %s", flowUUID) + + // gen, err := generator.NewWGPUGenerator(width, height, "kernels/v210_bars.wgsl") + gen, err := generator.NewWGPUGenerator(width, height, "kernels/v210_bars_move.wgsl") + if err != nil { + log.Fatalf("wgpu init failed: %v", err) + } + defer gen.Close() + + 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, + fpsNum, + fpsDen, + ) + 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() + + 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 + var tick uint32 // animation clock: small counter, not the huge grain index + 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) + } + if err := gen.GenerateFrame(gwa.Payload, int(tick)); err != nil { + log.Fatalf("GenerateFrame(%d): %v", idx, err) + } + if err := gwa.Commit(gwa.TotalSlices, 0); err != nil { + log.Fatalf("Commit(%d): %v", idx, err) + } + + written++ + idx++ + tick++ + 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)) + } +} diff --git a/cmd/wgpu-sample/main.go b/cmd/wgpu-sample/main.go new file mode 100644 index 0000000..785a8c8 --- /dev/null +++ b/cmd/wgpu-sample/main.go @@ -0,0 +1,282 @@ +// 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/go.mod b/go.mod new file mode 100644 index 0000000..73293f8 --- /dev/null +++ b/go.mod @@ -0,0 +1,17 @@ +module mxl-pattern-generator + +go 1.26.0 + +require ( + github.com/gogpu/gputypes v0.8.0 + github.com/gogpu/wgpu v0.34.3 + github.com/qvest-digital/go-mxl v1.1.0-rc.4 +) + +require ( + github.com/go-webgpu/goffi v0.6.3 // indirect + 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 +) diff --git a/go.sum b/go.sum new file mode 100644 index 0000000..465c687 --- /dev/null +++ b/go.sum @@ -0,0 +1,16 @@ +github.com/go-webgpu/goffi v0.6.3 h1:p4gKGikHBAQ/8iUiew9MV4C5M1ZGIsk8QGFGuJjMj+A= +github.com/go-webgpu/goffi v0.6.3/go.mod h1:wfoxNsJkU+5RFbV1kNN1kunhc1lFHuJKK3zpgx08/uM= +github.com/go-webgpu/webgpu v0.5.5 h1:pIrXzRg0LRlNjNmR+ZNo/ERN88YaObzbCY5oYhir5SA= +github.com/go-webgpu/webgpu v0.5.5/go.mod h1:vgIuNTa1UlZ4njCGY6Pmp/0c5T5o2Ml2fQ0znLc7B98= +github.com/gogpu/gpucontext v0.31.3 h1:SMnYOQgr+s5issi1R/Oezbo8lV5Lb6HWUZ4HujFgrZA= +github.com/gogpu/gpucontext v0.31.3/go.mod h1:pnyWQA9lpvESPAKoqzKlsg3pDMFSkxdjonu4gWvnNtA= +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/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= +golang.org/x/sys v0.47.0 h1:o7XGOvZQCADBQQ4Y7VNq2dRWQR7JmOUW8Kxx4ZsNgWs= +golang.org/x/sys v0.47.0/go.mod h1:4GL1E5IUh+htKOUEOaiffhrAeqysfVGipDYzABqnCmw= diff --git a/internal/flow-def/flow-def.go b/internal/flow-def/flow-def.go new file mode 100644 index 0000000..c0497d3 --- /dev/null +++ b/internal/flow-def/flow-def.go @@ -0,0 +1,111 @@ +// MXL Flow Definition helper +package flowdef + +import ( + "encoding/json" + "errors" +) + +const ( + TYPE_VIDEO = iota + TYPE_VIDEO_ALPHA + TYPE_AUDIO +) + +var ( + ErrUnknownFlowType = errors.New("Unknown flow type provided") + ErrIncorrectFlowUUID = errors.New("Incorrect flow uuid") +) + +type flowDef struct { + Description string `json:"description"` + Id string `json:"id"` + Tags map[string][]string `json:"tags"` + Format string `json:"format"` + Label string `json:"label"` + Parents []string `json:"parents"` + MediaType string `json:"media_type"` + GrainRate grainRate `json:"grain_rate"` + FrameWidth uint `json:"frame_width"` + FrameHeight uint `json:"frame_height"` + InterlaceMode string `json:"interlace_mode"` + ColorSpace string `json:"colorspace"` + Components [3]videoComponent `json:"components"` +} + +type grainRate struct { + Numerator uint `json:"numerator"` + Denominator uint `json:"denominator"` +} + +type videoComponent struct { + Name string `json:"name"` + Width uint `json:"width"` + Height uint `json:"height"` + BitDepth uint `json:"bit_depth"` +} + +func NewFlowDefJSON( + feedType int, + uuid string, + width uint, + height uint, + fpsNum uint, + fpsDen uint, +) (string, error) { + if feedType != TYPE_VIDEO && + feedType != TYPE_VIDEO_ALPHA && + feedType != TYPE_AUDIO { + return "", ErrUnknownFlowType + } + if uuid == "" { + return "", ErrIncorrectFlowUUID + } + tags := map[string][]string{ + "urn:x-nmos:tag:grouphint/v1.0": { + "___one day I will follow NMOS specs___:Video", + }, + } + flowDef := flowDef{ + Description: "go-mxl-pattern-gen generated feed", + Id: uuid, + Tags: tags, + Format: "urn:x-nmos:format:video", + Label: "go-mxl-pattern-gen generated feed", + Parents: nil, + MediaType: "video/v210", + GrainRate: grainRate{ + Numerator: fpsNum, + Denominator: fpsDen, + }, + FrameWidth: width, + FrameHeight: height, + InterlaceMode: "progressive", + ColorSpace: "BT709", + Components: [3]videoComponent{ + videoComponent{ + Name: "Y", + Width: width, + Height: height, + BitDepth: 10, + }, + videoComponent{ + Name: "Cb", + Width: width / 2, + Height: height, + BitDepth: 10, + }, + videoComponent{ + Name: "Cr", + Width: width / 2, + Height: height, + BitDepth: 10, + }, + }, + } + jsonBytes, err := json.Marshal(flowDef) + if err != nil { + return "", err + } + return string(jsonBytes), nil +} diff --git a/internal/generator/font.go b/internal/generator/font.go new file mode 100644 index 0000000..0b97725 --- /dev/null +++ b/internal/generator/font.go @@ -0,0 +1,102 @@ +package generator + +// 8x16 1-bpp bitmap font baked from DejaVu Sans Mono 13px (public-domain +// style console glyphs). ASCII 32..127, one byte per row, LSB = leftmost pixel. +var fontAtlas8x16 = [96 * 16]byte{ + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x00, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x24, 0x24, 0x24, 0x24, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x48, 0x68, 0x28, 0xFE, 0x24, 0x24, 0x7F, 0x16, 0x12, 0x12, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x3C, 0x16, 0x02, 0x1C, 0x38, 0x40, 0x62, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x0E, 0x09, 0x09, 0x6E, 0x18, 0x76, 0xD0, 0xD0, 0x70, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x06, 0x04, 0x0C, 0x9A, 0xD3, 0x63, 0x62, 0x5C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x18, 0x18, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x10, 0x10, 0x18, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x18, 0x10, 0x10, 0x00, 0x00, 0x00, + 0x00, 0x0C, 0x08, 0x18, 0x18, 0x10, 0x10, 0x10, 0x10, 0x10, 0x18, 0x08, 0x0C, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x08, 0x4A, 0x1C, 0x1C, 0x4A, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x18, 0x18, 0x18, 0x7F, 0x18, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x18, 0x18, 0x08, 0x08, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x3C, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x60, 0x20, 0x30, 0x10, 0x10, 0x08, 0x08, 0x0C, 0x04, 0x06, 0x02, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x26, 0x62, 0x42, 0x5A, 0x42, 0x62, 0x26, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x1E, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x7C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x22, 0x60, 0x20, 0x30, 0x18, 0x0C, 0x06, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x22, 0x60, 0x20, 0x3C, 0x60, 0x60, 0x62, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x30, 0x38, 0x28, 0x24, 0x26, 0x22, 0x7E, 0x20, 0x20, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3E, 0x06, 0x06, 0x3E, 0x20, 0x60, 0x60, 0x22, 0x1C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x06, 0x02, 0x3E, 0x66, 0x42, 0x42, 0x66, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x60, 0x20, 0x30, 0x10, 0x18, 0x18, 0x08, 0x0C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x66, 0x62, 0x26, 0x3C, 0x66, 0x42, 0x66, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x26, 0x62, 0x62, 0x66, 0x7C, 0x60, 0x20, 0x1C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x18, 0x18, 0x00, 0x00, 0x00, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x18, 0x18, 0x00, 0x00, 0x00, 0x18, 0x18, 0x08, 0x08, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x40, 0x38, 0x06, 0x06, 0x38, 0x40, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x7F, 0x00, 0x00, 0x7F, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x02, 0x1C, 0x70, 0x70, 0x1C, 0x02, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x60, 0x60, 0x30, 0x18, 0x08, 0x00, 0x08, 0x08, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x46, 0x42, 0xF3, 0xC9, 0xC9, 0xC9, 0xF3, 0x02, 0x06, 0x38, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x18, 0x18, 0x3C, 0x34, 0x24, 0x66, 0x7E, 0x42, 0xC3, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3E, 0x62, 0x62, 0x62, 0x3E, 0x62, 0x42, 0x62, 0x3E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x38, 0x44, 0x06, 0x02, 0x02, 0x02, 0x06, 0x44, 0x38, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x1E, 0x22, 0x62, 0x42, 0x42, 0x42, 0x62, 0x22, 0x1E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x06, 0x06, 0x06, 0x7E, 0x06, 0x06, 0x06, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x06, 0x06, 0x06, 0x7E, 0x06, 0x06, 0x06, 0x06, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x46, 0x02, 0x02, 0x72, 0x42, 0x42, 0x46, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x42, 0x42, 0x42, 0x42, 0x7E, 0x42, 0x42, 0x42, 0x42, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x32, 0x1E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x42, 0x22, 0x12, 0x0E, 0x1E, 0x12, 0x32, 0x62, 0x42, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x06, 0x06, 0x06, 0x06, 0x06, 0x06, 0x06, 0x06, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x67, 0x67, 0x67, 0x5F, 0x5B, 0x5B, 0x43, 0x43, 0x43, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x46, 0x46, 0x4E, 0x4A, 0x5A, 0x52, 0x72, 0x62, 0x62, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x66, 0x62, 0x42, 0x42, 0x42, 0x62, 0x66, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3E, 0x66, 0x46, 0x46, 0x66, 0x3E, 0x06, 0x06, 0x06, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x66, 0x62, 0x42, 0x42, 0x42, 0x62, 0x66, 0x3C, 0x30, 0x20, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3E, 0x62, 0x62, 0x62, 0x3E, 0x32, 0x62, 0x42, 0xC2, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x3C, 0x06, 0x02, 0x06, 0x3C, 0x60, 0x40, 0x62, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0xFF, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x42, 0x42, 0x42, 0x42, 0x42, 0x42, 0x42, 0x66, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x43, 0x42, 0x62, 0x26, 0x24, 0x24, 0x3C, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0xC1, 0xC3, 0xC3, 0x5B, 0x5A, 0x5A, 0x76, 0x66, 0x66, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x42, 0x66, 0x3C, 0x18, 0x18, 0x3C, 0x24, 0x66, 0x43, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x43, 0x66, 0x24, 0x3C, 0x18, 0x18, 0x18, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x60, 0x20, 0x30, 0x18, 0x08, 0x04, 0x06, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x38, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x38, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x02, 0x06, 0x04, 0x0C, 0x08, 0x08, 0x10, 0x10, 0x30, 0x20, 0x60, 0x00, 0x00, + 0x00, 0x1C, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x1C, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x18, 0x3C, 0x26, 0x42, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xFF, 0x00, + 0x00, 0x00, 0x0C, 0x08, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3C, 0x22, 0x60, 0x7C, 0x62, 0x62, 0x7C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x02, 0x02, 0x02, 0x02, 0x3E, 0x66, 0x46, 0x42, 0x46, 0x66, 0x3E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x38, 0x44, 0x06, 0x06, 0x06, 0x44, 0x38, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x60, 0x60, 0x60, 0x60, 0x7C, 0x66, 0x62, 0x62, 0x62, 0x66, 0x7C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3C, 0x66, 0x42, 0x7E, 0x02, 0x46, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x70, 0x18, 0x08, 0x08, 0x7E, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x7C, 0x66, 0x62, 0x62, 0x62, 0x66, 0x7C, 0x60, 0x20, 0x1C, 0x00, + 0x00, 0x02, 0x02, 0x02, 0x02, 0x3E, 0x66, 0x66, 0x62, 0x62, 0x62, 0x62, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x18, 0x00, 0x00, 0x00, 0x1C, 0x18, 0x18, 0x18, 0x18, 0x18, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x10, 0x00, 0x00, 0x00, 0x1C, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x10, 0x18, 0x0E, 0x00, + 0x00, 0x06, 0x06, 0x06, 0x06, 0x66, 0x36, 0x1E, 0x1E, 0x36, 0x66, 0x46, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x0E, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x08, 0x70, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x7E, 0x5A, 0x5A, 0x5A, 0x5A, 0x5A, 0x5A, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3E, 0x66, 0x66, 0x62, 0x62, 0x62, 0x62, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3C, 0x66, 0x42, 0x42, 0x42, 0x66, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3E, 0x66, 0x46, 0x42, 0x46, 0x66, 0x3E, 0x02, 0x02, 0x02, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x7C, 0x66, 0x62, 0x62, 0x62, 0x66, 0x7C, 0x60, 0x60, 0x60, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x7C, 0x0C, 0x0C, 0x0C, 0x0C, 0x0C, 0x0C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x3C, 0x26, 0x06, 0x3C, 0x20, 0x22, 0x3C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x08, 0x08, 0x7E, 0x08, 0x08, 0x08, 0x08, 0x08, 0x78, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x62, 0x62, 0x62, 0x62, 0x66, 0x66, 0x7C, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x42, 0x62, 0x26, 0x24, 0x34, 0x18, 0x18, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0xC1, 0xC3, 0x5A, 0x5A, 0x7E, 0x66, 0x26, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x62, 0x24, 0x18, 0x18, 0x3C, 0x24, 0x42, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x42, 0x66, 0x24, 0x24, 0x3C, 0x18, 0x18, 0x18, 0x08, 0x06, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x7E, 0x20, 0x10, 0x18, 0x0C, 0x04, 0x7E, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x30, 0x18, 0x18, 0x18, 0x08, 0x0E, 0x08, 0x18, 0x18, 0x18, 0x18, 0x30, 0x00, 0x00, 0x00, + 0x00, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x18, 0x00, 0x00, + 0x00, 0x0E, 0x08, 0x18, 0x18, 0x18, 0x30, 0x18, 0x18, 0x18, 0x08, 0x08, 0x0E, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x0E, 0x70, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, + 0x00, 0x00, 0x00, 0x7E, 0x40, 0x40, 0x40, 0x40, 0x40, 0x40, 0x40, 0x40, 0x40, 0x7E, 0x00, 0x00, +} diff --git a/internal/generator/generator.go b/internal/generator/generator.go new file mode 100644 index 0000000..64dd4cd --- /dev/null +++ b/internal/generator/generator.go @@ -0,0 +1,6 @@ +package generator + +type FrameGenerator interface { + GenerateFrame(dest []byte, frameIndex int) error + Close() error +} diff --git a/internal/generator/opencl.go b/internal/generator/opencl.go new file mode 100644 index 0000000..355d423 --- /dev/null +++ b/internal/generator/opencl.go @@ -0,0 +1,302 @@ +// OpenCL pattern gen +package generator + +/* +#cgo linux LDFLAGS: -lOpenCL +#cgo windows LDFLAGS: -lOpenCL + +#include +#include + +static cl_program build_cl_program(cl_context ctx, cl_device_id dev, const char* src, cl_int* err) { + cl_program prog = clCreateProgramWithSource(ctx, 1, &src, NULL, err); + if (*err != CL_SUCCESS) return NULL; + *err = clBuildProgram(prog, 1, &dev, NULL, NULL, NULL); + return prog; +} + +static const char* get_build_log(cl_program prog, cl_device_id dev) { + static char buf[4096]; + size_t n = 0; + if (clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, sizeof(buf) - 1, buf, &n) != CL_SUCCESS) return ""; + buf[n < sizeof(buf) - 1 ? n : sizeof(buf) - 1] = '\0'; + return buf; +} +*/ +import "C" +import ( + "fmt" + "os" + "unsafe" +) + +type OpenCLGenerator struct { + ctx C.cl_context + queue C.cl_command_queue + kernel C.cl_kernel + prog C.cl_program + dev C.cl_device_id + atlasBuf C.cl_mem + textBuf C.cl_mem + width int + height int +} + +// TextConfig describes a burn-in text overlay rendered by the kernel +// with the baked 8x16 bitmap font. FG/BG are 10-bit {Y, Cb, Cr}. +type TextConfig struct { + Str string + X, Y int + Scale int + FG [3]uint32 + BG [3]uint32 + Boxed bool +} + +func (g *OpenCLGenerator) setArg(idx int, size C.size_t, p unsafe.Pointer) error { + if err := C.clSetKernelArg(g.kernel, C.cl_uint(idx), size, p); err != C.CL_SUCCESS { + return fmt.Errorf("opencl: set arg %d failed: %d", idx, err) + } + return nil +} + +// setTextArgs pushes the overlay arguments (kernel args 5..16) to the kernel. +// buf is the text byte buffer; a placeholder is allowed when len(cfg.Str) == 0, +// because the kernel checks text_len before dereferencing the text pointer. +func (g *OpenCLGenerator) setTextArgs(buf C.cl_mem, cfg TextConfig) error { + textLen := C.int(len(cfg.Str)) + tx, ty := C.int(cfg.X), C.int(cfg.Y) + scale := C.int(cfg.Scale) + if scale < 1 { + scale = 1 + } + hasBG := C.int(0) + if cfg.Boxed { + hasBG = 1 + } + fgY, fgCb, fgCr := C.cl_uint(cfg.FG[0]), C.cl_uint(cfg.FG[1]), C.cl_uint(cfg.FG[2]) + bgY, bgCb, bgCr := C.cl_uint(cfg.BG[0]), C.cl_uint(cfg.BG[1]), C.cl_uint(cfg.BG[2]) + args := []struct { + idx int + p unsafe.Pointer + sz C.size_t + }{ + {5, unsafe.Pointer(&buf), C.sizeof_cl_mem}, + {6, unsafe.Pointer(&textLen), C.sizeof_int}, + {7, unsafe.Pointer(&tx), C.sizeof_int}, + {8, unsafe.Pointer(&ty), C.sizeof_int}, + {9, unsafe.Pointer(&scale), C.sizeof_int}, + {10, unsafe.Pointer(&hasBG), C.sizeof_int}, + {11, unsafe.Pointer(&fgY), C.sizeof_cl_uint}, + {12, unsafe.Pointer(&fgCb), C.sizeof_cl_uint}, + {13, unsafe.Pointer(&fgCr), C.sizeof_cl_uint}, + {14, unsafe.Pointer(&bgY), C.sizeof_cl_uint}, + {15, unsafe.Pointer(&bgCb), C.sizeof_cl_uint}, + {16, unsafe.Pointer(&bgCr), C.sizeof_cl_uint}, + } + for _, a := range args { + if err := g.setArg(a.idx, a.sz, a.p); err != nil { + return err + } + } + return nil +} + +// SetText installs (or, with an empty string, removes) the burn-in overlay. +func (g *OpenCLGenerator) SetText(cfg TextConfig) error { + if len(cfg.Str) > 0 && (cfg.Scale < 1 || cfg.Scale > 8) { + return fmt.Errorf("opencl: text scale must be 1..8, got %d", cfg.Scale) + } + if len(cfg.Str) > 1024 { + return fmt.Errorf("opencl: text too long: %d bytes", len(cfg.Str)) + } + + var err C.cl_int + buf := g.atlasBuf // placeholder: text_len = 0 keeps the kernel off it + if len(cfg.Str) > 0 { + cData := C.CBytes([]byte(cfg.Str)) + defer C.free(cData) + buf = C.clCreateBuffer(g.ctx, C.CL_MEM_READ_ONLY|C.CL_MEM_COPY_HOST_PTR, + C.size_t(len(cfg.Str)), cData, &err) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: text buffer upload failed: %d", err) + } + } + + if err := g.setTextArgs(buf, cfg); err != nil { + if len(cfg.Str) > 0 { + C.clReleaseMemObject(buf) + } + return err + } + + if g.textBuf != nil && g.textBuf != buf { + C.clReleaseMemObject(g.textBuf) + } + if len(cfg.Str) > 0 { + g.textBuf = buf + } else { + g.textBuf = nil + } + return nil +} + +func NewOpenCLGenerator(width, height uint, kernelPath string) (*OpenCLGenerator, error) { + var err C.cl_int + var platform C.cl_platform_id + + if err = C.clGetPlatformIDs(1, &platform, nil); err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: platform not found: %d", err) + } + + g := &OpenCLGenerator{width: int(width), height: int(height)} + if err = C.clGetDeviceIDs(platform, C.CL_DEVICE_TYPE_GPU, 1, &g.dev, nil); err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: no suitable gpu device found") + } + + g.ctx = C.clCreateContext(nil, 1, &g.dev, nil, nil, &err) + if err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: ctx creation error") + } + + g.queue = C.clCreateCommandQueueWithProperties(g.ctx, g.dev, nil, &err) + if err != C.CL_SUCCESS { + g.queue = C.clCreateCommandQueue(g.ctx, g.dev, 0, &err) + } + + src, errGo := os.ReadFile(kernelPath) + if errGo != nil { + g.Close() + return nil, fmt.Errorf("opencl: read kernel error: %w", errGo) + } + + cSrc := C.CString(string(src)) + defer C.free(unsafe.Pointer(cSrc)) + + g.prog = C.build_cl_program(g.ctx, g.dev, cSrc, &err) + if err != C.CL_SUCCESS { + msg := "clCreateProgramWithSource failed" + if g.prog != nil { + msg = C.GoString(C.get_build_log(g.prog, g.dev)) + } + g.Close() + return nil, fmt.Errorf("opencl: compilation failed: %s", msg) + } + + cKName := C.CString("generate_v210_pattern") + defer C.free(unsafe.Pointer(cKName)) + g.kernel = C.clCreateKernel(g.prog, cKName, &err) + if err != C.CL_SUCCESS { + g.Close() + return nil, fmt.Errorf("opencl: kernel init failed") + } + + g.atlasBuf = C.clCreateBuffer(g.ctx, C.CL_MEM_READ_ONLY|C.CL_MEM_COPY_HOST_PTR, + C.size_t(len(fontAtlas8x16)), unsafe.Pointer(&fontAtlas8x16[0]), &err) + if err != C.CL_SUCCESS { + g.Close() + return nil, fmt.Errorf("opencl: font atlas upload failed: %d", err) + } + if err := g.setArg(4, C.sizeof_cl_mem, unsafe.Pointer(&g.atlasBuf)); err != nil { + g.Close() + return nil, err + } + + // Overlay defaults: no text, white on black if SetText is ever called. + if err := g.setTextArgs(g.atlasBuf, TextConfig{ + Scale: 1, + FG: [3]uint32{940, 512, 512}, + BG: [3]uint32{64, 512, 512}, + }); err != nil { + g.Close() + return nil, err + } + + return g, nil +} + +func (g *OpenCLGenerator) GenerateFrame(dest []byte, frameIndex int) error { + var err C.cl_int + byteSize := C.size_t(len(dest)) + + // ВАЖНО: Мы маппим память MXL SDK ([]byte) напрямую в OpenCL буфер. + // GPU запишет данные прямо в общую память MXL без лишнего копирования (Zero-Copy). + if want := (g.width * g.height / 6) * 16; len(dest) < want { + return fmt.Errorf("opencl: dest too small: %d bytes, need %d", len(dest), want) + } + clMem := C.clCreateBuffer( + g.ctx, + C.CL_MEM_WRITE_ONLY|C.CL_MEM_USE_HOST_PTR, + byteSize, + unsafe.Pointer(&dest[0]), + &err, + ) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: mxl host pointer mapping failed: %d", err) + } + defer C.clReleaseMemObject(clMem) + + if err := g.setArg(0, C.sizeof_cl_mem, unsafe.Pointer(&clMem)); err != nil { + return err + } + w, h, f := C.int(g.width), C.int(g.height), C.int(frameIndex) + if err := g.setArg(1, C.sizeof_int, unsafe.Pointer(&w)); err != nil { + return err + } + if err := g.setArg(2, C.sizeof_int, unsafe.Pointer(&h)); err != nil { + return err + } + if err := g.setArg(3, C.sizeof_int, unsafe.Pointer(&f)); err != nil { + return err + } + + // Каждый поток GPU пакует блок из 6 пикселей (16 байт) + globalWorkSize := C.size_t((g.width * g.height) / 6) + err = C.clEnqueueNDRangeKernel(g.queue, g.kernel, 1, nil, &globalWorkSize, nil, 0, nil, nil) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: dispatch failed: %d", err) + } + + // Гарантированная синхронизация записи GPU в host-память: map/unmap. + // На discrete GPU драйвер держит копию в device memory; одного clFinish + // недостаточно — данные обязаны появиться в host_ptr только после unmap. + mapped := C.clEnqueueMapBuffer(g.queue, clMem, C.CL_TRUE, C.CL_MAP_READ, 0, byteSize, 0, nil, nil, &err) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: map failed: %d", err) + } + if err = C.clEnqueueUnmapMemObject(g.queue, clMem, mapped, 0, nil, nil); err != C.CL_SUCCESS { + return fmt.Errorf("opencl: unmap failed: %d", err) + } + if err = C.clFinish(g.queue); err != C.CL_SUCCESS { + return fmt.Errorf("opencl: finish failed: %d", err) + } + return nil +} + +func (g *OpenCLGenerator) Close() error { + if g.textBuf != nil { + C.clReleaseMemObject(g.textBuf) + g.textBuf = nil + } + if g.atlasBuf != nil { + C.clReleaseMemObject(g.atlasBuf) + g.atlasBuf = nil + } + if g.kernel != nil { + C.clReleaseKernel(g.kernel) + g.kernel = nil + } + if g.prog != nil { + C.clReleaseProgram(g.prog) + g.prog = nil + } + if g.queue != nil { + C.clReleaseCommandQueue(g.queue) + g.queue = nil + } + if g.ctx != nil { + C.clReleaseContext(g.ctx) + g.ctx = nil + } + return nil +} diff --git a/internal/generator/wgpu.go b/internal/generator/wgpu.go new file mode 100644 index 0000000..d1f0ea7 --- /dev/null +++ b/internal/generator/wgpu.go @@ -0,0 +1,236 @@ +// WGPU pattern gen: renders v210 frames with a wgpu (WebGPU/Vulkan) +// compute shader. Pure Go — no cgo in the GPU layer. +package generator + +import ( + "context" + "encoding/binary" + "fmt" + "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 wgpuWorkgroupSize = 64 + +// WGPUGenerator renders v210 frames on the GPU via wgpu and implements +// FrameGenerator. Per frame the kernel packs the pattern into a storage +// 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 +} + +var _ FrameGenerator = (*WGPUGenerator)(nil) + +func NewWGPUGenerator(width, height uint, kernelPath string) (*WGPUGenerator, error) { + g := &WGPUGenerator{ + width: int(width), + height: int(height), + blocks: int(width*height) / 6, + params: make([]byte, 16), + } + g.frameSize = uint64(g.blocks) * 16 + binary.LittleEndian.PutUint32(g.params[0:], uint32(width)) + binary.LittleEndian.PutUint32(g.params[4:], uint32(height)) + + var err error + if g.instance, err = wgpu.CreateInstance(nil); err != nil { + return nil, fmt.Errorf("wgpu: create instance: %w", err) + } + if g.adapter, err = g.instance.RequestAdapter(nil); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: request adapter: %w", err) + } + if g.device, err = g.adapter.RequestDevice(nil); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: request device: %w", err) + } + g.queue = g.device.Queue() + + wgsl, err := os.ReadFile(kernelPath) + if err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: read kernel: %w", err) + } + if g.shader, err = g.device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{ + Label: "v210-shader", WGSL: string(wgsl), + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: shader: %w", err) + } + if g.out, err = g.device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "v210-out", Size: g.frameSize, + 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, + Usage: wgpu.BufferUsageCopyDst | wgpu.BufferUsageMapRead, + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: host buffer: %w", err) + } + if g.uniform, err = g.device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "v210-params", Size: uint64(len(g.params)), + Usage: wgpu.BufferUsageUniform | wgpu.BufferUsageCopyDst, + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: params buffer: %w", err) + } + if err := g.queue.WriteBuffer(g.uniform, 0, g.params); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: write params: %w", err) + } + 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}}, + }, + }); 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))}, + }, + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: bind group: %w", err) + } + if g.pl, err = g.device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{ + Label: "v210-pl", BindGroupLayouts: []*wgpu.BindGroupLayout{g.bgl}, + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: pipeline layout: %w", err) + } + if g.pipeline, err = g.device.CreateComputePipeline(&wgpu.ComputePipelineDescriptor{ + Label: "v210-pipeline", Layout: g.pl, Module: g.shader, EntryPoint: "main", + }); err != nil { + g.Close() + return nil, fmt.Errorf("wgpu: pipeline: %w", err) + } + return g, nil +} + +func (g *WGPUGenerator) GenerateFrame(dest []byte, frameIndex int) error { + if uint64(len(dest)) < g.frameSize { + return fmt.Errorf("wgpu: dest too small: %d bytes, need %d", len(dest), g.frameSize) + } + + binary.LittleEndian.PutUint32(g.params[8:], uint32(frameIndex)) + if err := g.queue.WriteBuffer(g.uniform, 0, g.params); err != nil { + return fmt.Errorf("wgpu: write params: %w", err) + } + + encoder, err := g.device.CreateCommandEncoder(nil) + if err != nil { + return fmt.Errorf("wgpu: encoder: %w", err) + } + pass, err := encoder.BeginComputePass(nil) + if err != nil { + return fmt.Errorf("wgpu: compute pass: %w", err) + } + pass.SetPipeline(g.pipeline) + pass.SetBindGroup(0, g.bg, nil) + pass.Dispatch(uint32((g.blocks+wgpuWorkgroupSize-1)/wgpuWorkgroupSize), 1, 1) + if err := pass.End(); err != nil { + return fmt.Errorf("wgpu: end pass: %w", err) + } + encoder.CopyBufferToBuffer(g.out, 0, g.host, 0, g.frameSize) + cmd, err := encoder.Finish() + if err != nil { + return fmt.Errorf("wgpu: finish: %w", err) + } + if _, err := g.queue.Submit(cmd); err != nil { + return fmt.Errorf("wgpu: submit: %w", err) + } + + ctx, cancel := context.WithTimeout(context.Background(), 5*time.Second) + defer cancel() + if err := g.host.Map(ctx, wgpu.MapModeRead, 0, g.frameSize); err != nil { + return fmt.Errorf("wgpu: map: %w", err) + } + rng, err := g.host.MappedRange(0, g.frameSize) + if err != nil { + _ = g.host.Unmap() + return fmt.Errorf("wgpu: mapped range: %w", err) + } + copy(dest, rng.Bytes()) + return g.host.Unmap() +} + +func (g *WGPUGenerator) Close() error { + if g.pipeline != nil { + g.pipeline.Release() + g.pipeline = nil + } + if g.pl != nil { + g.pl.Release() + g.pl = nil + } + if g.bg != nil { + g.bg.Release() + g.bg = nil + } + if g.bgl != nil { + g.bgl.Release() + g.bgl = nil + } + if g.shader != nil { + g.shader.Release() + g.shader = nil + } + if g.uniform != nil { + g.uniform.Release() + g.uniform = nil + } + if g.host != nil { + g.host.Release() + g.host = nil + } + if g.out != nil { + g.out.Release() + g.out = nil + } + if g.device != nil { + g.device.Release() + g.device = nil + } + if g.adapter != nil { + g.adapter.Release() + g.adapter = nil + } + if g.instance != nil { + g.instance.Release() + g.instance = nil + } + return nil +} diff --git a/internal/generator/wgpu_move_test.go b/internal/generator/wgpu_move_test.go new file mode 100644 index 0000000..67f715e --- /dev/null +++ b/internal/generator/wgpu_move_test.go @@ -0,0 +1,74 @@ +package generator + +import ( + "encoding/binary" + "path/filepath" + "testing" +) + +func TestWGPUMoveSquare(t *testing.T) { + const width, height = 1920, 1080 + g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "v210_bars_move.wgsl")) + 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 + } + } + + // square center x=960 sits on the green bar (534/253/207): + // inverted -> Y=1004-534=470, Cb=1024-253=771, Cr=1024-207=817 + const invY, invCb, invCr = 470, 771, 817 + const greenY = 534 // bar 3 covers x in [823,1098) + + // tick 0: offset = 0, square centered at (960,540), half = 75 + if err := g.GenerateFrame(buf, 0); err != nil { + t.Fatalf("tick 0: %v", err) + } + if y, cb, cr := sample(960, 540); y != invY || cb != invCb || cr != invCr { + t.Fatalf("tick 0 center: got %d/%d/%d, want inverted green 470/771/817", y, cb, cr) + } + if y, _, _ := sample(860, 540); y != greenY { + t.Fatalf("tick 0 left of square: Y=%d, want green %d", y, greenY) + } + + // amplitude = centerX - half = 885: tick 79 (sin~1) puts the square at + // the far right, x in [1770,1920), over the blue bar (111/848/481): + // inverted -> Y=1004-111=893, Cb=1024-848=176, Cr=1024-481=543 + const invBlueY, invBlueCb, invBlueCr = 893, 176, 543 + if err := g.GenerateFrame(buf, 79); err != nil { + t.Fatalf("tick 79: %v", err) + } + if y, cb, cr := sample(1840, 540); y != invBlueY || cb != invBlueCb || cr != invBlueCr { + t.Fatalf("tick 79 shifted center: got %d/%d/%d, want inverted blue 893/176/543", y, cb, cr) + } + if y, _, _ := sample(960, 540); y != greenY { + t.Fatalf("tick 79 old center: Y=%d, want green %d (square moved away)", y, greenY) + } + // bars untouched far from the square + if y, _, _ := sample(100, 100); y != 721 { + t.Fatalf("tick 79 bars: Y=%d, want white 721", y) + } +} diff --git a/internal/generator/wgpu_test.go b/internal/generator/wgpu_test.go new file mode 100644 index 0000000..10db231 --- /dev/null +++ b/internal/generator/wgpu_test.go @@ -0,0 +1,67 @@ +package generator + +import ( + "encoding/binary" + "path/filepath" + "testing" +) + +func TestWGPUGenerator(t *testing.T) { + const width, height = 1920, 1080 + g, err := NewWGPUGenerator(width, height, filepath.Join("..", "..", "kernels", "v210_bars.wgsl")) + if err != nil { + t.Fatalf("init: %v", err) + } + defer g.Close() + + buf := make([]byte, width*height*8/3) + if err := g.GenerateFrame(buf, 0); err != nil { + t.Fatalf("GenerateFrame: %v", err) + } + if err := g.GenerateFrame(buf, 1); err != nil { + t.Fatalf("GenerateFrame 2: %v", err) + } + + want := [7][3]uint32{ + {721, 512, 512}, {674, 176, 543}, {581, 589, 176}, + {534, 253, 207}, {251, 771, 817}, {204, 435, 848}, {111, 848, 481}, + } + 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(buf[off:]) + w1 := binary.LittleEndian.Uint32(buf[off+4:]) + w2 := binary.LittleEndian.Uint32(buf[off+8:]) + w3 := binary.LittleEndian.Uint32(buf[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 := want[barOf(x)] + if yv != b[0] { + t.Fatalf("pixel (%d,%d): Y=%d want %d", x, y, yv, b[0]) + } + if x%2 == 0 && (cb != b[1] || cr != b[2]) { + t.Fatalf("pixel (%d,%d): Cb=%d Cr=%d want %d/%d", x, y, cb, cr, b[1], b[2]) + } + } + } +} diff --git a/kernels/v210_bars.wgsl b/kernels/v210_bars.wgsl new file mode 100644 index 0000000..08cdfac --- /dev/null +++ b/kernels/v210_bars.wgsl @@ -0,0 +1,50 @@ +// 75% SMPTE color bars, 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, + _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; + + // 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); + + var y: array; + var cb: array; + var cr: array; + for (var i = 0u; i < 6u; i++) { + let b = bar_index(x + i); + y[i] = y_tab[b]; + cb[i] = cb_tab[b]; + cr[i] = cr_tab[b]; + } + + // 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_bars_move.wgsl b/kernels/v210_bars_move.wgsl new file mode 100644 index 0000000..1a97270 --- /dev/null +++ b/kernels/v210_bars_move.wgsl @@ -0,0 +1,77 @@ +// 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); + let b = bar_index(x + i); + 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_tab[b]; + cb[i] = 1024u - cb_tab[b]; + cr[i] = 1024u - cr_tab[b]; + } else { + y[i] = y_tab[b]; + cb[i] = cb_tab[b]; + cr[i] = cr_tab[b]; + } + } + + // 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_pack.cl b/kernels/v210_pack.cl new file mode 100644 index 0000000..938f2ef --- /dev/null +++ b/kernels/v210_pack.cl @@ -0,0 +1,72 @@ +// 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/linkcheck b/linkcheck new file mode 100755 index 0000000..4e21740 Binary files /dev/null and b/linkcheck differ diff --git a/scripts/gen_font.py b/scripts/gen_font.py new file mode 100644 index 0000000..bcacd7c --- /dev/null +++ b/scripts/gen_font.py @@ -0,0 +1,53 @@ +#!/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") diff --git a/v210-mxl-architecture.md b/v210-mxl-architecture.md new file mode 100644 index 0000000..6c2368d --- /dev/null +++ b/v210-mxl-architecture.md @@ -0,0 +1,485 @@ +# Архитектура v210 Генератора для Media eXchange Layer (go-mxl) + +Профессиональная Zero-GC и Zero-Copy архитектура гибридного **CPU/GPU v210 тестового генератора**, оптимизированная под работу с официальным C SDK Media eXchange Layer (`libmxl`) через официальные Go-биндинги `go-mxl` от Qvest Digital. + +## 🏢 Финальная структура проекта + +```text +v210-generator/ +├── cmd/ +│ └── v210-gen/ +│ └── main.go # Cobra CLI & главный цикл записи в MXL Flow +├── internal/ +│ └── generator/ +│ ├── generator.go # Интерфейс генератора +│ ├── cpu.go # CPU битовый упаковщик (Пишет сразу в []byte) +│ └── opencl.go # OpenCL Cgo упаковщик (Маппит []byte из MXL в GPU) +└── kernels/ + └── v210_pack.cl # Шейдер упаковки v210 +``` + +--- + +## 1. Архитектура Генератора: `internal/generator/generator.go` + +Так как `go-mxl` предоставляет память в виде стандартных байтовых срезов (`[]byte`), мы адаптируем интерфейс. Генераторы больше не создают буферы, они принимают сырой `[]byte`, выделенный с помощью SDK. + +```go +package generator + +type FrameGenerator interface { + // GenerateFrame упаковывает данные напрямую в предоставленный байтовый срез. + // Этот срез будет являться физической памятью MXL Shared Memory. + GenerateFrame(dest []byte, frameIndex int) error + Close() error +} +``` + +--- + +## 2. GPU Движок под MXL: `internal/generator/opencl.go` + +Поскольку `go-mxl` дает нам прямой доступ к разделяемой памяти, мы берем этот срез `[]byte`, берем его адрес в RAM через `unsafe.Pointer` и отдаем драйверу OpenCL через флаг `CL_MEM_USE_HOST_PTR`. Видеокарта запишет v210 паттерн напрямую в MXL, минуя кэш Go. + +```go +package generator + +/* +#cgo linux LDFLAGS: -lOpenCL +#cgo windows LDFLAGS: -lOpenCL + +#include +#include + +static cl_program build_cl_program(cl_context ctx, cl_device_id dev, const char* src, cl_int* err) { + cl_program prog = clCreateProgramWithSource(ctx, 1, &src, NULL, err); + if (*err != CL_SUCCESS) return NULL; + *err = clBuildProgram(prog, 1, &dev, NULL, NULL, NULL); + return prog; +} +*/ +import "C" +import ( + "fmt" + "os" + "unsafe" +) + +type OpenCLGenerator struct { + ctx C.cl_context + queue C.cl_command_queue + kernel C.cl_kernel + prog C.cl_program + dev C.cl_device_id + width int + height int +} + +func NewOpenCLGenerator(width, height int, kernelPath string) (*OpenCLGenerator, error) { + var err C.cl_int + var platform C.cl_platform_id + + if err = C.clGetPlatformIDs(1, &platform, nil); err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: platform not found: %d", err) + } + + g := &OpenCLGenerator{width: width, height: height} + if err = C.clGetDeviceIDs(platform, C.CL_DEVICE_TYPE_GPU, 1, &g.dev, nil); err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: no suitable gpu device found") + } + + g.ctx = C.clCreateContext(nil, 1, &g.dev, nil, nil, &err) + if err != C.CL_SUCCESS { + return nil, fmt.Errorf("opencl: ctx creation error") + } + + g.queue = C.clCreateCommandQueueWithProperties(g.ctx, g.dev, nil, &err) + if err != C.CL_SUCCESS { + g.queue = C.clCreateCommandQueue(g.ctx, g.dev, 0, &err) + } + + src, errGo := os.ReadFile(kernelPath) + if errGo != nil { + g.Close() + return nil, fmt.Errorf("opencl: read kernel error: %w", errGo) + } + + cSrc := C.CString(string(src)) + defer C.free(unsafe.Pointer(cSrc)) + + g.prog = C.build_cl_program(g.ctx, g.dev, cSrc, &err) + if err != C.CL_SUCCESS { + g.Close() + return nil, fmt.Errorf("opencl: compilation failed") + } + + cKName := C.CString("generate_v210_pattern") + defer C.free(unsafe.Pointer(cKName)) + g.kernel = C.clCreateKernel(g.prog, cKName, &err) + if err != C.CL_SUCCESS { + g.Close() + return nil, fmt.Errorf("opencl: kernel init failed") + } + + return g, nil +} + +func (g *OpenCLGenerator) GenerateFrame(dest []byte, frameIndex int) error { + var err C.cl_int + byteSize := C.size_t(len(dest)) + + // ВАЖНО: Мы маппим память MXL SDK ([]byte) напрямую в OpenCL буфер. + // GPU запишет данные прямо в общую память MXL без лишнего копирования (Zero-Copy). + clMem := C.clCreateBuffer( + g.ctx, + C.CL_MEM_WRITE_ONLY|C.CL_MEM_USE_HOST_PTR, + byteSize, + unsafe.Pointer(&dest), // Указатель на первый элемент MXL-слайса + &err, + ) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: mxl host pointer mapping failed: %d", err) + } + defer C.clReleaseMemObject(clMem) + + C.clSetKernelArg(g.kernel, 0, C.sizeof_cl_mem, unsafe.Pointer(&clMem)) + w, h, f := C.int(g.width), C.int(g.height), C.int(frameIndex) + C.clSetKernelArg(g.kernel, 1, C.sizeof_int, unsafe.Pointer(&w)) + C.clSetKernelArg(g.kernel, 2, C.sizeof_int, unsafe.Pointer(&h)) + C.clSetKernelArg(g.kernel, 3, C.sizeof_int, unsafe.Pointer(&f)) + + // Каждый поток GPU пакует блок из 4 пикселей (16 байт) + globalWorkSize := C.size_t((g.width * g.height) / 4) + err = C.clEnqueueNDRangeKernel(g.queue, g.kernel, 1, nil, &globalWorkSize, nil, 0, nil, nil) + if err != C.CL_SUCCESS { + return fmt.Errorf("opencl: dispatch failed: %d", err) + } + + // Ждем окончания записи видеокарты в память MXL + C.clFinish(g.queue) + return nil +} + +func (g *OpenCLGenerator) Close() error { + if g.kernel != nil { C.clReleaseKernel(g.kernel) } + if g.prog != nil { C.clReleaseProgram(g.prog) } + if g.queue != nil { C.clReleaseCommandQueue(g.queue) } + if g.ctx != nil { C.clReleaseContext(g.ctx) } + return nil +} +``` + +--- + +## 3. CPU Паковщик: `internal/generator/cpu.go` + +CPU-генератор делает то же самое, но использует безопасную кастомную битовую маску для работы с `[]byte` напрямую вместо `[]uint32`, чтобы избежать преобразования типов данных. + +```go +package generator + +import "encoding/binary" + +type CPUGenerator struct { + width int + height int +} + +func NewCPUGenerator(width, height int) *CPUGenerator { + return &CPUGenerator{width: width, height: height} +} + +func (g *CPUGenerator) GenerateFrame(dest []byte, frameIndex int) error { + totalBlocks := (g.width * g.height) / 4 + + for blockID := 0; blockID < totalBlocks; blockID++ { + baseX := (blockID * 4) % g.width + + y0 := uint32((baseX * 1023) / g.width) + y1 := uint32(((baseX + 1) * 1023) / g.width) + y2 := uint32(((baseX + 2) * 1023) / g.width) + y3 := uint32(((baseX + 3) * 1023) / g.width) + u0, v0, u2, v2 := uint32(512), uint32(512), uint32(512), uint32(512) + + w0 := (u0 & 0x3FF) | ((y0 & 0x3FF) << 10) | ((v0 & 0x3FF) << 20) + w1 := (y1 & 0x3FF) | ((u2 & 0x3FF) << 10) | ((y2 & 0x3FF) << 20) + w2 := (v2 & 0x3FF) | ((y3 & 0x3FF) << 10) | ((u0 & 0x3FF) << 20) + w3 := uint32(0) + + // Записываем 4 слова (16 байт) прямо в байтовый слайс памяти MXL + byteIdx := blockID * 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) + } + return nil +} + +func (g *CPUGenerator) Close() error { return nil } +``` + +--- + +## 4. Интеграция и Главный Цикл MXL: `cmd/v210-gen/main.go` + +Здесь мы запускаем MXL инстанс и пишем фреймы. Обратите внимание, что мы **не используем `sync.Pool`**. Нам больше не нужно менеджить память в Go — MXL SDK сам выдает нам кусок разделяемой памяти для текущего индекса зерна (Grain). Наша задача — наполнить его и вызвать `Commit()`. + +```go +package main + +import ( + "fmt" + "log" + "time" + "v210-generator/internal/generator" + + "github.com/qvest-digital/go-mxl/mxl" +) + +func main() { + useGPU := true + width, height := 1920, 1080 + flowUUID := "your-video-flow-uuid-here" + + // 1. Инициализируем инстанс MXL (работает через /dev/shm/mxl) + inst, err := mxl.NewInstance("/dev/shm/mxl", "") + if err != nil { + log.Fatalf("MXL Init failed: %v", err) + } + defer inst.Close() + + // 2. Создаем Writer для отправки потока данных в MXL Fabric + writer, err := inst.NewWriter(flowUUID) + if err != nil { + log.Fatalf("Failed to create MXL Writer: %v", err) + } + defer writer.Close() + + // 3. Выбираем движок рендеринга + var engine generator.FrameGenerator + if useGPU { + engine, err = generator.NewOpenCLGenerator(width, height, "../../kernels/v210_pack.cl") + if err != nil { + log.Printf("⚠️ GPU OpenCL failed: %v. Falling back to CPU.", err) + engine = generator.NewCPUGenerator(width, height) + } + } else { + engine = generator.NewCPUGenerator(width, height) + } + defer engine.Close() + + // Получаем параметры таймингов из конфига потока + info, _ := writer.Info() + rate := info.Config.Common.GrainRate + idx := mxl.CurrentIndex(rate) + + fmt.Println("📺 Broadcasting v210 patterns directly into MXL Shared Memory...") + + for { + // Алоцируем Grain (зерно данных) в разделяемой памяти MXL. + // В этот момент SDK выделяет структуру под капот. + grain, err := writer.OpenGrain(idx, 100*time.Millisecond) + if err != nil { + log.Printf("Failed to open MXL grain for index %d: %v", idx, err) + time.Sleep(10 * time.Millisecond) + continue + } + + // grain.Payload — это []byte, ссылающийся напрямую на Linux Shared Memory + // Мы передаем его в наш генератор + err = engine.GenerateFrame(grain.Payload, int(idx)) + if err != nil { + log.Printf("Render error at index %d: %v", idx, err) + grain.Cancel() // Отменяем транзакцию для этого кадра, если произошла ошибка + idx++ + continue + } + + // Фиксируем (коммитим) кадр. С этого момента MXL рассылает его подписчикам (Readers) + grain.Commit() + + idx++ + } +} +``` + +## 🏎️ Почему эта связка идеальна? +* **Полный Zero-Copy**: Благодаря `go-mxl` от Qvest Digital, буферы памяти выделяются прямо в Linux Shared Memory (`/dev/shm/mxl`). Мы берем указатель на этот слайс и маппим его в OpenCL с помощью `CL_MEM_USE_HOST_PTR`. Видеокарта выполняет операции упаковки и складывает кадры прямо в оперативную память обмена. +* **0% нагрузки на GC в Go**: В главном цикле `for` отсутствуют динамические аллокации в куче (heap), рантайм Go работает без пауз на очистку памяти (STW), гарантируя идеальный frame pacing для Real-time Broadcast систем. + + +## BARS +__kernel void generate_v210_pattern(__global uint* output, int width, int height, int frame_index) { + int g_id = get_global_id(0); + int total_blocks = (width * height) / 4; + if (g_id >= total_blocks) return; + + int pixel_x = (g_id * 4) % width; + int pixel_y = (g_id * 4) / width; + + // В SMPTE RP 219 полосы занимают верхние 67% экрана. Ниже идут другие элементы. + // Для простоты примера сделаем полосы на весь экран, либо вы можете сделать split по pixel_y. + + // Вычисляем ширину одной полосы (всего их 7) + int bar_width = width / 7; + int bar_index = pixel_x / bar_width; + if (bar_index > 6) bar_index = 6; // Защита от выхода за границы из-за округления + + // Массивы констант для 75% SMPTE полос + ushort Y_table[7] = {721, 647, 569, 495, 315, 241, 163}; + ushort U_table[7] = {512, 176, 599, 263, 761, 425, 848}; + ushort V_table[7] = {512, 557, 176, 221, 803, 848, 467}; + + // Так как поток берет 4 пикселя подряд, определим цвета для каждого из них. + // (В 95% случаев они попадут в одну полосу, но на стыках полос пиксели будут разными) + ushort y0 = Y_table[(pixel_x + 0) / bar_width > 6 ? 6 : (pixel_x + 0) / bar_width]; + ushort y1 = Y_table[(pixel_x + 1) / bar_width > 6 ? 6 : (pixel_x + 1) / bar_width]; + ushort y2 = Y_table[(pixel_x + 2) / bar_width > 6 ? 6 : (pixel_x + 2) / bar_width]; + ushort y3 = Y_table[(pixel_x + 3) / bar_width > 6 ? 6 : (pixel_x + 3) / bar_width]; + + // v210 — это 4:2:2. Берем Хрому (U/V) от первого пикселя каждой пары (субдискретизация) + ushort u0 = U_table[(pixel_x + 0) / bar_width > 6 ? 6 : (pixel_x + 0) / bar_width]; + ushort v0 = V_table[(pixel_x + 0) / bar_width > 6 ? 6 : (pixel_x + 0) / bar_width]; + + ushort u2 = U_table[(pixel_x + 2) / bar_width > 6 ? 6 : (pixel_x + 2) / bar_width]; + ushort v2 = V_table[(pixel_x + 2) / bar_width > 6 ? 6 : (pixel_x + 2) / bar_width]; + + // Битовая упаковка v210 в четыре 32-битных слова + uint word0 = (u0 & 0x3FF) | ((y0 & 0x3FF) << 10) | ((v0 & 0x3FF) << 20); + uint word1 = (y1 & 0x3FF) | ((u2 & 0x3FF) << 10) | ((y2 & 0x3FF) << 20); + uint word2 = (v2 & 0x3FF) | ((y3 & 0x3FF) << 10) | ((u0 & 0x3FF) << 20); + uint word3 = 0; // 4-е слово в схеме v210 несет только выравнивание (высшие биты пусты) + + 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; +} + +## BARS CPU +package generator + +import "encoding/binary" + +type CPUGenerator struct { + width int + height int +} + +func NewCPUGenerator(width, height int) *CPUGenerator { + return &CPUGenerator{width: width, height: height} +} + +func (g *CPUGenerator) GenerateFrame(dest []byte, frameIndex int) error { + totalBlocks := (g.width * g.height) / 4 + barWidth := g.width / 7 + + // SMPTE Таблицы констант + Y_table := [7]uint32{721, 647, 569, 495, 315, 241, 163} + U_table := [7]uint32{512, 176, 599, 263, 761, 425, 848} + V_table := [7]uint32{512, 557, 176, 221, 803, 848, 467} + + getIdx := func(x int) int { + idx := x / barWidth + if idx > 6 { return 6 } + return idx + } + + for blockID := 0; blockID < totalBlocks; blockID++ { + baseX := (blockID * 4) % g.width + + // Запрашиваем индексы цветов для 4 пикселей в блоке + idx0 := getIdx(baseX + 0) + idx1 := getIdx(baseX + 1) + idx2 := getIdx(baseX + 2) + idx3 := getIdx(baseX + 3) + + y0, y1, y2, y3 := Y_table[idx0], Y_table[idx1], Y_table[idx2], Y_table[idx3] + u0, v0 := U_table[idx0], V_table[idx0] + u2, v2 := U_table[idx2], V_table[idx2] + + // Упаковка по спецификации v210 + w0 := (u0 & 0x3FF) | ((y0 & 0x3FF) << 10) | ((v0 & 0x3FF) << 20) + w1 := (y1 & 0x3FF) | ((u2 & 0x3FF) << 10) | ((y2 & 0x3FF) << 20) + w2 := (v2 & 0x3FF) | ((y3 & 0x3FF) << 10) | ((u0 & 0x3FF) << 20) + w3 := uint32(0) + + byteIdx := blockID * 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) + } + return nil +} + +func (g *CPUGenerator) Close() error { return nil } + + +struct Uniforms { + u_dt: f32, // Delta time for this frame (4 bytes) + u_frame: u32, // Current frame number (4 bytes) +}; + +@group(0) @binding(0) var uniforms: Uniforms; + +struct VertexOutput { + @builtin(position) position: vec4, + @location(0) color: vec4, +}; + +// We store the persistent position on the GPU using a global variable +// so we can incrementally add 'dt' to it frame over frame. +var squarePosition: vec2 = vec2(0.0, 0.0); +var velocity: vec2 = vec2(0.4, 0.2); // Moves diagonally + +@vertex +fn vs_main(@builtin(vertex_index) vertexIndex: u32) -> VertexOutput { + var pos = array, 6>( + vec2(-0.2, 0.2), + vec2(-0.2, -0.2), + vec2( 0.2, 0.2), + vec2( 0.2, 0.2), + vec2(-0.2, -0.2), + vec2( 0.2, -0.2) + ); + + var output: VertexOutput; + + // 1. Frame-based logic: Change color every 60 frames + var color = vec4(0.3, 0.6, 0.9, 1.0); // Light blue + if ((uniforms.u_frame / 60u) % 2u == 0u) { + color = vec4(0.9, 0.4, 0.3, 1.0); // Switch to Coral Red + } + + // 2. DT-based logic: Calculate movement using uniform velocity * dt + // Note: Because vertex shaders run per-vertex, complex state changes are usually calculated + // using total accumulated time or updated via JS. Here we use an explicit math function: + let speedMultiplier = 1.5; + let accumulatedGuess = f32(uniforms.u_frame) * uniforms.u_dt * speedMultiplier; + + let offsetX = sin(accumulatedGuess) * 0.5; + let offsetY = cos(accumulatedGuess * 0.5) * 0.3; + + let animatedPosition = pos[vertexIndex] + vec2(offsetX, offsetY); + + output.position = vec4(animatedPosition, 0.0, 1.0); + output.color = color; + return output; +} + +@fragment +fn fs_main(@location(0) color: vec4) -> @location(0) vec4 { + return color; +} + + +## font + +Option Kernel cost Quality Effort +1-bpp bitmap (current) 1 bit test blocky, aliased, integer scale only, 96 glyphs done +8-bpp AA atlas + 1 lerp smooth glyphs, fractional positioning, proper color blend ~1 hour — same architecture, atlas becomes grayscale, mix(bg, fg, a) instead of a bit test +MSDF atlas (Valve/msdfgen) + 1 fetch + smoothstep crisp at any scale from one atlas, outlines/shadows/glow nearly free ~1 day — offline bake step, shader grows ~40 lines +CPU rasterize → upload mask (x/image/font, any TTF) same as AA perfect (hinting, kerning, Unicode), arbitrary fonts at init ~half day + a dependency; CPU cost only on text change (timecode = once/sec, nothing)