perf(encode/nvenc-linux): LN2 v1 — stream-ordered submit via NvEncSetIOCudaStreams

Bind the session's IO streams to the encode thread's high-priority copy
stream: in sync-retrieve depth-1 use the per-frame input copy and cursor
blend now enqueue with NO cuStreamSynchronize and encode_picture orders
after them on the stream. Same stream both directions, so the encode's
completion is inserted into the stream and later work (the next frame's
copy into a reused ring slot) waits for it.

Soundness gate: the fast path engages only when pending is empty (true
depth-1 usage) — every prior encode was drained by a blocking poll, and
the caller holds the frame payload across the matching poll (contract now
documented on Encoder::submit; both host loops already comply). Pipelined
callers and PUNKTFUNK_NVENC_ASYNC mode keep the blocking copies.

True zero-copy input registration (registering the worker-owned IPC
buffer directly) stays the LN2 v2 follow-up — it needs a contiguous
worker-pool NV12 layout and a registration<->IPC-mapping lifetime tie.

PUNKTFUNK_NVENC_STREAM_ORDERED=0 restores the old blocking behavior.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

crates/pf-encode/src/enc/codec.rs

@@ -219,6 +219,11 @@ pub struct EncoderCaps {
 
 /// A hardware encoder. One per session; runs on the encode thread.
 pub trait Encoder: Send {
+    /// Submit one captured frame for encoding. Lifetime contract: the caller must keep `frame`
+    /// (and its GPU payload) alive until this frame's AU has been returned by
+    /// [`poll`](Self::poll) — a stream-ordered backend (Linux direct-NVENC's IO-stream binding)
+    /// may still be reading the payload asynchronously after `submit` returns. Both host encode
+    /// loops already hold the frame across their poll drain; new callers must do the same.
     fn submit(&mut self, frame: &CapturedFrame) -> Result<()>;
     /// [`submit`](Self::submit) with the **wire frame index** this frame's AU will carry — the
     /// number the packetizer stamps on it and the client's loss reports/RFI requests name. The

crates/pf-encode/src/enc/linux/mod.rs

@@ -826,7 +826,7 @@ impl NvencEncoder {
                         (*f).linesize[i] as usize,
                     )
                 });
-                pf_zerocopy::cuda::copy_yuv444_to_device(buf, dsts)
+                pf_zerocopy::cuda::copy_yuv444_to_device(buf, dsts, true)
             } else if self.want_444 {
                 ffi::av_frame_free(&mut f);
                 bail!(
@@ -839,11 +839,11 @@ impl NvencEncoder {
                 let y_pitch = (*f).linesize[0] as usize;
                 let uv_ptr = (*f).data[1] as pf_zerocopy::cuda::CUdeviceptr;
                 let uv_pitch = (*f).linesize[1] as usize;
-                pf_zerocopy::cuda::copy_nv12_to_device(buf, y_ptr, y_pitch, uv_ptr, uv_pitch)
+                pf_zerocopy::cuda::copy_nv12_to_device(buf, y_ptr, y_pitch, uv_ptr, uv_pitch, true)
             } else {
                 let dst_ptr = (*f).data[0] as pf_zerocopy::cuda::CUdeviceptr;
                 let dst_pitch = (*f).linesize[0] as usize;
-                pf_zerocopy::cuda::copy_device_to_device(buf, dst_ptr, dst_pitch)
+                pf_zerocopy::cuda::copy_device_to_device(buf, dst_ptr, dst_pitch, true)
             };
             if let Err(e) = copy_res {
                 ffi::av_frame_free(&mut f);

crates/pf-encode/src/enc/linux/nvenc_cuda.rs

@@ -16,9 +16,14 @@
 //!     ([`zerocopy::cuda::InputSurface`]): each captured `FramePayload::Cuda` `DeviceBuffer` is
 //!     device→device copied into the current ring slot (via the existing `copy_*_to_device`
 //!     helpers) before `encode_picture`. This mirrors the libav path's recycled-hwframe-pool copy
-//!     (NVENC rejects a null-`buf[0]` frame and its CUDADEVICEPTR registration cache is bounded +
+//!     (NVENC rejects a null-`buf[0]` frame; the captured buffer is worker-owned CUDA-IPC memory
-//!     pointer-keyed, so registering a fresh pool pointer each frame would thrash it) — so it is
+//!     recycled on drop, so registering it directly needs a contiguous worker-pool layout + a
-//!     zero regression versus today; true zero-copy input registration is a follow-up.
+//!     registration↔IPC-mapping lifetime tie — the true zero-copy follow-up, plan §7 LN2 v2).
+//!     **Stream-ordered submit** (default, `PUNKTFUNK_NVENC_STREAM_ORDERED=0` reverts): the
+//!     session's IO streams are bound to the encode thread's copy stream
+//!     (`NvEncSetIOCudaStreams`), so in sync-retrieve depth-1 use the copy + cursor blend enqueue
+//!     with NO per-frame `cuStreamSynchronize` and the encode orders after them on the stream —
+//!     the submit path's CPU stalls are gone even though the copy itself remains.
 //!
 //! **Two-thread retrieve** (`PUNKTFUNK_NVENC_ASYNC=1`, the same opt-in knob as the Windows
 //! backend — gpu-contention plan §5.B, latency plan T2.2): NVENC *async mode*
@@ -120,6 +125,14 @@ struct EncodeApi {
     encode_picture:
         unsafe extern "C" fn(*mut c_void, *mut nv::NV_ENC_PIC_PARAMS) -> nv::NVENCSTATUS,
     invalidate_ref_frames: unsafe extern "C" fn(*mut c_void, u64) -> nv::NVENCSTATUS,
+    /// `NvEncSetIOCudaStreams` — binds the session's input/output ordering to a CUDA stream so the
+    /// input copy + cursor blend can enqueue without a CPU sync (stream-ordered submit). The two
+    /// `NV_ENC_CUSTREAM_PTR` args are pointers TO `CUstream` values.
+    set_io_cuda_streams: unsafe extern "C" fn(
+        *mut c_void,
+        nv::NV_ENC_CUSTREAM_PTR,
+        nv::NV_ENC_CUSTREAM_PTR,
+    ) -> nv::NVENCSTATUS,
 }
 
 /// Resolve the table once per process. `Err` = NVENC genuinely unavailable (no NVIDIA driver/.so,
@@ -212,6 +225,7 @@ fn load_api() -> std::result::Result<EncodeApi, String> {
             unmap_input_resource: list.nvEncUnmapInputResource.ok_or(MISSING)?,
             encode_picture: list.nvEncEncodePicture.ok_or(MISSING)?,
             invalidate_ref_frames: list.nvEncInvalidateRefFrames.ok_or(MISSING)?,
+            set_io_cuda_streams: list.nvEncSetIOCudaStreams.ok_or(MISSING)?,
         };
         std::mem::forget(lib); // keep the .so mapped for the fn pointers' lifetime (process)
         Ok(api)
@@ -245,6 +259,18 @@ fn async_inflight_cap() -> usize {
         .clamp(2, POOL - 1)
 }
 
+/// Stream-ordered submit (`PUNKTFUNK_NVENC_STREAM_ORDERED`, default ON; `0` = the pre-existing
+/// blocking copies). With the session's IO streams bound to the encode thread's copy stream
+/// (`NvEncSetIOCudaStreams`), the input copy + cursor blend enqueue with NO CPU sync and
+/// `encode_picture` orders after them on the stream — deleting the 1–3 per-frame
+/// `cuStreamSynchronize` stalls from the submit path (latency plan §7 LN2). Sync-retrieve mode
+/// only, and only while nothing is in flight (see the gate in [`Encoder::submit`]).
+fn stream_ordered_requested() -> bool {
+    std::env::var("PUNKTFUNK_NVENC_STREAM_ORDERED")
+        .map(|v| v.trim() != "0")
+        .unwrap_or(true)
+}
+
 /// One in-flight encode handed to the retrieve thread: the output bitstream to (blocking-)lock.
 /// Raw pointer travels as `usize` (a process-global driver handle; the thread is joined before
 /// the session it belongs to is destroyed).
@@ -460,6 +486,14 @@ pub struct NvencCudaEncoder {
     /// The two-thread retrieve runtime (`PUNKTFUNK_NVENC_ASYNC`) — `None` in the default
     /// single-thread mode and between sessions. Exists only `init_session`→`teardown`.
     async_rt: Option<AsyncRetrieve>,
+    /// Boxed `CUstream` the session's IO-stream binding points at (`NvEncSetIOCudaStreams` takes
+    /// POINTERS to `CUstream`, and this struct moves — the pointee needs a stable heap address for
+    /// the session's lifetime). Null when stream-ordering is off; freed in `teardown` AFTER the
+    /// session is destroyed.
+    io_stream: *mut *mut c_void,
+    /// Stream-ordered submit armed for the live session (sync-retrieve mode only; see
+    /// [`stream_ordered_requested`]). The per-frame gate additionally requires `pending` empty.
+    stream_ordered: bool,
 }
 
 // SAFETY: the `!Send` fields are the raw NVENC session handle (`encoder`), the shared `CUcontext`
@@ -539,6 +573,8 @@ impl NvencCudaEncoder {
             split_mode: nv::NV_ENC_SPLIT_ENCODE_MODE::NV_ENC_SPLIT_DISABLE_MODE as u32,
             last_rfi_range: None,
             async_rt: None,
+            io_stream: ptr::null_mut(),
+            stream_ordered: false,
         })
     }
 
@@ -578,6 +614,14 @@ impl NvencCudaEncoder {
                  session's slot toward the concurrent-session cap"
             );
         }
+        // The boxed CUstream the IO-stream binding pointed at — freed only now, AFTER the session
+        // that referenced it is destroyed (created by `Box::into_raw` in `init_session`, freed
+        // exactly once here; `io_stream` is nulled so a re-init can't double-free).
+        if !self.io_stream.is_null() {
+            drop(Box::from_raw(self.io_stream));
+            self.io_stream = ptr::null_mut();
+        }
+        self.stream_ordered = false;
         self.ring.clear(); // drops the InputSurfaces, freeing their CUDA allocations
         self.bitstreams.clear();
         self.pending.clear();
@@ -834,7 +878,11 @@ impl NvencCudaEncoder {
         // `try_open_session` just returned (and `best` only when non-null). `create_bitstream_buffer`
         // and `register_resource` take `enc`, the chosen live session, and `&mut` locals whose
         // `version` is set and which outlive the synchronous call. `InputSurface::alloc_*` returns a
-        // live pitched CUDA allocation on the shared context. No handle escapes the encode thread.
+        // live pitched CUDA allocation on the shared context. `set_io_cuda_streams` takes `enc` plus
+        // two pointers to the boxed live `CUstream` (`Box::into_raw`), which outlives the session —
+        // freed exactly once: in `teardown` after `destroy_encoder` when armed, or via
+        // `Box::from_raw` right here on the rejection path (where `io_stream` is never set). No
+        // handle escapes the encode thread.
         unsafe {
             // Bind to the shared CUDA context; make it current on this (encode) thread for both the
             // session open and every subsequent device→device input copy.
@@ -994,6 +1042,45 @@ impl NvencCudaEncoder {
                     "NVENC two-thread retrieve enabled (submit thread + blocking-lock thread)"
                 );
             }
+            // Stream-ordered submit (latency plan §7 LN2): bind the session's IO streams to this
+            // thread's copy stream so the input copy + cursor blend enqueue with no CPU sync and
+            // `encode_picture` orders after them. Same stream both ways: input-stream semantics
+            // start the encode only after our enqueued copies, output-stream semantics insert the
+            // encode's completion INTO the stream — so later stream work (the next frame's copy
+            // into a reused ring slot) also waits for it. Sync-retrieve mode only: in two-thread
+            // mode the captured buffer may be recycled after `submit` returns while the stream
+            // still holds its copy (the blocking copies are the lifetime guarantee there).
+            if self.async_rt.is_none() && stream_ordered_requested() {
+                let stream = cuda::copy_stream_handle();
+                if !stream.is_null() {
+                    // The pointee must outlive the session (the driver takes CUstream POINTERS) —
+                    // box it; `teardown` frees it after `destroy_encoder`.
+                    let holder = Box::into_raw(Box::new(stream));
+                    match (api().set_io_cuda_streams)(
+                        enc,
+                        holder as nv::NV_ENC_CUSTREAM_PTR,
+                        holder as nv::NV_ENC_CUSTREAM_PTR,
+                    )
+                    .nv_ok()
+                    {
+                        Ok(()) => {
+                            self.io_stream = holder;
+                            self.stream_ordered = true;
+                            tracing::info!(
+                                "NVENC stream-ordered submit armed (IO streams bound — no CPU \
+                                 sync in the submit path)"
+                            );
+                        }
+                        Err(e) => {
+                            drop(Box::from_raw(holder));
+                            tracing::debug!(
+                                status = ?e,
+                                "NvEncSetIOCudaStreams rejected — keeping blocking copies"
+                            );
+                        }
+                    }
+                }
+            }
             tracing::info!(
                 mode = %format_args!("{}x{}@{}", self.width, self.height, self.fps),
                 bit_depth = self.bit_depth,
@@ -1007,8 +1094,10 @@ impl NvencCudaEncoder {
     }
 
     /// Copy the captured `DeviceBuffer` into the ring slot's registered input surface (device→device
-    /// on the shared context, synchronized by the copy helpers).
+    /// on the shared context). `sync` blocks until the copy completes (the pre-existing behavior);
-    fn copy_into_slot(&self, buf: &cuda::DeviceBuffer, slot: usize) -> Result<()> {
+    /// `!sync` enqueues on the encode thread's copy stream and leaves ordering to the session's
+    /// 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;
@@ -1023,16 +1112,16 @@ impl NvencCudaEncoder {
                     (base + pitch as u64 * hh, pitch),
                     (base + 2 * pitch as u64 * hh, pitch),
                 ];
-                cuda::copy_yuv444_to_device(buf, planes)
+                cuda::copy_yuv444_to_device(buf, planes, sync)
             }
             nv::NV_ENC_BUFFER_FORMAT::NV_ENC_BUFFER_FORMAT_NV12 => {
                 if !buf.is_nv12() {
                     bail!("NV12 session but the captured buffer has no chroma plane");
                 }
                 // Contiguous NV12: UV follows Y at base + pitch*height, same pitch.
-                cuda::copy_nv12_to_device(buf, base, pitch, base + pitch as u64 * hh, pitch)
+                cuda::copy_nv12_to_device(buf, base, pitch, base + pitch as u64 * hh, pitch, sync)
             }
-            _ => cuda::copy_device_to_device(buf, base, pitch),
+            _ => cuda::copy_device_to_device(buf, base, pitch, sync),
         }
     }
 
@@ -1149,10 +1238,20 @@ impl Encoder for NvencCudaEncoder {
         // loop's `submit_us` folds all four together; this is what splits them apart.
         let sample = pf_host_config::config().perf && self.frames % 120 == 0;
         self.frames += 1;
+
+        // Stream-ordered fast path (§7 LN2): enqueue the copy + blend with no CPU sync and let the
+        // IO-stream binding order `encode_picture` after them — but ONLY while nothing is in
+        // flight (true depth-1 usage). The gate is what makes this sound: with `pending` empty,
+        // every prior encode was drained by a blocking `poll`, so (a) the ring slot being reused
+        // was fully read, and (b) the caller still holds this frame's payload across the matching
+        // `poll` (both host loops do — see `Encoder::submit`'s doc), which blocks until the encode
+        // (and therefore the enqueued copy) completed. A pipelined caller (pending non-empty)
+        // falls back to the blocking copy so an early-recycled source can never be read late.
+        let ordered = self.stream_ordered && self.pending.is_empty();
         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)?;
+        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
@@ -1196,12 +1295,12 @@ impl Encoder for NvencCudaEncoder {
                 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)
+                        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)
+                        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),
+                    _ => cb.blend_argb(s.ptr, s.pitch, w, h, ov.w, ov.h, ov.x, ov.y, !ordered),
                 };
                 if let Err(e) = r {
                     if !self.cursor_blend_warned {
@@ -1227,8 +1326,11 @@ impl Encoder for NvencCudaEncoder {
         // `pic` (`NV_ENC_PIC_PARAMS`, version set) points `inputBuffer` at `mp.mappedResource` and
         // `outputBitstream` at the live pool bitstream `bitstreams[slot]`; the optional SEI scratch is
         // stack-local and outlives the synchronous `encode_picture`. The input surface for `slot` was
-        // just filled by the (synchronized) device→device copy and is not overwritten until this slot
+        // just filled by the device→device copy — either synchronized (blocking mode) or ordered
-        // is reused POOL submits later, by which time this encode was polled (POOL ≥ in-flight depth).
+        // before this encode by the session's IO-stream binding (`ordered` — same stream, see the
+        // gate above) — and is not overwritten until this slot is reused POOL submits later, by
+        // which time this encode was polled (POOL ≥ in-flight depth; in ordered mode the poll's
+        // blocking lock additionally proves the enqueued copy completed).
         unsafe {
             let reg = self.ring[slot].reg;
             let mut mp = nv::NV_ENC_MAP_INPUT_RESOURCE {

crates/pf-zerocopy/src/imp/cuda.rs

@@ -251,6 +251,32 @@ unsafe fn copy_blocking(copy: &CUDA_MEMCPY2D, what: &str) -> Result<()> {
     ck(cuStreamSynchronize(stream), "cuStreamSynchronize")
 }
 
+/// Issue `copy` on this thread's priority stream WITHOUT waiting — for stream-ordered consumers
+/// only (the direct-NVENC submit path with `NvEncSetIOCudaStreams` bound to this stream): the
+/// stream, not the CPU, orders completion, so the SOURCE must stay valid until the downstream
+/// stream work (the encode) has finished.
+unsafe fn copy_async(copy: &CUDA_MEMCPY2D, what: &str) -> Result<()> {
+    ck(cuMemcpy2DAsync_v2(copy, copy_stream()), what)
+}
+
+/// `copy_blocking` when `sync`, else `copy_async` — the shared tail of the public `copy_*_to_device`
+/// helpers, whose `sync: false` mode carries `copy_async`'s source-lifetime contract.
+unsafe fn copy_issue(copy: &CUDA_MEMCPY2D, what: &str, sync: bool) -> Result<()> {
+    if sync {
+        copy_blocking(copy, what)
+    } else {
+        copy_async(copy, what)
+    }
+}
+
+/// The calling thread's copy/launch stream as a raw handle, for binding external stream-ordering
+/// (the direct-NVENC `NvEncSetIOCudaStreams` hookup). Null = the NULL stream (priority-stream
+/// creation failed) — callers should treat null as "stream-ordering unavailable" and keep their
+/// blocking copies. The shared context must be current on this thread.
+pub fn copy_stream_handle() -> *mut c_void {
+    copy_stream() // CUstream IS *mut c_void (opaque CUstream_st*)
+}
+
 /// Max cursor-overlay bitmap edge (px) uploaded to the device blend buffer — matches the Vulkan path.
 pub const CURSOR_MAX: u32 = 256;
 
@@ -354,6 +380,7 @@ impl CursorBlend {
         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);
@@ -370,7 +397,7 @@ impl CursorBlend {
             &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)
+        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)`.
@@ -385,6 +412,7 @@ impl CursorBlend {
         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);
@@ -401,7 +429,7 @@ impl CursorBlend {
             &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)
+        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`).
@@ -416,6 +444,7 @@ impl CursorBlend {
         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);
@@ -441,16 +470,20 @@ impl CursorBlend {
             (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, then synchronize.
+    /// 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(());
@@ -458,9 +491,11 @@ impl CursorBlend {
         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); grid/block
+        // locals whose types match the kernel's C parameters (per the call site above) — CUDA
-        // dims are non-zero. Launched on the copy stream (ordered after the input-surface copy that
+        // copies the parameter values during `cuLaunchKernel` itself, so they need not outlive
-        // `copy_into_slot` already synchronized) then synchronized. Requires the context current.
+        // 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(
@@ -478,7 +513,10 @@ impl CursorBlend {
                 ),
                 "cuLaunchKernel(cursor)",
             )?;
-            ck(cuStreamSynchronize(stream), "cuStreamSynchronize(cursor)")
+            if sync {
+                ck(cuStreamSynchronize(stream), "cuStreamSynchronize(cursor)")?;
+            }
+            Ok(())
         }
     }
 }
@@ -1128,10 +1166,13 @@ pub fn copy_mapped_yuv444(
 /// Copy a pitched device buffer into another device region (device→device), e.g. our imported
 /// [`DeviceBuffer`] into a pooled CUDA surface NVENC owns. Both are 4-byte (BGRx) pixels.
 /// The caller must have the shared context current on this thread (see [`make_current`]).
+/// `sync: false` enqueues without a CPU wait (stream-ordered consumers only — `src` must stay
+/// valid until the downstream stream work completes; see [`copy_stream_handle`]).
 pub fn copy_device_to_device(
     src: &DeviceBuffer,
     dst_ptr: CUdeviceptr,
     dst_pitch: usize,
+    sync: bool,
 ) -> Result<()> {
     let copy = CUDA_MEMCPY2D {
         srcMemoryType: CU_MEMORYTYPE_DEVICE,
@@ -1144,23 +1185,27 @@ pub fn copy_device_to_device(
         Height: src.height as usize,
         ..Default::default()
     };
-    // SAFETY: `copy_blocking` is unsafe (issues a CUDA copy); the caller must have the shared
+    // SAFETY: `copy_issue` is unsafe (issues a CUDA copy); the caller must have the shared
     // context current (documented). `&copy` is a live local device→device `CUDA_MEMCPY2D` outliving
-    // the synchronous call: `srcDevice`/`srcPitch` are `src`'s live allocation, `dstDevice`/
+    // the enqueue: `srcDevice`/`srcPitch` are `src`'s live allocation, `dstDevice`/`dstPitch` the
-    // `dstPitch` the caller's live region, `width*4`×`height` within both. Wrapper → live table.
+    // caller's live region, `width*4`×`height` within both; `sync: false` shifts the source-
-    unsafe { copy_blocking(&copy, "cuMemcpy2DAsync_v2(dev->dev)") }
+    // lifetime obligation to the caller (documented above). Wrapper → live table.
+    unsafe { copy_issue(&copy, "cuMemcpy2DAsync_v2(dev->dev)", sync) }
 }
 
 /// Copy our imported NV12 [`DeviceBuffer`] (Y + UV planes) into NVENC's two-plane CUDA surface
 /// `(y_dst, y_pitch)` / `(uv_dst, uv_pitch)` (`av_hwframe_get_buffer`'s `data[0]`/`data[1]` +
 /// `linesize[0]`/`linesize[1]`). The Y plane is `width`×`height` bytes; the chroma plane is
 /// `(width/2)·2` bytes × `height/2` rows. The caller must have the shared context current.
+/// `sync: false` enqueues without a CPU wait (stream-ordered consumers only — `src` must stay
+/// valid until the downstream stream work completes; see [`copy_stream_handle`]).
 pub fn copy_nv12_to_device(
     src: &DeviceBuffer,
     y_dst: CUdeviceptr,
     y_pitch: usize,
     uv_dst: CUdeviceptr,
     uv_pitch: usize,
+    sync: bool,
 ) -> Result<()> {
     let (src_uv_ptr, src_uv_pitch) = src
         .uv
@@ -1189,15 +1234,16 @@ pub fn copy_nv12_to_device(
         Height: h / 2,
         ..Default::default()
     };
-    // SAFETY: two unsafe `copy_blocking` device→device copies; the caller must have the shared
+    // SAFETY: two unsafe `copy_issue` device→device copies; the caller must have the shared
     // context current (documented). `&y`/`&uv` are live local `CUDA_MEMCPY2D`s outliving each
-    // synchronous call. All four device pointers are valid: `src.ptr`/`src_uv_ptr` come from a live
+    // enqueue. All four device pointers are valid: `src.ptr`/`src_uv_ptr` come from a live
     // NV12 `DeviceBuffer` (its `.uv` presence was checked via `ok_or_else`), `y_dst`/`uv_dst` are
     // the caller's live NVENC surface planes; the luma copy is `w`×`h`, the chroma copy
-    // `(w/2)*2`×`h/2`, each within its planes. Wrappers → live table.
+    // `(w/2)*2`×`h/2`, each within its planes; `sync: false` shifts the source-lifetime obligation
+    // to the caller (documented above). Wrappers → live table.
     unsafe {
-        copy_blocking(&y, "cuMemcpy2DAsync_v2(nv12 Y dev->dev)")?;
+        copy_issue(&y, "cuMemcpy2DAsync_v2(nv12 Y dev->dev)", sync)?;
-        copy_blocking(&uv, "cuMemcpy2DAsync_v2(nv12 UV dev->dev)")
+        copy_issue(&uv, "cuMemcpy2DAsync_v2(nv12 UV dev->dev)", sync)
     }
 }
 
@@ -1205,7 +1251,13 @@ pub fn copy_nv12_to_device(
 /// (`av_hwframe_get_buffer`'s `data[0..3]` + `linesize[0..3]` for a `yuv444p` frames context).
 /// Each plane is `width`×`height` bytes; the source planes sit at row offsets `0/H/2H` of the
 /// single allocation. The caller must have the shared context current.
-pub fn copy_yuv444_to_device(src: &DeviceBuffer, dsts: [(CUdeviceptr, usize); 3]) -> Result<()> {
+/// `sync: false` enqueues without a CPU wait (stream-ordered consumers only — `src` must stay
+/// valid until the downstream stream work completes; see [`copy_stream_handle`]).
+pub fn copy_yuv444_to_device(
+    src: &DeviceBuffer,
+    dsts: [(CUdeviceptr, usize); 3],
+    sync: bool,
+) -> Result<()> {
     anyhow::ensure!(src.yuv444, "copy_yuv444_to_device on a non-YUV444 buffer");
     let w = src.width as usize;
     let h = src.height as usize;
@@ -1221,12 +1273,13 @@ pub fn copy_yuv444_to_device(src: &DeviceBuffer, dsts: [(CUdeviceptr, usize); 3]
             Height: h,
             ..Default::default()
         };
-        // SAFETY: unsafe `copy_blocking` device→device copy; the caller must have the shared
+        // SAFETY: unsafe `copy_issue` device→device copy; the caller must have the shared
-        // context current (documented). `&copy` is a live local outliving the synchronous call;
+        // context current (documented). `&copy` is a live local outliving the enqueue;
         // `src.ptr + pitch·h·i` stays within the live 3·H-row stacked allocation (`yuv444`
         // checked above), `dst_ptr`/`dst_pitch` is the caller's live NVENC plane; `w`×`h` fits
-        // both. Wrapper → live table.
+        // both; `sync: false` shifts the source-lifetime obligation to the caller (documented
-        unsafe { copy_blocking(&copy, "cuMemcpy2DAsync_v2(yuv444 plane dev->dev)")? };
+        // above). Wrapper → live table.
+        unsafe { copy_issue(&copy, "cuMemcpy2DAsync_v2(yuv444 plane dev->dev)", sync)? };
     }
     Ok(())
 }