From d2c46eaf3c67ef8eea46ad5b26768feaddc5e97d Mon Sep 17 00:00:00 2001 From: enricobuehler Date: Wed, 22 Jul 2026 19:56:56 +0200 Subject: [PATCH] =?UTF-8?q?feat(encode/nvenc):=20SPIR-V=20cursor=20blend?= =?UTF-8?q?=20over=20Vulkan-allocated=20input=20slots=20=E2=80=94=20retire?= =?UTF-8?q?=20the=20PTX=20kernels?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A vendored PTX blob is JIT'd against the driver's ISA ceiling, so the cursor-blend module silently dies on drivers older than the generating toolkit (CUDA_ERROR_UNSUPPORTED_PTX_VERSION/INVALID_PTX, 222/218 — the KWin leg's invisible composite cursor on driver 595/CUDA 13.2 vs a CUDA 13.3 blob). SPIR-V has no such coupling, and the in-tree precedent already exists twice (vulkan_video's CSC blend, VkBridge's exportable OPAQUE_FD → cuImportExternalMemory bridge). New pf_zerocopy::vkslot::VkSlotBlend: the direct-SDK NVENC encoder now allocates its input ring as exportable Vulkan buffers CUDA-imports (same contiguous InputSurface layouts, pitch = row bytes rounded to 256), and the cursor composite is a compute dispatch over the cursor's rectangle (cursor_blend.comp, vendored .spv; spec-constant selects ARGB/NV12/YUV444; BT.709 limited, matching the retired .cu). The surface SSBO is uint[] with every invocation owning whole words — no 8-bit-storage device dependency. Cursor-bearing frames force the existing CPU-synced submit path so the CUDA copy → Vulkan dispatch (fence-waited) → NVENC encode ordering is CPU-established; cursorless frames keep the stream-ordered fast path untouched. Any bring-up/alloc/registration failure falls back wholesale to plain pitched CUDA surfaces (never a mixed or short ring): sessions always encode, composite mode just loses the cursor, warned once. cursor_blend.cu / cursor_blend.ptx and the CursorBlend PTX loader are deleted. Co-Authored-By: Claude Fable 5 --- .../pf-encode/src/enc/linux/cursor_blend.cu | 105 --- .../pf-encode/src/enc/linux/cursor_blend.ptx | 584 -------------- crates/pf-encode/src/enc/linux/nvenc_cuda.rs | 314 +++++--- crates/pf-zerocopy/src/imp/cuda.rs | 261 +------ crates/pf-zerocopy/src/imp/cuda/ffi.rs | 76 +- crates/pf-zerocopy/src/imp/cursor_blend.comp | 170 +++++ crates/pf-zerocopy/src/imp/cursor_blend.spv | Bin 0 -> 20412 bytes crates/pf-zerocopy/src/imp/mod.rs | 1 + crates/pf-zerocopy/src/imp/vkslot.rs | 710 ++++++++++++++++++ 9 files changed, 1088 insertions(+), 1133 deletions(-) delete mode 100644 crates/pf-encode/src/enc/linux/cursor_blend.cu delete mode 100644 crates/pf-encode/src/enc/linux/cursor_blend.ptx create mode 100644 crates/pf-zerocopy/src/imp/cursor_blend.comp create mode 100644 crates/pf-zerocopy/src/imp/cursor_blend.spv create mode 100644 crates/pf-zerocopy/src/imp/vkslot.rs diff --git a/crates/pf-encode/src/enc/linux/cursor_blend.cu b/crates/pf-encode/src/enc/linux/cursor_blend.cu deleted file mode 100644 index acd0464f..00000000 --- a/crates/pf-encode/src/enc/linux/cursor_blend.cu +++ /dev/null @@ -1,105 +0,0 @@ -// Cursor-overlay blend kernels for the CUDA/NVENC path (cursor-as-metadata). The cursor bitmap is -// straight-alpha RGBA, row-packed (stride = curW*4). Blended into the encoder-OWNED NVENC input -// surface — never the compositor's dmabuf. One thread per cursor pixel (ARGB / YUV444) or per 2x2 -// chroma block (NV12). Coefficients are BT.709 limited, matching rgb2yuv.comp so the cursor colour -// matches the rest of the frame regardless of which zero-copy backend encodes it. -// -// Build (regenerate cursor_blend.ptx after editing): -// nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx -// PTX is JIT'd by the driver forward to the actual GPU, so a compute_75 (Turing) baseline runs on -// every Turing-or-newer NVENC GPU. (CUDA 13's nvcc no longer targets pre-Turing archs.) - -typedef unsigned char u8; - -__device__ __forceinline__ u8 blend8(int dst, int src, int a) { - return (u8)((src * a + dst * (255 - a)) / 255); -} - -// Packed 4-byte surface. NVENC's ARGB format stores bytes B,G,R,A in memory; the cursor is R,G,B,A. -extern "C" __global__ void blend_argb( - u8* surf, int pitch, int surfW, int surfH, - const u8* cur, int curW, int curH, int ox, int oy) -{ - int cx = blockIdx.x * blockDim.x + threadIdx.x; - int cy = blockIdx.y * blockDim.y + threadIdx.y; - if (cx >= curW || cy >= curH) return; - int px = ox + cx, py = oy + cy; - if (px < 0 || py < 0 || px >= surfW || py >= surfH) return; - const u8* s = cur + (size_t)(cy * curW + cx) * 4; - int a = s[3]; - if (a == 0) return; - u8* d = surf + (size_t)py * pitch + (size_t)px * 4; - d[0] = blend8(d[0], s[2], a); // B <- cursor B - d[1] = blend8(d[1], s[1], a); // G <- cursor G - d[2] = blend8(d[2], s[0], a); // R <- cursor R -} - -// Planar YUV444: three full-res planes stacked at base, base+plane, base+2*plane (plane=pitch*surfH). -extern "C" __global__ void blend_yuv444( - u8* base, int pitch, int surfW, int surfH, - const u8* cur, int curW, int curH, int ox, int oy) -{ - int cx = blockIdx.x * blockDim.x + threadIdx.x; - int cy = blockIdx.y * blockDim.y + threadIdx.y; - if (cx >= curW || cy >= curH) return; - int px = ox + cx, py = oy + cy; - if (px < 0 || py < 0 || px >= surfW || py >= surfH) return; - const u8* s = cur + (size_t)(cy * curW + cx) * 4; - int a = s[3]; - if (a == 0) return; - float R = s[0], G = s[1], B = s[2]; - int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f); - int U = (int)(128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B + 0.5f); - int V = (int)(128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B + 0.5f); - size_t plane = (size_t)pitch * surfH; - u8* yp = base + (size_t)py * pitch + px; - u8* up = base + plane + (size_t)py * pitch + px; - u8* vp = base + 2 * plane + (size_t)py * pitch + px; - *yp = blend8(*yp, Y, a); - *up = blend8(*up, U, a); - *vp = blend8(*vp, V, a); -} - -// NV12: full-res Y plane + interleaved half-res UV plane. One thread per 2x2 luma block; each blends -// up to four Y samples and one (alpha-weighted) UV sample. -extern "C" __global__ void blend_nv12( - u8* yb, int yPitch, u8* uvb, int uvPitch, int surfW, int surfH, - const u8* cur, int curW, int curH, int ox, int oy) -{ - int bx = blockIdx.x * blockDim.x + threadIdx.x; - int by = blockIdx.y * blockDim.y + threadIdx.y; - int base_cx = bx * 2, base_cy = by * 2; - if (base_cx >= curW || base_cy >= curH) return; - float ua = 0.0f, va = 0.0f, wa = 0.0f; - int cnt = 0; - for (int j = 0; j < 2; j++) { - for (int i = 0; i < 2; i++) { - int cx = base_cx + i, cy = base_cy + j; - if (cx >= curW || cy >= curH) continue; - int px = ox + cx, py = oy + cy; - if (px < 0 || py < 0 || px >= surfW || py >= surfH) continue; - const u8* s = cur + (size_t)(cy * curW + cx) * 4; - int a = s[3]; - if (a == 0) continue; - float R = s[0], G = s[1], B = s[2]; - int Y = (int)(16.0f + 0.1826f * R + 0.6142f * G + 0.0620f * B + 0.5f); - u8* yp = yb + (size_t)py * yPitch + px; - *yp = blend8(*yp, Y, a); - ua += (128.0f - 0.1006f * R - 0.3386f * G + 0.4392f * B) * a; - va += (128.0f + 0.4392f * R - 0.3989f * G - 0.0403f * B) * a; - wa += a; - cnt++; - } - } - if (wa <= 0.0f || cnt == 0) return; - // The chroma sample covering this block's top-left surface pixel. - int uvx = (ox + base_cx) / 2; - int uvy = (oy + base_cy) / 2; - if (uvx < 0 || uvy < 0 || uvx * 2 >= surfW || uvy * 2 >= surfH) return; - int U = (int)(ua / wa + 0.5f); - int V = (int)(va / wa + 0.5f); - int amean = (int)(wa / cnt + 0.5f); - u8* uv = uvb + (size_t)uvy * uvPitch + (size_t)uvx * 2; - uv[0] = blend8(uv[0], U, amean); - uv[1] = blend8(uv[1], V, amean); -} diff --git a/crates/pf-encode/src/enc/linux/cursor_blend.ptx b/crates/pf-encode/src/enc/linux/cursor_blend.ptx deleted file mode 100644 index 1c218508..00000000 --- a/crates/pf-encode/src/enc/linux/cursor_blend.ptx +++ /dev/null @@ -1,584 +0,0 @@ -// -// Generated by NVIDIA NVVM Compiler -// -// Compiler Build ID: CL-38244171 -// Cuda compilation tools, release 13.3, V13.3.73 -// Based on NVVM 7.0.1 -// - -// .version lowered from the generating toolkit's 9.3 (CUDA 13.3) by hand: a driver refuses PTX -// whose ISA version is newer than its JIT (CUDA_ERROR_UNSUPPORTED_PTX_VERSION, 222 — driver -// 595.58 = CUDA 13.2 rejects a 9.3 stamp), while a stamp OLDER than the body's actual syntax is -// rejected as CUDA_ERROR_INVALID_PTX (218 — 8.0 was tried and failed against this body). 9.2 is -// the highest ISA a CUDA-13.2 driver JITs and the lowest this 13.3-emitted body accepts. -// TODO(portability): drivers older than the 13.2 era will refuse 9.2 too — the real fix is -// regenerating cursor_blend.cu with an older toolkit (CUDA 12.x) so the body itself is old, then -// stamping that toolkit's native version. Re-test on the oldest driver box after any regen. -.version 9.2 -.target sm_75 -.address_size 64 - - // .globl blend_argb - -.visible .entry blend_argb( - .param .u64 blend_argb_param_0, - .param .u32 blend_argb_param_1, - .param .u32 blend_argb_param_2, - .param .u32 blend_argb_param_3, - .param .u64 blend_argb_param_4, - .param .u32 blend_argb_param_5, - .param .u32 blend_argb_param_6, - .param .u32 blend_argb_param_7, - .param .u32 blend_argb_param_8 -) -{ - .reg .pred %p<10>; - .reg .b16 %rs<2>; - .reg .b32 %r<34>; - .reg .b64 %rd<17>; - - - ld.param.u64 %rd2, [blend_argb_param_0]; - ld.param.u32 %r5, [blend_argb_param_1]; - ld.param.u32 %r6, [blend_argb_param_2]; - ld.param.u32 %r7, [blend_argb_param_3]; - ld.param.u64 %rd3, [blend_argb_param_4]; - ld.param.u32 %r8, [blend_argb_param_5]; - ld.param.u32 %r11, [blend_argb_param_6]; - ld.param.u32 %r9, [blend_argb_param_7]; - ld.param.u32 %r10, [blend_argb_param_8]; - mov.u32 %r12, %ntid.x; - mov.u32 %r13, %ctaid.x; - mov.u32 %r14, %tid.x; - mad.lo.s32 %r1, %r13, %r12, %r14; - mov.u32 %r15, %ntid.y; - mov.u32 %r16, %ctaid.y; - mov.u32 %r17, %tid.y; - mad.lo.s32 %r2, %r16, %r15, %r17; - setp.ge.s32 %p1, %r1, %r8; - setp.ge.s32 %p2, %r2, %r11; - or.pred %p3, %p1, %p2; - @%p3 bra $L__BB0_4; - - add.s32 %r3, %r1, %r9; - add.s32 %r4, %r2, %r10; - or.b32 %r18, %r4, %r3; - setp.lt.s32 %p4, %r18, 0; - setp.ge.s32 %p5, %r3, %r6; - or.pred %p6, %p5, %p4; - setp.ge.s32 %p7, %r4, %r7; - or.pred %p8, %p7, %p6; - @%p8 bra $L__BB0_4; - - mad.lo.s32 %r19, %r2, %r8, %r1; - mul.wide.s32 %rd4, %r19, 4; - cvta.to.global.u64 %rd5, %rd3; - add.s64 %rd1, %rd5, %rd4; - ld.global.u8 %rs1, [%rd1+3]; - setp.eq.s16 %p9, %rs1, 0; - @%p9 bra $L__BB0_4; - - cvt.u32.u16 %r20, %rs1; - mul.wide.s32 %rd6, %r4, %r5; - mul.wide.s32 %rd7, %r3, 4; - add.s64 %rd8, %rd6, %rd7; - cvta.to.global.u64 %rd9, %rd2; - add.s64 %rd10, %rd9, %rd8; - ld.global.u8 %r21, [%rd10]; - ld.global.u8 %r22, [%rd1+2]; - xor.b32 %r23, %r20, 255; - mul.lo.s32 %r24, %r23, %r21; - mad.lo.s32 %r25, %r22, %r20, %r24; - mul.wide.u32 %rd11, %r25, -2139062143; - shr.u64 %rd12, %rd11, 39; - st.global.u8 [%rd10], %rd12; - ld.global.u8 %r26, [%rd10+1]; - ld.global.u8 %r27, [%rd1+1]; - mul.lo.s32 %r28, %r23, %r26; - mad.lo.s32 %r29, %r27, %r20, %r28; - mul.wide.u32 %rd13, %r29, -2139062143; - shr.u64 %rd14, %rd13, 39; - st.global.u8 [%rd10+1], %rd14; - ld.global.u8 %r30, [%rd10+2]; - ld.global.u8 %r31, [%rd1]; - mul.lo.s32 %r32, %r23, %r30; - mad.lo.s32 %r33, %r31, %r20, %r32; - mul.wide.u32 %rd15, %r33, -2139062143; - shr.u64 %rd16, %rd15, 39; - st.global.u8 [%rd10+2], %rd16; - -$L__BB0_4: - ret; - -} - // .globl blend_yuv444 -.visible .entry blend_yuv444( - .param .u64 blend_yuv444_param_0, - .param .u32 blend_yuv444_param_1, - .param .u32 blend_yuv444_param_2, - .param .u32 blend_yuv444_param_3, - .param .u64 blend_yuv444_param_4, - .param .u32 blend_yuv444_param_5, - .param .u32 blend_yuv444_param_6, - .param .u32 blend_yuv444_param_7, - .param .u32 blend_yuv444_param_8 -) -{ - .reg .pred %p<10>; - .reg .b16 %rs<5>; - .reg .f32 %f<16>; - .reg .b32 %r<49>; - .reg .b64 %rd<14>; - - - ld.param.u64 %rd2, [blend_yuv444_param_0]; - ld.param.u32 %r5, [blend_yuv444_param_1]; - ld.param.u32 %r6, [blend_yuv444_param_2]; - ld.param.u32 %r7, [blend_yuv444_param_3]; - ld.param.u64 %rd3, [blend_yuv444_param_4]; - ld.param.u32 %r8, [blend_yuv444_param_5]; - ld.param.u32 %r11, [blend_yuv444_param_6]; - ld.param.u32 %r9, [blend_yuv444_param_7]; - ld.param.u32 %r10, [blend_yuv444_param_8]; - mov.u32 %r12, %ntid.x; - mov.u32 %r13, %ctaid.x; - mov.u32 %r14, %tid.x; - mad.lo.s32 %r1, %r13, %r12, %r14; - mov.u32 %r15, %ntid.y; - mov.u32 %r16, %ctaid.y; - mov.u32 %r17, %tid.y; - mad.lo.s32 %r2, %r16, %r15, %r17; - setp.ge.s32 %p1, %r1, %r8; - setp.ge.s32 %p2, %r2, %r11; - or.pred %p3, %p1, %p2; - @%p3 bra $L__BB1_4; - - add.s32 %r3, %r1, %r9; - add.s32 %r4, %r2, %r10; - or.b32 %r18, %r4, %r3; - setp.lt.s32 %p4, %r18, 0; - setp.ge.s32 %p5, %r3, %r6; - or.pred %p6, %p5, %p4; - setp.ge.s32 %p7, %r4, %r7; - or.pred %p8, %p7, %p6; - @%p8 bra $L__BB1_4; - - mad.lo.s32 %r19, %r2, %r8, %r1; - mul.wide.s32 %rd4, %r19, 4; - cvta.to.global.u64 %rd5, %rd3; - add.s64 %rd1, %rd5, %rd4; - ld.global.u8 %rs1, [%rd1+3]; - setp.eq.s16 %p9, %rs1, 0; - @%p9 bra $L__BB1_4; - - cvt.u32.u16 %r20, %rs1; - ld.global.u8 %rs2, [%rd1]; - cvt.rn.f32.u16 %f1, %rs2; - ld.global.u8 %rs3, [%rd1+1]; - cvt.rn.f32.u16 %f2, %rs3; - ld.global.u8 %rs4, [%rd1+2]; - cvt.rn.f32.u16 %f3, %rs4; - fma.rn.f32 %f4, %f1, 0f3E3AFB7F, 0f41800000; - fma.rn.f32 %f5, %f2, 0f3F1D3C36, %f4; - fma.rn.f32 %f6, %f3, 0f3D7DF3B6, %f5; - add.f32 %f7, %f6, 0f3F000000; - cvt.rzi.s32.f32 %r21, %f7; - fma.rn.f32 %f8, %f1, 0fBDCE075F, 0f43000000; - fma.rn.f32 %f9, %f2, 0fBEAD5CFB, %f8; - fma.rn.f32 %f10, %f3, 0f3EE0DED3, %f9; - add.f32 %f11, %f10, 0f3F000000; - cvt.rzi.s32.f32 %r22, %f11; - fma.rn.f32 %f12, %f1, 0f3EE0DED3, 0f43000000; - fma.rn.f32 %f13, %f2, 0fBECC3C9F, %f12; - fma.rn.f32 %f14, %f3, 0fBD25119D, %f13; - add.f32 %f15, %f14, 0f3F000000; - cvt.rzi.s32.f32 %r23, %f15; - mul.wide.s32 %rd6, %r4, %r5; - cvt.s64.s32 %rd7, %r3; - add.s64 %rd8, %rd6, %rd7; - cvta.to.global.u64 %rd9, %rd2; - add.s64 %rd10, %rd9, %rd8; - ld.global.u8 %r24, [%rd10]; - mul.lo.s32 %r25, %r21, %r20; - xor.b32 %r26, %r20, 255; - mad.lo.s32 %r27, %r26, %r24, %r25; - mul.hi.s32 %r28, %r27, -2139062143; - add.s32 %r29, %r28, %r27; - shr.u32 %r30, %r29, 31; - shr.u32 %r31, %r29, 7; - add.s32 %r32, %r31, %r30; - st.global.u8 [%rd10], %r32; - mul.wide.s32 %rd11, %r7, %r5; - add.s64 %rd12, %rd10, %rd11; - ld.global.u8 %r33, [%rd12]; - mul.lo.s32 %r34, %r22, %r20; - mad.lo.s32 %r35, %r26, %r33, %r34; - mul.hi.s32 %r36, %r35, -2139062143; - add.s32 %r37, %r36, %r35; - shr.u32 %r38, %r37, 31; - shr.u32 %r39, %r37, 7; - add.s32 %r40, %r39, %r38; - st.global.u8 [%rd12], %r40; - add.s64 %rd13, %rd12, %rd11; - ld.global.u8 %r41, [%rd13]; - mul.lo.s32 %r42, %r23, %r20; - mad.lo.s32 %r43, %r26, %r41, %r42; - mul.hi.s32 %r44, %r43, -2139062143; - add.s32 %r45, %r44, %r43; - shr.u32 %r46, %r45, 31; - shr.u32 %r47, %r45, 7; - add.s32 %r48, %r47, %r46; - st.global.u8 [%rd13], %r48; - -$L__BB1_4: - ret; - -} - // .globl blend_nv12 -.visible .entry blend_nv12( - .param .u64 blend_nv12_param_0, - .param .u32 blend_nv12_param_1, - .param .u64 blend_nv12_param_2, - .param .u32 blend_nv12_param_3, - .param .u32 blend_nv12_param_4, - .param .u32 blend_nv12_param_5, - .param .u64 blend_nv12_param_6, - .param .u32 blend_nv12_param_7, - .param .u32 blend_nv12_param_8, - .param .u32 blend_nv12_param_9, - .param .u32 blend_nv12_param_10 -) -{ - .reg .pred %p<43>; - .reg .b16 %rs<17>; - .reg .f32 %f<108>; - .reg .b32 %r<123>; - .reg .b64 %rd<35>; - - - ld.param.u64 %rd11, [blend_nv12_param_0]; - ld.param.u32 %r21, [blend_nv12_param_1]; - ld.param.u64 %rd10, [blend_nv12_param_2]; - ld.param.u32 %r22, [blend_nv12_param_3]; - ld.param.u32 %r23, [blend_nv12_param_4]; - ld.param.u32 %r24, [blend_nv12_param_5]; - ld.param.u64 %rd12, [blend_nv12_param_6]; - ld.param.u32 %r25, [blend_nv12_param_7]; - ld.param.u32 %r26, [blend_nv12_param_8]; - ld.param.u32 %r27, [blend_nv12_param_9]; - ld.param.u32 %r28, [blend_nv12_param_10]; - cvta.to.global.u64 %rd1, %rd11; - cvta.to.global.u64 %rd2, %rd12; - mov.u32 %r29, %ntid.x; - mov.u32 %r30, %ctaid.x; - mov.u32 %r31, %tid.x; - mad.lo.s32 %r32, %r30, %r29, %r31; - mov.u32 %r33, %ntid.y; - mov.u32 %r34, %ctaid.y; - mov.u32 %r35, %tid.y; - mad.lo.s32 %r36, %r34, %r33, %r35; - shl.b32 %r1, %r32, 1; - shl.b32 %r2, %r36, 1; - setp.ge.s32 %p1, %r1, %r25; - setp.ge.s32 %p2, %r2, %r26; - or.pred %p3, %p1, %p2; - mov.f32 %f102, 0f00000000; - mov.f32 %f103, 0f00000000; - mov.f32 %f104, 0f00000000; - @%p3 bra $L__BB2_19; - - cvt.s64.s32 %rd3, %r21; - add.s32 %r3, %r2, %r28; - setp.ge.s32 %p4, %r3, %r24; - mul.lo.s32 %r4, %r2, %r25; - mul.wide.s32 %rd4, %r3, %r21; - add.s32 %r5, %r1, %r27; - or.b32 %r38, %r5, %r3; - setp.lt.s32 %p5, %r38, 0; - mov.u32 %r121, 0; - setp.ge.s32 %p6, %r5, %r23; - or.pred %p7, %p6, %p5; - or.pred %p8, %p4, %p7; - @%p8 bra $L__BB2_4; - - add.s32 %r40, %r1, %r4; - mul.wide.s32 %rd13, %r40, 4; - add.s64 %rd5, %rd2, %rd13; - ld.global.u8 %rs1, [%rd5+3]; - setp.eq.s16 %p9, %rs1, 0; - @%p9 bra $L__BB2_4; - - cvt.u32.u16 %r42, %rs1; - ld.global.u8 %rs5, [%rd5]; - cvt.rn.f32.u16 %f31, %rs5; - ld.global.u8 %rs6, [%rd5+1]; - cvt.rn.f32.u16 %f32, %rs6; - ld.global.u8 %rs7, [%rd5+2]; - cvt.rn.f32.u16 %f33, %rs7; - fma.rn.f32 %f34, %f31, 0f3E3AFB7F, 0f41800000; - fma.rn.f32 %f35, %f32, 0f3F1D3C36, %f34; - fma.rn.f32 %f36, %f33, 0f3D7DF3B6, %f35; - add.f32 %f37, %f36, 0f3F000000; - cvt.rzi.s32.f32 %r43, %f37; - cvt.s64.s32 %rd14, %r5; - add.s64 %rd15, %rd4, %rd14; - add.s64 %rd16, %rd1, %rd15; - ld.global.u8 %r44, [%rd16]; - mul.lo.s32 %r45, %r43, %r42; - xor.b32 %r46, %r42, 255; - mad.lo.s32 %r47, %r46, %r44, %r45; - mul.hi.s32 %r48, %r47, -2139062143; - add.s32 %r49, %r48, %r47; - shr.u32 %r50, %r49, 31; - shr.u32 %r51, %r49, 7; - add.s32 %r52, %r51, %r50; - st.global.u8 [%rd16], %r52; - fma.rn.f32 %f38, %f31, 0fBDCE075F, 0f43000000; - fma.rn.f32 %f39, %f32, 0fBEAD5CFB, %f38; - fma.rn.f32 %f40, %f33, 0f3EE0DED3, %f39; - cvt.rn.f32.u16 %f104, %rs1; - fma.rn.f32 %f102, %f40, %f104, 0f00000000; - fma.rn.f32 %f41, %f31, 0f3EE0DED3, 0f43000000; - fma.rn.f32 %f42, %f32, 0fBECC3C9F, %f41; - fma.rn.f32 %f43, %f33, 0fBD25119D, %f42; - fma.rn.f32 %f103, %f43, %f104, 0f00000000; - mov.u32 %r121, 1; - -$L__BB2_4: - add.s32 %r7, %r1, 1; - setp.ge.s32 %p10, %r7, %r25; - @%p10 bra $L__BB2_8; - - add.s32 %r8, %r7, %r27; - or.b32 %r53, %r8, %r3; - setp.lt.s32 %p12, %r53, 0; - setp.ge.s32 %p13, %r8, %r23; - or.pred %p14, %p13, %p12; - or.pred %p15, %p4, %p14; - @%p15 bra $L__BB2_8; - - add.s32 %r54, %r7, %r4; - mul.wide.s32 %rd17, %r54, 4; - add.s64 %rd6, %rd2, %rd17; - ld.global.u8 %rs2, [%rd6+3]; - setp.eq.s16 %p16, %rs2, 0; - @%p16 bra $L__BB2_8; - - cvt.u32.u16 %r55, %rs2; - ld.global.u8 %rs8, [%rd6]; - cvt.rn.f32.u16 %f44, %rs8; - ld.global.u8 %rs9, [%rd6+1]; - cvt.rn.f32.u16 %f45, %rs9; - ld.global.u8 %rs10, [%rd6+2]; - cvt.rn.f32.u16 %f46, %rs10; - fma.rn.f32 %f47, %f44, 0f3E3AFB7F, 0f41800000; - fma.rn.f32 %f48, %f45, 0f3F1D3C36, %f47; - fma.rn.f32 %f49, %f46, 0f3D7DF3B6, %f48; - add.f32 %f50, %f49, 0f3F000000; - cvt.rzi.s32.f32 %r56, %f50; - cvt.s64.s32 %rd18, %r8; - add.s64 %rd19, %rd4, %rd18; - add.s64 %rd20, %rd1, %rd19; - ld.global.u8 %r57, [%rd20]; - mul.lo.s32 %r58, %r56, %r55; - xor.b32 %r59, %r55, 255; - mad.lo.s32 %r60, %r59, %r57, %r58; - mul.hi.s32 %r61, %r60, -2139062143; - add.s32 %r62, %r61, %r60; - shr.u32 %r63, %r62, 31; - shr.u32 %r64, %r62, 7; - add.s32 %r65, %r64, %r63; - st.global.u8 [%rd20], %r65; - fma.rn.f32 %f51, %f44, 0fBDCE075F, 0f43000000; - fma.rn.f32 %f52, %f45, 0fBEAD5CFB, %f51; - fma.rn.f32 %f53, %f46, 0f3EE0DED3, %f52; - cvt.rn.f32.u16 %f54, %rs2; - fma.rn.f32 %f102, %f53, %f54, %f102; - fma.rn.f32 %f55, %f44, 0f3EE0DED3, 0f43000000; - fma.rn.f32 %f56, %f45, 0fBECC3C9F, %f55; - fma.rn.f32 %f57, %f46, 0fBD25119D, %f56; - fma.rn.f32 %f103, %f57, %f54, %f103; - add.f32 %f104, %f104, %f54; - add.s32 %r121, %r121, 1; - -$L__BB2_8: - add.s32 %r11, %r2, 1; - setp.ge.s32 %p17, %r11, %r26; - add.s32 %r12, %r11, %r28; - add.s32 %r13, %r4, %r25; - cvt.s64.s32 %rd21, %r12; - mul.lo.s64 %rd7, %rd21, %rd3; - @%p17 bra $L__BB2_12; - - setp.ge.s32 %p18, %r12, %r24; - or.b32 %r66, %r5, %r12; - setp.lt.s32 %p19, %r66, 0; - or.pred %p21, %p6, %p19; - or.pred %p22, %p18, %p21; - @%p22 bra $L__BB2_12; - - add.s32 %r67, %r1, %r13; - mul.wide.s32 %rd22, %r67, 4; - add.s64 %rd8, %rd2, %rd22; - ld.global.u8 %rs3, [%rd8+3]; - setp.eq.s16 %p23, %rs3, 0; - @%p23 bra $L__BB2_12; - - cvt.u32.u16 %r68, %rs3; - ld.global.u8 %rs11, [%rd8]; - cvt.rn.f32.u16 %f58, %rs11; - ld.global.u8 %rs12, [%rd8+1]; - cvt.rn.f32.u16 %f59, %rs12; - ld.global.u8 %rs13, [%rd8+2]; - cvt.rn.f32.u16 %f60, %rs13; - fma.rn.f32 %f61, %f58, 0f3E3AFB7F, 0f41800000; - fma.rn.f32 %f62, %f59, 0f3F1D3C36, %f61; - fma.rn.f32 %f63, %f60, 0f3D7DF3B6, %f62; - add.f32 %f64, %f63, 0f3F000000; - cvt.rzi.s32.f32 %r69, %f64; - cvt.s64.s32 %rd23, %r5; - add.s64 %rd24, %rd7, %rd23; - add.s64 %rd25, %rd1, %rd24; - ld.global.u8 %r70, [%rd25]; - mul.lo.s32 %r71, %r69, %r68; - xor.b32 %r72, %r68, 255; - mad.lo.s32 %r73, %r72, %r70, %r71; - mul.hi.s32 %r74, %r73, -2139062143; - add.s32 %r75, %r74, %r73; - shr.u32 %r76, %r75, 31; - shr.u32 %r77, %r75, 7; - add.s32 %r78, %r77, %r76; - st.global.u8 [%rd25], %r78; - fma.rn.f32 %f65, %f58, 0fBDCE075F, 0f43000000; - fma.rn.f32 %f66, %f59, 0fBEAD5CFB, %f65; - fma.rn.f32 %f67, %f60, 0f3EE0DED3, %f66; - cvt.rn.f32.u16 %f68, %rs3; - fma.rn.f32 %f102, %f67, %f68, %f102; - fma.rn.f32 %f69, %f58, 0f3EE0DED3, 0f43000000; - fma.rn.f32 %f70, %f59, 0fBECC3C9F, %f69; - fma.rn.f32 %f71, %f60, 0fBD25119D, %f70; - fma.rn.f32 %f103, %f71, %f68, %f103; - add.f32 %f104, %f104, %f68; - add.s32 %r121, %r121, 1; - -$L__BB2_12: - or.pred %p26, %p17, %p10; - @%p26 bra $L__BB2_16; - - setp.ge.s32 %p27, %r12, %r24; - add.s32 %r16, %r7, %r27; - or.b32 %r79, %r16, %r12; - setp.lt.s32 %p28, %r79, 0; - setp.ge.s32 %p29, %r16, %r23; - or.pred %p30, %p29, %p28; - or.pred %p31, %p27, %p30; - @%p31 bra $L__BB2_16; - - add.s32 %r80, %r7, %r13; - mul.wide.s32 %rd26, %r80, 4; - add.s64 %rd9, %rd2, %rd26; - ld.global.u8 %rs4, [%rd9+3]; - setp.eq.s16 %p32, %rs4, 0; - @%p32 bra $L__BB2_16; - - cvt.u32.u16 %r81, %rs4; - ld.global.u8 %rs14, [%rd9]; - cvt.rn.f32.u16 %f72, %rs14; - ld.global.u8 %rs15, [%rd9+1]; - cvt.rn.f32.u16 %f73, %rs15; - ld.global.u8 %rs16, [%rd9+2]; - cvt.rn.f32.u16 %f74, %rs16; - fma.rn.f32 %f75, %f72, 0f3E3AFB7F, 0f41800000; - fma.rn.f32 %f76, %f73, 0f3F1D3C36, %f75; - fma.rn.f32 %f77, %f74, 0f3D7DF3B6, %f76; - add.f32 %f78, %f77, 0f3F000000; - cvt.rzi.s32.f32 %r82, %f78; - cvt.s64.s32 %rd27, %r16; - add.s64 %rd28, %rd7, %rd27; - add.s64 %rd29, %rd1, %rd28; - ld.global.u8 %r83, [%rd29]; - mul.lo.s32 %r84, %r82, %r81; - xor.b32 %r85, %r81, 255; - mad.lo.s32 %r86, %r85, %r83, %r84; - mul.hi.s32 %r87, %r86, -2139062143; - add.s32 %r88, %r87, %r86; - shr.u32 %r89, %r88, 31; - shr.u32 %r90, %r88, 7; - add.s32 %r91, %r90, %r89; - st.global.u8 [%rd29], %r91; - fma.rn.f32 %f79, %f72, 0fBDCE075F, 0f43000000; - fma.rn.f32 %f80, %f73, 0fBEAD5CFB, %f79; - fma.rn.f32 %f81, %f74, 0f3EE0DED3, %f80; - cvt.rn.f32.u16 %f82, %rs4; - fma.rn.f32 %f102, %f81, %f82, %f102; - fma.rn.f32 %f83, %f72, 0f3EE0DED3, 0f43000000; - fma.rn.f32 %f84, %f73, 0fBECC3C9F, %f83; - fma.rn.f32 %f85, %f74, 0fBD25119D, %f84; - fma.rn.f32 %f103, %f85, %f82, %f103; - add.f32 %f104, %f104, %f82; - add.s32 %r121, %r121, 1; - -$L__BB2_16: - setp.eq.s32 %p33, %r121, 0; - setp.le.f32 %p34, %f104, 0f00000000; - or.pred %p35, %p34, %p33; - @%p35 bra $L__BB2_19; - - shr.u32 %r92, %r5, 31; - add.s32 %r93, %r5, %r92; - shr.s32 %r19, %r93, 1; - setp.lt.s32 %p36, %r3, -1; - setp.lt.s32 %p37, %r5, -1; - or.pred %p38, %p37, %p36; - and.b32 %r94, %r93, -2; - setp.ge.s32 %p39, %r94, %r23; - or.pred %p40, %p38, %p39; - shr.u32 %r95, %r3, 31; - add.s32 %r96, %r3, %r95; - shr.s32 %r20, %r96, 1; - and.b32 %r97, %r96, -2; - setp.ge.s32 %p41, %r97, %r24; - or.pred %p42, %p40, %p41; - @%p42 bra $L__BB2_19; - - div.rn.f32 %f86, %f102, %f104; - add.f32 %f87, %f86, 0f3F000000; - cvt.rzi.s32.f32 %r98, %f87; - div.rn.f32 %f88, %f103, %f104; - add.f32 %f89, %f88, 0f3F000000; - cvt.rzi.s32.f32 %r99, %f89; - cvt.rn.f32.s32 %f90, %r121; - div.rn.f32 %f91, %f104, %f90; - add.f32 %f92, %f91, 0f3F000000; - cvt.rzi.s32.f32 %r100, %f92; - mul.wide.s32 %rd30, %r20, %r22; - mul.wide.s32 %rd31, %r19, 2; - add.s64 %rd32, %rd30, %rd31; - cvta.to.global.u64 %rd33, %rd10; - add.s64 %rd34, %rd33, %rd32; - ld.global.u8 %r101, [%rd34]; - mul.lo.s32 %r102, %r100, %r98; - mov.u32 %r103, 255; - sub.s32 %r104, %r103, %r100; - mad.lo.s32 %r105, %r104, %r101, %r102; - mul.hi.s32 %r106, %r105, -2139062143; - add.s32 %r107, %r106, %r105; - shr.u32 %r108, %r107, 31; - shr.u32 %r109, %r107, 7; - add.s32 %r110, %r109, %r108; - st.global.u8 [%rd34], %r110; - ld.global.u8 %r111, [%rd34+1]; - mul.lo.s32 %r112, %r100, %r99; - mad.lo.s32 %r113, %r104, %r111, %r112; - mul.hi.s32 %r114, %r113, -2139062143; - add.s32 %r115, %r114, %r113; - shr.u32 %r116, %r115, 31; - shr.u32 %r117, %r115, 7; - add.s32 %r118, %r117, %r116; - st.global.u8 [%rd34+1], %r118; - -$L__BB2_19: - ret; - -} - diff --git a/crates/pf-encode/src/enc/linux/nvenc_cuda.rs b/crates/pf-encode/src/enc/linux/nvenc_cuda.rs index 38bbfed1..119a0ee2 100644 --- a/crates/pf-encode/src/enc/linux/nvenc_cuda.rs +++ b/crates/pf-encode/src/enc/linux/nvenc_cuda.rs @@ -69,6 +69,7 @@ use super::{AuChunk, ChromaFormat, Codec, EncodedFrame, Encoder, EncoderCaps}; use anyhow::{anyhow, bail, Context, Result}; use pf_frame::{CapturedFrame, FramePayload}; use pf_zerocopy::cuda::{self, InputSurface}; +use pf_zerocopy::vkslot::{SlotFormat, VkSlotBlend, VkSlotRef}; use std::collections::VecDeque; use std::ffi::c_void; use std::ptr; @@ -76,16 +77,6 @@ use std::sync::mpsc; use nvidia_video_codec_sdk::sys::nvEncodeAPI as nv; -/// Prebuilt PTX for the cursor-overlay blend kernels (cursor-as-metadata). Source is -/// `cursor_blend.cu` beside this file; regenerate with -/// `nvcc -ptx -arch=compute_75 cursor_blend.cu -o cursor_blend.ptx` after editing. JIT'd by the -/// driver, so it runs on any Turing-or-newer GPU. ⚠️ The `.version` stamp is load-bearing (see -/// the comment in the .ptx): a driver refuses PTX with an ISA newer than its JIT (error 222), -/// and a stamp older than the body's real syntax is INVALID_PTX (218) — so a new toolkit's -/// output silently kills cursor compositing on older-driver boxes. Prefer regenerating with the -/// OLDEST toolchain that compiles the .cu, and re-test on the oldest driver box. -const CURSOR_PTX: &[u8] = include_bytes!("cursor_blend.ptx"); - // --------------------------------------------------------------------------------------------- // Runtime-loaded NVENC entry table (Linux). Same shape as the Windows backend's `EncodeApi`, minus // the async-event entry points (Windows-only). Resolved once from `libnvidia-encode.so.1` — the two @@ -432,13 +423,54 @@ fn buffer_format(buf: &cuda::DeviceBuffer) -> nv::NV_ENC_BUFFER_FORMAT { } } -/// One encoder-owned CUDA input surface + its NVENC registration. The surface is copied into each +/// One encoder-owned input surface + its NVENC registration. The surface is copied into each /// use (device→device) and the registration is created once at session init, unregistered at teardown. struct RingSlot { - surface: InputSurface, + surface: SlotSurface, reg: nv::NV_ENC_REGISTERED_PTR, } +/// The ring slot's backing allocation: Vulkan external memory CUDA-imported (the normal case — +/// blendable by the SPIR-V cursor pass, see `vkslot.rs`) or a plain pitched CUDA allocation (the +/// fallback when Vulkan bring-up fails: sessions still encode, composite mode just has no +/// cursor). Both present the same `(ptr, pitch, height)` NVENC-registration vocabulary. +enum SlotSurface { + Cuda(InputSurface), + /// Backing objects live in the encoder's [`VkSlotBlend`] (freed by its `free_slots`); the + /// ref itself is Copy and carries the registered geometry. + Vk(VkSlotRef), +} + +/// The [`SlotFormat`] for an NVENC buffer format (the ring-build + blend vocabulary). +fn slot_fmt_of(fmt: nv::NV_ENC_BUFFER_FORMAT) -> SlotFormat { + match fmt { + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => SlotFormat::Yuv444, + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => SlotFormat::Nv12, + _ => SlotFormat::Argb, + } +} + +impl SlotSurface { + fn ptr(&self) -> pf_zerocopy::cuda::CUdeviceptr { + match self { + SlotSurface::Cuda(s) => s.ptr, + SlotSurface::Vk(r) => r.ptr, + } + } + fn pitch(&self) -> usize { + match self { + SlotSurface::Cuda(s) => s.pitch, + SlotSurface::Vk(r) => r.pitch, + } + } + fn height(&self) -> u32 { + match self { + SlotSurface::Cuda(s) => s.height, + SlotSurface::Vk(r) => r.height, + } + } +} + /// `doNotWait` sampling cadence inside [`Encoder::poll_chunk`] — the probe measured ~200 µs /// between slice completions on the 5070 Ti, so 50 µs keeps the added per-chunk delivery delay /// well under one slice time without hammering the driver. @@ -534,16 +566,18 @@ pub struct NvencCudaEncoder { split_mode: u32, /// The last reference-frame range we invalidated — dedupes repeated RFI requests for one loss. last_rfi_range: Option<(i64, i64)>, - /// Cursor-as-metadata GPU blend (loaded lazily on the first frame that carries a cursor, once the - /// CUDA context is current). `None` until then or if the module load fails; `cursor_tried` stops - /// re-attempting a failed load every frame. `cursor_serial` tracks the uploaded bitmap. - cursor: Option, + /// Cursor-as-metadata GPU blend: the Vulkan device + SPIR-V compute pass the ring's + /// external-memory slots are allocated through (`vkslot.rs`) — the driver-portable + /// replacement for the retired PTX kernels. Brought up once at session init (`cursor_tried` + /// stops re-attempts); `None` = bring-up failed, the ring fell back to plain CUDA + /// allocations and composite mode degrades to no cursor. `cursor_serial` tracks the + /// uploaded bitmap. + vk_blend: Option, cursor_tried: bool, cursor_serial: u64, - /// Suppress-until-success latches for the per-frame cursor upload/blend warns: a persistent - /// failure sits in the submit() hot path, so warn once per failure streak (reset on success) - /// rather than on every cursor-bearing frame, which would evict the log ring. - cursor_upload_warned: bool, + /// Suppress-until-success latch for the per-frame blend warn: a persistent failure sits in + /// the submit() hot path, so warn once per failure streak (reset on success) rather than on + /// every cursor-bearing frame, which would evict the log ring. cursor_blend_warned: bool, /// One-shot latch for [`diagnose_failed_open`](Self::diagnose_failed_open) so a rebuild-retry /// burst (the session loop's bounded encoder resets) logs the diagnosis once, not per attempt. @@ -649,10 +683,9 @@ impl NvencCudaEncoder { frame_idx: 0, force_kf: false, pending_anchor: false, - cursor: None, + vk_blend: None, cursor_tried: false, cursor_serial: u64::MAX, - cursor_upload_warned: false, cursor_blend_warned: false, diagnosed: false, inited: false, @@ -743,7 +776,12 @@ impl NvencCudaEncoder { // (the forfeit contract), and the next session re-latches the arming at init. self.subframe_chunks = false; self.chunk = None; - self.ring.clear(); // drops the InputSurfaces, freeing their CUDA allocations + self.ring.clear(); // drops the CUDA InputSurfaces; Vk slots are freed just below + if let Some(vk) = &mut self.vk_blend { + // The Vulkan-backed slots' memory (and its CUDA mapping) — the device itself stays + // up for the next session's ring (`cursor_tried` keeps bring-up one-shot). + vk.free_slots(); + } self.bitstreams.clear(); self.pending.clear(); self.encoder = ptr::null_mut(); @@ -1132,36 +1170,102 @@ impl NvencCudaEncoder { // Encoder-owned input-surface ring: allocate + register POOL surfaces in the negotiated // format. Registered once here, mapped per submit, unregistered at teardown. - for _ in 0..POOL { - let surface = match self.buffer_fmt { - nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => { - InputSurface::alloc_yuv444(self.width, self.height) - } - nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => { - InputSurface::alloc_nv12(self.width, self.height) - } - _ => InputSurface::alloc_rgb(self.width, self.height), + // Preferred backing = Vulkan external memory CUDA-imported (`vkslot.rs`), so the + // SPIR-V cursor blend can composite into the very bytes NVENC encodes; any bring-up + // or per-slot failure falls back to plain pitched CUDA allocations (sessions always + // encode — composite mode just loses the cursor, warned below). + if !self.cursor_tried { + self.cursor_tried = true; + match VkSlotBlend::new() { + Ok(v) => self.vk_blend = Some(v), + Err(e) => tracing::warn!( + error = %format!("{e:#}"), + "NVENC (Linux): Vulkan slot-blend bring-up failed — plain CUDA input \ + surfaces, cursor compositing unavailable" + ), } - .context("alloc NVENC input surface")?; - let mut rr = nv::NV_ENC_REGISTER_RESOURCE { - version: nv::NV_ENC_REGISTER_RESOURCE_VER, - resourceType: - nv::NV_ENC_INPUT_RESOURCE_TYPE::NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR, - width: self.width, - height: self.height, - pitch: surface.pitch as u32, - resourceToRegister: surface.ptr as *mut c_void, - bufferFormat: self.buffer_fmt, - bufferUsage: nv::NV_ENC_BUFFER_USAGE::NV_ENC_INPUT_IMAGE, - ..Default::default() - }; - (api().register_resource)(self.encoder, &mut rr) - .nv_ok() - .map_err(|e| nvenc_status::call_err("register_resource (CUDADEVICEPTR)", e))?; - self.ring.push(RingSlot { - surface, - reg: rr.registeredResource, - }); + } + let slot_fmt = slot_fmt_of(self.buffer_fmt); + // Two attempts: the full ring on Vulkan slots, else (any failure) the full ring on + // plain CUDA — never a mixed ring (it would blend on some slots only: a flickering + // cursor) and never a short one. + 'ring: for use_vk in [self.vk_blend.is_some(), false] { + if !use_vk && self.vk_blend.is_some() { + // Second attempt: retire the Vulkan side wholesale first. + for s in self.ring.drain(..) { + let _ = (api().unregister_resource)(self.encoder, s.reg); + } + if let Some(vk) = &mut self.vk_blend { + vk.free_slots(); + } + self.vk_blend = None; + } + for _ in 0..POOL { + let surface = if use_vk { + let vk = self.vk_blend.as_mut().expect("use_vk implies Some"); + match vk.alloc_slot(slot_fmt, self.width, self.height) { + Ok(r) => SlotSurface::Vk(r), + Err(e) => { + tracing::warn!( + error = %format!("{e:#}"), + "NVENC (Linux): Vulkan slot alloc failed — rebuilding the \ + ring on plain CUDA surfaces (cursor compositing \ + unavailable)" + ); + continue 'ring; + } + } + } else { + SlotSurface::Cuda( + match self.buffer_fmt { + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => { + InputSurface::alloc_yuv444(self.width, self.height) + } + nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => { + InputSurface::alloc_nv12(self.width, self.height) + } + _ => InputSurface::alloc_rgb(self.width, self.height), + } + .context("alloc NVENC input surface")?, + ) + }; + let mut rr = nv::NV_ENC_REGISTER_RESOURCE { + version: nv::NV_ENC_REGISTER_RESOURCE_VER, + resourceType: + nv::NV_ENC_INPUT_RESOURCE_TYPE::NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR, + width: self.width, + height: self.height, + pitch: surface.pitch() as u32, + resourceToRegister: surface.ptr() as *mut c_void, + bufferFormat: self.buffer_fmt, + bufferUsage: nv::NV_ENC_BUFFER_USAGE::NV_ENC_INPUT_IMAGE, + ..Default::default() + }; + match (api().register_resource)(self.encoder, &mut rr).nv_ok() { + Ok(()) => {} + Err(e) if use_vk => { + // NVENC refusing the imported pointer is a Vulkan-side condition + // too — same wholesale fallback. + tracing::warn!( + error = ?e, + "NVENC (Linux): registering a Vulkan-imported slot failed — \ + rebuilding the ring on plain CUDA surfaces" + ); + continue 'ring; + } + Err(e) => { + return Err(nvenc_status::call_err( + "register_resource (CUDADEVICEPTR)", + e, + )) + } + } + self.ring.push(RingSlot { + surface, + reg: rr.registeredResource, + }); + } + break 'ring; // full ring built } self.inited = true; @@ -1246,9 +1350,9 @@ impl NvencCudaEncoder { /// IO-stream binding (stream-ordered submit — see the gate in [`Encoder::submit`]). fn copy_into_slot(&self, buf: &cuda::DeviceBuffer, slot: usize, sync: bool) -> Result<()> { let s = &self.ring[slot].surface; - let base = s.ptr; - let pitch = s.pitch; - let hh = s.height as u64; + let base = s.ptr(); + let pitch = s.pitch(); + let hh = s.height() as u64; match self.buffer_fmt { nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => { if !buf.yuv444 { @@ -1400,79 +1504,57 @@ impl Encoder for NvencCudaEncoder { // `async_rt` must be absent too: in two-thread mode the frame may be recycled right after // submit returns while the stream still holds its copy (belt-and-braces — an escalated // session was rebuilt without the binding, so `stream_ordered` is false there anyway). - let ordered = self.stream_ordered && self.async_rt.is_none() && self.pending.is_empty(); + // Cursor-bearing frames additionally force the CPU-synced path: the Vulkan blend sits + // between the CUDA copy and the encode, and its cross-API ordering is fence/CPU- + // established, not stream-ordered. Frames without a cursor (games hide it; client-draws + // sessions strip it) keep the stream-ordered fast path untouched. + let ordered = self.stream_ordered + && self.async_rt.is_none() + && self.pending.is_empty() + && captured.cursor.is_none(); let t0 = std::time::Instant::now(); // Copy the captured buffer into this slot's input surface before encoding it. self.copy_into_slot(buf, slot, !ordered)?; let t_copy = t0.elapsed(); - // Cursor-as-metadata: blend the overlay into this slot's OWNED input surface (a tiny kernel - // over the cursor's rect — never the compositor's dmabuf). The PTX module loads lazily on the - // first cursor frame now that the CUDA context is current; any failure degrades to no cursor, - // never a dropped frame. + // Cursor-as-metadata: blend the overlay into this slot's OWNED input surface via the + // SPIR-V compute pass (a dispatch over the cursor's rect — never the compositor's + // dmabuf). Cursor-bearing frames forced `ordered = false` above, so the CUDA copy has + // completed before the Vulkan dispatch and the fence-waited dispatch completes before + // the encode below — the cross-API ordering is CPU-established. Any failure degrades to + // no cursor, never a dropped frame. if let Some(ov) = &captured.cursor { - if !self.cursor_tried { - self.cursor_tried = true; - match cuda::CursorBlend::new(CURSOR_PTX) { - Ok(cb) => { - // Success is as diagnosis-critical as failure: a silent no-cursor session - // must be attributable to "overlay never arrived", not "module never said". - tracing::info!("NVENC (Linux): cursor blend module loaded"); - self.cursor = Some(cb); - } - Err(e) => tracing::warn!( - error = %format!("{e:#}"), - "NVENC (Linux): cursor blend module load failed — cursor not composited" - ), - } - } - if let Some(cb) = &self.cursor { + if let (Some(vk), SlotSurface::Vk(vref)) = + (self.vk_blend.as_mut(), &self.ring[slot].surface) + { if self.cursor_serial != ov.serial { - match cb.upload(ov.rgba.as_slice(), ov.w, ov.h) { - Ok(()) => { - self.cursor_serial = ov.serial; - self.cursor_upload_warned = false; - } - Err(e) => { - if !self.cursor_upload_warned { - self.cursor_upload_warned = true; - tracing::warn!( - error = %format!("{e:#}"), - serial = ov.serial, - "NVENC (Linux): cursor upload failed — cursor not composited" - ); - } - } - } + vk.upload_cursor(ov.rgba.as_slice(), ov.w, ov.h); + self.cursor_serial = ov.serial; } - let s = &self.ring[slot].surface; - // surfW = content width; surfH = the surface's allocated height (matches - // `copy_into_slot`'s plane math). Cursor pixels past the content are in cropped - // padding rows — harmless. - let (w, h) = (self.width, s.height); - let r = match self.buffer_fmt { - nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_YUV444 => { - cb.blend_yuv444(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered) - } - nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => { - cb.blend_nv12(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered) - } - _ => cb.blend_argb(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered), - }; + // surfW = content width; the blend derives plane strides from the slot's luma + // height. Cursor pixels past the content land in cropped padding rows — harmless. + let r = vk.blend_ref( + vref, + slot_fmt_of(self.buffer_fmt), + self.width, + ov.w, + ov.h, + ov.x, + ov.y, + ); if let Err(e) = r { if !self.cursor_blend_warned { self.cursor_blend_warned = true; tracing::warn!( error = %format!("{e:#}"), - "NVENC (Linux): cursor blend launch failed — cursor not composited" + "NVENC (Linux): cursor blend dispatch failed — cursor not composited" ); } } else { self.cursor_blend_warned = false; - // TEMP KWin composite probe (DROP BEFORE MERGE): prove the blend launches and - // with what geometry — the on-glass symptom is a loaded module and a silent, - // invisible cursor. + // TEMP KWin composite probe (DROP BEFORE MERGE): prove the blend dispatches + // and with what geometry. { use std::sync::atomic::{AtomicU64, Ordering as ProbeOrd}; static PROBE_BLEND: AtomicU64 = AtomicU64::new(0); @@ -1485,14 +1567,18 @@ impl Encoder for NvencCudaEncoder { ov_y = ov.y, ov_w = ov.w, ov_h = ov.h, - surf_w = w, - surf_h = h, visible = ov.visible, - "cursor blend probe: kernel launched" + "cursor blend probe: vulkan dispatch submitted" ); } } } + } else if !self.cursor_blend_warned { + self.cursor_blend_warned = true; + tracing::warn!( + "NVENC (Linux): cursor overlay present but no Vulkan blend (bring-up failed \ + earlier) — cursor not composited" + ); } } @@ -1540,7 +1626,7 @@ impl Encoder for NvencCudaEncoder { version: nv::NV_ENC_PIC_PARAMS_VER, inputWidth: self.width, inputHeight: self.height, - inputPitch: self.ring[slot].surface.pitch as u32, + inputPitch: self.ring[slot].surface.pitch() as u32, inputBuffer: mp.mappedResource, bufferFmt: mp.mappedBufferFmt, outputBitstream: self.bitstreams[slot], diff --git a/crates/pf-zerocopy/src/imp/cuda.rs b/crates/pf-zerocopy/src/imp/cuda.rs index 97779d4c..0daef3c9 100644 --- a/crates/pf-zerocopy/src/imp/cuda.rs +++ b/crates/pf-zerocopy/src/imp/cuda.rs @@ -7,8 +7,12 @@ //! and ffmpeg's `hevc_nvenc` (encode thread) — each thread makes it current before use; //! * device memory: pitched allocations, the reusable `BufferPool`/`DeviceBuffer`, IPC //! export/import, host readback, and the plane copies; -//! * GL / external-memory interop (`RegisteredTexture`, `ExternalDmabuf`); and -//! * the CUDA cursor-blend kernel (`CursorBlend`). +//! * GL / external-memory interop (`RegisteredTexture`, `ExternalDmabuf`). +//! +//! (The CUDA cursor-blend PTX kernel that used to live here is retired: vendored PTX is JIT'd +//! against the driver's ISA ceiling and silently dies on older drivers. The NVENC cursor blend +//! is now the SPIR-V compute pass in [`super::vkslot`], dispatched over Vulkan-allocated, +//! CUDA-imported input slots.) //! //! (We use GL interop, not EGL interop: `cuGraphicsEGLRegisterImage` is Tegra-only on the desktop //! driver — see [`super::egl`].) @@ -18,7 +22,6 @@ #![deny(clippy::undocumented_unsafe_blocks)] use anyhow::{bail, Result}; -use std::ffi::CStr; use std::os::raw::{c_uint, c_void}; use std::sync::{Arc, Mutex, OnceLock}; @@ -280,258 +283,6 @@ pub fn copy_stream_handle() -> *mut c_void { /// Max cursor-overlay bitmap edge (px) uploaded to the device blend buffer — matches the Vulkan path. pub const CURSOR_MAX: u32 = 256; -/// GPU cursor-overlay compositor for the NVENC path (cursor-as-metadata): loads the `cursor_blend` -/// PTX module once and blends a straight-alpha RGBA cursor into an encoder-OWNED NVENC input surface -/// (ARGB / NV12 / YUV444) with a small kernel launched over the cursor's rectangle — no full-frame -/// pass, and the compositor's dmabuf is never touched. The cursor bitmap lives in a device buffer -/// re-uploaded only when it changes. Requires `context()` to have succeeded (driver present). -pub struct CursorBlend { - module: CUmodule, - f_argb: CUfunction, - f_nv12: CUfunction, - f_yuv444: CUfunction, - cur_buf: CUdeviceptr, // device RGBA staging (CURSOR_MAX²·4, tight rows) -} - -// SAFETY: process-lifetime driver handles used only from the encode thread with the shared context -// current — like [`DeviceBuffer`], moving the struct between threads cannot dangle or race. -unsafe impl Send for CursorBlend {} - -impl CursorBlend { - /// Load the embedded PTX image and resolve the three blend kernels + a device cursor buffer. - pub fn new(ptx: &[u8]) -> Result { - // cuModuleLoadData reads a PTX image as a NUL-terminated string; the embedded .ptx is not, - // so append a terminator. - let mut image = ptx.to_vec(); - image.push(0); - let mut module: CUmodule = std::ptr::null_mut(); - // SAFETY: `&mut module` is a live out-param the driver fills; `image` is a NUL-terminated PTX - // byte image that outlives the synchronous load. `ck` bails on error before `module` is used. - unsafe { - ck( - cuModuleLoadData(&mut module, image.as_ptr() as *const c_void), - "cuModuleLoadData(cursor_blend)", - )?; - } - let getf = |name: &CStr| -> Result { - let mut f: CUfunction = std::ptr::null_mut(); - // SAFETY: `module` loaded above; each name is a valid NUL-terminated symbol present in - // the module (verified in the .ptx `.entry` list); `&mut f` is a live out-param. - unsafe { - ck( - cuModuleGetFunction(&mut f, module, name.as_ptr()), - "cuModuleGetFunction", - )?; - } - Ok(f) - }; - let f_argb = getf(c"blend_argb")?; - let f_nv12 = getf(c"blend_nv12")?; - let f_yuv444 = getf(c"blend_yuv444")?; - let mut cur_buf: CUdeviceptr = 0; - // SAFETY: `&mut cur_buf` is a live out-param; the size fits the CURSOR_MAX² RGBA buffer. - unsafe { - ck( - cuMemAlloc_v2(&mut cur_buf, (CURSOR_MAX * CURSOR_MAX * 4) as usize), - "cuMemAlloc(cursor)", - )?; - } - Ok(CursorBlend { - module, - f_argb, - f_nv12, - f_yuv444, - cur_buf, - }) - } - - /// Upload the cursor RGBA (`cw*ch*4`, tight rows) into the device blend buffer. Call only when - /// the bitmap changes; position moves are just kernel args. - pub fn upload(&self, rgba: &[u8], cw: u32, ch: u32) -> Result<()> { - let cw = cw.min(CURSOR_MAX); - let ch = ch.min(CURSOR_MAX); - let row = cw as usize * 4; - let copy = CUDA_MEMCPY2D { - srcMemoryType: 1, // HOST - srcHost: rgba.as_ptr() as *const c_void, - srcPitch: row, - dstMemoryType: CU_MEMORYTYPE_DEVICE, - dstDevice: self.cur_buf, - dstPitch: row, - WidthInBytes: row, - Height: ch as usize, - ..Default::default() - }; - // SAFETY: HOST→DEVICE 2D copy of `row*ch` bytes; `rgba` covers at least that (caller passes - // `cw*ch*4`), `cur_buf` is the CURSOR_MAX²·4 device alloc (row ≤ CURSOR_MAX·4, ch ≤ CURSOR_MAX). - // Synchronous via `copy_blocking`. Requires the context current (caller's contract). - unsafe { copy_blocking(©, "cursor HtoD") } - } - - /// Blend into a packed 4-byte (NVENC ARGB) owned surface at `(ox,oy)`. - #[allow(clippy::too_many_arguments)] // surface geometry + cursor size + offset — a struct would just be unpacked at the call - pub fn blend_argb( - &self, - surf: CUdeviceptr, - pitch: usize, - w: u32, - h: u32, - cw: u32, - ch: u32, - ox: i32, - oy: i32, - sync: bool, - ) -> Result<()> { - let (mut a_surf, mut a_cur) = (surf, self.cur_buf); - let (mut a_pitch, mut a_w, mut a_h) = (pitch as i32, w as i32, h as i32); - let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); - let (mut a_ox, mut a_oy) = (ox, oy); - let mut args: [*mut c_void; 9] = [ - &mut a_surf as *mut _ as *mut c_void, - &mut a_pitch as *mut _ as *mut c_void, - &mut a_w as *mut _ as *mut c_void, - &mut a_h as *mut _ as *mut c_void, - &mut a_cur as *mut _ as *mut c_void, - &mut a_cw as *mut _ as *mut c_void, - &mut a_ch as *mut _ as *mut c_void, - &mut a_ox as *mut _ as *mut c_void, - &mut a_oy as *mut _ as *mut c_void, - ]; - self.launch(self.f_argb, a_cw as u32, a_ch as u32, &mut args, sync) - } - - /// Blend into an owned planar YUV444 surface (3 stacked full-res planes) at `(ox,oy)`. - #[allow(clippy::too_many_arguments)] // surface geometry + cursor size + offset — a struct would just be unpacked at the call - pub fn blend_yuv444( - &self, - base: CUdeviceptr, - pitch: usize, - w: u32, - h: u32, - cw: u32, - ch: u32, - ox: i32, - oy: i32, - sync: bool, - ) -> Result<()> { - let (mut a_base, mut a_cur) = (base, self.cur_buf); - let (mut a_pitch, mut a_w, mut a_h) = (pitch as i32, w as i32, h as i32); - let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); - let (mut a_ox, mut a_oy) = (ox, oy); - let mut args: [*mut c_void; 9] = [ - &mut a_base as *mut _ as *mut c_void, - &mut a_pitch as *mut _ as *mut c_void, - &mut a_w as *mut _ as *mut c_void, - &mut a_h as *mut _ as *mut c_void, - &mut a_cur as *mut _ as *mut c_void, - &mut a_cw as *mut _ as *mut c_void, - &mut a_ch as *mut _ as *mut c_void, - &mut a_ox as *mut _ as *mut c_void, - &mut a_oy as *mut _ as *mut c_void, - ]; - self.launch(self.f_yuv444, a_cw as u32, a_ch as u32, &mut args, sync) - } - - /// Blend into an owned NV12 surface (Y plane at `base`, interleaved UV at `base + pitch*h`). - #[allow(clippy::too_many_arguments)] // surface geometry + cursor size + offset — a struct would just be unpacked at the call - pub fn blend_nv12( - &self, - base: CUdeviceptr, - pitch: usize, - w: u32, - h: u32, - cw: u32, - ch: u32, - ox: i32, - oy: i32, - sync: bool, - ) -> Result<()> { - let (mut a_yb, mut a_uvb, mut a_cur) = (base, base + pitch as u64 * h as u64, self.cur_buf); - let (mut a_yp, mut a_uvp) = (pitch as i32, pitch as i32); - let (mut a_w, mut a_h) = (w as i32, h as i32); - let (mut a_cw, mut a_ch) = (cw.min(CURSOR_MAX) as i32, ch.min(CURSOR_MAX) as i32); - let (mut a_ox, mut a_oy) = (ox, oy); - let mut args: [*mut c_void; 11] = [ - &mut a_yb as *mut _ as *mut c_void, - &mut a_yp as *mut _ as *mut c_void, - &mut a_uvb as *mut _ as *mut c_void, - &mut a_uvp as *mut _ as *mut c_void, - &mut a_w as *mut _ as *mut c_void, - &mut a_h as *mut _ as *mut c_void, - &mut a_cur as *mut _ as *mut c_void, - &mut a_cw as *mut _ as *mut c_void, - &mut a_ch as *mut _ as *mut c_void, - &mut a_ox as *mut _ as *mut c_void, - &mut a_oy as *mut _ as *mut c_void, - ]; - // One thread per 2x2 luma block → grid over ceil(cw/2) × ceil(ch/2). - self.launch( - self.f_nv12, - (a_cw as u32).div_ceil(2), - (a_ch as u32).div_ceil(2), - &mut args, - sync, - ) - } - - /// Launch `f` over a `work_w × work_h` grid (16×16 blocks) on the copy stream; `sync` waits - /// for it, `!sync` leaves completion to the stream (stream-ordered consumers only — the - /// kernel PARAMETERS are copied at launch time, so the arg locals need not outlive the call). - fn launch( - &self, - f: CUfunction, - work_w: u32, - work_h: u32, - args: &mut [*mut c_void], - sync: bool, - ) -> Result<()> { - if work_w == 0 || work_h == 0 { - return Ok(()); - } - const B: u32 = 16; - let stream = copy_stream(); - // SAFETY: `f` is a resolved kernel from our loaded module; `args` holds pointers to live - // locals whose types match the kernel's C parameters (per the call site above) — CUDA - // copies the parameter values during `cuLaunchKernel` itself, so they need not outlive - // the call. Grid/block dims are non-zero. Launched on the copy stream (ordered after the - // input-surface copy issued on the same stream); `sync` waits, `!sync` leaves ordering to - // the stream (the NVENC IO-stream binding). Requires the context current. - unsafe { - ck( - cuLaunchKernel( - f, - work_w.div_ceil(B), - work_h.div_ceil(B), - 1, - B, - B, - 1, - 0, - stream, - args.as_mut_ptr(), - std::ptr::null_mut(), - ), - "cuLaunchKernel(cursor)", - )?; - if sync { - ck(cuStreamSynchronize(stream), "cuStreamSynchronize(cursor)")?; - } - Ok(()) - } - } -} - -impl Drop for CursorBlend { - fn drop(&mut self) { - // SAFETY: `cur_buf`/`module` are our own handles, freed exactly once here; the context is - // current on the encode thread that drops the encoder. Errors are ignored on teardown. - unsafe { - let _ = cuMemFree_v2(self.cur_buf); - let _ = cuModuleUnload(self.module); - } - } -} - /// Allocate one pitched device buffer for `width`x`height` 4-byte pixels; returns `(ptr, pitch)`. fn alloc_pitched(width: u32, height: u32) -> Result<(CUdeviceptr, usize)> { let mut ptr: CUdeviceptr = 0; diff --git a/crates/pf-zerocopy/src/imp/cuda/ffi.rs b/crates/pf-zerocopy/src/imp/cuda/ffi.rs index 236abda1..046ad04e 100644 --- a/crates/pf-zerocopy/src/imp/cuda/ffi.rs +++ b/crates/pf-zerocopy/src/imp/cuda/ffi.rs @@ -9,7 +9,7 @@ #![deny(clippy::undocumented_unsafe_blocks)] use anyhow::{bail, Result}; -use std::os::raw::{c_char, c_int, c_uint, c_void}; +use std::os::raw::{c_int, c_uint, c_void}; use std::sync::OnceLock; pub type CUresult = c_uint; // CUDA_SUCCESS == 0 @@ -20,8 +20,6 @@ pub type CUdeviceptr = u64; pub type CUgraphicsResource = *mut c_void; pub type CUarray = *mut c_void; pub type CUexternalMemory = *mut c_void; // opaque CUextMemory_st* -pub type CUmodule = *mut c_void; // opaque CUmod_st* -pub type CUfunction = *mut c_void; // opaque CUfunc_st* /// `CUmemorytype` (cuda.h): HOST=1, DEVICE=2, ARRAY=3, UNIFIED=4. pub const CU_MEMORYTYPE_DEVICE: c_uint = 2; @@ -144,26 +142,6 @@ pub(crate) struct CudaApi { cuIpcGetMemHandle: unsafe extern "C" fn(*mut CUipcMemHandle, CUdeviceptr) -> CUresult, cuIpcOpenMemHandle: unsafe extern "C" fn(*mut CUdeviceptr, CUipcMemHandle, c_uint) -> CUresult, cuIpcCloseMemHandle: unsafe extern "C" fn(CUdeviceptr) -> CUresult, - // Cursor-overlay blend: a linear device alloc + a PTX module with the blend kernels launched - // over the cursor's small rectangle (see [`CursorBlend`]). - cuMemAlloc_v2: unsafe extern "C" fn(*mut CUdeviceptr, usize) -> CUresult, - cuModuleLoadData: unsafe extern "C" fn(*mut CUmodule, *const c_void) -> CUresult, - cuModuleUnload: unsafe extern "C" fn(CUmodule) -> CUresult, - cuModuleGetFunction: unsafe extern "C" fn(*mut CUfunction, CUmodule, *const c_char) -> CUresult, - #[allow(clippy::type_complexity)] - cuLaunchKernel: unsafe extern "C" fn( - CUfunction, - c_uint, - c_uint, - c_uint, - c_uint, - c_uint, - c_uint, - c_uint, - CUstream, - *mut *mut c_void, - *mut *mut c_void, - ) -> CUresult, } // SAFETY: every field is a bare `extern "C" fn` address into the leaked, process-lifetime // `libcuda` mapping (`cuda_api` `forget`s the `Library`, so it is never unloaded) — an immutable @@ -236,11 +214,6 @@ pub(crate) fn cuda_api() -> Option<&'static CudaApi> { .or_else(|_| lib.get(b"cuIpcOpenMemHandle\0")) .ok()?, cuIpcCloseMemHandle: *lib.get(b"cuIpcCloseMemHandle\0").ok()?, - cuMemAlloc_v2: *lib.get(b"cuMemAlloc_v2\0").ok()?, - cuModuleLoadData: *lib.get(b"cuModuleLoadData\0").ok()?, - cuModuleUnload: *lib.get(b"cuModuleUnload\0").ok()?, - cuModuleGetFunction: *lib.get(b"cuModuleGetFunction\0").ok()?, - cuLaunchKernel: *lib.get(b"cuLaunchKernel\0").ok()?, }; std::mem::forget(lib); // keep libcuda mapped for the fn pointers' lifetime (process) Some(api) @@ -304,53 +277,6 @@ pub(crate) unsafe fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult { None => CU_ERROR_NOT_LOADED, } } -pub(crate) unsafe fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, size: usize) -> CUresult { - match cuda_api() { - Some(a) => (a.cuMemAlloc_v2)(dptr, size), - None => CU_ERROR_NOT_LOADED, - } -} -pub(crate) unsafe fn cuModuleLoadData(m: *mut CUmodule, image: *const c_void) -> CUresult { - match cuda_api() { - Some(a) => (a.cuModuleLoadData)(m, image), - None => CU_ERROR_NOT_LOADED, - } -} -pub(crate) unsafe fn cuModuleUnload(m: CUmodule) -> CUresult { - match cuda_api() { - Some(a) => (a.cuModuleUnload)(m), - None => CU_ERROR_NOT_LOADED, - } -} -pub(crate) unsafe fn cuModuleGetFunction( - f: *mut CUfunction, - m: CUmodule, - name: *const c_char, -) -> CUresult { - match cuda_api() { - Some(a) => (a.cuModuleGetFunction)(f, m, name), - None => CU_ERROR_NOT_LOADED, - } -} -#[allow(clippy::too_many_arguments)] -pub(crate) unsafe fn cuLaunchKernel( - f: CUfunction, - gx: c_uint, - gy: c_uint, - gz: c_uint, - bx: c_uint, - by: c_uint, - bz: c_uint, - shmem: c_uint, - stream: CUstream, - params: *mut *mut c_void, - extra: *mut *mut c_void, -) -> CUresult { - match cuda_api() { - Some(a) => (a.cuLaunchKernel)(f, gx, gy, gz, bx, by, bz, shmem, stream, params, extra), - None => CU_ERROR_NOT_LOADED, - } -} pub(crate) unsafe fn cuMemcpy2DAsync_v2(copy: *const CUDA_MEMCPY2D, stream: CUstream) -> CUresult { match cuda_api() { Some(a) => (a.cuMemcpy2DAsync_v2)(copy, stream), diff --git a/crates/pf-zerocopy/src/imp/cursor_blend.comp b/crates/pf-zerocopy/src/imp/cursor_blend.comp new file mode 100644 index 00000000..0d940361 --- /dev/null +++ b/crates/pf-zerocopy/src/imp/cursor_blend.comp @@ -0,0 +1,170 @@ +#version 450 +// Cursor-overlay blend for the direct-SDK NVENC path (cursor-as-metadata), dispatched over the +// cursor's rectangle only — the Vulkan replacement for the retired cursor_blend.cu PTX kernels +// (PTX is JIT'd against the driver's ISA ceiling, so a vendored blob silently dies on older +// drivers: CUDA errors 222/218 on-glass; SPIR-V has no such coupling). The NVENC input surface is +// Vulkan-allocated, CUDA-imported external memory (see vkslot.rs), so this shader writes the very +// bytes NVENC encodes. +// +// MODE (spec constant): 0 = packed 4-byte ARGB (NVENC byte order B,G,R,A), 1 = NV12 (Y plane + +// interleaved half-res UV at row surfH), 2 = planar YUV444 (3 full-res planes stacked at +// pitch*surfH). BT.709 limited-range coefficients — identical to rgb2nv12_buf.comp and the +// retired .cu, so the cursor colour matches the frame regardless of backend. +// +// The surface SSBO is uint[] (no 8-bit storage dependency — maximum driver reach): every +// invocation exclusively owns the 32-bit words it read-modify-writes. ARGB: one invocation per +// cursor pixel = one word. NV12/YUV444: one invocation per WORD-ALIGNED 4-px luma span (per two +// rows for NV12, whose 2 chroma bytes-pairs land in one exclusive word). Spans are aligned to the +// SURFACE, not the cursor, so neighbouring invocations never share a word even at odd `ox`. +// +// Rebuild: glslc cursor_blend.comp -o cursor_blend.spv (vendored beside this file) + +layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in; + +layout(constant_id = 0) const uint MODE = 0; + +layout(std430, binding = 0) buffer Surf { uint surf[]; }; +layout(std430, binding = 1) readonly buffer Cur { uint cur[]; }; + +layout(push_constant) uniform Push { + uint pitch; // surface row stride, bytes (4-aligned by construction) + uint surfW; // content width, px + uint surfH; // luma rows (plane stride multiplier) + uint curW; // cursor bitmap width, px + uint curH; // cursor bitmap height, px + int ox; // cursor top-left on the surface, px (may be negative) + int oy; +} pc; + +// Cursor texel (straight-alpha RGBA, tight rows) or (0,0,0,0) outside the bitmap. +uvec4 cursor_px(int cx, int cy) { + if (cx < 0 || cy < 0 || cx >= int(pc.curW) || cy >= int(pc.curH)) return uvec4(0); + uint w = cur[uint(cy) * pc.curW + uint(cx)]; + return uvec4(w & 0xFFu, (w >> 8) & 0xFFu, (w >> 16) & 0xFFu, (w >> 24) & 0xFFu); // R,G,B,A +} + +uint blend8(uint dst, uint src, uint a) { + return (src * a + dst * (255u - a)) / 255u; +} + +// BT.709 limited RGB→Y/U/V (matches the retired .cu / rgb2nv12_buf.comp). +uint y_of(uvec4 s) { + return uint(clamp(16.0 + 0.1826 * float(s.r) + 0.6142 * float(s.g) + 0.0620 * float(s.b) + 0.5, 0.0, 255.0)); +} +float u_of(uvec4 s) { return 128.0 - 0.1006 * float(s.r) - 0.3386 * float(s.g) + 0.4392 * float(s.b); } +float v_of(uvec4 s) { return 128.0 + 0.4392 * float(s.r) - 0.3989 * float(s.g) - 0.0403 * float(s.b); } + +// Read-modify-write one byte lane of a word index. +void rmw_byte(uint word_idx, uint lane, uint val8, uint a) { + uint w = surf[word_idx]; + uint shift = lane * 8u; + uint d = (w >> shift) & 0xFFu; + uint b = blend8(d, val8, a); + surf[word_idx] = (w & ~(0xFFu << shift)) | (b << shift); +} + +void main() { + if (MODE == 0u) { + // ARGB: one invocation per cursor pixel; each surface pixel is one exclusive word. + int cx = int(gl_GlobalInvocationID.x); + int cy = int(gl_GlobalInvocationID.y); + if (cx >= int(pc.curW) || cy >= int(pc.curH)) return; + int px = pc.ox + cx, py = pc.oy + cy; + if (px < 0 || py < 0 || px >= int(pc.surfW) || py >= int(pc.surfH)) return; + uvec4 s = cursor_px(cx, cy); + if (s.a == 0u) return; + uint idx = (uint(py) * pc.pitch + uint(px) * 4u) / 4u; + uint w = surf[idx]; + uint b = blend8(w & 0xFFu, s.b, s.a); // B lane + uint g = blend8((w >> 8) & 0xFFu, s.g, s.a); // G lane + uint r = blend8((w >> 16) & 0xFFu, s.r, s.a); // R lane + surf[idx] = (w & 0xFF000000u) | (r << 16) | (g << 8) | b; + return; + } + + // NV12 / YUV444: one invocation per SURFACE-word-aligned 4-px luma span. Span origin: + // x0 = floor(ox/4)*4 + span*4 (surface px), rows walk the cursor rect. + int span = int(gl_GlobalInvocationID.x); + int row = int(gl_GlobalInvocationID.y); + int x0 = (pc.ox >> 2) << 2; // word-aligned start at/left-of ox (ox may be negative) + int px0 = x0 + span * 4; + + if (MODE == 2u) { + // YUV444: rows walk cursor rows one at a time. + if (row >= int(pc.curH)) return; + int py = pc.oy + row; + if (py < 0 || py >= int(pc.surfH)) return; + uint plane = pc.pitch * pc.surfH; + for (int i = 0; i < 4; i++) { + int px = px0 + i; + int cx = px - pc.ox; + if (px < 0 || px >= int(pc.surfW)) continue; + uvec4 s = cursor_px(cx, row); + if (s.a == 0u) continue; + uint off = uint(py) * pc.pitch + uint(px); + uint U = uint(clamp(u_of(s) + 0.5, 0.0, 255.0)); + uint V = uint(clamp(v_of(s) + 0.5, 0.0, 255.0)); + rmw_byte(off / 4u, off % 4u, y_of(s), s.a); + rmw_byte((plane + off) / 4u, (plane + off) % 4u, U, s.a); + rmw_byte((2u * plane + off) / 4u, (2u * plane + off) % 4u, V, s.a); + } + return; + } + + // NV12: rows walk 2-row luma blocks (row = block row). The span's 4 luma px × 2 rows are + // exclusive words; its 2 chroma samples (4 bytes) are one exclusive word. + int base_cy = row * 2; + if (base_cy >= int(pc.curH)) return; + // Luma: 4 px × 2 rows. + for (int j = 0; j < 2; j++) { + int cy = base_cy + j; + int py = pc.oy + cy; + if (cy >= int(pc.curH) || py < 0 || py >= int(pc.surfH)) continue; + for (int i = 0; i < 4; i++) { + int px = px0 + i; + int cx = px - pc.ox; + if (px < 0 || px >= int(pc.surfW)) continue; + uvec4 s = cursor_px(cx, cy); + if (s.a == 0u) continue; + uint off = uint(py) * pc.pitch + uint(px); + rmw_byte(off / 4u, off % 4u, y_of(s), s.a); + } + } + // Chroma: two UV samples covering the span's 2x2 blocks, alpha-weighted like the .cu kernel. + // The UV plane starts at row surfH; sample (uvx, uvy) lives at uv_base + uvy*pitch + uvx*2. + // Guard: only spans whose px0 is 4-aligned own their chroma word (px0 is by construction). + int py_top = pc.oy + base_cy; + int uvy = py_top >> 1; + if (py_top < 0 || uvy < 0 || uvy * 2 >= int(pc.surfH)) return; + uint uv_base = pc.pitch * pc.surfH; + for (int hf = 0; hf < 2; hf++) { + // Each hf = one 2x2 luma block = one UV sample (2 bytes). + int bx = px0 + hf * 2; + if (bx < 0 || bx >= int(pc.surfW)) continue; + int uvx = bx >> 1; + float ua = 0.0, va = 0.0, wa = 0.0; + int cnt = 0; + for (int j = 0; j < 2; j++) { + for (int i = 0; i < 2; i++) { + int px = bx + i; + int py = py_top + j; + int cx = px - pc.ox; + int cy = base_cy + j; + if (px < 0 || py < 0 || px >= int(pc.surfW) || py >= int(pc.surfH)) continue; + uvec4 s = cursor_px(cx, cy); + if (s.a == 0u) continue; + ua += u_of(s) * float(s.a); + va += v_of(s) * float(s.a); + wa += float(s.a); + cnt++; + } + } + if (wa <= 0.0 || cnt == 0) continue; + uint U = uint(clamp(ua / wa + 0.5, 0.0, 255.0)); + uint V = uint(clamp(va / wa + 0.5, 0.0, 255.0)); + uint amean = uint(clamp(wa / float(cnt) + 0.5, 0.0, 255.0)); + uint off = uv_base + uint(uvy) * pc.pitch + uint(uvx) * 2u; + rmw_byte(off / 4u, off % 4u, U, amean); + rmw_byte((off + 1u) / 4u, (off + 1u) % 4u, V, amean); + } +} diff --git a/crates/pf-zerocopy/src/imp/cursor_blend.spv b/crates/pf-zerocopy/src/imp/cursor_blend.spv new file mode 100644 index 0000000000000000000000000000000000000000..046d149a9f5d3933aba0aa8821d4dd2ec76410f0 GIT binary patch literal 20412 zcmZ9U2Y_BxwS|8&lR_tSgg^oWLXjFu=pjG?1VZmHOd1TN3^OT2S_&SSlY7m1d#!c$*=L_}?>%>BcrCrw>)U87(O9~% zY-4^)qj~jj^g?N%^=aCXqYoavS@*Q@n+@G&3mukg^f>CXJU+d#t(3{_9aAukS1Bu? ztM#Y!tBU?D{I3`Gen@L$sm93BZ6n8w8992-wsD=EZQavmO`6a)sbk87w(%WZ6UI&J zm@z@WeuZDhlyQ@$k1rj1<5BkCyU~}?G#cZkcXdzgYU@08NXM2t@(+nzTNoF0r}p4w z#IwYG8!LmyPMR=f{PshpZ@FX7Kcr?G$1%QpS`)7b?(Q1b#4CZ@n?Ai8YhY)!O`SMo z#`K|fW?!pgyQ})OvD0hy1B&`V;2E|0fs5+ z{@|Gu@d4No)4Q5 qX0)*aMCn|J}Zr%h|)e(=HMXtiUJdzrV^xC;eb&fUYAx7NV1RGw)@@mA)CcXV0}ZHdill+j249 z(%2F1b}`0CaOQVkaK^X)Vr^O*V`}`M#du5OF!%}0d|MkwEmqgsIJ(AG1B(_U1Fl@yv#IO`X{s+u7h#d3veOw;e=$g}6Cyt&JPR)y-nv zE%1(N-tTI*>1ls2IM0vA!I|qP^eN{0sl|9p<5~FBi4&Xsz63tB8t2P3zt`c1RdsJH zR@d5izs5hR@sDf#lN$fD#+T}~c+N}L_%b!VT#c_);{$4Z?HV6g;~T=qws%ixV^cfl zTftAN=6nZm_WCZitM>Qd;H*oVdUX$s)4#$qz*(1>`V{LjYcbx^m<#V&m)6Ebi`BI@ zE~)WLYy7erzr4n;l+S9LHnp>7?_3R^KEtuLHtuNRHizFF`#tazc@jAf4}!;5V|k|O z*K?2lRovd=^DKBqg`We@tnl;TaZ`HURWE=upRa&3Z*PJ#Z*PmMxqTm;Is8Oi>1*;Z z&wQ>0F6VRY8sD(Sw^)p~G)BXF_I^v_U~u;OQQ-E;6Ph=$b1)X3b?E|^eRkLQv>KmL z<7d_Q>>8g_tR8`V~)jARdu!K&HbJP*Jce-w(?OJCcE+ulom-<63a9z<~*#+I*3@si(% zE&IA3-qY9T@6yj>w>ICU#`CiMLTnlDMYw%=S?{^gv)Nh|KcC?8H7F}F`_APcdbCgR zDzq`5ffVBmgmai}7Io{L_kk3z)!<8F)ts+E%*j$%wX`$F@)g%_5IwjiV)YdLR)yQn z_JiQ7Vf9mAl`Pl94y0J;nB>+CqLuUH7!ogbWz>z)H=&JV*p%Y6Hods-j9)MM^~^(S zV_UHK=({5^yNH*a?rL>KBFY3vEbA@)`=9?CG|ez+{bd~`7}zSF?)FU>9Dgt z+F_UQx$tdw8nCO}xbrAWP*Oi1eEVBx?40@w;O3k9i@{vY_V#y4lYU-*X_KEdBL3g( zsb9pefVcI1tQp_-S5ZDX=Eg<&T?2pexfd4kTlCrK(M9}D_?7(!@9eza2X+k3`~6tC z^ZEeRoYkCnx%XhkVVg%N+O4X2@SgM<%8qg${*!V&McwuG9{Ly7HZ2tU`8Re6in{&C zU#&FjU#qZfoR8*yrGd{k>-&P;SL&%>7OmuK!uhwTelVmQ$56QW*xr6dz{@?{X=Aoj z#&sb)lCp65&MnZD$n=}YdJ zzT}?iOFps2J=d4@p6kQ4p6g3KyT(1+m-U|SOYRvz-1+jHAFlPBUvkg-CHK4^?)rG< z57&C`FS+ObaQpMzU-FjP;$Q!O71s8xc15#_ZuPg#`ilRTzgNA z`>jyc`@InEGtF;?l0RGHemkUI>-R&+U#M}vBg%TeB}(r1M9KZ8DETJ^cRu~LDC_;c zD7oJk;rjc1QF6a8!ksU_FG}t=M#=roD7oJmCHGsSj!OzcoreyT<3# zxZfK3zsK9;WvthSyi0udzC=;;9U~sjo5uI)>lELyzLR&x*LU6<6!p{LzB~U@X}b{D zcal8q{|mPLJWAUC53HW{Z-Lcp?>kJcud&~z7}I#3A>W~R&h%`Vc<)wv;=Kn~yNNa% zW8bG3NBuhDs(nCFcP>7{dN~*BA5zqu3vqJ!7;JvYEjG!c!Pr>TR#j~Yaa`A9s z9Q6-~pMJFOusY_p1U}~GI_Pf>=Gg-FGEaSaQPj*+{0YutHFNa5(~n|~eXtqFl9e95 z6x=zUO&jNYX}CV>u8HRewZvQwY)sEiiLpFbJ$wbQbCsC=;rgf>)3beMWvev7C)t$pNuwKrg`sx%l=S`fxz5q6#to2%8 z$L#q~AMb+!VD+r^+F-S;^*UhVsJqtbN1L@C2sSs@On-Ao-h;r_oA-KHFY{Jkm!f9g z;^e(P*f~nx8-UGwaOGp(8-mr7_YkmJ^4@&K^}u&Rc-pf9?yiUnt%Hw}9xW1N$psDMB5LP~n=cw~=DA=EA*{`vz*Kc2}b?NhPuzBBK<$VNLZ7gjw zmLuWn8H?X0YRT;wu(_4*xntq#_IWf`9{)CQeJsbLsq23nR-Uo6gB^?g8q0e94#QfP zv5W(kV;K)un^=rx0$e>~nFv-(ZYP4xt$gQqz}4;Z1gt#%Cxh!_nS`dU|4CSR#xfb~ zSnStW*6TMGYhA`N1?-*-p9*$tS(8)1YMqq$bb*btuv*J*u$tq4u+pZ%ZEIcHO$R%d zX*UC`<{n79nQ+@$mv*Ou-B0nE1=deJ_tKZZxtAQ1V^z~P&l1mG&gW~`GtfOVmGgfl zntGnuXMxp>bvjm__H)3tFZbeHH1)Kf2UfHFY^>b&eX-|Id{*_rI!@07=T`ay_{h(L zs|}_&XKLqz)zyr9HdZY$E&!{Ae;I7A=I1i7THcA5gN>tZ-`~J`*|&9Hr>NPt zIP0pGdHojHewW22w{L^h!@mPI*Np#4us-UJ{|c;H{J#rU%bDmZu$s^9oFl#mS2ypg zv0mn_{(Xv?d5aVG2Vmo7&VC3sch^=QpRd<|)iY<;g4Hr-KLQ&^-E&0x(dHa+9oXEQ zC;iRAJg>)knWsKKrl^^xIP*e^-(w1yRd5U|0P&0`*8u-`RY^n+y_>lTit*6gKej7%zLqFiTME7m>JJP zuv)qQ9)zpAt`A|o9IN`TC~A&XoVX8zldJdaui@spX655LJpxwG{`(DBE&K0LuyNE~ zkMyI>{`)Q1+?;3q&A~i>hxIZ~eIBEznWxxzuC-h}c|Qquj-0C}u=3RX0X(~^dkQPh zS@vnLXIa}=XZy#o>aLe(**}8ShEj~_96SS7&sp|QU^U+zj^i(2FUMispDAkQEiT8Q zo?M>;UqNwh?dN&0TJED4!1a6WZ)obdkNys}oqFQ_1MIpc?u%fxoL&D3ww=0j@GMr% z_MZ7)0XqlQy^NI)V}918?q6X4_DOwCrTrVMelhwoj6}az;r!!g_>G43{!Yc;qfD;G z`Wl)(_VE%{p8o#>PXBL!<>~)T@F3!-r~m(g)yw|>2j^e2f9vc0zlElc{lAWt+rR7Z zHrP4Ix&9rn+6YQ~-UaJ3pEt?nhu}4^>gMnf)_%kvV~r>N z1gmfM=VxGJmiv=!)=<~~Q>=U#y63BvqW*@0x4<1&_DU;UAN9;f zZ?JmC>UjFV%Q$`EWt@I+ebh5&OM=ynWA6U!+Sp=$Zf$+?To&v;vCf}i%QGj-gX?p$ z0-C!1{)}6mIq45}PVCQqjHgffS`qA6a*wS9_xnjbJ}ZOE@vQ>aPd)RrDp#&)oTQd;4>|1F+U7-a26SxOM*gUY@xd1g_8Bx@hY9 z55&qdck6+jJNvUAgeb#S* z_3~M7-R2ZE^A;y>HFNU!A34ixUHNG}%WQ+)mg433x2x*48N>EqW8~dG46Npu%|Skt zVk~nIr|phl#~;2E*jQPMox%F3o2$Q@QFFff5`R~)^R+BEF?K^!4<8P;U1IJI)<@l# z{(eU-{v*KpXHE71*VklEH1({>USPGX$w;tq)Dw3U*uFAPdxO=S`)e7GKl9xOZd>be zZr>NI?%M8$^>SX-M^n_CS8@8!WV2ZCJIm9}iZ`zMKGdzWWq0 zC!(q6Onm~_cIw6)hgD0=6T!yJcsjsp<(c{jZ?tJQFlGkk2ZUw3v6!Av;O8_p3|^i=BZCNMa?|L#&fOZ z>dE_5uyf>G&BV%6_a(6RtaY=n@|>wp1AC^njdl9Wz^c1mo~cgyv5}>)RXHx@GcbR*10|#td@KA9I)%*I%hr3MN`lF{yZ^7J#o(m zyRL~lAFSpcwx0{Ywo`Ww=3>>-$CtrsIU9ThtX7^4zFPU@9e*KQEk2$pFNQmp_Hhwb zo-^gwz@919t-l0oKjKTV#uHzL)i>+&b#Q%szJaE$|K(VD&H͓}998`pUHlxyN1 z%iLW7KAe*K{aawQoB^y;v%SwG*Z8Q4Z(DI~;(Z&u6(#Y$16E5s>(q?5HP(3BVDtR= zF4#R#zSpioQ%@h?18+!4AKwS-qwbtsiB*gL55Q{q+1(GpYRTzpuv+}D0qdXV&$VE+ z>`Tv1KZ4sXKG%WO^vOHz$8g(Pmv`FrV0E7dH(Uoyk3~p-GJNFj2vD96U^sUXa>{hV(x<>k&m$~1D^)h#TenwF< zcX7^IYRU0VaQ*)NIhuNW?y7vsar^>JJ29D=mu=!|xmoC8G zNAa@${;FP^xqbj_jI8TIu$uWeR{1X}#xk}zZGQzW-=z=1jg@tK7_4Tj?A>3(ZD*bN zKZsRJ%tyd#$>%p1~Brk?%t z1XwNa!Y9GTQFlJmk2d?~DX_UYPx_mKc|MKxGEaT}Kv6SKvGH6hxq9;c6WF|SFaH^= zmV5abu$puESFAr<5I zm%(b;JMzC%jAd-G{l0?Dp8Z!y; z_P=1~GyH#GV`YE71=dI1T;Ig1#s3}fl?8tn>>Q;1dtiOkZT~h_E$u%58`FL0cg=@j z_hET{{0L1w=f{u1wo^Ce`&hNa{1m)l!9N3+Inkt{kGkzY!K$VG5@2)6eD;E?xkmA6 zfjgG?w8GW&$@9NA+_u)`UhM-`_a5kn_3|E2?@Lj0jl}6|NpR-KbL&#@`n|d|ntJZl zWx#5=SC<7FN8SCKezdt)mjj!d`$KIcL3eNXH*)gStpj$ySm)o%$#dQw2==^f8|(D(Z|>Bc1JB!oz-r}rdtJDC&fDvO z)pFh*4EAyy)~!!bGjDM@4s~<&JURqyU*_T8^2zhF+Ks^etXMs{Z49pGwh5YgCo%om z@upz4a~O~Oi(q4^+ppSYV4uBnDxb~aYWkXsb!y3TOK|e^Zw%$hb1U#5VyY+4t-;JWI9x65ty8mo-Zj1#^KCba-tUP@l0(NhO9||vXI1H|jdU7}%tX}4D1UxxdU(ex4G=0qB zV5~fS9u02py=vbb19uIk!o7!&g{x&=jsx39Jvp_3)ytfYhbJfN>p8Wf>0?gK{|*ss zPQDAqfj!I1y%)yAuZH_>uzmtuZ8&S?dvGGy_etK%`ka8Ko^OQ{!M0O(59`wbwlC|_ z*GXW{j_K=UxLV>&0;~D^fW(;$Hb2HP*}Lr}|bDHS-fEpP^v$$vw9N*n6)0`|};)>N&&h1Xjx#ZfCG@ z)Sbihqs(zWeVfUu&LorYBr`XXIS5MvtgUvh7 zp+mrG8ONdE`urb;rk-cj;b7aTC+-no*D-OA1gqs9KMHI+_59u5G2nwJ?knpvZ^weI z^S*Td%2Rh7SS^2}csy9mzqO1{J6NCmz3Dix+E_~3jR&ij?Iys>b|=8qCQ{PwM6i0< zt^;1SI~lHa5+&^>fz@qy0C(h)*vXXRDdu8aeOx#7-1qLIl__2;P`vNGulpBx#R9Jc zb`PwA&3tr%^$9-(tY6l%3#^ZNe7eE*5k3uU8^4jwX*$KRs3*=0u=(x>ucpGb(*wZ1Oy_(sN!u3$B)P`fTv&6!kg8 z@hmcjavEheC9%&2+vhw=VxI$6PwaESYKeUwcrHafvCnT}p6}Si{xaCu7f=%WD`54+ z{wi24u`dMAr>H0PMPTDw7wv0c^@}O_dz{O_OHekUSbquSQcC_F=W?(CdInM{vO!aWTeBu<@-+>}$csF4y8maP`E#4y=~gKL%ezQO~?z4>rDa&ZE4X z$3bwfavlf5oyT>t$>&C}dO43jfvYEDfJl+LYPwZcS)e`$|@SPO(#J&e?eCwP?dFHVf zcnHO7JxZR1>lb)%fj21dhG6Gt(XYOynYL8I#cMNLr|1Eey!5;(blXv6q!1}1? zUG;mgdK<-<&VyQFJ_%N{&N+~$?hoJ%sygRGp1MDRT?gAaH}d>93(tU8rWnt5*6Zh- zTbCGr0y~G{e+Dlo_+P+|Zz0cje9sUBijQcmZ dKI#uvef=G5EbFqy{{Ux=&D}An>6`QY{{xWd?;-#I literal 0 HcmV?d00001 diff --git a/crates/pf-zerocopy/src/imp/mod.rs b/crates/pf-zerocopy/src/imp/mod.rs index a0374cc0..19c24036 100644 --- a/crates/pf-zerocopy/src/imp/mod.rs +++ b/crates/pf-zerocopy/src/imp/mod.rs @@ -15,6 +15,7 @@ pub mod client; pub mod cuda; pub mod egl; pub mod proto; +pub mod vkslot; pub mod vulkan; pub mod worker; diff --git a/crates/pf-zerocopy/src/imp/vkslot.rs b/crates/pf-zerocopy/src/imp/vkslot.rs new file mode 100644 index 00000000..92968d45 --- /dev/null +++ b/crates/pf-zerocopy/src/imp/vkslot.rs @@ -0,0 +1,710 @@ +//! Vulkan-allocated NVENC input slots + the Vulkan compute cursor blend — the driver-portable +//! replacement for the retired `cursor_blend.cu` PTX kernels (design: remote-desktop-sweep §8, +//! Phase A). A vendored PTX blob is JIT'd against the driver's ISA ceiling, so it silently dies +//! on drivers older than the generating toolkit (CUDA errors 222/218 on-glass — the KWin leg's +//! invisible composite cursor); SPIR-V has no such coupling. +//! +//! ```text +//! exportable VkBuffer ──vkGetMemoryFdKHR(OPAQUE_FD)──▶ cuImportExternalMemory ──▶ CUdeviceptr +//! ▲ │ NVENC registers + encodes +//! └── cursor_blend.comp dispatch (cursor rect only) ◀───┘ CUDA copies frames in +//! ``` +//! +//! The direct-SDK NVENC encoder allocates its input ring through [`VkSlotBlend::alloc_slot`] +//! instead of `cuMemAllocPitch`: same contiguous layouts (`InputSurface` docs), but the memory is +//! Vulkan external memory both APIs address. Per cursor-bearing frame the encoder CPU-syncs its +//! CUDA copy, then [`VkSlotBlend::blend`] dispatches the compute blend over the cursor's +//! rectangle and fence-waits — the same coherence ceremony [`super::vulkan::VkBridge`] ships for +//! its CSC (fence-ordered cross-API access on NVIDIA, no queue-family transfer needed). Frames +//! without a cursor never touch Vulkan, keeping the stream-ordered fast path intact. +//! +//! Falls back cleanly: if bring-up fails the encoder allocates plain CUDA surfaces and composite +//! mode degrades to no cursor (warned once) — never a failed session. + +// Every `unsafe` block in this file carries a `// SAFETY:` proof; enforce it (unsafe-proof program). +#![deny(clippy::undocumented_unsafe_blocks)] + +use super::cuda::{self, CUdeviceptr}; +use anyhow::{anyhow, Context as _, Result}; +use ash::vk; + +/// Max cursor-overlay bitmap edge (px) — matches [`cuda::CURSOR_MAX`] and the capture-side clamp. +pub const CURSOR_MAX: u32 = cuda::CURSOR_MAX; + +/// The vendored SPIR-V for `cursor_blend.comp` (beside this file; rebuild with +/// `glslc cursor_blend.comp -o cursor_blend.spv`). +const CURSOR_SPV: &[u8] = include_bytes!("cursor_blend.spv"); + +/// NVENC input-surface layout — selects the spec-constant `MODE` pipeline and the allocation +/// arithmetic (mirroring `InputSurface`'s contiguous layouts). +#[derive(Clone, Copy, PartialEq, Eq, Debug)] +pub enum SlotFormat { + /// Packed 4-byte ARGB (NVENC byte order B,G,R,A): `pitch × height`. + Argb, + /// NV12: Y rows `[0, H)` + interleaved UV rows `[H, 3H/2)` under one pitch. + Nv12, + /// Planar YUV444: three full-res planes stacked at `pitch × height` intervals. + Yuv444, +} + +impl SlotFormat { + fn mode(self) -> u32 { + match self { + SlotFormat::Argb => 0, + SlotFormat::Nv12 => 1, + SlotFormat::Yuv444 => 2, + } + } + fn row_bytes(self, width: u32) -> u64 { + match self { + SlotFormat::Argb => width as u64 * 4, + SlotFormat::Nv12 | SlotFormat::Yuv444 => width as u64, + } + } + fn rows(self, height: u32) -> u64 { + match self { + SlotFormat::Argb => height as u64, + SlotFormat::Nv12 => height as u64 + (height as u64 / 2).max(1), + SlotFormat::Yuv444 => height as u64 * 3, + } + } +} + +/// What the encoder holds per ring slot: the CUDA view it registers with NVENC plus the id it +/// hands back to [`VkSlotBlend::blend`]. The backing Vulkan objects + CUDA mapping live in the +/// [`VkSlotBlend`] (freed by [`free_slots`](VkSlotBlend::free_slots) / drop), so this is Copy — +/// the encoder's ring keeps its existing shape. +#[derive(Clone, Copy)] +pub struct VkSlotRef { + /// Device pointer NVENC registers (CUDA's mapping of the Vulkan memory). + pub ptr: CUdeviceptr, + /// Row stride in bytes (ours: row bytes rounded up to 256). + pub pitch: usize, + /// Luma rows (the plane-stride multiplier, as in `InputSurface`). + pub height: u32, + /// Index into the blend's slot table. + pub id: usize, +} + +/// One allocated slot's backing objects, freed together in reverse order (CUDA mapping first). +struct SlotAlloc { + buffer: vk::Buffer, + memory: vk::DeviceMemory, + /// CUDA's import of the exported OPAQUE_FD — must drop BEFORE the Vulkan memory is freed. + cuda: cuda::ExternalDmabuf, + size: u64, +} + +/// 28-byte push-constant block matching `cursor_blend.comp`'s `Push`. +#[repr(C)] +struct Push { + pitch: u32, + surf_w: u32, + surf_h: u32, + cur_w: u32, + cur_h: u32, + ox: i32, + oy: i32, +} + +pub struct VkSlotBlend { + _entry: ash::Entry, + instance: ash::Instance, + device: ash::Device, + ext_fd: ash::khr::external_memory_fd::Device, + queue: vk::Queue, + cmd_pool: vk::CommandPool, + cmd: vk::CommandBuffer, + fence: vk::Fence, + mem_props: vk::PhysicalDeviceMemoryProperties, + shader: vk::ShaderModule, + desc_layout: vk::DescriptorSetLayout, + pipe_layout: vk::PipelineLayout, + desc_pool: vk::DescriptorPool, + desc_set: vk::DescriptorSet, + /// One pipeline per [`SlotFormat`], indexed by `mode()` (spec constant). + pipelines: [vk::Pipeline; 3], + /// Host-visible cursor bitmap staging (CURSOR_MAX²·4, tight rows), persistently mapped. + cur_buf: vk::Buffer, + cur_mem: vk::DeviceMemory, + cur_map: *mut u8, + slots: Vec, +} + +// SAFETY: raw Vulkan handles + a persistently-mapped pointer, all uniquely owned by this struct +// and destroyed exactly once in `Drop`; used from the encoder thread but moved with it. `Send` +// only (not `Sync`), matching the single-thread use — transferring opaque handles cannot dangle. +unsafe impl Send for VkSlotBlend {} + +impl VkSlotBlend { + /// Bring up the device + blend pipelines. Requires the CUDA shared context (the encoder's) to + /// be established; picks the NVIDIA physical device (the NVENC path is NVIDIA by definition). + pub fn new() -> Result { + // SAFETY: standard ash bring-up, same shape as `VkBridge::new` — every call is `unsafe` + // only because ash cannot statically verify handle/CreateInfo validity. Every + // `*CreateInfo`/`AllocateInfo` is built from locals that live for the duration of the + // synchronous call reading them; every handle passed was created and `?`-checked in this + // same function. Shares nothing across threads. + unsafe { + let entry = ash::Entry::load().context("load libvulkan")?; + let app = vk::ApplicationInfo::default().api_version(vk::API_VERSION_1_1); + let instance = entry + .create_instance( + &vk::InstanceCreateInfo::default().application_info(&app), + None, + ) + .context("vkCreateInstance")?; + let phys = match instance + .enumerate_physical_devices() + .context("enumerate GPUs")? + .into_iter() + .find(|&p| instance.get_physical_device_properties(p).vendor_id == 0x10DE) + { + Some(p) => p, + None => { + instance.destroy_instance(None); + return Err(anyhow!("no NVIDIA Vulkan device")); + } + }; + let mem_props = instance.get_physical_device_memory_properties(phys); + let qf = match instance + .get_physical_device_queue_family_properties(phys) + .iter() + .position(|q| q.queue_flags.contains(vk::QueueFlags::COMPUTE)) + { + Some(i) => i as u32, + None => { + instance.destroy_instance(None); + return Err(anyhow!("no compute-capable queue family")); + } + }; + let prio = [1.0f32]; + let qci = [vk::DeviceQueueCreateInfo::default() + .queue_family_index(qf) + .queue_priorities(&prio)]; + let exts = [ash::khr::external_memory_fd::NAME.as_ptr()]; + let device = match instance.create_device( + phys, + &vk::DeviceCreateInfo::default() + .queue_create_infos(&qci) + .enabled_extension_names(&exts), + None, + ) { + Ok(d) => d, + Err(e) => { + instance.destroy_instance(None); + return Err(e).context("vkCreateDevice (external_memory_fd supported?)"); + } + }; + // From here teardown-on-error goes through `destroy_partial`, which tolerates null + // handles — build everything into an incrementally-filled struct. + let ext_fd = ash::khr::external_memory_fd::Device::new(&instance, &device); + let queue = device.get_device_queue(qf, 0); + let mut me = VkSlotBlend { + _entry: entry, + instance, + device, + ext_fd, + queue, + cmd_pool: vk::CommandPool::null(), + cmd: vk::CommandBuffer::null(), + fence: vk::Fence::null(), + mem_props, + shader: vk::ShaderModule::null(), + desc_layout: vk::DescriptorSetLayout::null(), + pipe_layout: vk::PipelineLayout::null(), + desc_pool: vk::DescriptorPool::null(), + desc_set: vk::DescriptorSet::null(), + pipelines: [vk::Pipeline::null(); 3], + cur_buf: vk::Buffer::null(), + cur_mem: vk::DeviceMemory::null(), + cur_map: std::ptr::null_mut(), + slots: Vec::new(), + }; + me.init_objects(qf).inspect_err(|_| { + // `Drop` runs the same teardown and tolerates the nulls left by a partial init. + })?; + tracing::info!( + "Vulkan slot blend ready (exportable NVENC inputs + SPIR-V cursor blend)" + ); + Ok(me) + } + } + + /// The non-device objects: command machinery, cursor staging, descriptor + pipelines. + fn init_objects(&mut self, qf: u32) -> Result<()> { + // SAFETY: same contract as `new` — ash calls on the live `self.device` with builder infos + // from locals outliving each synchronous call; created handles are stored into `self` + // immediately so the caller's `Drop` frees them on any later failure. + unsafe { + let d = &self.device; + self.cmd_pool = d + .create_command_pool( + &vk::CommandPoolCreateInfo::default() + .queue_family_index(qf) + .flags(vk::CommandPoolCreateFlags::RESET_COMMAND_BUFFER), + None, + ) + .context("create command pool")?; + self.cmd = d + .allocate_command_buffers( + &vk::CommandBufferAllocateInfo::default() + .command_pool(self.cmd_pool) + .level(vk::CommandBufferLevel::PRIMARY) + .command_buffer_count(1), + ) + .context("allocate command buffer")?[0]; + self.fence = d + .create_fence(&vk::FenceCreateInfo::default(), None) + .context("create fence")?; + + // Cursor staging: host-visible+coherent SSBO, persistently mapped. + let cur_size = (CURSOR_MAX * CURSOR_MAX * 4) as u64; + self.cur_buf = d + .create_buffer( + &vk::BufferCreateInfo::default() + .size(cur_size) + .usage(vk::BufferUsageFlags::STORAGE_BUFFER), + None, + ) + .context("create cursor buffer")?; + let reqs = d.get_buffer_memory_requirements(self.cur_buf); + let mem_type = self + .memory_type( + reqs.memory_type_bits, + vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT, + ) + .context("cursor buffer memory type")?; + self.cur_mem = d + .allocate_memory( + &vk::MemoryAllocateInfo::default() + .allocation_size(reqs.size) + .memory_type_index(mem_type), + None, + ) + .context("allocate cursor memory")?; + d.bind_buffer_memory(self.cur_buf, self.cur_mem, 0) + .context("bind cursor memory")?; + self.cur_map = d + .map_memory(self.cur_mem, 0, cur_size, vk::MemoryMapFlags::empty()) + .context("map cursor memory")? as *mut u8; + + // Descriptor set: binding 0 = surface SSBO (rebound per blend), 1 = cursor SSBO. + let bindings = [ + vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::COMPUTE), + vk::DescriptorSetLayoutBinding::default() + .binding(1) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::COMPUTE), + ]; + self.desc_layout = d + .create_descriptor_set_layout( + &vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings), + None, + ) + .context("create descriptor layout")?; + let pc = [vk::PushConstantRange::default() + .stage_flags(vk::ShaderStageFlags::COMPUTE) + .size(std::mem::size_of::() as u32)]; + let dl = [self.desc_layout]; + self.pipe_layout = d + .create_pipeline_layout( + &vk::PipelineLayoutCreateInfo::default() + .set_layouts(&dl) + .push_constant_ranges(&pc), + None, + ) + .context("create pipeline layout")?; + let pool_sizes = [vk::DescriptorPoolSize::default() + .ty(vk::DescriptorType::STORAGE_BUFFER) + .descriptor_count(2)]; + self.desc_pool = d + .create_descriptor_pool( + &vk::DescriptorPoolCreateInfo::default() + .max_sets(1) + .pool_sizes(&pool_sizes), + None, + ) + .context("create descriptor pool")?; + let dls = [self.desc_layout]; + self.desc_set = d + .allocate_descriptor_sets( + &vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(self.desc_pool) + .set_layouts(&dls), + ) + .context("allocate descriptor set")?[0]; + + // The shader + one pipeline per MODE (spec constant 0). + if CURSOR_SPV.len() % 4 != 0 { + anyhow::bail!("cursor_blend.spv is not word-aligned"); + } + let words: Vec = CURSOR_SPV + .chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect(); + self.shader = d + .create_shader_module(&vk::ShaderModuleCreateInfo::default().code(&words), None) + .context("create blend shader module")?; + for mode in 0u32..3 { + let entries = [vk::SpecializationMapEntry::default() + .constant_id(0) + .offset(0) + .size(4)]; + let data = mode.to_le_bytes(); + let spec = vk::SpecializationInfo::default() + .map_entries(&entries) + .data(&data); + let stage = vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::COMPUTE) + .module(self.shader) + .name(c"main") + .specialization_info(&spec); + let info = [vk::ComputePipelineCreateInfo::default() + .stage(stage) + .layout(self.pipe_layout)]; + let p = d + .create_compute_pipelines(vk::PipelineCache::null(), &info, None) + .map_err(|(_, e)| e) + .context("create blend pipeline")?[0]; + self.pipelines[mode as usize] = p; + } + } + Ok(()) + } + + fn memory_type(&self, type_bits: u32, flags: vk::MemoryPropertyFlags) -> Result { + (0..self.mem_props.memory_type_count) + .find(|&i| { + type_bits & (1 << i) != 0 + && self.mem_props.memory_types[i as usize] + .property_flags + .contains(flags) + }) + .ok_or_else(|| anyhow!("no memory type for flags {flags:?}")) + } + + /// Allocate one NVENC input slot as exportable Vulkan memory mapped into CUDA. Layout matches + /// `InputSurface` (contiguous planes under one pitch); pitch = row bytes rounded to 256. + pub fn alloc_slot(&mut self, fmt: SlotFormat, width: u32, height: u32) -> Result { + let pitch = (fmt.row_bytes(width) + 255) & !255; + let size = pitch * fmt.rows(height); + // SAFETY: exportable-buffer allocation, the exact `VkBridge::ensure_dst` incantation: + // `ExternalMemoryBufferCreateInfo`/`ExportMemoryAllocateInfo` declare OPAQUE_FD, + // `MemoryDedicatedAllocateInfo` ties the memory to the buffer; every info is a local + // outliving its synchronous call and every failure path destroys the objects created so + // far exactly once. `get_memory_fd` hands us an fd that `import_owned_fd` either adopts + // (driver owns it) or closes on failure. + unsafe { + let d = &self.device; + let mut ext_info = vk::ExternalMemoryBufferCreateInfo::default() + .handle_types(vk::ExternalMemoryHandleTypeFlags::OPAQUE_FD); + let buffer = d + .create_buffer( + &vk::BufferCreateInfo::default() + .size(size) + .usage(vk::BufferUsageFlags::STORAGE_BUFFER) + .push_next(&mut ext_info), + None, + ) + .context("create slot buffer")?; + let reqs = d.get_buffer_memory_requirements(buffer); + let mem_type = match self + .memory_type(reqs.memory_type_bits, vk::MemoryPropertyFlags::DEVICE_LOCAL) + { + Ok(t) => t, + Err(e) => { + d.destroy_buffer(buffer, None); + return Err(e); + } + }; + let mut export = vk::ExportMemoryAllocateInfo::default() + .handle_types(vk::ExternalMemoryHandleTypeFlags::OPAQUE_FD); + let mut dedicated = vk::MemoryDedicatedAllocateInfo::default().buffer(buffer); + let memory = match d.allocate_memory( + &vk::MemoryAllocateInfo::default() + .allocation_size(reqs.size) + .memory_type_index(mem_type) + .push_next(&mut export) + .push_next(&mut dedicated), + None, + ) { + Ok(m) => m, + Err(e) => { + d.destroy_buffer(buffer, None); + return Err(e).context("allocate exportable slot memory"); + } + }; + if let Err(e) = d.bind_buffer_memory(buffer, memory, 0) { + d.free_memory(memory, None); + d.destroy_buffer(buffer, None); + return Err(e).context("bind slot memory"); + } + let fd = match self.ext_fd.get_memory_fd( + &vk::MemoryGetFdInfoKHR::default() + .memory(memory) + .handle_type(vk::ExternalMemoryHandleTypeFlags::OPAQUE_FD), + ) { + Ok(f) => f, + Err(e) => { + d.free_memory(memory, None); + d.destroy_buffer(buffer, None); + return Err(e).context("vkGetMemoryFdKHR(slot)"); + } + }; + let ext = match cuda::ExternalDmabuf::import_owned_fd(fd, reqs.size) { + Ok(c) => c, + Err(e) => { + d.free_memory(memory, None); + d.destroy_buffer(buffer, None); + return Err(e).context("cuImportExternalMemory(slot OPAQUE_FD)"); + } + }; + let r = VkSlotRef { + ptr: ext.ptr, + pitch: pitch as usize, + height, + id: self.slots.len(), + }; + self.slots.push(SlotAlloc { + buffer, + memory, + cuda: ext, + size, + }); + Ok(r) + } + } + + /// Free every allocated slot (encoder teardown, alongside its ring clear). CUDA mappings drop + /// first (field order in [`SlotAlloc`] frees `cuda` via its own `Drop` before we free the VK + /// objects explicitly here). + pub fn free_slots(&mut self) { + for s in self.slots.drain(..) { + drop(s.cuda); // CUDA's view of the memory goes first + // SAFETY: `buffer`/`memory` were created in `alloc_slot`, are uniquely owned by the + // drained `SlotAlloc`, and are destroyed exactly once here. No queue work can be + // in flight: every `blend` fence-waits before returning. + unsafe { + self.device.destroy_buffer(s.buffer, None); + self.device.free_memory(s.memory, None); + } + } + } + + /// Upload the cursor RGBA (`cw*ch*4`, tight rows) into the mapped staging buffer. Call only + /// when the bitmap changes; position moves are push constants. + pub fn upload_cursor(&mut self, rgba: &[u8], cw: u32, ch: u32) { + let cw = cw.min(CURSOR_MAX); + let ch = ch.min(CURSOR_MAX); + let len = (cw * ch * 4) as usize; + let len = len.min(rgba.len()); + // SAFETY: `cur_map` is the live persistent mapping of the CURSOR_MAX²·4 host-coherent + // allocation (created in `init_objects`, unmapped only in `Drop`); `len` is clamped to + // both the source slice and the buffer capacity. No blend is in flight (every `blend` + // fence-waits before returning), so no GPU read races this host write. + unsafe { + std::ptr::copy_nonoverlapping(rgba.as_ptr(), self.cur_map, len); + } + } + + /// Blend the uploaded cursor into `slot` at `(ox, oy)`: record, submit, fence-wait. The + /// caller has CPU-synced its CUDA frame copy first; the fence wait makes the shader's writes + /// visible to the subsequent NVENC encode (the `VkBridge` precedent: fence-ordered cross-API + /// access, no queue-family transfer — NVIDIA-only path). + #[allow(clippy::too_many_arguments)] // surface geometry + cursor rect — unpacked kernel args + pub fn blend_ref( + &mut self, + slot: &VkSlotRef, + fmt: SlotFormat, + surf_w: u32, + cw: u32, + ch: u32, + ox: i32, + oy: i32, + ) -> Result<()> { + let alloc = self + .slots + .get(slot.id) + .ok_or_else(|| anyhow!("bad slot id {}", slot.id))?; + let cw = cw.min(CURSOR_MAX); + let ch = ch.min(CURSOR_MAX); + if cw == 0 || ch == 0 { + return Ok(()); + } + let push = Push { + pitch: slot.pitch as u32, + surf_w, + surf_h: slot.height, + cur_w: cw, + cur_h: ch, + ox, + oy, + }; + // Dispatch geometry (must match cursor_blend.comp): ARGB = per cursor px; NV12/YUV444 = + // per word-aligned 4-px span × (2-row blocks | rows). + let (gx, gy) = match fmt { + SlotFormat::Argb => (cw.div_ceil(8), ch.div_ceil(8)), + _ => { + let x0 = (ox >> 2) << 2; + let spans = ((ox + cw as i32) - x0 + 3).div_euclid(4).max(1) as u32; + let rows = match fmt { + SlotFormat::Nv12 => ch.div_ceil(2), + _ => ch, + }; + (spans.div_ceil(8), rows.div_ceil(8)) + } + }; + // SAFETY: single-threaded record/submit/wait on handles this struct owns. The descriptor + // update is safe because no prior submission is in flight (every blend fence-waits and + // the fence is reset before reuse). Buffer infos and barrier structs are locals outliving + // their synchronous calls. The dispatch's shader accesses stay in-bounds by the shader's + // own guards (surfW/surfH/curW/curH from `push`) against the slot allocation sized in + // `alloc_slot` for exactly that geometry. + unsafe { + let d = &self.device; + let surf_info = [vk::DescriptorBufferInfo::default() + .buffer(alloc.buffer) + .offset(0) + .range(alloc.size)]; + let cur_info = [vk::DescriptorBufferInfo::default() + .buffer(self.cur_buf) + .offset(0) + .range((CURSOR_MAX * CURSOR_MAX * 4) as u64)]; + let writes = [ + vk::WriteDescriptorSet::default() + .dst_set(self.desc_set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .buffer_info(&surf_info), + vk::WriteDescriptorSet::default() + .dst_set(self.desc_set) + .dst_binding(1) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .buffer_info(&cur_info), + ]; + d.update_descriptor_sets(&writes, &[]); + + d.begin_command_buffer( + self.cmd, + &vk::CommandBufferBeginInfo::default() + .flags(vk::CommandBufferUsageFlags::ONE_TIME_SUBMIT), + ) + .context("begin blend cmd")?; + // CUDA wrote the frame into this memory outside Vulkan's view — make it visible to + // the shader (external-memory coherence ceremony; NVIDIA honors this with the fence + // ordering alone, the barrier is the spec-shaped belt-and-braces). + let acquire = [vk::MemoryBarrier::default() + .src_access_mask(vk::AccessFlags::MEMORY_WRITE) + .dst_access_mask(vk::AccessFlags::SHADER_READ | vk::AccessFlags::SHADER_WRITE)]; + d.cmd_pipeline_barrier( + self.cmd, + vk::PipelineStageFlags::TOP_OF_PIPE, + vk::PipelineStageFlags::COMPUTE_SHADER, + vk::DependencyFlags::empty(), + &acquire, + &[], + &[], + ); + d.cmd_bind_pipeline( + self.cmd, + vk::PipelineBindPoint::COMPUTE, + self.pipelines[fmt.mode() as usize], + ); + d.cmd_bind_descriptor_sets( + self.cmd, + vk::PipelineBindPoint::COMPUTE, + self.pipe_layout, + 0, + &[self.desc_set], + &[], + ); + let bytes = std::slice::from_raw_parts( + (&push as *const Push) as *const u8, + std::mem::size_of::(), + ); + d.cmd_push_constants( + self.cmd, + self.pipe_layout, + vk::ShaderStageFlags::COMPUTE, + 0, + bytes, + ); + d.cmd_dispatch(self.cmd, gx.max(1), gy.max(1), 1); + // Release the shader's writes so the post-fence CUDA/NVENC reads see them. + let release = [vk::MemoryBarrier::default() + .src_access_mask(vk::AccessFlags::SHADER_WRITE) + .dst_access_mask(vk::AccessFlags::MEMORY_READ)]; + d.cmd_pipeline_barrier( + self.cmd, + vk::PipelineStageFlags::COMPUTE_SHADER, + vk::PipelineStageFlags::BOTTOM_OF_PIPE, + vk::DependencyFlags::empty(), + &release, + &[], + &[], + ); + d.end_command_buffer(self.cmd).context("end blend cmd")?; + let cmds = [self.cmd]; + let submit = [vk::SubmitInfo::default().command_buffers(&cmds)]; + d.queue_submit(self.queue, &submit, self.fence) + .context("submit blend")?; + let r = d.wait_for_fences(&[self.fence], true, 1_000_000_000); + d.reset_fences(&[self.fence]).ok(); + r.context("blend fence wait")?; + } + Ok(()) + } +} + +impl Drop for VkSlotBlend { + fn drop(&mut self) { + self.free_slots(); + // SAFETY: every handle below was created in `new`/`init_objects` (or is null from a + // partial init — Vulkan destroy/free calls are defined no-ops on null handles) and is + // uniquely owned; each is destroyed exactly once here, pipelines/layouts/pools before the + // device, the device before the instance. No work is in flight (`blend` fence-waits). + unsafe { + let d = &self.device; + for p in self.pipelines { + if p != vk::Pipeline::null() { + d.destroy_pipeline(p, None); + } + } + if self.shader != vk::ShaderModule::null() { + d.destroy_shader_module(self.shader, None); + } + if self.desc_pool != vk::DescriptorPool::null() { + d.destroy_descriptor_pool(self.desc_pool, None); + } + if self.pipe_layout != vk::PipelineLayout::null() { + d.destroy_pipeline_layout(self.pipe_layout, None); + } + if self.desc_layout != vk::DescriptorSetLayout::null() { + d.destroy_descriptor_set_layout(self.desc_layout, None); + } + if !self.cur_map.is_null() { + d.unmap_memory(self.cur_mem); + } + if self.cur_buf != vk::Buffer::null() { + d.destroy_buffer(self.cur_buf, None); + } + if self.cur_mem != vk::DeviceMemory::null() { + d.free_memory(self.cur_mem, None); + } + if self.fence != vk::Fence::null() { + d.destroy_fence(self.fence, None); + } + if self.cmd_pool != vk::CommandPool::null() { + d.destroy_command_pool(self.cmd_pool, None); + } + d.destroy_device(None); + self.instance.destroy_instance(None); + } + } +}