gx
ad543054fc
Architectural review (2026-05-15) указал что cudaStreamSynchronize-only на producer-side не достаточен для cross-process visibility — NVIDIA Programming Guide §3.2.8 требует cudaIpcEventHandle_t. Phase 0 PoC v1 не проверял этот случай из-за cudaMemcpy который имеет implicit barriers. spike-v2 воспроизводит правильный сценарий: consumer запускает verify_kernel на ОТДЕЛЬНОМ stream'е (real-world use case — PyTorch / OpenCV CUDA), pattern включает row-based component для отлова partial-frame torn. Запуск 4 scenarios × 1500/600 frames: A-fhd60 (stream sync, FHD@60): 0 torn, p99=267µs, max=14.7ms B-fhd60 (event sync, FHD@60): 0 torn, p99=344µs, max=5.2ms A-4k30 (stream sync, 4K@30): 0 torn, p99=606µs, max=4.4ms B-4k30 (event sync, 4K@30): 0 torn, p99=437µs, max=3.7ms Все 4 показали 0 torn frames. R1 на single-host single-GPU фактически не воспроизводится — но NVIDIA contractually не гарантирует это. Decision: events as default (R1/R2 resolved). Architecture.md §6.6 закрыт. Tradeoff: mean latency +20µs, max latency в 3× ниже (predictable tail) + future-proof для multi-GPU. Также Dockerfile.dev — апдейт CUDA до 13.0.3 (12.4 не существует с devel-ubuntu24.04). Связано с PR review: R1, R2, R3 (R3, R4 — в следующих коммитах).
Phase 0 spike-v2 — kernel-on-consumer-stream torn-frame test
Архитектор (review 2026-05-15) указал на R1 / R2 в docs/architecture.md:
оригинальный PoC (tools/spike/) использовал cudaMemcpy на consumer-стороне,
который имеет implicit cross-context barriers. Это маскирует фундаментальную
проблему: cudaStreamSynchronize на producer не достаточен для гарантии
visibility GPU-writes для consumer'а на отдельном stream'е.
NVIDIA Programming Guide §3.2.8 явно говорит: «Synchronization between producer
and consumer processes is required and the responsibility of the application» —
то есть нужен cudaIpcEventHandle_t (cross-process CUDA event), не stream sync.
Что проверяем
Два сценария:
Сценарий A: stream sync only (текущий PoC v1 дизайн)
Producer:
- Fill_Y_kernel записывает
seq + rowв каждый пиксель строкиrowY-plane cudaStreamSynchronize(producer_stream)- Publish seq через atomic в SHM
Consumer:
- Read seq atomic ACQUIRE
- Verify_kernel на consumer_stream'е проверяет что Y-plane имеет ожидаемое значение во всех строках. Если есть строка с другим значением → torn.
cudaDeviceSynchronizeчтобы kernel завершился, проверить result.
Ожидаем: torn frames на high-fps (60+) на FullHD или 4K.
Сценарий B: event handle sync (предлагаемый fix)
Producer:
cudaEventCreateWithFlags(cudaEventInterprocess | cudaEventDisableTiming)cudaIpcGetEventHandle→ handle в SHM- На каждый publish:
cudaEventRecord(event, producer_stream)(вместо sync) - Publish seq
Consumer:
- На subscribe:
cudaIpcOpenEventHandle→ local event - На каждый next:
cudaStreamWaitEvent(consumer_stream, event, 0)— GPU-side wait - Verify_kernel на том же consumer_stream'е
- Result должен быть без torn frames даже на high resolution + high fps
Acceptance
- A показал torn frames > 0 → R1 confirmed → нужны events
- B показал torn frames == 0 → R2 fix validated
- Latency B vs A — изменения не должно быть значимого (events = GPU-side wait, no CPU sync)
Запуск (внутри cuframes-dev container)
cd /workspace
cmake -B build-v2 -S tools/spike-v2 -G Ninja
cmake --build build-v2
./tools/spike-v2/bench.sh