Files

T

Dmitry Sergeev f48195a8bf WGPU is coming

2026-09-07 17:18:03 +03:00

19 KiB

Raw Blame History

Архитектура v210 Генератора для Media eXchange Layer (go-mxl)

Профессиональная Zero-GC и Zero-Copy архитектура гибридного CPU/GPU v210 тестового генератора, оптимизированная под работу с официальным C SDK Media eXchange Layer (libmxl) через официальные Go-биндинги go-mxl от Qvest Digital.

🏢 Финальная структура проекта

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.

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.

package generator

/*
#cgo linux LDFLAGS: -lOpenCL
#cgo windows LDFLAGS: -lOpenCL

#include <CL/cl.h>
#include <stdlib.h>

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, чтобы избежать преобразования типов данных.

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().

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<vec2, 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<f32>(0.3, 0.6, 0.9, 1.0); // Light blue
if ((uniforms.u_frame / 60u) % 2u == 0u) {
    color = vec4<f32>(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<f32>(offsetX, offsetY);

output.position = vec4<f32>(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)

19 KiB Raw Blame History Unescape Escape