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<cuda::CursorBlend>,
+    /// 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<VkSlotBlend>,
     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<CursorBlend> {
-        // 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<CUfunction> {
-            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(&copy, "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 00000000..046d149a
Binary files /dev/null and b/crates/pf-zerocopy/src/imp/cursor_blend.spv differ
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<SlotAlloc>,
+}
+
+// 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<VkSlotBlend> {
+        // 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::<Push>() 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<u32> = 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<u32> {
+        (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<VkSlotRef> {
+        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::<Push>(),
+            );
+            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);
+        }
+    }
+}