# Архитектура 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)