Remove excessive
This commit is contained in:
Dmitry Sergeev
@@ -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)
|
|
||||||
}
|
|
||||||
}
|
|
||||||
@@ -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
|
|
||||||
}
|
|
||||||
@@ -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;
|
|
||||||
}
|
|
||||||
@@ -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")
|
|
||||||
Reference in New Issue
Block a user