Dmitry Sergeev
486 lines
19 KiB
Markdown
486 lines
19 KiB
Markdown
# Архитектура 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 <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`, чтобы избежать преобразования типов данных.
|
||
|
||
```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<uniform> uniforms: Uniforms;
|
||
|
||
struct VertexOutput {
|
||
@builtin(position) position: vec4<f32>,
|
||
@location(0) color: vec4<f32>,
|
||
};
|
||
|
||
// We store the persistent position on the GPU using a global variable
|
||
// so we can incrementally add 'dt' to it frame over frame.
|
||
var<private> squarePosition: vec2<f32> = vec2<f32>(0.0, 0.0);
|
||
var<private> velocity: vec2<f32> = vec2<f32>(0.4, 0.2); // Moves diagonally
|
||
|
||
@vertex
|
||
fn vs_main(@builtin(vertex_index) vertexIndex: u32) -> VertexOutput {
|
||
var pos = array<vec2<f32>, 6>(
|
||
vec2<f32>(-0.2, 0.2),
|
||
vec2<f32>(-0.2, -0.2),
|
||
vec2<f32>( 0.2, 0.2),
|
||
vec2<f32>( 0.2, 0.2),
|
||
vec2<f32>(-0.2, -0.2),
|
||
vec2<f32>( 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<f32>) -> @location(0) vec4<f32> {
|
||
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)
|