diff --git a/Makefile b/Makefile index 0083467..334a71b 100644 --- a/Makefile +++ b/Makefile @@ -4,19 +4,17 @@ # Harmless for binaries that don't need it. LD = -ldflags=-linkmode=internal -.PHONY: build test vet wgpu-gen clean +.PHONY: build run test clean build: - go build $(LD) ./... + go build -o ./build/mxl-gen $(LD) ./cmd/mxl-pattern + +run: + go run $(LD) ./cmd/mxl-pattern $(ARGS) test: - go test $(LD) ./... - -vet: - go vet ./... - -wgpu-gen: - go run $(LD) ./cmd/wgpu-gen + go test ./... clean: go clean + fm -f ./build/mxl-gen diff --git a/cmd/wgpu-gen/main.go b/cmd/mxl-pattern/main.go similarity index 100% rename from cmd/wgpu-gen/main.go rename to cmd/mxl-pattern/main.go diff --git a/internal/generator/opencl.go b/internal/generator/opencl.go deleted file mode 100644 index 355d423..0000000 --- a/internal/generator/opencl.go +++ /dev/null @@ -1,302 +0,0 @@ -// 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 -}