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 — в следующих коммитах).
68 lines
2.2 KiB
C++
68 lines
2.2 KiB
C++
// Phase 0 spike-v2 — общие типы. Расширенный относительно spike v1: добавлена
|
|
// поддержка cuda IPC event handle (для scenarios B) и pattern-fill per-row
|
|
// (для verify внутри кадра).
|
|
|
|
#pragma once
|
|
|
|
#include <cuda_runtime.h>
|
|
#include <cstdint>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <ctime>
|
|
|
|
namespace cuframes_spike_v2 {
|
|
|
|
constexpr int RING_SIZE = 2;
|
|
|
|
// Pattern: pixel [row][col] = (seq * 31 + row * 7) & 0xFF
|
|
// Использует разные значения по строкам — позволяет verify обнаружить если часть
|
|
// кадра ещё имеет старый seq.
|
|
__host__ __device__ inline uint8_t pattern_value(uint64_t seq, int row) {
|
|
return static_cast<uint8_t>((seq * 31u + row * 7u) & 0xFF);
|
|
}
|
|
|
|
struct FrameMeta {
|
|
int32_t width;
|
|
int32_t height;
|
|
int32_t pitch_y;
|
|
};
|
|
|
|
struct SlotDescriptor {
|
|
cudaIpcMemHandle_t mem_handle;
|
|
uint64_t producer_seq;
|
|
int64_t pts_ns;
|
|
};
|
|
|
|
struct SharedHeader {
|
|
uint32_t magic;
|
|
uint32_t version;
|
|
int32_t use_events; // 1 = sync mode B (events), 0 = sync mode A
|
|
cudaIpcEventHandle_t event_handle; // valid only if use_events == 1
|
|
FrameMeta meta;
|
|
SlotDescriptor slots[RING_SIZE];
|
|
uint64_t global_seq;
|
|
// Diagnostics
|
|
uint64_t torn_frame_count; // consumer записывает; producer читает для лога
|
|
};
|
|
|
|
constexpr uint32_t MAGIC = 0xCC7C2D02u;
|
|
constexpr uint32_t VERSION = 2;
|
|
|
|
#define CHECK_CUDA(call) do { \
|
|
cudaError_t _err = (call); \
|
|
if (_err != cudaSuccess) { \
|
|
fprintf(stderr, "CUDA error at %s:%d: %s\n", \
|
|
__FILE__, __LINE__, cudaGetErrorString(_err)); \
|
|
std::exit(1); \
|
|
} \
|
|
} while (0)
|
|
|
|
static inline int64_t now_ns() {
|
|
timespec ts;
|
|
clock_gettime(CLOCK_MONOTONIC, &ts);
|
|
return static_cast<int64_t>(ts.tv_sec) * 1000000000LL + ts.tv_nsec;
|
|
}
|
|
|
|
} // namespace cuframes_spike_v2
|