refactor(host/W6.2): extract the video encode backends into the pf-encode crate

encode.rs + encode/* (NVENC, VAAPI, native AMF, AMF/QSV ffmpeg, direct-SDK
NVENC/CUDA, raw Vulkan-Video, PyroWave, openh264) move into crates/pf-encode
behind one Encoder trait + open_video selector (plan §W6). The crate speaks the
shared frame vocabulary (pf-frame: CapturedFrame/PixelFormat + the DXGI identity
D3d11Frame/make_device) and pf-zerocopy (CUDA context/buffers), and NEVER
pf-capture — the capture→encode edge is one-way (ZeroCopyPolicy, prior commit).

Dep moves: the heavy encoder deps (ffmpeg-next, the NVENC SDK, openh264,
pyrowave-sys) move from the host to pf-encode; the host's
nvenc/amf-qsv/vulkan-encode/pyrowave features now FORWARD to pf-encode/*. The
host keeps a mod-encode shim (pub use pf_encode) so every crate::encode::* path
(negotiator + GameStream/native/mgmt planes) is unchanged.

resolve_render_adapter_luid moves from the host's windows/win_adapter.rs into
pf-gpu (both pf-encode and pf-capture need it as a peer of GPU selection); its 5
call sites (encode amf/nvenc, capture idd_push/synthetic_nv12, vdisplay manager)
rewire to pf_gpu::resolve_render_adapter_luid and win_adapter.rs is deleted.
pf-frame's make_device gains a # Safety section (public-unsafe-fn lint, latent
since the pf-frame carve — a full-workspace -D warnings clippy catches it).

Verified: Linux clippy -D warnings (pf-encode + host nvenc,vulkan-encode,pyrowave
--all-targets) + 13/13 pf-encode + 299/299 host tests; Windows clippy -D warnings
(pf-encode nvenc,amf-qsv --all-targets + host nvenc,amf-qsv --all-targets)
Finished exit 0.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

crates/pf-encode/src/enc/linux/cursor_blend.cu

@@ -0,0 +1,105 @@
+// 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);
+}