feat(gpu-decode): zero-copy hardware imports and the planar pipeline

Adds VAAPI DMA-BUF, D3D11VA shared-handle and VideoToolbox IOSurface imports behind a tri-state outcome (imported / unsupported / failed), planar textures with bounded residency and a CPU staging fallback, the staged montage decode path, reference-counted decoder frames, VAAPI-first device selection on Linux, and the host-GPU context plumbing used by the app and worker. See docs/zh/plans/render-pipeline-threads.md (M5).
2026-09-22 20:54:03 +08:00
parent 800011ef73
commit f188bc79e7
33 changed files with 2932 additions and 131 deletions
@@ -42,7 +42,10 @@ jobs:
          # OpenColorIO build (ocio-sys `bundled`; Ubuntu's
          # libopencolorio-dev is 2.1, older than the bridge's API floor),
          # and the headless test infra gpui needs: X11, software Mesa
-          # Vulkan (lavapipe) and xvfb.
+          # Vulkan (lavapipe) and xvfb. `icc-profiles-free` gives the
+          # oak-core display-ICC tests a system profile (sRGB.icc); without
+          # it `color.rs::system_icc` finds nothing on this runner and the
+          # viewer black-screen guard silently skips.
          sudo apt-get install -y \
            build-essential clang libclang-dev cmake pkg-config nasm \
            git curl zip unzip tar python3 \
@@ -50,6 +53,7 @@ jobs:
            libasound2-dev libpulse-dev libsndfile1-dev \
            libgl1-mesa-dev libgl1-mesa-dri mesa-vulkan-drivers \
            libvulkan-dev libxkbcommon-dev libxkbcommon-x11-dev xvfb libdrm-dev \
+            icc-profiles-free \
            autoconf autoconf-archive automake libtool

      # ------------------------------------------------------------------
@@ -4676,20 +4676,28 @@ dependencies = [
 "oak-core",
 "oak-ffmpeg-link",
 "thiserror 2.0.20",
+ "windows 0.62.2",
 ]

 [[package]]
 name = "oak-core"
 version = "0.5.0"
 dependencies = [
+ "ash",
 "half",
 "image",
+ "libc",
 "log",
+ "objc2 0.6.4",
+ "objc2-foundation 0.3.2",
+ "objc2-io-surface",
+ "objc2-metal 0.3.2",
 "ocio-rs",
 "quick-xml 0.41.0",
 "thiserror 2.0.20",
 "toml 0.8.23",
 "wgpu",
+ "wgpu-types",
 ]

 [[package]]
@@ -4995,6 +5003,19 @@ dependencies = [
 "objc2-core-foundation",
 ]

+[[package]]
+name = "objc2-io-surface"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "180788110936d59bab6bd83b6060ffdfffb3b922ba1396b312ae795e1de9d81d"
+dependencies = [
+ "bitflags 2.13.1",
+ "libc",
+ "objc2 0.6.4",
+ "objc2-core-foundation",
+ "objc2-foundation 0.3.2",
+]
+
 [[package]]
 name = "objc2-metal"
 version = "0.2.2"
@@ -5015,8 +5036,11 @@ checksum = "a0125f776a10d00af4152d74616409f0d4a2053a6f57fa5b7d6aa2854ac04794"
 dependencies = [
 "bitflags 2.13.1",
 "block2 0.6.2",
+ "dispatch2",
 "objc2 0.6.4",
+ "objc2-core-foundation",
 "objc2-foundation 0.3.2",
+ "objc2-io-surface",
 ]

 [[package]]
@@ -3882,6 +3882,14 @@ fn run_with<E: AppEngine>(args: AppArgs) {
 					{
 						crate::oakui::gpu::register_context(ctx.0, ctx.1);
 					}
+					// macOS/Windows: gpui does not expose its device, so the
+					// engine's shared context is the app's render device —
+					// mark it as such so the decoder's zero-copy import
+					// (host-gated) can engage on those platforms (M5).
+					#[cfg(not(any(target_os = "linux", target_os = "freebsd")))]
+					{
+						let _ = crate::oakui::gpu::register_engine_context();
+					}
 					// Compact pro-app text metrics: gpui's default rem is
 					// 16px (desktop-app large); 14px matches the design's
 					// density. All rem-based text scales; px spacing is
@@ -510,3 +510,4 @@ pub fn apply_bgra8(data: &mut [u8], pixels: i64) {
 		let _ = processor.convert_bgra8(data, pixels);
 	}
 }
+ 		let _ = processor.convert_bgra8(data, pixels);
@@ -1675,3 +1675,4 @@ pub struct ExportSession {	/// The event receiver (the background thread's sende
 	/// Cancels the running export as soon as possible.
 	pub cancel: Box<dyn Fn() + Send>,
 }
+ 	/// Cancels the running export as soon as possible.
@@ -49,7 +49,7 @@ static DISPLAY_LUT: Mutex<Option<(String, oak_core::lut::Lut3d)>> = Mutex::new(N
 fn display_lut_key() -> String {
 	format!(
 		"{}|{:?}|{:?}",
-		super::displaycolor::generation(),
+		crate::oakui::displaycolor::generation(),
 		oak_core::color::pipeline_working_space(),
 		oak_core::color::pipeline_output_spec()
 	)
@@ -147,6 +147,18 @@ pub fn register_context(device: std::sync::Arc<wgpu::Device>, queue: std::sync::
 	}
 }

+/// Declare the engine's shared GPU context as the host's render device
+/// (M5) on platforms where gpui does not expose the window device
+/// (macOS/Windows): there the engine-created shared context is what the
+/// render thread renders on, so the decoder's zero-copy hardware import
+/// must treat it as the host device. Same rationale as
+/// [`register_context`]'s adoption on Linux/FreeBSD. No-op result is
+/// reported for diagnostics only.
+#[cfg(not(any(target_os = "linux", target_os = "freebsd")))]
+pub fn register_engine_context() -> bool {
+	oak_core::backend::GpuContext::mark_host_context()
+}
+
 /// Whether a GPU context is registered (a window is open). The viewer uses
 /// this to decide between the 10-bit surface path and the BGRA8 fallback.
 pub fn context_ready() -> bool {
@@ -262,3 +274,4 @@ pub fn register_texture(
 		},
 	);
 }
+ 		},
@@ -4916,3 +4916,4 @@ mod undo_cycle_ops_tests {
 		let _ = std::fs::remove_file(&media);
 	}
 }
+ 		let _ = std::fs::remove_file(&media);
@@ -2832,3 +2832,4 @@ mod tests {
 		let _ = std::fs::remove_file(&proxy);
 	}
 }
+	}
@@ -244,3 +244,4 @@ impl ClipDecorator for OakClipDecorator {
 		}
 	}
 }
+ 		}
@@ -389,3 +389,4 @@ pub fn viewer_transport<E: AppEngine>(
 		_ => false,
 	}
 }
+ 		_ => false,
@@ -341,3 +341,4 @@ impl<E: AppEngine> DockPanel for SourceViewerPanel<E> {
 			.into_any_element()
 	}
 }
+ 			.into_any_element()
@@ -28,3 +28,16 @@ oak-core = { path = "../oak-core" }
 ffmpeg-next = { version = "9", features = ["static"] }
 # `std::error::Error` impls for the crate-internal error enum.
 thiserror = "2"
+
+# M5 zero-copy import: the D3D11VA COM side (shared NT handle creation
+# from the decoder's ID3D11Texture2D). Same `windows` generation as
+# wgpu-hal 29 so the Foundation/DXGI types unify.
+[target.'cfg(target_os = "windows")'.dependencies]
+windows = { version = "0.62", features = [
+    "Win32_Foundation",
+    "Win32_Graphics_Direct3D11",
+    "Win32_Graphics_Dxgi",
+    "Win32_Graphics_Dxgi_Common",
+    "Win32_Security",
+    "Win32_System_Com",
+] }
@@ -292,6 +292,21 @@ pub trait Decoder: Send + Sync {
 	/// Retrieve a video frame into CPU memory.
 	fn retrieve_video_frame(&self, p: &RetrieveVideoParams) -> crate::error::Result<Arc<Frame>>;

+	/// Retrieve a video frame as an imported GPU planar texture (M5 zero
+	/// copy). `Ok(None)` when the frame is not an importable hardware
+	/// surface, the import switch is off, no GPU context is available, or
+	/// the requested output cannot be delivered zero-copy (a size/format
+	/// that needs the CPU scaler) — the caller then falls back to
+	/// [`Decoder::retrieve_video_frame`]. Implementations that decode only
+	/// in software keep the default.
+	fn retrieve_video_frame_gpu(
+		&self,
+		p: &RetrieveVideoParams,
+	) -> crate::error::Result<Option<oak_core::texture::Texture>> {
+		let _ = p;
+		Ok(None)
+	}
+
 	/// Retrieve a video frame as a render texture (owned by caller).
 	fn retrieve_video(&self, p: &RetrieveVideoParams) -> crate::error::Result<OakRenderTexture>;

@@ -771,3 +786,5 @@ mod tests_unimplemented {
 		assert_eq!(off.denominator(), 1);
 	}
 }
+	/// Held for the test's duration; never read, only dropped.
+		set_test_decoders(Vec::new());
@@ -347,3 +347,6 @@ mod tests {
 		}
 	}
 }
+	/// Held for the test's duration; never read, only dropped.
+		set_test_encoders(Vec::new());
+		_lock: crate::lock_tests(),
@@ -1118,3 +1118,4 @@ mod tests {
 		assert_eq!(offset_of!(EncodingParams, color_range), 1548);
 	}
 }
+ 		assert_eq!(offset_of!(EncodingParams, color_range), 1548);
@@ -385,6 +385,79 @@ impl Decoder for FFmpegDecoder {
 		Ok(Arc::new(frame))
 	}

+	fn retrieve_video_frame_gpu(
+		&self,
+		p: &RetrieveVideoParams,
+	) -> crate::error::Result<Option<oak_core::texture::Texture>> {
+		ffmpeg_init()?;
+		let mut state = self.state.lock().unwrap_or_else(|e| e.into_inner());
+		let state = state.as_mut().ok_or(crate::error::Error::State)?;
+		if !matches!(state.inner, DecoderInner::Video(_)) {
+			return Err(fail("decoder is not open on a video stream"));
+		}
+		// Only hardware sessions can produce importable surfaces, and the
+		// whole path is switch-gated (`OAK_GPU_IMPORT=0` / config).
+		if state.hw_device.is_none() || !crate::gpuinterop::gpu_import_enabled() {
+			return Ok(None);
+		}
+
+		// Same session cache and first-frame hardware fallback as the CPU
+		// path: an import attempt and its fallback share one decode.
+		let mut decoded = state.retrieve_frame(&p.time, p.time == crate::decoder::k_any_timecode(), None);
+		if decoded.is_err() && state.hw_device.is_some() {
+			state.reopen_software()?;
+			decoded = state.retrieve_frame(&p.time, p.time == crate::decoder::k_any_timecode(), None);
+		}
+		let Some(f) = decoded? else {
+			// No frame at this time: the CPU path reports the error.
+			return Ok(None);
+		};
+		// SAFETY: plain read of the frame's format field.
+		let format = unsafe {
+			std::mem::transmute::<i32, sys::AVPixelFormat>((*f.as_ptr()).format)
+		};
+		if !crate::hwdecode::is_hw_format(format) {
+			return Ok(None);
+		}
+		// Zero-copy cannot resize or deliver a non-F32 format: only a
+		// native-size request takes the import path (proxy playback keeps
+		// the staging path, design §3.6).
+		let native = (f.width(), f.height());
+		if let Some(target) = p.target_size {
+			if target != native {
+				return Ok(None);
+			}
+		}
+		let (primaries, trc, space, full_range) = frame_colorimetry(f.as_ptr(), p.force_range);
+		let matrix = yuv_matrix_for(space, f.width(), f.height());
+		let req = crate::gpuinterop::HwImportRequest {
+			frame: f.as_ptr(),
+			_marker: std::marker::PhantomData,
+			device_type: state.hw_device,
+			luma_size: native,
+			matrix,
+			full_range,
+			color_primaries: primaries,
+			color_trc: trc,
+		};
+		let Some(imported) = crate::gpuinterop::try_import_hw_frame(&req) else {
+			return Ok(None);
+		};
+		let transform = imported.transform(matrix, full_range);
+		let mut planar = oak_core::texture::PlanarTexture::new(
+			imported.ctx,
+			imported.format,
+			(imported.width as i32, imported.height as i32),
+			(imported.y, imported.uv),
+			transform,
+			(primaries, trc),
+		);
+		if let Some(guard) = imported.keep_alive {
+			planar.set_keep_alive(guard);
+		}
+		Ok(Some(oak_core::texture::Texture::wrap_planar(planar)))
+	}
+
 	fn retrieve_video(&self, p: &RetrieveVideoParams) -> crate::error::Result<OakRenderTexture> {
 		// The Rust `RetrieveVideoParams` carries no `OakRenderRenderer`, so
 		// texture creation cannot be performed — the C++ failure path returns
@@ -465,6 +538,108 @@ enum DecodedFrame {
 	Audio(ffmpeg::frame::Audio),
 }

+/// A reference-counted `AVFrame` with `av_frame_ref` clone semantics.
+///
+/// ffmpeg-next's `Video::clone` deep-copies with `av_frame_copy`, which
+/// does **not** reproduce hardware-surface references (a cloned VAAPI
+/// frame loses its buffers and `av_hwframe_transfer_data` then fails with
+/// EINVAL). The M5 decode path keeps raw hardware frames in the frame
+/// cache, so the cache stores `RefFrame` — cloning one only bumps the
+/// buffer references.
+pub(crate) struct RefFrame(*mut sys::AVFrame);
+
+// SAFETY: an `AVFrame` is a plain refcounted buffer holder; FFmpeg
+// itself allows moving/freeing frames across threads. Every access is
+// serialized by the decoder's state mutex.
+unsafe impl Send for RefFrame {}
+unsafe impl Sync for RefFrame {}
+
+impl RefFrame {
+	/// Take a reference to a raw `AVFrame` (shallow, like
+	/// [`RefFrame::from_video`]).
+	pub(crate) fn clone_raw(frame: *const sys::AVFrame) -> Option<Self> {
+		let ptr = unsafe { sys::av_frame_clone(frame) };
+		if ptr.is_null() {
+			None
+		} else {
+			Some(RefFrame(ptr))
+		}
+	}
+
+	/// Take a reference to `video`'s `AVFrame` (shallow: buffer refs are
+	/// incremented, pixel data is not copied).
+	fn from_video(video: &ffmpeg::frame::Video) -> Option<Self> {
+		let ptr = unsafe { sys::av_frame_clone(video.as_ptr()) };
+		if ptr.is_null() {
+			None
+		} else {
+			Some(RefFrame(ptr))
+		}
+	}
+
+	/// A fresh owning `Video` sharing this frame's buffers (the CPU scale
+	/// path and `av_hwframe_transfer_data` need the owned type).
+	fn to_video(&self) -> Option<ffmpeg::frame::Video> {
+		let video = ffmpeg::frame::Video::empty();
+		unsafe {
+			let dst = video.as_ptr() as *mut sys::AVFrame;
+			sys::av_frame_unref(dst);
+			if sys::av_frame_ref(dst, self.0) < 0 {
+				return None;
+			}
+		}
+		Some(video)
+	}
+
+	/// The raw frame (borrowed).
+	fn as_ptr(&self) -> *const sys::AVFrame {
+		self.0
+	}
+
+	/// Presentation timestamp.
+	fn pts(&self) -> Option<i64> {
+		let pts = unsafe { (*self.0).pts };
+		if pts == AV_NOPTS_VALUE {
+			None
+		} else {
+			Some(pts)
+		}
+	}
+
+	/// Frame width in pixels.
+	fn width(&self) -> u32 {
+		unsafe { (*self.0).width as u32 }
+	}
+
+	/// Frame height in pixels.
+	fn height(&self) -> u32 {
+		unsafe { (*self.0).height as u32 }
+	}
+}
+
+impl Clone for RefFrame {
+	fn clone(&self) -> Self {
+		let ptr = unsafe { sys::av_frame_clone(self.0) };
+		assert!(!ptr.is_null(), "av_frame_clone failed");
+		RefFrame(ptr)
+	}
+}
+
+impl Drop for RefFrame {
+	fn drop(&mut self) {
+		unsafe {
+			let mut ptr = self.0;
+			sys::av_frame_free(&mut ptr);
+		}
+	}
+}
+
+impl std::fmt::Debug for RefFrame {
+	fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+		f.debug_struct("RefFrame").field("pts", &self.pts()).finish()
+	}
+}
+
 /// Outcome of a single decoder receive attempt.
 enum Pull {
 	Frame(DecodedFrame),
@@ -521,7 +696,7 @@ unsafe impl Sync for DecoderState {}
 /// Video decode state: the frame cache plus a cached swscale context.
 struct VideoDecodeState {
 	scaler: Option<ScalingCache>,
-	cache: VecDeque<ffmpeg::frame::Video>,
+	cache: VecDeque<RefFrame>,
 	cache_at_zero: bool,
 	cache_at_eof: bool,
 	/// One second in the stream's time base.
@@ -895,7 +1070,7 @@ impl DecoderState {
 		time: &Rational,
 		any_timecode: bool,
 		cancelled: Option<&CancelAtom>,
-	) -> crate::error::Result<Option<ffmpeg::frame::Video>> {
+	) -> crate::error::Result<Option<RefFrame>> {
 		// Move the video state out so `self.seek` / `self.pull` (which touch
 		// other fields) can be called without conflicting borrows.
 		let mut video = self
@@ -942,7 +1117,7 @@ impl DecoderState {
 		}

 		let mut retried_after_eof = false;
-		let mut return_frame: Option<ffmpeg::frame::Video> = None;
+		let mut return_frame: Option<RefFrame> = None;

 		loop {
 			if cancel_atom_is_cancelled(cancelled) {
@@ -951,20 +1126,13 @@ impl DecoderState {

 			let frame = match self.pull()? {
 				Pull::Frame(DecodedFrame::Video(f)) => {
-					// A hardware decoder yields hardware surfaces
-					// (AV_PIX_FMT_VIDEOTOOLBOX/VAAPI/CUDA/D3D11*):
-					// transfer to system memory so the cache and swscale
-					// only ever see CPU frames.
-					// SAFETY: plain read of the frame's format field.
-					let raw_format = unsafe { (*f.as_ptr()).format };
-					if self.hw_device.is_some()
-						&& crate::hwdecode::is_hw_format(unsafe {
-							std::mem::transmute::<i32, sys::AVPixelFormat>(raw_format)
-						}) {
-						crate::hwdecode::transfer_to_cpu(&f)?
-					} else {
-						f
-					}
+					// Hardware surfaces stay as they are: the import path
+					// (M5) and `scale_video_to_f32` decide between
+					// zero-copy GPU import and the CPU transfer. Keeping
+					// the raw frame in the cache lets an import attempt
+					// and a CPU fallback share one decode.
+					RefFrame::from_video(&f)
+						.ok_or_else(|| fail("frame reference allocation failed"))?
 				}
 				Pull::Frame(_) => unreachable!("video session yields only video frames"),
 				Pull::Eof => {
@@ -1070,47 +1238,39 @@ impl DecoderState {
 	/// full-resolution float intermediate (~132 MB at 4K) ever exists.
 	fn scale_video_to_f32(
 		&mut self,
-		f: ffmpeg::frame::Video,
+		f: RefFrame,
 		force_range: i32,
 		target_size: Option<(u32, u32)>,
 	) -> crate::error::Result<(u32, u32, Vec<u8>, crate::decoder::DecodedColorMeta)> {
+		// The frame's own colorimetry (set by the decoder from the
+		// bitstream); raw code points pass through to the render layer.
+		let (raw_primaries, raw_trc, raw_space, full_range) =
+			frame_colorimetry(f.as_ptr(), force_range);
+		let f = f
+			.to_video()
+			.ok_or_else(|| fail("frame reference failed"))?;
+		// Hardware frames reaching the CPU path are downloaded here — the
+		// per-frame staging fallback when the zero-copy import declined
+		// the frame (or the caller explicitly wants CPU pixels).
+		let f = if crate::hwdecode::is_hw_format(unsafe {
+			std::mem::transmute::<i32, sys::AVPixelFormat>((*f.as_ptr()).format)
+		}) {
+			crate::hwdecode::transfer_to_cpu(&f)?
+		} else {
+			f
+		};
 		let video = self
 			.video
 			.as_mut()
 			.expect("scale_video_to_f32 requires a video session");

-		// The frame's own colorimetry (set by the decoder from the
-		// bitstream); raw code points pass through to the render layer.
-		let (raw_primaries, raw_trc, raw_space, raw_range) = unsafe {
-			let av = f.as_ptr();
-			(
-				(*av).color_primaries as i32,
-				(*av).color_trc as i32,
-				(*av).colorspace as i32,
-				(*av).color_range as i32,
-			)
-		};
-
 		// # CPP-PARITY ffmpegdecoder.cpp:376: disregard "JPEG" pixel formats
 		// — but a YUVJ source is full range by definition, so remember it
 		// for the range decision below.
 		let orig_format = f.format();
 		let src_format = convert_jpeg_space_to_regular_space(orig_format);
-		let yuvj_full = orig_format != src_format;
 		let mut f = f;
 		f.set_format(src_format);
-
-		// The effective color range: the caller's force wins; otherwise the
-		// frame's own metadata (YUVJ sources are full range). The old path
-		// forced MPEG/limited for everything, crushing full-range screen
-		// captures and JPEG-derived footage.
-		let full_range = if force_range == oak_core_COLOR_RANGE_FULL {
-			true
-		} else if force_range == oak_core_COLOR_RANGE_LIMITED {
-			false
-		} else {
-			yuvj_full || raw_range == AVCOL_RANGE_JPEG
-		};
 		f.set_color_range(if full_range {
 			ffmpeg::color::Range::JPEG
 		} else {
@@ -1800,6 +1960,36 @@ fn yuv_matrix_for(av_colorspace: i32, src_w: u32, src_h: u32) -> YuvMatrix {
 	}
 }

+/// The frame's raw colorimetry and effective range: `(color_primaries,
+/// color_trc, colorspace, full_range)`. Shared by the CPU scale path and
+/// the GPU import path so both hand the render layer identical
+/// colorimetry. `force_range` wins when it is one of the explicit
+/// `oak_core_COLOR_RANGE_*` values; otherwise the frame's own metadata
+/// decides (YUVJ sources are full range by definition — the caller
+/// resolves the format before this call for the CPU path, but the raw
+/// AVFrame format is also checked here).
+fn frame_colorimetry(f: *const sys::AVFrame, force_range: i32) -> (i32, i32, i32, bool) {
+	// SAFETY: plain reads of the frame's color fields.
+	let (primaries, trc, space, range) = unsafe {
+		(
+			(*f).color_primaries as i32,
+			(*f).color_trc as i32,
+			(*f).colorspace as i32,
+			(*f).color_range as i32,
+		)
+	};
+	let format = unsafe { Pixel::from(std::mem::transmute::<i32, sys::AVPixelFormat>((*f).format)) };
+	let yuvj_full = convert_jpeg_space_to_regular_space(format) != format;
+	let full_range = if force_range == oak_core_COLOR_RANGE_FULL {
+		true
+	} else if force_range == oak_core_COLOR_RANGE_LIMITED {
+		false
+	} else {
+		yuvj_full || range == AVCOL_RANGE_JPEG
+	};
+	(primaries, trc, space, full_range)
+}
+
 /// Bit depth (bits per component) and YUV-ness of a pixel format, from its
 /// `AVPixFmtDescriptor` (8 and false for formats without one — none in
 /// practice for decoder output).
@@ -1845,7 +2035,7 @@ fn apply_sws_colorspace(
 }

 /// The frame's presentation timestamp (NOPTS when unset).
-fn pts_of(f: Option<&ffmpeg::frame::Video>) -> Option<i64> {
+fn pts_of(f: Option<&RefFrame>) -> Option<i64> {
 	f.and_then(|f| f.pts())
 }

@@ -1853,7 +2043,7 @@ fn pts_of(f: Option<&ffmpeg::frame::Video>) -> Option<i64> {
 ///
 /// # CPP-PARITY
 /// `FFmpegDecoder::get_frame_from_cache`.
-fn get_frame_from_cache(video: &VideoDecodeState, t: i64) -> Option<ffmpeg::frame::Video> {
+fn get_frame_from_cache(video: &VideoDecodeState, t: i64) -> Option<RefFrame> {
 	let front = pts_of(video.cache.front()).unwrap_or(AV_NOPTS_VALUE);
 	let back = pts_of(video.cache.back()).unwrap_or(AV_NOPTS_VALUE);
 	if t < front {
@@ -0,0 +1,715 @@
+// Oak Video Editor - Non-Linear Video Editor
+// Copyright (C) 2026 Oak Team
+//
+// This program is free software: you can redistribute it and/or modify
+// it under the terms of the GNU General Public License as published by
+// the Free Software Foundation, either version 3 of the License, or
+// (at your option) any later version.
+//
+// This program is distributed in the hope that it will be useful,
+// but WITHOUT ANY WARRANTY; without even the implied warranty of
+// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+// GNU General Public License for more details.
+//
+// You should have received a copy of the GNU General Public License
+// along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+//! M5: zero-copy hardware-decode import.
+//!
+//! A decoded hardware surface (VAAPI/NVDEC/D3D11VA/VideoToolbox) is
+//! imported into the engine's wgpu device as a pair of planar textures
+//! (luma + interleaved chroma) instead of being downloaded with
+//! `av_hwframe_transfer_data`. The render side turns the planes into
+//! working-space RGBA with the M2 YUV→RGB pass; the CPU staging path
+//! stays the per-frame fallback.
+//!
+//! Platform surface extraction lives here (FFmpeg types); the raw
+//! Vulkan/Metal import lives in `oak_core::backend::external`.
+//!
+//! ## Counters
+//!
+//! - [`HW_IMPORTS`] — frames that took the zero-copy path.
+//! - [`HW_IMPORT_FALLBACKS`] — attempts that were tried and failed.
+//! - [`HW_IMPORT_UNSUPPORTED`] — frames that skipped the import because
+//!   no import path exists (NVDEC, a platform switch off, an
+//!   unimportable layout); not failures.
+//! - `hwdecode::HW_TRANSFERS` — frames downloaded to system memory; a
+//!   zero-copy hardware frame must not move it.
+
+use std::sync::Arc;
+use std::sync::atomic::{AtomicU32, AtomicU64, Ordering};
+
+use ffmpeg::ffi as sys;
+use ffmpeg_next as ffmpeg;
+use oak_core::backend::{GpuContext, YuvTransform};
+#[cfg(target_os = "linux")]
+use oak_core::backend::ImportGuard;
+use oak_core::colormath::YuvMatrix;
+use oak_core::texture::PlanarFormat;
+
+/// Frames that were imported to the GPU instead of downloaded (M5
+/// acceptance: the hardware path keeps `hwdecode::HW_TRANSFERS` at 0).
+pub static HW_IMPORTS: AtomicU64 = AtomicU64::new(0);
+
+/// Import attempts that were actually attempted and failed (per-frame
+/// fallback; feeds the consecutive-failure breaker). "Unsupported"
+/// outcomes are deliberately NOT counted here — a missing import path is
+/// not a failure (M5 audit).
+pub static HW_IMPORT_FALLBACKS: AtomicU64 = AtomicU64::new(0);
+
+/// Frames that took the staging path because no import path exists for
+/// this device/machine/configuration: NVDEC/CUDA has no public handle
+/// export, a per-platform switch is off, or the surface layout is not
+/// importable. Observability only: neither counted as a fallback nor fed
+/// to the breaker.
+pub static HW_IMPORT_UNSUPPORTED: AtomicU64 = AtomicU64::new(0);
+
+/// Consecutive import failures after which the process stops attempting
+/// (the platform combination is evidently unsupported on this machine).
+/// A single failure does not disable anything — the next frame is
+/// attempted afresh, as the per-frame fallback contract requires.
+const STICKY_FAILURE_LIMIT: u32 = 8;
+
+static CONSECUTIVE_FAILURES: AtomicU32 = AtomicU32::new(0);
+
+/// The config key of the zero-copy import switch (1 = import, 0 =
+/// staging). Default ON; `OAK_GPU_IMPORT=0` overrides both.
+pub const CONFIG_KEY_GPU_DECODE_IMPORT: &str = "GpuDecodeImport";
+
+/// The per-platform import switches (each platform's import is an
+/// independent, revertable row of design §3.6).
+pub const CONFIG_KEY_VAAPI_IMPORT: &str = "GpuDecodeImportVaapi";
+/// Windows D3D11VA import switch.
+pub const CONFIG_KEY_D3D11_IMPORT: &str = "GpuDecodeImportD3d11";
+/// macOS VideoToolbox import switch.
+pub const CONFIG_KEY_VIDEOTOOLBOX_IMPORT: &str = "GpuDecodeImportVideotoolbox";
+
+/// Whether the zero-copy hardware import is preferred. `OAK_GPU_IMPORT=0`
+/// force-disables it (diagnostics/CI without importable surfaces), same
+/// convention as `OAK_HWACCEL`.
+pub fn gpu_import_enabled() -> bool {
+	if let Ok(v) = std::env::var("OAK_GPU_IMPORT") {
+		return v != "0";
+	}
+	switch_from_config(CONFIG_KEY_GPU_DECODE_IMPORT)
+}
+
+/// One platform's import switch (global import switch AND platform
+/// switch). `env` overrides the config key, like the global switch.
+fn platform_import_enabled(env: &str, config: &str) -> bool {
+	if let Ok(v) = std::env::var(env) {
+		return v != "0";
+	}
+	switch_from_config(config)
+}
+
+fn switch_from_config(key: &str) -> bool {
+	match oak_core::configstore::ConfigStore::instance().get(None, key) {
+		Ok(value) => value != "false",
+		Err(_) => true,
+	}
+}
+
+/// Reset the import counters (tests).
+pub fn reset_import_counters() {
+	HW_IMPORTS.store(0, Ordering::Relaxed);
+	HW_IMPORT_FALLBACKS.store(0, Ordering::Relaxed);
+	HW_IMPORT_UNSUPPORTED.store(0, Ordering::Relaxed);
+	CONSECUTIVE_FAILURES.store(0, Ordering::Relaxed);
+}
+
+/// Everything the import needs from the decoder's frame and colorimetry.
+pub struct HwImportRequest<'a> {
+	/// The decoded frame (a hardware surface when import is possible).
+	/// Borrowed for the duration of the attempt; the import guard takes
+	/// its own reference to the underlying surface.
+	pub frame: *const sys::AVFrame,
+	/// Lifetime marker for the borrowed frame.
+	pub _marker: std::marker::PhantomData<&'a ()>,
+	/// The session's hardware device type (`None` = software decode).
+	pub device_type: Option<sys::AVHWDeviceType>,
+	/// Frame dimensions (luma).
+	pub luma_size: (u32, u32),
+	/// The frame's YUV→RGB luma matrix.
+	pub matrix: YuvMatrix,
+	/// Full/limited range.
+	pub full_range: bool,
+	/// Source color-primaries code point.
+	pub color_primaries: i32,
+	/// Source transfer-characteristic code point.
+	pub color_trc: i32,
+}
+
+/// A successfully imported hardware frame: the plane tokens plus the
+/// colorimetry the render side needs to resolve them to working-space
+/// RGBA.
+pub struct ImportedHwFrame {
+	/// Luma plane token.
+	pub y: u64,
+	/// Interleaved chroma plane token.
+	pub uv: u64,
+	/// Plane layout (bit depth / interleave).
+	pub format: PlanarFormat,
+	/// Luma dimensions.
+	pub width: u32,
+	/// Frame height.
+	pub height: u32,
+	/// The GPU device that owns the planes.
+	pub ctx: Arc<GpuContext>,
+	/// Extra lifetime guard for platforms whose GPU texture types cannot
+	/// carry a drop callback (macOS): the decoder's pixel buffer.
+	pub keep_alive: Option<Arc<dyn std::any::Any + Send + Sync>>,
+}
+
+impl ImportedHwFrame {
+	/// The YUV→RGB transform for the planes (the bit depth picks the
+	/// exact limited-range expansion: 8 for NV12, 16 for P010).
+	pub fn transform(&self, matrix: YuvMatrix, full_range: bool) -> YuvTransform {
+		YuvTransform::from_matrix_depth(matrix, full_range, self.format.bit_depth())
+	}
+}
+
+/// Try to import `req.frame` as a pair of planar GPU textures. `None`
+/// means "use the staging path for this frame" — the caller must then
+/// call the CPU retrieval (the decoder session cache still holds the
+/// decoded surface, so no re-decode happens).
+pub fn try_import_hw_frame(req: &HwImportRequest<'_>) -> Option<ImportedHwFrame> {
+	if !gpu_import_enabled() {
+		return None;
+	}
+	if CONSECUTIVE_FAILURES.load(Ordering::Relaxed) >= STICKY_FAILURE_LIMIT {
+		return None;
+	}
+	let device_type = req.device_type?;
+	// Only a host-installed device (the app's UI/render context) may take
+	// the frame: a lazily created worker/CLI context must keep decoding on
+	// the CPU path. Checked before touching the frame so the common
+	// "no import here" case costs nothing.
+	if !GpuContext::host_gpu_installed() {
+		return None;
+	}
+	if !crate::hwdecode::is_hw_format(frame_format(req.frame)) {
+		return None;
+	}
+	let Some(ctx) = GpuContext::shared() else {
+		// The host slot was cleared between the two checks (shutdown):
+		// nothing was attempted, so this is not a failure.
+		HW_IMPORT_UNSUPPORTED.fetch_add(1, Ordering::Relaxed);
+		return None;
+	};
+	match platform_import(req, &ctx, device_type) {
+		ImportOutcome::Imported {
+			y,
+			uv,
+			format,
+			keep_alive,
+		} => {
+			CONSECUTIVE_FAILURES.store(0, Ordering::Relaxed);
+			HW_IMPORTS.fetch_add(1, Ordering::Relaxed);
+			Some(ImportedHwFrame {
+				y,
+				uv,
+				format,
+				width: req.luma_size.0,
+				height: req.luma_size.1,
+				ctx,
+				keep_alive,
+			})
+		}
+		ImportOutcome::Unsupported => {
+			HW_IMPORT_UNSUPPORTED.fetch_add(1, Ordering::Relaxed);
+			None
+		}
+		ImportOutcome::Failed => {
+			bump_fallback();
+			None
+		}
+	}
+}
+
+/// Outcome of one platform import attempt.
+enum ImportOutcome {
+	/// The planes were imported; the tokens and keep-alive guard.
+	Imported {
+		y: u64,
+		uv: u64,
+		format: PlanarFormat,
+		keep_alive: Option<Arc<dyn std::any::Any + Send + Sync>>,
+	},
+	/// No import path exists for this device/machine/configuration (the
+	/// frame takes the staging path without being a failure).
+	Unsupported,
+	/// The import path exists and the attempt failed: per-frame fallback,
+	/// counted and fed to the consecutive-failure breaker.
+	Failed,
+}
+
+fn bump_fallback() {
+	HW_IMPORT_FALLBACKS.fetch_add(1, Ordering::Relaxed);
+	CONSECUTIVE_FAILURES.fetch_add(1, Ordering::Relaxed);
+}
+
+/// The frame's raw pixel format (hardware variants included).
+fn frame_format(frame: *const sys::AVFrame) -> sys::AVPixelFormat {
+	// SAFETY: plain read of the frame's format field.
+	unsafe { std::mem::transmute::<i32, sys::AVPixelFormat>((*frame).format) }
+}
+
+#[cfg(target_os = "linux")]
+fn platform_import(
+	req: &HwImportRequest<'_>,
+	ctx: &Arc<GpuContext>,
+	device_type: sys::AVHWDeviceType,
+) -> ImportOutcome {
+	use sys::AVHWDeviceType;
+	if device_type != AVHWDeviceType::AV_HWDEVICE_TYPE_VAAPI {
+		// CUDA/NVDEC surfaces have no public fd/handle export; the
+		// staging path is the honest fallback (design §3.6) — NOT a
+		// failure.
+		return ImportOutcome::Unsupported;
+	}
+	if !platform_import_enabled("OAK_GPU_IMPORT_VAAPI", CONFIG_KEY_VAAPI_IMPORT) {
+		return ImportOutcome::Unsupported;
+	}
+	vaapi_import(req, ctx)
+}
+
+#[cfg(target_os = "windows")]
+fn platform_import(
+	req: &HwImportRequest<'_>,
+	ctx: &Arc<GpuContext>,
+	device_type: sys::AVHWDeviceType,
+) -> ImportOutcome {
+	if device_type != sys::AVHWDeviceType::AV_HWDEVICE_TYPE_D3D11VA {
+		return ImportOutcome::Unsupported;
+	}
+	if !platform_import_enabled("OAK_GPU_IMPORT_D3D11", CONFIG_KEY_D3D11_IMPORT) {
+		return ImportOutcome::Unsupported;
+	}
+	d3d11_import(req, ctx)
+}
+
+#[cfg(target_os = "macos")]
+fn platform_import(
+	req: &HwImportRequest<'_>,
+	ctx: &Arc<GpuContext>,
+	device_type: sys::AVHWDeviceType,
+) -> ImportOutcome {
+	if device_type != sys::AVHWDeviceType::AV_HWDEVICE_TYPE_VIDEOTOOLBOX {
+		return ImportOutcome::Unsupported;
+	}
+	if !platform_import_enabled("OAK_GPU_IMPORT_VIDEOTOOLBOX", CONFIG_KEY_VIDEOTOOLBOX_IMPORT) {
+		return ImportOutcome::Unsupported;
+	}
+	videotoolbox_import(req, ctx)
+}
+
+#[cfg(not(any(target_os = "linux", target_os = "windows", target_os = "macos")))]
+fn platform_import(
+	_req: &HwImportRequest<'_>,
+	_ctx: &Arc<GpuContext>,
+	_device_type: sys::AVHWDeviceType,
+) -> ImportOutcome {
+	ImportOutcome::Unsupported
+}
+
+// ---------------------------------------------------------------------------
+// Linux: VAAPI surface -> DRM PRIME descriptor -> DMA-BUF plane imports
+// ---------------------------------------------------------------------------
+
+/// DRM format fourcc values (little-endian `fourcc_code`).
+#[cfg(target_os = "linux")]
+mod drm_fourcc {
+	pub const R8: u32 = 0x2020_3852; // 'R','8',' ',' '
+	pub const R16: u32 = 0x2036_3152; // 'R','1','6',' '
+	pub const GR88: u32 = 0x3838_5252; // 'R','R','8','8'
+	pub const RG88: u32 = 0x3838_4752; // 'R','G','8','8'
+	pub const RG1616: u32 = 0x3631_4752; // 'R','G','1','6'
+}
+
+
+/// Map a VAAPI frame to its DRM PRIME descriptor (FFmpeg calls
+/// `vaExportSurfaceHandle` internally; no CPU download happens).
+#[cfg(target_os = "linux")]
+fn vaapi_import(req: &HwImportRequest<'_>, ctx: &Arc<GpuContext>) -> ImportOutcome {
+	let mut drm_frame = ffmpeg::frame::Video::empty();
+	// SAFETY: `drm_frame` is a valid AVFrame; the map only writes into it
+	// and takes references to the source frame's hwframe mapping.
+	let rc = unsafe {
+		let raw = drm_frame.as_mut_ptr();
+		(*raw).format = sys::AVPixelFormat::AV_PIX_FMT_DRM_PRIME as i32;
+		sys::av_hwframe_map(raw, req.frame, sys::AV_HWFRAME_MAP_READ as i32)
+	};
+	if rc < 0 {
+		// The VAAPI session exists but the surface could not be exported
+		// as DRM PRIME: a real failure (the breaker bounds the retries).
+		return ImportOutcome::Failed;
+	}
+	// SAFETY: on success `data[0]` is the descriptor FFmpeg allocated for
+	// this mapping; it stays valid until `drm_frame` is unref'd.
+	let desc = unsafe { &*((*drm_frame.as_ptr()).data[0] as *const sys::AVDRMFrameDescriptor) };
+	if desc.nb_layers != 2 {
+		return ImportOutcome::Unsupported;
+	}
+	let layer_format = |i: usize| desc.layers[i].format;
+	let (format, y_fourcc, uv_fourcc) = match (layer_format(0), layer_format(1)) {
+		(a, b) if a == drm_fourcc::R8 && (b == drm_fourcc::GR88 || b == drm_fourcc::RG88) => {
+			(PlanarFormat::Nv12, a, b)
+		}
+		(a, b) if a == drm_fourcc::R16 && b == drm_fourcc::RG1616 => {
+			(PlanarFormat::P010, a, b)
+		}
+		_ => return ImportOutcome::Unsupported,
+	};
+	let _ = (y_fourcc, uv_fourcc);
+	for layer in &desc.layers[..desc.nb_layers as usize] {
+		if layer.nb_planes != 1 {
+			return ImportOutcome::Unsupported;
+		}
+	}
+	let (lw, lh) = req.luma_size;
+	if lw == 0 || lh == 0 {
+		return ImportOutcome::Unsupported;
+	}
+	let (cw, ch) = (lw.div_ceil(2), lh.div_ceil(2));
+	let plane = |layer: usize| {
+		let l = &desc.layers[layer];
+		let p = &l.planes[0];
+		let obj = &desc.objects[p.object_index as usize];
+		oak_core::backend::external::DmaBufPlane {
+			fd: obj.fd,
+			offset: p.offset as u64,
+			pitch: p.pitch as u64,
+			modifier: obj.format_modifier,
+		}
+	};
+	// Shared keep-alive: the source frame pins the VAAPI surface until
+	// the last plane texture is destroyed (the dmabuf fds Vulkan imported
+	// keep the buffer itself alive even after the mapping drops).
+	let Some(keep) = crate::ffmpeg::RefFrame::clone_raw(req.frame) else {
+		return ImportOutcome::Failed;
+	};
+	let keep = Arc::new(keep);
+	let y_guard = keep.clone();
+	let y = match ctx.import_dmabuf_plane(lw, lh, format.y_format(), plane(0), guard(y_guard)) {
+		Ok(y) => y,
+		Err(oak_core::error::Error::State) | Err(oak_core::error::Error::Invalid) => {
+			return ImportOutcome::Unsupported
+		}
+		Err(_) => return ImportOutcome::Failed,
+	};
+	let uv_guard = keep.clone();
+	match ctx.import_dmabuf_plane(cw, ch, format.uv_format(), plane(1), guard(uv_guard)) {
+		Ok(uv) => ImportOutcome::Imported {
+			y,
+			uv,
+			format,
+			keep_alive: None,
+		},
+		Err(_) => {
+			ctx.destroy_texture(y);
+			ImportOutcome::Failed
+		}
+	}
+}
+
+/// A keep-alive import guard holding one `Arc` of the decoder frame
+/// (Linux/VAAPI path; Windows attaches its guard inline and macOS keeps
+/// the frame on the planar texture).
+#[cfg(target_os = "linux")]
+fn guard(value: Arc<crate::ffmpeg::RefFrame>) -> Option<ImportGuard> {
+	Some(Box::new(move || drop(value)))
+}
+
+// ---------------------------------------------------------------------------
+// Windows: D3D11VA texture -> shared NT handle -> Vulkan import
+// ---------------------------------------------------------------------------
+
+/// Import a D3D11VA texture (Windows): one NV12/P010 texture shared as an
+/// NT handle, imported as a single multi-planar image whose two plane
+/// views feed the planar YUV→RGB pass. The decoder's frame reference is
+/// kept alive by the texture's drop guard.
+#[cfg(target_os = "windows")]
+fn d3d11_import(req: &HwImportRequest<'_>, ctx: &Arc<GpuContext>) -> ImportOutcome {
+	use windows::core::Interface;
+	use windows::Win32::Foundation::CloseHandle;
+	use windows::Win32::Graphics::Direct3D11::{ID3D11Texture2D, D3D11_TEXTURE2D_DESC};
+	use windows::Win32::Graphics::Dxgi::Common::{DXGI_FORMAT_NV12, DXGI_FORMAT_P010};
+	use windows::Win32::Graphics::Dxgi::{IDXGIResource1, DXGI_SHARED_RESOURCE_READ};
+
+	// AV_PIX_FMT_D3D11: data[0] = ID3D11Texture2D*, data[1] = subresource
+	// index. D3D11VA frames cycle through the decoder's surface ARRAY, so
+	// the slice index is part of the frame identity — the import selects
+	// it via an array-layer view (M5 audit).
+	let texture_ptr = unsafe { (*req.frame).data[0] } as *mut std::ffi::c_void;
+	if texture_ptr.is_null() {
+		return ImportOutcome::Unsupported;
+	}
+	let slice = unsafe { (*req.frame).data[1] } as usize as u32;
+	// SAFETY: FFmpeg owns the texture reference for the frame's lifetime.
+	let texture: &ID3D11Texture2D = unsafe { &*(texture_ptr as *const ID3D11Texture2D) };
+	let mut desc = D3D11_TEXTURE2D_DESC::default();
+	unsafe { texture.GetDesc(&mut desc) };
+	let array_size = desc.ArraySize.max(1);
+	if slice >= array_size {
+		return ImportOutcome::Unsupported;
+	}
+	let (texture_format, planar) = match desc.Format {
+		DXGI_FORMAT_NV12 => (oak_core::wgpu::TextureFormat::NV12, PlanarFormat::Nv12),
+		DXGI_FORMAT_P010 => (oak_core::wgpu::TextureFormat::P010, PlanarFormat::P010),
+		_ => return ImportOutcome::Unsupported,
+	};
+	// SAFETY: the D3D11 texture supports IDXGIResource1 (every shared
+	// texture does); a cast failure means it is not a shareable resource.
+	let Ok(resource) = texture.cast::<IDXGIResource1>() else {
+		return ImportOutcome::Unsupported;
+	};
+	// Read-only NT handle; Vulkan references the resource, so the handle
+	// is closed again right after the import.
+	let handle = unsafe {
+		resource.CreateSharedHandle(
+			None,
+			DXGI_SHARED_RESOURCE_READ.0,
+			None::<&windows::core::PCWSTR>,
+		)
+	};
+	let Ok(handle) = handle else {
+		return ImportOutcome::Failed;
+	};
+	let Some(keep) = crate::ffmpeg::RefFrame::clone_raw(req.frame) else {
+		unsafe {
+			let _ = CloseHandle(handle);
+		}
+		return ImportOutcome::Failed;
+	};
+	let (w, h) = req.luma_size;
+	let guard: Arc<crate::ffmpeg::RefFrame> = Arc::new(keep);
+	let result = ctx.import_d3d11_shared_texture(
+		handle.0 as isize,
+		w,
+		h,
+		texture_format,
+		array_size,
+		slice,
+		Some(Box::new(move || drop(guard))),
+	);
+	unsafe {
+		let _ = CloseHandle(handle);
+	}
+	match result {
+		Ok((y, uv)) => ImportOutcome::Imported {
+			y,
+			uv,
+			format: planar,
+			keep_alive: None,
+		},
+		// State/Invalid are capability gaps (missing device feature, wrong
+		// format), not failed attempts.
+		Err(oak_core::error::Error::State) | Err(oak_core::error::Error::Invalid) => {
+			ImportOutcome::Unsupported
+		}
+		Err(_) => ImportOutcome::Failed,
+	}
+}
+
+// ---------------------------------------------------------------------------
+// macOS: VideoToolbox CVPixelBuffer -> IOSurface -> Metal textures
+// ---------------------------------------------------------------------------
+
+/// `CVPixelBufferGetIOSurface` (CoreVideo).
+#[cfg(target_os = "macos")]
+#[link(name = "CoreVideo", kind = "framework")]
+extern "C" {
+	fn CVPixelBufferGetIOSurface(pixel_buffer: *mut std::ffi::c_void) -> *mut std::ffi::c_void;
+}
+
+/// Import a VideoToolbox frame (macOS): its CVPixelBuffer's IOSurface is
+/// wrapped in one Metal texture per plane. Metal textures cannot carry a
+/// drop callback, so the decoder's frame reference travels back in
+/// `ImportedHwFrame::keep_alive` and is held by the planar texture.
+#[cfg(target_os = "macos")]
+fn videotoolbox_import(req: &HwImportRequest<'_>, ctx: &Arc<GpuContext>) -> ImportOutcome {
+	// AV_PIX_FMT_VIDEOTOOLBOX keeps the CVPixelBuffer in data[3].
+	let pixel_buffer = unsafe { (*req.frame).data[3] } as *mut std::ffi::c_void;
+	if pixel_buffer.is_null() {
+		return ImportOutcome::Unsupported;
+	}
+	let surface = unsafe { CVPixelBufferGetIOSurface(pixel_buffer) };
+	if surface.is_null() {
+		return ImportOutcome::Unsupported;
+	}
+	let frames = unsafe { (*req.frame).hw_frames_ctx };
+	if frames.is_null() {
+		return ImportOutcome::Unsupported;
+	}
+	// SAFETY: `hw_frames_ctx` is an AVHWFramesContext for this frame.
+	let sw_format = unsafe { (*(*frames).data as *const sys::AVHWFramesContext).sw_format };
+	let format = match sw_format {
+		sys::AVPixelFormat::AV_PIX_FMT_NV12 => PlanarFormat::Nv12,
+		sys::AVPixelFormat::AV_PIX_FMT_P010LE | sys::AVPixelFormat::AV_PIX_FMT_P010BE => {
+			PlanarFormat::P010
+		}
+		_ => return ImportOutcome::Unsupported,
+	};
+	let (w, h) = req.luma_size;
+	let (cw, ch) = (w.div_ceil(2), h.div_ceil(2));
+	// SAFETY: `surface` is the CVPixelBuffer's IOSurface and the decoder
+	// frame reference is kept alive by the import guard (planar texture).
+	let Ok(y) = (unsafe { ctx.import_iosurface_plane(surface, 0, w, h, format.y_format()) }) else {
+		return ImportOutcome::Failed;
+	};
+	let uv = match unsafe { ctx.import_iosurface_plane(surface, 1, cw, ch, format.uv_format()) } {
+		Ok(uv) => uv,
+		Err(_) => {
+			ctx.destroy_texture(y);
+			return ImportOutcome::Failed;
+		}
+	};
+	let Some(keep) = crate::ffmpeg::RefFrame::clone_raw(req.frame) else {
+		ctx.destroy_texture(y);
+		ctx.destroy_texture(uv);
+		return ImportOutcome::Failed;
+	};
+	ImportOutcome::Imported {
+		y,
+		uv,
+		format,
+		keep_alive: Some(Arc::new(keep)),
+	}
+}
+
+#[cfg(test)]
+mod tests {
+	use super::*;
+
+	fn empty_frame() -> ffmpeg::frame::Video {
+		ffmpeg::frame::Video::empty()
+	}
+
+	fn request(frame: &ffmpeg::frame::Video, device_type: sys::AVHWDeviceType) -> HwImportRequest<'_> {
+		HwImportRequest {
+			frame: unsafe { frame.as_ptr() },
+			_marker: std::marker::PhantomData,
+			device_type: Some(device_type),
+			luma_size: (64, 64),
+			matrix: YuvMatrix::Bt709,
+			full_range: false,
+			color_primaries: 1,
+			color_trc: 1,
+		}
+	}
+
+	/// Test-only RAII override for a process-global environment variable:
+	/// restores the previous value (or absence) on drop, so a panicking
+	/// assertion cannot leak the override into the next serialized test and
+	/// an externally preset value survives the test.
+	struct EnvVarGuard {
+		name: &'static str,
+		previous: Option<std::ffi::OsString>,
+	}
+
+	impl EnvVarGuard {
+		/// Set `name=value` for the guard's lifetime.
+		fn set(name: &'static str, value: &str) -> Self {
+			let previous = std::env::var_os(name);
+			std::env::set_var(name, value);
+			Self { name, previous }
+		}
+
+		/// Remove `name` for the guard's lifetime.
+		fn remove(name: &'static str) -> Self {
+			let previous = std::env::var_os(name);
+			std::env::remove_var(name);
+			Self { name, previous }
+		}
+	}
+
+	impl Drop for EnvVarGuard {
+		fn drop(&mut self) {
+			match self.previous.take() {
+				Some(value) => std::env::set_var(self.name, value),
+				None => std::env::remove_var(self.name),
+			}
+		}
+	}
+
+	/// Without a host-installed device the import declines before any
+	/// frame work (and is not counted as a fallback: nothing failed).
+	#[test]
+	fn import_requires_a_host_installed_context() {
+		let _g = crate::lock_tests();
+		reset_import_counters();
+		assert!(!GpuContext::host_gpu_installed());
+		let frame = empty_frame();
+		let req = request(&frame, sys::AVHWDeviceType::AV_HWDEVICE_TYPE_VAAPI);
+		assert!(try_import_hw_frame(&req).is_none());
+		assert_eq!(HW_IMPORTS.load(Ordering::Relaxed), 0);
+		assert_eq!(HW_IMPORT_FALLBACKS.load(Ordering::Relaxed), 0);
+		assert_eq!(HW_IMPORT_UNSUPPORTED.load(Ordering::Relaxed), 0);
+	}
+
+	/// The counter classification (M5 audit): a device type with no import
+	/// path (NVDEC/CUDA) is "unsupported", never a failure — it must not
+	/// feed the consecutive-failure breaker or `HW_IMPORT_FALLBACKS`.
+	/// Exercised at the platform-function level so no real GPU is needed.
+	///
+	/// (`platform_import` is private; this test lives in the same module.)
+	#[cfg(target_os = "linux")]
+	#[test]
+	fn non_importable_device_type_is_unsupported_not_failed() {
+		let _g = crate::lock_tests();
+		// A context is required by the platform function's signature;
+		// the non-VAAPI device check happens before any of its fields are
+		// touched, so a lazily-created local context is never used.
+		let Some(ctx) = GpuContext::create(oak_core::backend::BackendKind::Auto) else {
+			eprintln!("no GPU adapter; skipping classification test");
+			return;
+		};
+		let frame = empty_frame();
+		let req = request(&frame, sys::AVHWDeviceType::AV_HWDEVICE_TYPE_CUDA);
+		assert!(
+			matches!(
+				platform_import(&req, &ctx, sys::AVHWDeviceType::AV_HWDEVICE_TYPE_CUDA),
+				ImportOutcome::Unsupported
+			),
+			"CUDA/NVDEC has no export path and must classify as unsupported"
+		);
+	}
+
+	/// The `OAK_GPU_IMPORT=0` escape hatch disables the path outright, and
+	/// removing the override falls back to the default-on config. The
+	/// override is RAII-scoped, so an externally preset value survives the
+	/// test instead of being clobbered by `remove_var`.
+	#[test]
+	fn import_switch_disables_the_path() {
+		let _g = crate::lock_tests();
+		let previous = std::env::var_os("OAK_GPU_IMPORT");
+		{
+			let _off = EnvVarGuard::set("OAK_GPU_IMPORT", "0");
+			assert!(!gpu_import_enabled(), "explicit 0 disables the import path");
+		}
+		assert_eq!(
+			std::env::var_os("OAK_GPU_IMPORT"),
+			previous,
+			"the guard restores the externally preset value"
+		);
+		if previous.is_none() {
+			// No override: the config-store default keeps the import path
+			// on. (With an external value the caller's setting governs, so
+			// asserting the default would be wrong.)
+			let _unset = EnvVarGuard::remove("OAK_GPU_IMPORT");
+			assert!(
+				gpu_import_enabled(),
+				"unset falls back to the default-on config"
+			);
+		}
+	}
+
+	/// Plane layouts carry the expansion depth the YUV transform needs.
+	#[test]
+	fn planar_formats_carry_their_bit_depth() {
+		assert_eq!(PlanarFormat::Nv12.bit_depth(), 8);
+		assert_eq!(PlanarFormat::P010.bit_depth(), 16);
+		assert_eq!(PlanarFormat::Nv12.name(), "nv12");
+	}
+}
@@ -31,14 +31,15 @@
 //! hardware surface.
 //!
 //! - **macOS**: `AV_HWDEVICE_TYPE_VIDEOTOOLBOX`
-//! - **Linux**: `CUDA` (NVDEC, the discrete-GPU path), then `VAAPI`
+//! - **Linux**: `VAAPI` (DMA-BUF, the M5 import path), then `CUDA` (NVDEC)
 //! - **Windows**: `D3D11VA`, then `CUDA` (NVDEC)
 //!
 //! Device creation can fail on machines without the device/driver (a
 //! headless Linux box, no NVIDIA GPU) — the candidate is skipped and
 //! the next one (or the software decoder) is used. Hardware frames
 //! (`AV_PIX_FMT_VIDEOTOOLBOX` / `VAAPI` / `CUDA` / `D3D11VA_VLD` /
-//! `D3D11`) are transferred to system memory with
+//! `D3D11`) either go through the M5 zero-copy import
+//! (`crate::gpuinterop`) or are transferred to system memory with
 //! `av_hwframe_transfer_data` before the swscale conversion.

 use ffmpeg::ffi as sys;
@@ -131,16 +132,17 @@ pub fn device_type_candidates() -> &'static [sys::AVHWDeviceType] {
 	}
 	#[cfg(all(unix, not(target_os = "macos")))]
 	{
-		// CUDA (native NVDEC) first: it is the discrete-GPU path, and on
-		// multi-GPU boxes VAAPI's default render node can point at a
-		// device with no VA driver (e.g. NVIDIA without
-		// libva-nvidia-driver) while the AMD/Intel node would work —
-		// trying NVDEC first sidesteps that misdirection. VAAPI stays as
-		// the fallback for AMD/Intel-only machines (CUDA device creation
-		// fails fast without libcuda).
+		// VAAPI first: it is the DMA-BUF path the M5 zero-copy import
+		// consumes (`VK_EXT_external_memory_dma_buf`), and on NVIDIA boxes
+		// with the VA-API driver it is just as reachable as NVDEC — while
+		// a CUDA (NVDEC) surface has no public handle export, so frames
+		// decoded through it can only take the staging fallback. CUDA
+		// stays as the fallback for machines without libva; it also
+		// sidesteps VAAPI's default render node pointing at a device with
+		// no VA driver (CUDA device creation fails fast without libcuda).
 		&[
-			sys::AVHWDeviceType::AV_HWDEVICE_TYPE_CUDA,
 			sys::AVHWDeviceType::AV_HWDEVICE_TYPE_VAAPI,
+			sys::AVHWDeviceType::AV_HWDEVICE_TYPE_CUDA,
 		]
 	}
 }
@@ -363,3 +365,6 @@ mod tests {
 		assert!(!device_unavailable(dev));
 	}
 }
+	/// would let their attempts perturb the negative-cache assertion on
+#[cfg(test)]
+#[cfg(test)]
@@ -46,6 +46,7 @@ pub mod ffmpeg;
 pub mod footagedescription;
 pub mod frame;
 pub mod framemanager;
+pub mod gpuinterop;
 pub mod hwdecode;
 pub mod oiio;
 pub mod oiioframebridge;
@@ -769,3 +769,4 @@ mod tests_extra {
 		);
 	}
 }
+ 		);
@@ -15,6 +15,11 @@ ocio-rs = { version = "0.2", features = ["bundled"] }
 # generation (29, matching the gpui_wgpu device the app presents with, so
 # the render thread's textures are directly sampleable by the UI).
 wgpu = "29"
+# wgpu-types: the HAL descriptor types (`TextureUses`) that wgpu 29 does
+# re-export through `wgpu::hal`'s descriptors but not by name (M5 external
+# image import builds a `wgpu_hal::TextureDescriptor` by hand). Pinned to
+# the same generation as wgpu 29.
+wgpu-types = "29"
 log = "0.4.34"
 # half: f16↔f32 conversion for format-aware texture downloads (the M2
 # present target is Rgba16Float; the rest of the pipeline is Rgba32Float).
@@ -26,3 +31,35 @@ quick-xml = "0.41.0"
 # Pure-Rust image I/O for oiioutils.rs (TIFF only — the only format current
 # callers need).
 image = { version = "0.25", default-features = false, features = ["tiff"] }
+
+# --- M5 zero-copy hardware-decode import (platform natives) ---------------
+# Linux: raw Vulkan through wgpu-hal's own `ash` generation (the version is
+# pinned by wgpu-hal 29's dependency, so the `vk` types unify) + `dup(2)`
+# for the DMA-BUF fds handed to `VkImportMemoryFdInfoKHR`.
+[target.'cfg(target_os = "linux")'.dependencies]
+libc = "0.2"
+
+# The raw Vulkan import path is shared by Linux (DMA-BUF) and Windows
+# (D3D11 shared handle).
+[target.'cfg(any(target_os = "linux", target_os = "windows"))'.dependencies]
+ash = "0.38"
+
+# Windows: D3D11VA surfaces arrive as an `ID3D11Texture2D`; the shared NT
+# handle is handed to Vulkan's `VK_EXTERNAL_MEMORY_HANDLE_TYPE_D3D11_TEXTURE`
+# path. Only the raw ash structs are needed inside oak-core (the COM side
+# lives in oak-codec), so no `windows` dependency here.
+#
+# macOS: VideoToolbox surfaces are IOSurfaces; the Metal texture wrapping
+# them is built with the same objc2 generation wgpu-hal 29 uses.
+[target.'cfg(target_os = "macos")'.dependencies]
+objc2 = "0.6"
+objc2-foundation = "0.3"
+objc2-metal = { version = "0.3", features = [
+    "MTLAllocation",
+    "MTLDevice",
+    "MTLPixelFormat",
+    "MTLResource",
+    "MTLTexture",
+    "objc2-io-surface",
+] }
+objc2-io-surface = "0.3"
@@ -44,6 +44,12 @@ use crate::error::{Error, Result};
 use crate::frame::VideoParamsPod;
 use crate::texture::{Frame, Texture};

+pub mod external;
+
+#[cfg(target_os = "linux")]
+pub use external::DmaBufPlane;
+pub use external::ImportGuard;
+
 /// Backend selection preference (mapped onto wgpu backends).
 ///
 /// The choice is user-visible: the settings panel exposes a renderer
@@ -222,6 +228,15 @@ struct GpuTexture {
 	width: u32,
 	height: u32,
 	format: wgpu::TextureFormat,
+	/// The view aspect to use for sampling. Plane-imported textures
+	/// (Windows NV12: one texture, two plane views) register one token
+	/// per plane with the matching aspect; everything else is `All`.
+	aspect: wgpu::TextureAspect,
+	/// Selected array layer range for array-imported textures (Windows
+	/// D3D11VA frames are slices of one array texture): `Some((base,
+	/// count))` when this token aliases one slice. `None` = the whole
+	/// (single-layer) texture.
+	layer: Option<(u32, u32)>,
 }

 fn lock<T>(m: &Mutex<T>) -> MutexGuard<'_, T> {
@@ -295,6 +310,9 @@ pub struct GpuContext {
 	present: Mutex<Vec<(wgpu::TextureFormat, PresentPipeline)>>,
 	/// The compiled YUV→RGB pass (M5 decode import dependency).
 	yuv: Mutex<Option<PresentPipeline>>,
+	/// The compiled planar (imported hardware NV12/P010) YUV→RGB pass
+	/// (M5): luma + interleaved chroma plane bindings.
+	yuv_planar: Mutex<Option<PresentPipeline>>,
 	/// Caller-keyed 3D LUT textures (per-node color transforms; M2).
 	color_luts: Mutex<Vec<(String, u64)>>,
 	/// Set once this context has created a GPU resource (texture or
@@ -337,6 +355,12 @@ pub fn reset_gpu_transfer_counters() {
 /// The process-wide shared-context slot (`shared` / `install_shared`).
 struct SharedSlot {
 	decided: bool,
+	/// The HOST (app/UI) explicitly installed the context, as opposed to
+	/// the engine lazily creating one from the user config. External
+	/// hardware-surface imports (M5) only run on a host-installed device:
+	/// a lazily created private context in a worker/CLI/test must not
+	/// silently switch decode onto the GPU.
+	host_installed: bool,
 	ctx: Option<Arc<GpuContext>>,
 }

@@ -345,29 +369,46 @@ fn shared_slot() -> &'static Mutex<SharedSlot> {
 	SLOT.get_or_init(|| {
 		Mutex::new(SharedSlot {
 			decided: false,
+			host_installed: false,
 			ctx: None,
 		})
 	})
 }

+/// Parse an `OAK_REQUIRE_GPU` value: unset means "skipping is allowed";
+/// only `0`/`false` (case-insensitive) disable the requirement. Kept pure
+/// so the policy test needs no process-wide environment mutation (which
+/// would race the GPU acceptance tests reading the variable in parallel).
+fn require_gpu_from_value(value: Option<&str>) -> bool {
+	match value {
+		Some(v) => !(v == "0" || v.eq_ignore_ascii_case("false")),
+		None => false,
+	}
+}
+
 /// Whether GPU-dependent tests must hard-fail when no adapter is
 /// available. CI sets `OAK_REQUIRE_GPU=1` on the software-Vulkan runner
 /// (lavapipe is present), so a degraded environment fails the suite
 /// instead of silently losing the GPU acceptance signal.
 pub fn require_gpu_adapter() -> bool {
-	match std::env::var("OAK_REQUIRE_GPU") {
-		Ok(v) => !(v == "0" || v.eq_ignore_ascii_case("false")),
-		Err(_) => false,
+	require_gpu_from_value(std::env::var("OAK_REQUIRE_GPU").ok().as_deref())
+}
+
+/// Handle a missing adapter: panic when `required`, otherwise log the
+/// skip (the caller returns). Taking the flag as a parameter lets the
+/// policy test drive the hard-fail arm without flipping the real
+/// environment variable of a running test process.
+fn skip_or_fail_gpu_with(required: bool, what: &str) {
+	if required {
+		panic!("no GPU adapter available for {what} (OAK_REQUIRE_GPU is set)");
 	}
+	eprintln!("no GPU adapter; skipping {what}");
 }

 /// Handle a missing adapter in a GPU acceptance test: panic when
 /// `OAK_REQUIRE_GPU` is set, otherwise log the skip (the caller returns).
 pub fn skip_or_fail_gpu(what: &str) {
-	if require_gpu_adapter() {
-		panic!("no GPU adapter available for {what} (OAK_REQUIRE_GPU is set)");
-	}
-	eprintln!("no GPU adapter; skipping {what}");
+	skip_or_fail_gpu_with(require_gpu_adapter(), what);
 }

 /// A fresh [`GpuContext`] for a test, or `None` when no adapter exists
@@ -438,6 +479,18 @@ impl GpuContext {
 			if filterable {
 				required_features |= wgpu::Features::FLOAT32_FILTERABLE;
 			}
+			// M5 (Windows): the D3D11VA import wraps a multi-planar
+			// NV12/P010 texture; those formats need their features
+			// enabled at device creation (no-op on adapters without
+			// them — the import then falls back per frame).
+			for feature in [
+				wgpu::Features::TEXTURE_FORMAT_NV12,
+				wgpu::Features::TEXTURE_FORMAT_P010,
+			] {
+				if adapter.features().contains(feature) {
+					required_features |= feature;
+				}
+			}
 			let (device, queue) =
 				match pollster_block_on(adapter.request_device(&wgpu::DeviceDescriptor {
 					label: Some("oakrender"),
@@ -469,6 +522,7 @@ impl GpuContext {
 				display_lut: Mutex::new(None),
 				present: Mutex::new(Vec::new()),
 				yuv: Mutex::new(None),
+				yuv_planar: Mutex::new(None),
 				color_luts: Mutex::new(Vec::new()),
 				used: AtomicBool::new(false),
 			}));
@@ -508,6 +562,7 @@ impl GpuContext {
 			display_lut: Mutex::new(None),
 			present: Mutex::new(Vec::new()),
 			yuv: Mutex::new(None),
+			yuv_planar: Mutex::new(None),
 			color_luts: Mutex::new(Vec::new()),
 			used: AtomicBool::new(false),
 		})
@@ -600,6 +655,8 @@ impl GpuContext {
 				width: width as u32,
 				height: height as u32,
 				format,
+				aspect: wgpu::TextureAspect::All,
+				layer: None,
 			},
 		);
 		Ok(token)
@@ -1245,6 +1302,202 @@ impl GpuContext {
 		Ok(program)
 	}

+	/// Run the planar YUV→RGB pass (M5 zero-copy decode): the imported
+	/// luma plane (R8/R16, bind 0) and interleaved chroma plane
+	/// (R8G8/R16G16, bind 1) → one `Rgba32Float` destination. The
+	/// matrix/range live in `transform`, exactly like
+	/// [`GpuContext::run_yuv_to_rgb`].
+	pub fn run_planar_yuv_to_rgb(
+		&self,
+		y: u64,
+		uv: u64,
+		dst: u64,
+		transform: &YuvTransform,
+	) -> Result<()> {
+		let y_view = self.texture_view(y)?;
+		let uv_view = self.texture_view(uv)?;
+		let dst_tex = lock(&self.textures)
+			.get(&dst)
+			.cloned()
+			.ok_or(Error::NotFound)?;
+		let (y_size, uv_size) = {
+			let textures = lock(&self.textures);
+			let y = textures.get(&y).ok_or(Error::NotFound)?;
+			let uv = textures.get(&uv).ok_or(Error::NotFound)?;
+			((y.width, y.height), (uv.width, uv.height))
+		};
+		let pipeline = self.planar_yuv_pipeline()?;
+		let m = transform.matrix;
+		let b = transform.offset;
+		let params: [f32; 16] = [
+			m[0][0],
+			m[0][1],
+			m[0][2],
+			b[0],
+			m[1][0],
+			m[1][1],
+			m[1][2],
+			b[1],
+			m[2][0],
+			m[2][1],
+			m[2][2],
+			b[2],
+			uv_size.0 as f32 / y_size.0.max(1) as f32,
+			uv_size.1 as f32 / y_size.1.max(1) as f32,
+			0.0,
+			0.0,
+		];
+		let uniform = self.device.create_buffer(&wgpu::BufferDescriptor {
+			label: Some("oakrender-yuv-planar-params"),
+			size: 64,
+			usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+			mapped_at_creation: false,
+		});
+		self.queue
+			.write_buffer(&uniform, 0, &f32_uniform_bytes(&params));
+		let dst_view = dst_tex
+			.texture
+			.create_view(&wgpu::TextureViewDescriptor::default());
+		let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
+			label: Some("oakrender-yuv-planar-bg"),
+			layout: &pipeline.layout,
+			entries: &[
+				wgpu::BindGroupEntry {
+					binding: 0,
+					resource: wgpu::BindingResource::TextureView(&y_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 1,
+					resource: wgpu::BindingResource::TextureView(&uv_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 2,
+					resource: uniform.as_entire_binding(),
+				},
+			],
+		});
+		let mut encoder = self
+			.device
+			.create_command_encoder(&wgpu::CommandEncoderDescriptor {
+				label: Some("oakrender-yuv-planar"),
+			});
+		{
+			let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+				label: Some("oakrender-yuv-planar-pass"),
+				color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+					view: &dst_view,
+					depth_slice: None,
+					resolve_target: None,
+					ops: wgpu::Operations {
+						load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT),
+						store: wgpu::StoreOp::Store,
+					},
+				})],
+				depth_stencil_attachment: None,
+				timestamp_writes: None,
+				occlusion_query_set: None,
+				multiview_mask: None,
+			});
+			pass.set_pipeline(&pipeline.pipeline);
+			pass.set_bind_group(0, &bind_group, &[]);
+			pass.draw(0..3, 0..1);
+		}
+		self.queue.submit(Some(encoder.finish()));
+		Ok(())
+	}
+
+	/// The planar YUV→RGB pass pipeline (built once).
+	fn planar_yuv_pipeline(&self) -> Result<PresentPipeline> {
+		let mut cache = lock(&self.yuv_planar);
+		if let Some(p) = cache.as_ref() {
+			return Ok(p.clone());
+		}
+		let plane = |binding: u32| wgpu::BindGroupLayoutEntry {
+			binding,
+			visibility: wgpu::ShaderStages::FRAGMENT,
+			ty: wgpu::BindingType::Texture {
+				sample_type: wgpu::TextureSampleType::Float { filterable: false },
+				view_dimension: wgpu::TextureViewDimension::D2,
+				multisampled: false,
+			},
+			count: None,
+		};
+		let layout = self
+			.device
+			.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+				label: Some("oakrender-yuv-planar-layout"),
+				entries: &[
+					plane(0),
+					plane(1),
+					wgpu::BindGroupLayoutEntry {
+						binding: 2,
+						visibility: wgpu::ShaderStages::FRAGMENT,
+						ty: wgpu::BindingType::Buffer {
+							ty: wgpu::BufferBindingType::Uniform,
+							has_dynamic_offset: false,
+							min_binding_size: None,
+						},
+						count: None,
+					},
+				],
+			});
+		let vs = self
+			.device
+			.create_shader_module(wgpu::ShaderModuleDescriptor {
+				label: Some("oakrender-yuv-planar-vs"),
+				source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(EFFECT_VS_WGSL)),
+			});
+		let fs = self
+			.device
+			.create_shader_module(wgpu::ShaderModuleDescriptor {
+				label: Some("oakrender-yuv-planar-fs"),
+				source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(YUV_PLANAR_WGSL)),
+			});
+		let pipeline_layout = self
+			.device
+			.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+				label: Some("oakrender-yuv-planar-pipeline-layout"),
+				bind_group_layouts: &[Some(&layout)],
+				immediate_size: 0,
+			});
+		let scope = self.device.push_error_scope(wgpu::ErrorFilter::Validation);
+		let pipeline = self
+			.device
+			.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+				label: Some("oakrender-yuv-planar"),
+				layout: Some(&pipeline_layout),
+				vertex: wgpu::VertexState {
+					module: &vs,
+					entry_point: Some("vs_main"),
+					compilation_options: Default::default(),
+					buffers: &[],
+				},
+				primitive: wgpu::PrimitiveState::default(),
+				depth_stencil: None,
+				multisample: wgpu::MultisampleState::default(),
+				fragment: Some(wgpu::FragmentState {
+					module: &fs,
+					entry_point: Some("main"),
+					compilation_options: Default::default(),
+					targets: &[Some(wgpu::ColorTargetState {
+						format: wgpu::TextureFormat::Rgba32Float,
+						blend: None,
+						write_mask: wgpu::ColorWrites::ALL,
+					})],
+				}),
+				multiview_mask: None,
+				cache: None,
+			});
+		if let Some(err) = pollster_block_on(scope.pop()) {
+			return Err(Error::Failed(format!(
+				"planar YUV pipeline validation failed: {err}"
+			)));
+		}
+		let program = PresentPipeline { pipeline, layout };
+		*cache = Some(program.clone());
+		Ok(program)
+	}
+
 	/// Clear a texture to transparent black on the GPU (no CPU transfer):
 	/// the graph compositor's starting accumulator and single-sided
 	/// transition sides.
@@ -1605,11 +1858,40 @@ impl GpuContext {
 				return false;
 			}
 		}
+		slot.host_installed = ctx.is_some();
 		slot.ctx = ctx;
 		slot.decided = true;
 		true
 	}

+	/// True when the host installed the shared device (the app's UI
+	/// context). The M5 hardware-surface import requires this: a lazily
+	/// created engine context belongs to a worker/CLI/test and must not
+	/// pull decode onto the GPU.
+	pub fn host_gpu_installed() -> bool {
+		lock(shared_slot()).host_installed
+	}
+
+	/// Declare the process's shared context as the HOST's render device
+	/// (M5), creating it on first use if the engine has not yet.
+	///
+	/// On Linux/FreeBSD the app adopts the window's wgpu device via
+	/// [`GpuContext::install_shared`]; on macOS/Windows gpui does not
+	/// expose that device, so the engine-created shared context *is* the
+	/// app's render device and must be marked here — otherwise the decode
+	/// import's host gate would make the D3D11VA/VideoToolbox rows dead
+	/// code (M5 audit). Worker, CLI and test processes never call this,
+	/// which is what keeps their decode on the CPU staging path.
+	pub fn mark_host_context() -> bool {
+		let mut slot = lock(shared_slot());
+		if !slot.decided {
+			slot.ctx = Self::create(BackendKind::from_user_config());
+			slot.decided = true;
+		}
+		slot.host_installed = slot.ctx.is_some();
+		slot.ctx.is_some()
+	}
+
 	/// True when the app installed the shared context (as opposed to the
 	/// engine lazily creating one from user config). Tests/UI use this to
 	/// tell "the presenter's device" from "a private device".
@@ -1931,22 +2213,46 @@ pub struct YuvTransform {
 }

 impl YuvTransform {
-	/// Build the transform for a luma matrix and range.
+	/// Build the transform for a luma matrix and range. Inputs are
+	/// normalized 16-bit code values (code / 65535).
 	pub fn from_matrix(matrix: crate::colormath::YuvMatrix, full_range: bool) -> Self {
+		Self::from_matrix_depth(matrix, full_range, 16)
+	}
+
+	/// Build the transform for a luma matrix, range and source bit depth.
+	/// Plane inputs are normalized code values (code / (2^bits - 1));
+	/// the limited-range expansion uses the depth's own code limits
+	/// (16..235 luma, 16..240 chroma scaled to the depth).
+	pub fn from_matrix_depth(
+		matrix: crate::colormath::YuvMatrix,
+		full_range: bool,
+		bit_depth: u32,
+	) -> Self {
 		let (kr, kb) = matrix.kr_kb();
 		let kg = 1.0 - kr - kb;
-		// CPU reference: Y spans 219<<8 for limited, chroma 224<<8, both
-		// divided per code value; inputs here are code/65535.
-		let (ay, by, ac, bc) = if full_range {
-			(1.0, 0.0, 1.0, 32768.0 / 65535.0)
+		let max = ((1u32 << bit_depth) - 1) as f32;
+		// Code limits for this depth. `scale` is the 8-bit-code shift the
+		// CPU reference uses (depth 16 = 8-bit codes << 8, so 16 stays
+		// bit-identical with `colormath::yuv444p16_to_rgb_f32`).
+		let scale = if bit_depth > 8 {
+			1u32 << (bit_depth - 8)
+		} else {
+			1
+		};
+		let (y_black, y_white, c_mid, c_span) = if full_range {
+			(0.0, max, (1u32 << (bit_depth - 1)) as f32, max)
 		} else {
 			(
-				65535.0 / 56064.0,
-				4096.0 / 56064.0,
-				65535.0 / 57344.0,
-				32768.0 / 57344.0,
+				(16 * scale) as f32,
+				(235 * scale) as f32,
+				(128 * scale) as f32,
+				(224 * scale) as f32,
 			)
 		};
+		let ay = max / (y_white - y_black);
+		let by = y_black / (y_white - y_black);
+		let ac = max / c_span;
+		let bc = c_mid / c_span;
 		let rv = 2.0 * (1.0 - kr) * ac;
 		let bu = 2.0 * (1.0 - kb) * ac;
 		let gu = -(2.0 * kb * (1.0 - kb) / kg) * ac;
@@ -2096,6 +2402,46 @@ fn main(@builtin(position) frag: vec4<f32>) -> @location(0) vec4<f32> {
 }
 "#;

+/// The planar (hardware-import) YUV→RGB pass (M5): luma + interleaved
+/// chroma plane inputs (R8/R8G8 or R16/R16G16, normalized 0..1), one
+/// `Rgba32Float` output. Same matrix/range math as [`YUV_WGSL`].
+const YUV_PLANAR_WGSL: &str = r#"
+struct Params {
+    m0: vec4<f32>,
+    m1: vec4<f32>,
+    m2: vec4<f32>,
+    uv_scale: vec4<f32>,
+};
+@group(0) @binding(0) var y_tex: texture_2d<f32>;
+@group(0) @binding(1) var uv_tex: texture_2d<f32>;
+@group(0) @binding(2) var<uniform> p: Params;
+
+@fragment
+fn main(@builtin(position) frag: vec4<f32>) -> @location(0) vec4<f32> {
+    let yd = textureDimensions(y_tex);
+    let coord = clamp(
+        vec2<i32>(i32(frag.x), i32(frag.y)),
+        vec2<i32>(0, 0),
+        vec2<i32>(i32(yd.x), i32(yd.y)) - vec2<i32>(1, 1),
+    );
+    let yv = textureLoad(y_tex, coord, 0).r;
+    let ud = textureDimensions(uv_tex);
+    let uc = clamp(
+        vec2<i32>(vec2<f32>(coord) * p.uv_scale.xy),
+        vec2<i32>(0, 0),
+        vec2<i32>(i32(ud.x), i32(ud.y)) - vec2<i32>(1, 1),
+    );
+    let uv = textureLoad(uv_tex, uc, 0);
+    let yuv = vec3<f32>(yv, uv.r, uv.g);
+    let rgb = vec3<f32>(
+        dot(p.m0.xyz, yuv) + p.m0.w,
+        dot(p.m1.xyz, yuv) + p.m1.w,
+        dot(p.m2.xyz, yuv) + p.m2.w,
+    );
+    return vec4<f32>(rgb, 1.0);
+}
+"#;
+
 /// The fixed vertex stage for effect passes: a fullscreen triangle
 /// emitting `ove_texcoord`-convention UVs at location 0. UV v=0 is the
 /// first texture data row (the upload/download row order), so effect
@@ -2368,6 +2714,9 @@ impl DisplayRenderer {
 				let frame = unsafe { frame_from_pixels_for_upload(size, pixels, linesize) }?;
 				ctx.upload(*token, &frame)
 			}
+			// Imported hardware planes are immutable views of decoder
+			// memory (M5); uploading into them is never valid.
+			Texture::Planar(_) => Err(Error::State),
 			Texture::Cpu(frame) => {
 				let stride = frame.linesize_bytes();
 				if linesize != stride {
@@ -2477,7 +2826,9 @@ impl DisplayRenderer {
 pub fn texture_id_of(t: &Texture) -> i32 {
 	match t {
 		Texture::Gpu { token, .. } => *token as i32,
-		Texture::Cpu(_) => 0,
+		// Imported planar frames have no single displayable texture id;
+		// the footage path resolves them before presentation.
+		Texture::Planar(_) | Texture::Cpu(_) => 0,
 	}
 }

@@ -0,0 +1,883 @@
+// Oak Video Editor - Non-Linear Video Editor
+// Copyright (C) 2026 Oak Team
+//
+// This program is free software: you can redistribute it and/or modify
+// it under the terms of the GNU General Public License as published by
+// the Free Software Foundation, either version 3 of the License, or
+// (at your option) any later version.
+//
+// This program is distributed in the hope that it will be useful,
+// but WITHOUT ANY WARRANTY; without even the implied warranty of
+// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+// GNU General Public License for more details.
+//
+// You should have received a copy of the GNU General Public License
+// along with this program.  If not, see <http://www.gnu.org/licenses/>.
+
+//! M5: importing platform hardware surfaces as wgpu textures.
+//!
+//! The video decoder produces hardware surfaces (VAAPI/NVDEC/D3D11VA/
+//! VideoToolbox). Instead of downloading each frame to system memory and
+//! re-uploading it, the surface's platform handle is wrapped into a
+//! `wgpu::Texture` on the engine's own device:
+//!
+//! - **Linux**: the VAAPI surface's DMA-BUF plane (fd/offset/pitch/
+//!   modifier) is imported into a raw Vulkan image
+//!   (`VK_EXT_external_memory_dma_buf` + `VK_EXT_image_drm_format_modifier`)
+//!   and wrapped through `wgpu_hal::vulkan::Device::texture_from_raw`.
+//! - **Windows**: the D3D11VA texture's shared NT handle is wrapped via
+//!   `wgpu_hal::vulkan::Device::texture_from_d3d11_shared_handle`.
+//! - **macOS**: the VideoToolbox CVPixelBuffer's IOSurface is wrapped in a
+//!   Metal texture (`newTextureWithDescriptor:iosurface:plane:`) and
+//!   handed to `wgpu_hal::metal::Device::texture_from_raw`.
+//!
+//! Everything here is `unsafe` platform glue; the *only* consumer is the
+//! decode import in `oak-codec::gpuinterop`, and any failure must be
+//! treated as "fall back to the staging path for this frame" — never as a
+//! hard error that poisons later frames.
+
+use std::sync::Arc;
+
+use super::*;
+
+/// A DMA-BUF plane of a hardware frame (Linux/Vulkan import).
+///
+/// The `fd` is borrowed: the import duplicates it before handing it to
+/// Vulkan, so the caller keeps ownership of the original.
+#[cfg(target_os = "linux")]
+#[derive(Clone, Copy, Debug)]
+pub struct DmaBufPlane {
+	/// DMA-BUF file descriptor.
+	pub fd: i32,
+	/// Byte offset of the plane inside the buffer object.
+	pub offset: u64,
+	/// Row pitch in bytes.
+	pub pitch: u64,
+	/// DRM format modifier (`DRM_FORMAT_MOD_*`).
+	pub modifier: u64,
+}
+
+/// The Metal pixel format for an IOSurface plane (macOS).
+#[cfg(target_os = "macos")]
+fn iosurface_pixel_format(format: wgpu::TextureFormat) -> Option<objc2_metal::MTLPixelFormat> {
+	use objc2_metal::MTLPixelFormat;
+	Some(match format {
+		wgpu::TextureFormat::R8Unorm => MTLPixelFormat::R8Unorm,
+		wgpu::TextureFormat::Rg8Unorm => MTLPixelFormat::RG8Unorm,
+		wgpu::TextureFormat::R16Unorm => MTLPixelFormat::R16Unorm,
+		wgpu::TextureFormat::Rg16Unorm => MTLPixelFormat::RG16Unorm,
+		_ => return None,
+	})
+}
+
+/// Registration parameters for one imported texture alias: dimensions,
+/// array size, layout, plane aspect and the selected array slice (M5).
+#[derive(Clone, Copy)]
+struct ImportedTextureDesc {
+	width: u32,
+	height: u32,
+	array_layers: u32,
+	format: wgpu::TextureFormat,
+	aspect: wgpu::TextureAspect,
+	layer: Option<(u32, u32)>,
+}
+
+/// Caller-owned resource tied to an imported texture's lifetime (M5):
+/// the decoder's hardware `AVFrame` stays alive until the GPU texture is
+/// destroyed, so the decoder's surface pool cannot recycle the memory
+/// while the GPU still samples it. The closure runs exactly once, right
+/// before the platform resources are released.
+pub type ImportGuard = Box<dyn FnOnce() + Send + Sync>;
+
+/// The plane formats the import accepts. Kept small on purpose: luma and
+/// interleaved chroma for 8-bit NV12 and 16-bit-container P010.
+#[cfg(target_os = "linux")]
+fn import_vk_format(format: wgpu::TextureFormat) -> Option<ash::vk::Format> {
+	use ash::vk::Format;
+	Some(match format {
+		wgpu::TextureFormat::R8Unorm => Format::R8_UNORM,
+		wgpu::TextureFormat::Rg8Unorm => Format::R8G8_UNORM,
+		wgpu::TextureFormat::R16Unorm => Format::R16_UNORM,
+		wgpu::TextureFormat::Rg16Unorm => Format::R16G16_UNORM,
+		_ => return None,
+	})
+}
+
+impl GpuContext {
+	/// Import one DMA-BUF plane as a GPU texture (Linux/Vulkan).
+	///
+	/// Returns the registry token. `keep_alive` (the decoder's mapped
+	/// `AVFrame`) runs when the GPU texture is destroyed, after the GPU is
+	/// done with it.
+	#[cfg(target_os = "linux")]
+	pub fn import_dmabuf_plane(
+		&self,
+		width: u32,
+		height: u32,
+		format: wgpu::TextureFormat,
+		plane: DmaBufPlane,
+		keep_alive: Option<ImportGuard>,
+	) -> Result<u64> {
+		use ash::vk;
+		use wgpu::hal;
+
+		if self.kind != BackendKind::Vulkan || width == 0 || height == 0 {
+			return Err(Error::State);
+		}
+		let Some(vk_format) = import_vk_format(format) else {
+			return Err(Error::Invalid);
+		};
+		let guard = unsafe { self.device.as_hal::<hal::api::Vulkan>() }.ok_or(Error::State)?;
+		let hal_device: &hal::vulkan::Device = &guard;
+		let raw = hal_device.raw_device().clone();
+		let _ = hal_device.raw_physical_device();
+
+		let layouts = [vk::SubresourceLayout {
+			offset: plane.offset,
+			size: 0,
+			row_pitch: plane.pitch,
+			array_pitch: 0,
+			depth_pitch: 0,
+		}];
+		let external = vk::ExternalMemoryImageCreateInfo::default()
+			.handle_types(vk::ExternalMemoryHandleTypeFlags::DMA_BUF_EXT);
+		let mut modifier = vk::ImageDrmFormatModifierExplicitCreateInfoEXT::default()
+			.drm_format_modifier(plane.modifier)
+			.plane_layouts(&layouts);
+		modifier.p_next = (&external as *const vk::ExternalMemoryImageCreateInfo).cast();
+
+		let info = vk::ImageCreateInfo::default()
+			.image_type(vk::ImageType::TYPE_2D)
+			.format(vk_format)
+			.extent(vk::Extent3D {
+				width,
+				height,
+				depth: 1,
+			})
+			.mip_levels(1)
+			.array_layers(1)
+			.samples(vk::SampleCountFlags::TYPE_1)
+			.tiling(vk::ImageTiling::DRM_FORMAT_MODIFIER_EXT)
+			.usage(vk::ImageUsageFlags::SAMPLED)
+			.sharing_mode(vk::SharingMode::EXCLUSIVE)
+			.initial_layout(vk::ImageLayout::UNDEFINED)
+			.push_next(&mut modifier);
+		let image = unsafe { raw.create_image(&info, None) }
+			.map_err(|e| Error::Failed(format!("vkCreateImage(dmabuf): {e:?}")))?;
+
+		let req = unsafe { raw.get_image_memory_requirements(image) };
+
+		// Duplicate the fd: Vulkan takes ownership of the imported fd on a
+		// successful `vkAllocateMemory` and closes it when the memory is
+		// freed; the decoder keeps its own.
+		let dup_fd = unsafe { libc::dup(plane.fd) };
+		if dup_fd < 0 {
+			unsafe { raw.destroy_image(image, None) };
+			return Err(Error::Failed("dup(dmabuf fd) failed".into()));
+		}
+		// Pick the memory type by trying the compatible ones in order —
+		// `vkGetPhysicalDeviceMemoryProperties` is instance-level and
+		// wgpu-hal keeps its raw instance private. The compatible set is
+		// usually tiny (often one DEVICE_LOCAL type on discrete GPUs).
+		let mut import = vk::ImportMemoryFdInfoKHR::default()
+			.handle_type(vk::ExternalMemoryHandleTypeFlags::DMA_BUF_EXT)
+			.fd(dup_fd);
+		let mut dedicated = vk::MemoryDedicatedAllocateInfo::default().image(image);
+		let mut last_err = None;
+		let mut memory = None;
+		for type_index in 0..32u32 {
+			if req.memory_type_bits & (1 << type_index) == 0 {
+				continue;
+			}
+			let alloc = vk::MemoryAllocateInfo::default()
+				.allocation_size(req.size)
+				.memory_type_index(type_index)
+				.push_next(&mut dedicated)
+				.push_next(&mut import);
+			match unsafe { raw.allocate_memory(&alloc, None) } {
+				Ok(m) => {
+					memory = Some(m);
+					break;
+				}
+				Err(e) => last_err = Some((type_index, e)),
+			}
+		}
+		let Some(memory) = memory else {
+			unsafe {
+				libc::close(dup_fd);
+				raw.destroy_image(image, None);
+			}
+			return Err(Error::Failed(format!(
+				"vkAllocateMemory(import) failed for every compatible type: {last_err:?}"
+			)));
+		};
+		if let Err(e) = unsafe { raw.bind_image_memory(image, memory, 0) } {
+			unsafe {
+				raw.free_memory(memory, None);
+				raw.destroy_image(image, None);
+			}
+			return Err(Error::Failed(format!("vkBindImageMemory: {e:?}")));
+		}
+
+		let hal_desc = hal::TextureDescriptor {
+			label: Some("oakrender-imported-dmabuf"),
+			size: wgpu::Extent3d {
+				width,
+				height,
+				depth_or_array_layers: 1,
+			},
+			mip_level_count: 1,
+			sample_count: 1,
+			dimension: wgpu::TextureDimension::D2,
+			format,
+			usage: wgpu_types::TextureUses::RESOURCE,
+			memory_flags: hal::MemoryFlags::empty(),
+			view_formats: Vec::new(),
+		};
+		// wgpu-hal does not destroy external images (it cannot know how);
+		// the drop guard owns both the image and the imported memory, and
+		// carries the decoder's frame-keep-alive with it.
+		let raw_guard = raw.clone();
+		let drop_guard: hal::DropCallback = Box::new(move || unsafe {
+			if let Some(keep) = keep_alive {
+				keep();
+			}
+			raw_guard.destroy_image(image, None);
+			raw_guard.free_memory(memory, None);
+		});
+		let hal_texture = unsafe {
+			hal_device.texture_from_raw(
+				image,
+				&hal_desc,
+				Some(drop_guard),
+				hal::vulkan::TextureMemory::External,
+			)
+		};
+		self.register_imported_texture::<hal::api::Vulkan>(
+			hal_texture,
+			ImportedTextureDesc {
+				width,
+				height,
+				array_layers: 1,
+				format,
+				aspect: wgpu::TextureAspect::All,
+				layer: None,
+			},
+		)
+	}
+
+	/// Import one slice of a D3D11VA texture array shared as an NT handle
+	/// (Windows/Vulkan): one multi-planar image (`NV12`/`P010`) with
+	/// `array_size` layers, registered twice (planes 0/1) for `slice`.
+	///
+	/// D3D11VA frames cycle through the decoder's surface array
+	/// (`AVFrame::data[1]` is the slice index), so the imported VkImage
+	/// must carry the whole array and the tokens must select their slice —
+	/// otherwise only every N-th frame would be importable (M5 audit).
+	///
+	/// `keep_alive` runs when the image is destroyed; the handle itself is
+	/// owned by the caller (Vulkan takes a resource reference, not the
+	/// handle).
+	#[cfg(target_os = "windows")]
+	pub fn import_d3d11_shared_texture(
+		&self,
+		shared_handle: isize,
+		width: u32,
+		height: u32,
+		format: wgpu::TextureFormat,
+		array_size: u32,
+		slice: u32,
+		keep_alive: Option<ImportGuard>,
+	) -> Result<(u64, u64)> {
+		use ash::vk;
+		use wgpu::hal;
+
+		if self.kind != BackendKind::Vulkan || width == 0 || height == 0 {
+			return Err(Error::State);
+		}
+		let array_size = array_size.max(1);
+		if slice >= array_size {
+			return Err(Error::Invalid);
+		}
+		let (vk_format, required_feature) = match format {
+			wgpu::TextureFormat::NV12 => (
+				vk::Format::G8_B8R8_2PLANE_420_UNORM,
+				wgpu::Features::TEXTURE_FORMAT_NV12,
+			),
+			wgpu::TextureFormat::P010 => (
+				vk::Format::G10X6_B10X6R10X6_2PLANE_420_UNORM_3PACK16,
+				wgpu::Features::TEXTURE_FORMAT_P010,
+			),
+			_ => return Err(Error::Invalid),
+		};
+		// Multi-planar formats need their device feature; without it the
+		// hal texture would be created but every view invalid.
+		if !self.device.features().contains(required_feature) {
+			return Err(Error::State);
+		}
+		let guard = unsafe { self.device.as_hal::<hal::api::Vulkan>() }.ok_or(Error::State)?;
+		let hal_device: &hal::vulkan::Device = &guard;
+		let raw = hal_device.raw_device().clone();
+
+		let mut external = vk::ExternalMemoryImageCreateInfo::default()
+			.handle_types(vk::ExternalMemoryHandleTypeFlags::D3D11_TEXTURE);
+		let info = vk::ImageCreateInfo::default()
+			.image_type(vk::ImageType::TYPE_2D)
+			.format(vk_format)
+			.extent(vk::Extent3D {
+				width,
+				height,
+				depth: 1,
+			})
+			.mip_levels(1)
+			.array_layers(array_size)
+			.samples(vk::SampleCountFlags::TYPE_1)
+			.tiling(vk::ImageTiling::OPTIMAL)
+			.usage(vk::ImageUsageFlags::SAMPLED)
+			.sharing_mode(vk::SharingMode::EXCLUSIVE)
+			.initial_layout(vk::ImageLayout::UNDEFINED)
+			.push_next(&mut external);
+		let image = unsafe { raw.create_image(&info, None) }
+			.map_err(|e| Error::Failed(format!("vkCreateImage(d3d11): {e:?}")))?;
+
+		let req = unsafe { raw.get_image_memory_requirements(image) };
+		// Dedicated allocation is required for imported D3D11 textures.
+		let mut import = vk::ImportMemoryWin32HandleInfoKHR::default()
+			.handle_type(vk::ExternalMemoryHandleTypeFlags::D3D11_TEXTURE)
+			.handle(shared_handle);
+		let mut dedicated = vk::MemoryDedicatedAllocateInfo::default().image(image);
+		let mut last_err = None;
+		let mut memory = None;
+		for type_index in 0..32u32 {
+			if req.memory_type_bits & (1 << type_index) == 0 {
+				continue;
+			}
+			let alloc = vk::MemoryAllocateInfo::default()
+				.allocation_size(req.size)
+				.memory_type_index(type_index)
+				.push_next(&mut dedicated)
+				.push_next(&mut import);
+			match unsafe { raw.allocate_memory(&alloc, None) } {
+				Ok(m) => {
+					memory = Some(m);
+					break;
+				}
+				Err(e) => last_err = Some((type_index, e)),
+			}
+		}
+		let Some(memory) = memory else {
+			unsafe { raw.destroy_image(image, None) };
+			return Err(Error::Failed(format!(
+				"vkAllocateMemory(import d3d11) failed: {last_err:?}"
+			)));
+		};
+		if let Err(e) = unsafe { raw.bind_image_memory(image, memory, 0) } {
+			unsafe {
+				raw.free_memory(memory, None);
+				raw.destroy_image(image, None);
+			}
+			return Err(Error::Failed(format!("vkBindImageMemory: {e:?}")));
+		}
+
+		let hal_desc = hal::TextureDescriptor {
+			label: Some("oakrender-imported-d3d11"),
+			size: wgpu::Extent3d {
+				width,
+				height,
+				depth_or_array_layers: array_size,
+			},
+			mip_level_count: 1,
+			sample_count: 1,
+			dimension: wgpu::TextureDimension::D2,
+			format,
+			usage: wgpu_types::TextureUses::RESOURCE,
+			memory_flags: hal::MemoryFlags::empty(),
+			view_formats: Vec::new(),
+		};
+		let raw_guard = raw.clone();
+		let drop_guard: hal::DropCallback = Box::new(move || unsafe {
+			if let Some(keep) = keep_alive {
+				keep();
+			}
+			raw_guard.destroy_image(image, None);
+			raw_guard.free_memory(memory, None);
+		});
+		let hal_texture = unsafe {
+			hal_device.texture_from_raw(
+				image,
+				&hal_desc,
+				Some(drop_guard),
+				hal::vulkan::TextureMemory::External,
+			)
+		};
+		// Two plane views of the chosen slice of the multi-planar array:
+		// Y at plane 0, interleaved chroma at plane 1.
+		let desc = ImportedTextureDesc {
+			width,
+			height,
+			array_layers: array_size,
+			format,
+			aspect: wgpu::TextureAspect::Plane0,
+			layer: Some((slice, 1)),
+		};
+		let (y, texture) = self.register_imported_texture_shared::<hal::api::Vulkan>(
+			hal_texture, desc,
+		)?;
+		let uv = self.register_texture_entry(
+			texture,
+			ImportedTextureDesc {
+				// The chroma plane's view is half-size; the planar pass
+				// derives the sampling scale from the registered sizes.
+				width: width.div_ceil(2),
+				height: height.div_ceil(2),
+				aspect: wgpu::TextureAspect::Plane1,
+				..desc
+			},
+		);
+		Ok((y, uv))
+	}
+
+	/// Import one IOSurface plane as a Metal-backed texture (macOS): the
+	/// VideoToolbox pixel buffer's surface stays owned by the decoder; the
+	/// caller keeps it alive via the planar texture's guard.
+	///
+	/// # Safety
+	///
+	/// `iosurface` must be a live `IOSurfaceRef` that stays valid until
+	/// the imported texture is destroyed (the decoder's frame reference is
+	/// held by the planar texture's keep-alive guard).
+	#[cfg(target_os = "macos")]
+	pub unsafe fn import_iosurface_plane(
+		&self,
+		iosurface: *mut std::ffi::c_void,
+		plane: u32,
+		width: u32,
+		height: u32,
+		format: wgpu::TextureFormat,
+	) -> Result<u64> {
+		use objc2::rc::Retained;
+		use objc2::runtime::ProtocolObject;
+		use objc2_io_surface::IOSurfaceRef;
+		use objc2_metal::{
+			MTLDevice as _, MTLPixelFormat, MTLStorageMode, MTLTextureDescriptor,
+			MTLTextureType, MTLTextureUsage,
+		};
+		use wgpu::hal;
+
+		if self.kind != BackendKind::Metal || width == 0 || height == 0 || iosurface.is_null() {
+			return Err(Error::State);
+		}
+		let Some(pixel_format) = iosurface_pixel_format(format) else {
+			return Err(Error::Invalid);
+		};
+		let guard = unsafe { self.device.as_hal::<hal::api::Metal>() }.ok_or(Error::State)?;
+		let hal_device: &hal::metal::Device = &guard;
+		let device: &Retained<ProtocolObject<dyn objc2_metal::MTLDevice>> =
+			hal_device.raw_device();
+
+		let descriptor = MTLTextureDescriptor::new();
+		descriptor.setTextureType(MTLTextureType::Type2D);
+		descriptor.setPixelFormat(pixel_format);
+		descriptor.setStorageMode(MTLStorageMode::Shared);
+		descriptor.setUsage(MTLTextureUsage::ShaderRead);
+		// SAFETY: plain descriptor field setters.
+		unsafe {
+			descriptor.setWidth(width as usize);
+			descriptor.setHeight(height as usize);
+			descriptor.setMipmapLevelCount(1);
+		}
+		// SAFETY: the caller guarantees a live IOSurface for the texture's
+		// lifetime (the planar texture holds its keep-alive guard).
+		let surface: &IOSurfaceRef = unsafe { &*(iosurface as *const IOSurfaceRef) };
+		// SAFETY: `surface` is a live IOSurface for the texture's lifetime
+		// (the caller's planar keep-alive holds the pixel buffer).
+		let texture = device
+			.newTextureWithDescriptor_iosurface_plane(&descriptor, surface, plane as usize)
+			.ok_or_else(|| Error::Failed("newTextureWithDescriptor:iosurface: returned nil".into()))?;
+
+		let raw_type = MTLTextureType::Type2D;
+		let hal_texture = unsafe {
+			hal::metal::Device::texture_from_raw(
+				texture,
+				format,
+				raw_type,
+				1,
+				1,
+				hal::CopyExtent {
+					width,
+					height,
+					depth: 1,
+				},
+			)
+		};
+		self.register_imported_texture::<hal::api::Metal>(
+			hal_texture,
+			ImportedTextureDesc {
+				width,
+				height,
+				array_layers: 1,
+				format,
+				aspect: wgpu::TextureAspect::All,
+				layer: None,
+			},
+		)
+	}
+
+	/// Register a texture created through `create_texture_from_hal` in the
+	/// registry (M5 external imports; also the shared tail of every
+	/// platform import).
+	fn register_imported_texture<A: wgpu::hal::Api>(
+		&self,
+		hal_texture: A::Texture,
+		desc: ImportedTextureDesc,
+	) -> Result<u64> {
+		self.register_imported_texture_shared::<A>(hal_texture, desc)
+			.map(|(token, _)| token)
+	}
+
+	/// The shared registration body; returns the token plus the owning
+	/// `wgpu::Texture` so multi-plane/array imports (one texture, several
+	/// plane aspects or slices) can register aliases without recreating
+	/// the texture.
+	fn register_imported_texture_shared<A: wgpu::hal::Api>(
+		&self,
+		hal_texture: A::Texture,
+		desc: ImportedTextureDesc,
+	) -> Result<(u64, Arc<wgpu::Texture>)> {
+		let texture_desc = wgpu::TextureDescriptor {
+			label: Some("oakrender-imported"),
+			size: wgpu::Extent3d {
+				width: desc.width,
+				height: desc.height,
+				depth_or_array_layers: desc.array_layers.max(1),
+			},
+			mip_level_count: 1,
+			sample_count: 1,
+			dimension: wgpu::TextureDimension::D2,
+			format: desc.format,
+			usage: wgpu::TextureUsages::TEXTURE_BINDING,
+			view_formats: &[],
+		};
+		let texture = unsafe {
+			self.device
+				.create_texture_from_hal::<A>(hal_texture, &texture_desc)
+		};
+		self.used.store(true, Ordering::Release);
+		let texture = Arc::new(texture);
+		let token = self.register_texture_entry(texture.clone(), desc);
+		Ok((token, texture))
+	}
+
+	/// Insert one registry entry for `texture`.
+	fn register_texture_entry(&self, texture: Arc<wgpu::Texture>, desc: ImportedTextureDesc) -> u64 {
+		let token = self.next_token.fetch_add(1, Ordering::Relaxed);
+		lock(&self.textures).insert(
+			token,
+			GpuTexture {
+				texture,
+				width: desc.width,
+				height: desc.height,
+				format: desc.format,
+				aspect: desc.aspect,
+				layer: desc.layer,
+			},
+		);
+		token
+	}
+
+	/// The view for `token`, honoring plane aspects and array slices
+	/// (M5 imports).
+	pub(crate) fn texture_view(&self, token: u64) -> Result<wgpu::TextureView> {
+		let entry = lock(&self.textures)
+			.get(&token)
+			.cloned()
+			.ok_or(Error::NotFound)?;
+		let (base_array_layer, array_layer_count) = match entry.layer {
+			Some((base, count)) => (base, Some(count)),
+			None => (0, None),
+		};
+		Ok(entry.texture.create_view(&wgpu::TextureViewDescriptor {
+			aspect: entry.aspect,
+			base_array_layer,
+			array_layer_count,
+			// wgpu-core derives `D2Array` from the TEXTURE's layer count,
+			// not from the view range: an array-imported slice must ask
+			// for `D2` explicitly or the planar pass's `D2` layout
+			// validation rejects the bind group (M5 desktop review).
+			dimension: (array_layer_count == Some(1)).then_some(wgpu::TextureViewDimension::D2),
+			..Default::default()
+		}))
+	}
+}
+
+#[cfg(test)]
+mod tests {
+	//! Unit coverage for the safe parts of the import module. The unsafe
+	//! platform bodies are deliberately **not** exercised here: a real
+	//! import needs a hardware decoder surface (a VAAPI DMA-BUF, a D3D11VA
+	//! shared NT handle, a VideoToolbox IOSurface), which the
+	//! software-Vulkan CI runners do not have, and feeding them fake
+	//! handles would only assert that a driver rejects garbage.
+	//!
+	//! What covers the hardware paths instead:
+	//! - `import_dmabuf_plane` (Linux): a VAAPI machine running
+	//!   `oak-render/tests/footage_import_test.rs` (DMA-BUF + DRM format
+	//!   modifier import); the M5 real-hardware acceptance is recorded in
+	//!   `docs/zh/plans/render-pipeline-threads-m5-branch-coverage.txt`.
+	//! - `import_d3d11_shared_texture` (Windows): the Windows job on a host
+	//!   with D3D11VA; a lavapipe-only runner has no shared texture to
+	//!   import and the decode side falls back to staging.
+	//! - `import_iosurface_plane` (macOS): the macOS job with VideoToolbox
+	//!   active.
+	//! - `register_imported_texture*`/`texture_view`: the alias registry
+	//!   and view helper are unit-tested below on a locally created array
+	//!   texture (no platform import); the `Plane0`/`Plane1` aspects only
+	//!   become valid with the multi-planar formats of a real import.
+
+	use super::*;
+	use crate::backend::gpu_or_skip;
+	use crate::error::Error;
+
+	/// The Linux plane format map accepts exactly the luma/chroma planes
+	/// the planar import feeds it and rejects everything else.
+	#[cfg(target_os = "linux")]
+	#[test]
+	fn import_vk_format_maps_luma_and_chroma_planes() {
+		use ash::vk::Format;
+		assert_eq!(
+			import_vk_format(wgpu::TextureFormat::R8Unorm),
+			Some(Format::R8_UNORM)
+		);
+		assert_eq!(
+			import_vk_format(wgpu::TextureFormat::Rg8Unorm),
+			Some(Format::R8G8_UNORM)
+		);
+		assert_eq!(
+			import_vk_format(wgpu::TextureFormat::R16Unorm),
+			Some(Format::R16_UNORM)
+		);
+		assert_eq!(
+			import_vk_format(wgpu::TextureFormat::Rg16Unorm),
+			Some(Format::R16G16_UNORM)
+		);
+		for rejected in [
+			wgpu::TextureFormat::Rgba8Unorm,
+			wgpu::TextureFormat::Rgba32Float,
+			wgpu::TextureFormat::NV12,
+			wgpu::TextureFormat::P010,
+		] {
+			assert_eq!(
+				import_vk_format(rejected),
+				None,
+				"{rejected:?} is not a single-plane sample format"
+			);
+		}
+	}
+
+	/// The macOS plane map accepts the interleaved 8/16-bit layouts the
+	/// VideoToolbox import feeds and rejects everything else.
+	#[cfg(target_os = "macos")]
+	#[test]
+	fn iosurface_pixel_format_maps_supported_planes() {
+		use objc2_metal::MTLPixelFormat;
+		assert_eq!(
+			iosurface_pixel_format(wgpu::TextureFormat::R8Unorm),
+			Some(MTLPixelFormat::R8Unorm)
+		);
+		assert_eq!(
+			iosurface_pixel_format(wgpu::TextureFormat::Rg8Unorm),
+			Some(MTLPixelFormat::RG8Unorm)
+		);
+		assert_eq!(
+			iosurface_pixel_format(wgpu::TextureFormat::R16Unorm),
+			Some(MTLPixelFormat::R16Unorm)
+		);
+		assert_eq!(
+			iosurface_pixel_format(wgpu::TextureFormat::Rg16Unorm),
+			Some(MTLPixelFormat::RG16Unorm)
+		);
+		assert_eq!(
+			iosurface_pixel_format(wgpu::TextureFormat::Rgba8Unorm),
+			None
+		);
+	}
+
+	/// `import_dmabuf_plane`'s argument validation must reject zero-sized
+	/// imports and non-plane formats *before* it touches Vulkan, so this
+	/// runs without any DMA-BUF: the plane fd (here invalid on purpose)
+	/// is never reached.
+	#[cfg(target_os = "linux")]
+	#[test]
+	fn dmabuf_import_rejects_invalid_arguments_without_importing() {
+		let Some(ctx) = gpu_or_skip("the DMA-BUF argument validation test") else {
+			return;
+		};
+		if ctx.kind() != crate::backend::BackendKind::Vulkan {
+			eprintln!("SKIP: the DMA-BUF validation test needs a Vulkan context");
+			return;
+		}
+		let plane = DmaBufPlane {
+			fd: -1,
+			offset: 0,
+			pitch: 64,
+			modifier: 0,
+		};
+		let err = ctx
+			.import_dmabuf_plane(0, 64, wgpu::TextureFormat::R8Unorm, plane, None)
+			.unwrap_err();
+		assert!(matches!(err, Error::State), "zero width: {err:?}");
+		let err = ctx
+			.import_dmabuf_plane(64, 0, wgpu::TextureFormat::R8Unorm, plane, None)
+			.unwrap_err();
+		assert!(matches!(err, Error::State), "zero height: {err:?}");
+		let err = ctx
+			.import_dmabuf_plane(64, 64, wgpu::TextureFormat::Bgra8Unorm, plane, None)
+			.unwrap_err();
+		assert!(matches!(err, Error::Invalid), "non-plane format: {err:?}");
+	}
+
+	/// The registry/view helper every import shares: one array texture,
+	/// several aliases. A single-slice alias must produce the explicit `D2`
+	/// view the M5 desktop review fix requests (wgpu-core would otherwise
+	/// derive `D2Array` from the texture's layer count) **and bind into the
+	/// planar pass's `D2` bind-group layout** — the actual failure point of
+	/// that regression: `create_view` succeeds on a derived `D2Array`, the
+	/// bind group does not. Broader aliases keep the array dimension (and
+	/// are shown to be rejected by the `D2` layout, so the guard is
+	/// falsifiable), and unknown tokens are `NotFound`.
+	#[test]
+	fn imported_texture_aliases_create_valid_views_and_bind() {
+		let Some(ctx) = gpu_or_skip("the imported-texture alias test") else {
+			return;
+		};
+		// The shape a D3D11VA import registers: one 4-layer D2 texture.
+		let texture = Arc::new(ctx.device.create_texture(&wgpu::TextureDescriptor {
+			label: Some("oakrender-imported-test"),
+			size: wgpu::Extent3d {
+				width: 8,
+				height: 8,
+				depth_or_array_layers: 4,
+			},
+			mip_level_count: 1,
+			sample_count: 1,
+			dimension: wgpu::TextureDimension::D2,
+			format: wgpu::TextureFormat::Rgba8Unorm,
+			usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST,
+			view_formats: &[],
+		}));
+		let desc = ImportedTextureDesc {
+			width: 8,
+			height: 8,
+			array_layers: 4,
+			format: wgpu::TextureFormat::Rgba8Unorm,
+			aspect: wgpu::TextureAspect::All,
+			layer: Some((2, 1)),
+		};
+		let single = ctx.register_texture_entry(texture.clone(), desc);
+		let single_view = ctx.texture_view(single).expect("single-slice alias view");
+
+		let pair = ctx.register_texture_entry(
+			texture.clone(),
+			ImportedTextureDesc {
+				layer: Some((0, 2)),
+				..desc
+			},
+		);
+		let pair_view = ctx.texture_view(pair).expect("two-slice alias view");
+
+		let whole = ctx.register_texture_entry(
+			texture.clone(),
+			ImportedTextureDesc {
+				layer: None,
+				..desc
+			},
+		);
+		ctx.texture_view(whole).expect("whole-array alias view");
+
+		assert_ne!(single, pair, "aliases get distinct tokens");
+		assert_ne!(pair, whole, "aliases get distinct tokens");
+		assert!(
+			matches!(ctx.texture_view(u64::MAX), Err(Error::NotFound)),
+			"an unregistered token must be NotFound"
+		);
+
+		// The regression guard: the planar pass creates its bind group
+		// with a `D2` layout. Without `dimension: Some(D2)` the
+		// single-slice alias is a derived `D2Array` and this bind group
+		// fails validation, so the test must exercise bind-group creation
+		// (not just `create_view`). Use the production planar layout so the
+		// guard follows it if the layout changes.
+		let pipeline = ctx
+			.planar_yuv_pipeline()
+			.expect("the planar pass pipeline builds");
+		// The chroma slot must be a single-slice `D2` alias too: slice 0 of
+		// the same texture serves.
+		let chroma = ctx.register_texture_entry(
+			texture,
+			ImportedTextureDesc {
+				layer: Some((0, 1)),
+				..desc
+			},
+		);
+		let chroma_view = ctx.texture_view(chroma).expect("chroma alias view");
+		let uniform = ctx.device.create_buffer(&wgpu::BufferDescriptor {
+			label: Some("oakrender-imported-test-planar-params"),
+			size: 64,
+			usage: wgpu::BufferUsages::UNIFORM,
+			mapped_at_creation: false,
+		});
+		let scope = ctx.device.push_error_scope(wgpu::ErrorFilter::Validation);
+		let _bind_group = ctx.device.create_bind_group(&wgpu::BindGroupDescriptor {
+			label: Some("oakrender-imported-test-planar-bg"),
+			layout: &pipeline.layout,
+			entries: &[
+				wgpu::BindGroupEntry {
+					binding: 0,
+					resource: wgpu::BindingResource::TextureView(&single_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 1,
+					resource: wgpu::BindingResource::TextureView(&chroma_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 2,
+					resource: uniform.as_entire_binding(),
+				},
+			],
+		});
+		let error = pollster_block_on(scope.pop());
+		assert!(
+			error.is_none(),
+			"single-slice aliases must bind into the planar pass's D2 layout \
+			 (a D2Array view is rejected here): {error:?}"
+		);
+
+		// Negative control: the broader (D2Array) alias is rejected by the
+		// same layout, proving mismatch in the assertion above is
+		// observable through the error scope.
+		let scope = ctx.device.push_error_scope(wgpu::ErrorFilter::Validation);
+		let _rejected = ctx.device.create_bind_group(&wgpu::BindGroupDescriptor {
+			label: Some("oakrender-imported-test-planar-bg-rejected"),
+			layout: &pipeline.layout,
+			entries: &[
+				wgpu::BindGroupEntry {
+					binding: 0,
+					resource: wgpu::BindingResource::TextureView(&pair_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 1,
+					resource: wgpu::BindingResource::TextureView(&chroma_view),
+				},
+				wgpu::BindGroupEntry {
+					binding: 2,
+					resource: uniform.as_entire_binding(),
+				},
+			],
+		});
+		let error = pollster_block_on(scope.pop());
+		assert!(
+			error.is_some(),
+			"a D2Array alias must be rejected by the D2 planar layout"
+		);
+	}
+}
@@ -78,4 +78,9 @@ pub mod lut;
 pub use handle::CHandle;
 pub use rational::Rational;
 pub use samplefmt::{PixelFormat, SampleFormat};
+
+/// The wgpu generation the engine and the UI share (M2/M5). Re-exported
+/// so consumers of the external-import API can name texture formats
+/// without another dependency on the same major version.
+pub use wgpu;
 pub use timerange::{TimeRange, TimeRangeList};
@@ -20,8 +20,8 @@ use crate::PixelFormat;
 use crate::Rational;
 use std::sync::Arc;

-use crate::backend::{BackendKind, GpuContextLike};
-use crate::error::Result;
+use crate::backend::{BackendKind, GpuContextLike, YuvTransform};
+use crate::error::{Error, Result};
 use crate::frame::VideoParamsPod;

 /// A CPU frame (the payload oakcodec frames bridge into, and the value
@@ -199,6 +199,146 @@ impl Drop for GpuLease {
 	}
 }

+/// A hardware surface's chroma layout, as imported for the planar GPU
+/// textures (M5).
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub enum PlanarFormat {
+	/// 8-bit NV12: R8 luma + interleaved R8G8 chroma.
+	Nv12,
+	/// 10-bit P010: 16-bit left-aligned luma + interleaved 16-bit chroma.
+	P010,
+}
+
+impl PlanarFormat {
+	/// The single-channel luma plane format.
+	pub fn y_format(self) -> wgpu::TextureFormat {
+		match self {
+			PlanarFormat::Nv12 => wgpu::TextureFormat::R8Unorm,
+			PlanarFormat::P010 => wgpu::TextureFormat::R16Unorm,
+		}
+	}
+
+	/// The interleaved two-channel chroma plane format.
+	pub fn uv_format(self) -> wgpu::TextureFormat {
+		match self {
+			PlanarFormat::Nv12 => wgpu::TextureFormat::Rg8Unorm,
+			PlanarFormat::P010 => wgpu::TextureFormat::Rg16Unorm,
+		}
+	}
+
+	/// Display name (diagnostics).
+	pub fn name(self) -> &'static str {
+		match self {
+			PlanarFormat::Nv12 => "nv12",
+			PlanarFormat::P010 => "p010",
+		}
+	}
+
+	/// The source bit depth (the limited-range expansion depth for
+	/// [`crate::backend::YuvTransform::from_matrix_depth`]).
+	pub fn bit_depth(self) -> u32 {
+		match self {
+			PlanarFormat::Nv12 => 8,
+			PlanarFormat::P010 => 16,
+		}
+	}
+}
+
+/// A zero-copy hardware-decode frame (M5): the decoder's YUV hardware
+/// surface imported as two GPU plane textures (luma + interleaved
+/// chroma), plus everything the render thread needs to turn it into a
+/// working-space RGBA texture — [`GpuContext::run_planar_yuv_to_rgb`]
+/// for the matrix/range and the raw source colorimetry for the
+/// source→working transform.
+///
+/// This is a *transient* texture: footage decode produces it and the
+/// footage path resolves it to `Texture::Gpu` before any node samples
+/// it. Every other consumer sees only `Gpu`/`Cpu`.
+///
+/// [`GpuContext::run_planar_yuv_to_rgb`]: crate::backend::GpuContext::run_planar_yuv_to_rgb
+#[derive(Clone)]
+pub struct PlanarTexture {
+	/// Luma plane token (registry entry, `R8Unorm`/`R16Unorm`).
+	pub y: u64,
+	/// Interleaved chroma plane token (`R8G8Unorm`/`R16G16Unorm`).
+	pub uv: u64,
+	/// Plane layout (bit depth and chroma interleave).
+	pub format: PlanarFormat,
+	/// Frame width (luma resolution).
+	pub width: i32,
+	/// Frame height (luma resolution).
+	pub height: i32,
+	/// The YUV→RGB matrix/range for the planes.
+	pub transform: YuvTransform,
+	/// Source color-primaries code point (`AVCOL_PRI_*`).
+	pub color_primaries: i32,
+	/// Source transfer-characteristic code point (`AVCOL_TRC_*`).
+	pub color_trc: i32,
+	/// The context owning both plane textures.
+	pub ctx: Arc<dyn GpuContextLike>,
+	/// Plane-token ownership: dropping the last `PlanarTexture` clone
+	/// releases both registry tokens (and with them the imported
+	/// surfaces). Never read directly.
+	#[allow(dead_code)]
+	y_lease: Arc<GpuLease>,
+	#[allow(dead_code)]
+	uv_lease: Arc<GpuLease>,
+	/// Extra lifetime guard for platforms whose GPU texture types cannot
+	/// carry a drop callback (macOS/Metal): the decoder's pixel-buffer
+	/// reference is held here and released with the planar texture.
+	#[allow(dead_code)]
+	pub keep_alive: Option<Arc<dyn std::any::Any + Send + Sync>>,
+}
+
+impl PlanarTexture {
+	/// Build a planar texture from imported plane tokens: `planes` is
+	/// `(luma, chroma)`, `size` the luma dimensions.
+	pub fn new(
+		ctx: Arc<dyn GpuContextLike>,
+		format: PlanarFormat,
+		size: (i32, i32),
+		planes: (u64, u64),
+		transform: YuvTransform,
+		color: (i32, i32),
+	) -> Self {
+		let (width, height) = size;
+		let (y, uv) = planes;
+		let (color_primaries, color_trc) = color;
+		Self {
+			y,
+			uv,
+			format,
+			width,
+			height,
+			transform,
+			color_primaries,
+			color_trc,
+			y_lease: GpuLease::new(ctx.clone(), y),
+			uv_lease: GpuLease::new(ctx.clone(), uv),
+			keep_alive: None,
+			ctx,
+		}
+	}
+
+	/// Attach a non-CPU lifetime guard (macOS: the decoder's pixel
+	/// buffer), released with this texture.
+	pub fn set_keep_alive(&mut self, guard: Arc<dyn std::any::Any + Send + Sync>) {
+		self.keep_alive = Some(guard);
+	}
+}
+
+impl std::fmt::Debug for PlanarTexture {
+	fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+		f.debug_struct("PlanarTexture")
+			.field("y", &self.y)
+			.field("uv", &self.uv)
+			.field("format", &self.format)
+			.field("width", &self.width)
+			.field("height", &self.height)
+			.finish_non_exhaustive()
+	}
+}
+
 /// A texture: either backend-resident (GPU) or a CPU-frame wrapper.
 ///
 /// GPU textures carry an `Arc` to their [`GpuContext`](crate::backend::GpuContext)
@@ -226,6 +366,9 @@ pub enum Texture {
 		/// Shared token ownership (destroyed with the last clone).
 		lease: Arc<GpuLease>,
 	},
+	/// Imported hardware-decode YUV frame (M5); resolved to `Gpu` by the
+	/// footage path before consumers run.
+	Planar(PlanarTexture),
 	/// CPU-frame wrapper (uploaded lazily by the backend).
 	Cpu(Frame),
 }
@@ -248,6 +391,7 @@ impl std::fmt::Debug for Texture {
 				.field("height", height)
 				.field("format", format)
 				.finish(),
+			Texture::Planar(p) => f.debug_tuple("Texture::Planar").field(p).finish(),
 			Texture::Cpu(frame) => f.debug_tuple("Texture::Cpu").field(frame).finish(),
 		}
 	}
@@ -284,7 +428,7 @@ impl Texture {
 	/// True for dummy textures.
 	pub fn is_dummy(&self) -> bool {
 		match self {
-			Texture::Gpu { .. } => false,
+			Texture::Gpu { .. } | Texture::Planar(_) => false,
 			Texture::Cpu(f) => f.is_dummy(),
 		}
 	}
@@ -294,11 +438,33 @@ impl Texture {
 		Texture::Cpu(frame)
 	}

-	/// Read back into a CPU frame (downloads for GPU textures).
+	/// Wrap an imported hardware YUV frame (M5).
+	pub fn wrap_planar(planar: PlanarTexture) -> Self {
+		Texture::Planar(planar)
+	}
+
+	/// The imported planar frame, when this is one.
+	pub fn as_planar(&self) -> Option<&PlanarTexture> {
+		match self {
+			Texture::Planar(p) => Some(p),
+			_ => None,
+		}
+	}
+
+	/// True for imported hardware YUV frames (M5) — the footage path's
+	/// resolve trigger.
+	pub fn is_planar(&self) -> bool {
+		matches!(self, Texture::Planar(_))
+	}
+
+	/// Read back into a CPU frame (downloads for GPU textures). An
+	/// unresolved planar texture is an error: the caller must resolve
+	/// (YUV→RGB) through the footage path first.
 	pub fn to_frame(&self) -> Result<Frame> {
 		match self {
 			Texture::Cpu(f) => Ok(f.clone()),
 			Texture::Gpu { token, ctx, .. } => ctx.download(*token),
+			Texture::Planar(_) => Err(Error::State),
 		}
 	}

@@ -306,14 +472,18 @@ impl Texture {
 	pub fn size(&self) -> (i32, i32) {
 		match self {
 			Texture::Gpu { width, height, .. } => (*width, *height),
+			Texture::Planar(p) => (p.width, p.height),
 			Texture::Cpu(f) => (f.width, f.height),
 		}
 	}

-	/// The texture's pixel format.
+	/// The texture's pixel format. A planar texture reports `F32`: it is
+	/// transient, and the footage path always resolves it to an F32 RGBA
+	/// GPU texture before any consumer sees it.
 	pub fn format(&self) -> PixelFormat {
 		match self {
 			Texture::Gpu { format, .. } => *format,
+			Texture::Planar(_) => PixelFormat::F32,
 			Texture::Cpu(f) => f.format,
 		}
 	}
@@ -322,6 +492,7 @@ impl Texture {
 	pub fn backend(&self) -> BackendKind {
 		match self {
 			Texture::Gpu { backend, .. } => *backend,
+			Texture::Planar(p) => p.ctx.kind(),
 			Texture::Cpu(_) => BackendKind::Cpu,
 		}
 	}
@@ -1216,3 +1216,4 @@ pub struct ClipPreferences {
 	/// field 处理模式（kOfxImageField*）。
 	pub field: String,
 }
+ 	/// field 处理模式（kOfxImageField*）。
@@ -623,3 +623,8 @@ pub(crate) fn write_output_frame(dst: &mut Texture, image: &Image) -> crate::err
 	}
 	Ok(())
 }
+ 			ctx.upload(*token, &frame)
+ 				.map_err(|e| Error::Failed(format!("输出纹理上传失败：{e}")))?;
+ 		}
+		// 未解析的平面纹理不会成为插件输出目标（解码路径会先解析）。
+		Texture::Planar(_) => {
@@ -769,3 +769,4 @@ pub fn suite_v1() -> &'static PropertySuiteV1 {
 		get_dimension: prop_get_dimension,
 	})
 }
+ 		get_dimension: prop_get_dimension,
@@ -345,7 +345,9 @@ impl RenderEvalHooks {
                height,
                ..
            } => (*token, ctx.clone(), *width, *height),
-            Texture::Cpu(_) => unreachable!("handled above"),
+            // Imported planar frames never reach node processing (the
+            // footage path resolves them first); pass through defensively.
+            Texture::Cpu(_) | Texture::Planar(_) => return Ok(tex),
        };
        let concrete = ctx
            .as_any()
@@ -869,6 +871,11 @@ impl RenderEvalHooks {
                    scratch.push(token);
                    (token, (frame.width, frame.height))
                }
+                // Resolved by the footage path; never a shader input.
+                Texture::Planar(_) => {
+                    warn("planar texture reached a shader job input (unresolved)");
+                    return None;
+                }
            };
            // The pass size follows the effect input's texture (C++ the
            // job's video params = the main input size); any other bound
@@ -1125,7 +1132,7 @@ const MAX_CACHED_DECODERS: usize = 6;
 /// LRU tick source for [`DECODERS`].
 static DECODER_TICK: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(1);

-/// Count of [`render_footage_frame_inner`] entries, i.e. of real codec work
+/// Count of [`render_footage_frame_inner_opts`] entries, i.e. of real codec work
 /// on the footage path (frame-cache hits and decode-service LRU hits do not
 /// count). Visible for the decode-service tests, which use it to tell a
 /// cached frame from a re-decode.
@@ -1151,28 +1158,76 @@ pub fn reset_decode_invocations() {
 /// on every pre-render restart and on repeat plays).
 const MAX_CACHED_FRAMES: usize = 24;

-/// Decoded-frame LRU: `(filename, stream, time, w, h)` -> F32 CPU frame.
-/// Frame data is the expensive part (a 1080p frame ≈ 31 MB); the decoder
-/// session cache alone still re-decodes every `render_footage_frame`.
-/// The size is part of the key: the same media at a different target
-/// resolution is a different frame (an interleaved source-monitor/proxy
-/// request must not reuse a wrongly-sized pixel buffer).
-static DECODED_FRAMES: std::sync::OnceLock<
-    std::sync::Mutex<
-        std::collections::HashMap<(String, i32, (i64, i64), i32, i32), (Frame, u64)>,
-    >,
-> = std::sync::OnceLock::new();
+/// Cap on cached PLANAR hardware frames (M5 audit): each planar entry
+/// pins a decoder surface through its keep-alive guard, and the driver's
+/// surface pool is finite. Planar textures only need to bridge the decode
+/// thread → resolve, so they get a much smaller budget than the general
+/// frame LRU. Eviction drops only the import; the decoder session cache
+/// still holds the raw frame, so a re-request re-imports (no re-decode).
+const MAX_CACHED_PLANAR_FRAMES: usize = 4;

-fn decoded_frames() -> std::sync::MutexGuard<
-    'static,
-    std::collections::HashMap<(String, i32, (i64, i64), i32, i32), (Frame, u64)>,
-> {
+/// Decoded-frame LRU key: `(filename, stream, time, w, h)`. The size is
+/// part of the key: the same media at a different target resolution is a
+/// different frame (an interleaved source-monitor/proxy request must not
+/// reuse a wrongly-sized pixel buffer).
+type DecodedFrameKey = (String, i32, (i64, i64), i32, i32);
+
+/// Decoded-frame LRU map. Values are decoded textures (CPU frame, or the
+/// M5 imported planar texture when the hardware import took the frame)
+/// plus their LRU tick. Frame data is the expensive part (a 1080p F32
+/// frame ≈ 31 MB); the decoder session cache alone still re-decodes every
+/// `render_footage_frame`.
+type DecodedFrameCache = std::collections::HashMap<DecodedFrameKey, (Texture, u64)>;
+
+static DECODED_FRAMES: std::sync::OnceLock<std::sync::Mutex<DecodedFrameCache>> =
+    std::sync::OnceLock::new();
+
+fn decoded_frames() -> std::sync::MutexGuard<'static, DecodedFrameCache> {
    DECODED_FRAMES
        .get_or_init(|| std::sync::Mutex::new(std::collections::HashMap::new()))
        .lock()
        .unwrap_or_else(|e| e.into_inner())
 }

+/// Insert `texture` under `key` with the general and planar caps applied.
+fn insert_cached_frame(cache: &mut DecodedFrameCache, key: DecodedFrameKey, texture: Texture, tick: u64) {
+    if cache.contains_key(&key) {
+        return;
+    }
+    while cache.len() >= MAX_CACHED_FRAMES {
+        let Some(victim) = cache
+            .iter()
+            .filter(|(_, (_, t))| *t > 0)
+            .min_by_key(|(_, (_, t))| *t)
+            .map(|(k, _)| k.clone())
+        else {
+            break;
+        };
+        cache.remove(&victim);
+    }
+    if texture.is_planar() {
+        prune_planar_frames(cache, MAX_CACHED_PLANAR_FRAMES.saturating_sub(1));
+    }
+    cache.insert(key, (texture, tick));
+}
+
+/// Evict the least-recently-used planar entries until at most `keep`
+/// remain (hardware-surface residency bound, M5 audit).
+fn prune_planar_frames(cache: &mut DecodedFrameCache, keep: usize) {
+    let mut planar: Vec<(DecodedFrameKey, u64)> = cache
+        .iter()
+        .filter(|(_, (t, _))| t.is_planar())
+        .map(|(k, (_, tick))| (k.clone(), *tick))
+        .collect();
+    if planar.len() <= keep {
+        return;
+    }
+    planar.sort_by_key(|(_, tick)| *tick);
+    for (key, _) in planar.drain(..planar.len() - keep) {
+        cache.remove(&key);
+    }
+}
+
 /// Time-key rational (the frame-LRU's deterministic key half).
 fn time_key(t: &Rational) -> (i64, i64) {
    (t.numerator(), t.denominator())
@@ -1272,24 +1327,67 @@ pub fn render_footage_frame(
    format: PixelFormat,
 ) -> Result<Texture> {
    let started = std::time::Instant::now();
-    let result = match crate::pipeline::decode_service() {
-        Some(service) => {
-            let request = crate::pipeline::DecodeRequest {
-                filename: filename.to_string(),
+    let decode = |allow_import: bool| -> Result<Texture> {
+        match crate::pipeline::decode_service() {
+            Some(service) => {
+                let request = crate::pipeline::DecodeRequest {
+                    filename: filename.to_string(),
+                    stream_index,
+                    time,
+                    size,
+                    format,
+                    allow_import: true,
+                };
+                match service.request(request) {
+                    Some(result) => result,
+                    // The service went away (shutdown) mid-flight: decode here
+                    // rather than failing the frame.
+                    None => render_footage_frame_inner_opts(
+                        filename,
+                        stream_index,
+                        time,
+                        size,
+                        format,
+                        allow_import,
+                    ),
+                }
+            }
+            None => render_footage_frame_inner_opts(
+                filename,
                stream_index,
                time,
                size,
                format,
-            };
-            match service.request(request) {
-                Some(result) => result,
-                // The service went away (shutdown) mid-flight: decode here
-                // rather than failing the frame.
-                None => render_footage_frame_inner(filename, stream_index, time, size, format),
-            }
+                allow_import,
+            ),
        }
-        None => render_footage_frame_inner(filename, stream_index, time, size, format),
    };
+    let mut result = decode(true);
+    // M5: an imported hardware frame is planar YUV; resolve it to
+    // working-space RGBA on the GPU before it leaves the footage path.
+    // A resolution failure is a per-frame fallback: re-decode through the
+    // CPU scaler with the import disabled (the decoder session cache makes
+    // this a transfer, not a second decode).
+    if let Ok(texture) = &result {
+        if texture.is_planar() {
+            result = match resolve_planar_footage(texture) {
+                Ok(resolved) => Ok(resolved),
+                Err(err) => {
+                    eprintln!("planar footage resolve failed, staging fallback: {err:#}");
+                    // Bypass the decode service: its cache still holds the
+                    // planar entry that just failed to resolve.
+                    render_footage_frame_inner_opts(
+                        filename,
+                        stream_index,
+                        time,
+                        size,
+                        format,
+                        false,
+                    )
+                }
+            };
+        }
+    }
    if std::env::var_os("OAK_PERF").is_some() {
        eprintln!(
            "[decode] {:?} s{} time {}/{} ({:.3}s) size {:?} -> {:.3}s {:?}",
@@ -1306,15 +1404,84 @@ pub fn render_footage_frame(
    result
 }

-/// The synchronous decode behind [`render_footage_frame`] (frame-LRU →
-/// decoder session → codec → F32 frame). `pub(crate)` because the decode
-/// service runs exactly this on its own thread.
-pub(crate) fn render_footage_frame_inner(
+/// The CPU-staging decode for consumers that composite on the CPU (the
+/// montage compositor): same service/FRU semantics as
+/// [`render_footage_frame`], but the M5 zero-copy import is disabled and
+/// the result is always a CPU frame.
+///
+/// This is not just an optimization: the montage compositor runs on the
+/// CPU, so an imported `Texture::Gpu` would have to be downloaded again —
+/// and before this entry point existed it was silently SKIPPED (the M5
+/// audit's all-black montage bug). A planar texture must never reach a
+/// CPU consumer here.
+pub(crate) fn render_footage_frame_staged(
    filename: &str,
    stream_index: i32,
    time: Rational,
    size: (i32, i32),
    format: PixelFormat,
+) -> Result<Texture> {
+    let result = match crate::pipeline::decode_service() {
+        Some(service) => {
+            let request = crate::pipeline::DecodeRequest {
+                filename: filename.to_string(),
+                stream_index,
+                time,
+                size,
+                format,
+                allow_import: false,
+            };
+            match service.request(request) {
+                Some(result) => result,
+                None => render_footage_frame_inner_opts(
+                    filename,
+                    stream_index,
+                    time,
+                    size,
+                    format,
+                    false,
+                ),
+            }
+        }
+        None => render_footage_frame_inner_opts(
+            filename,
+            stream_index,
+            time,
+            size,
+            format,
+            false,
+        ),
+    };
+    // Defensive: a planar texture (a stale service entry from an older
+    // key layout) can never satisfy a staging request — re-decode CPU.
+    match result {
+        Ok(texture) if texture.is_planar() => render_footage_frame_inner_opts(
+            filename,
+            stream_index,
+            time,
+            size,
+            format,
+            false,
+        ),
+        other => other,
+    }
+}
+
+/// The synchronous decode behind [`render_footage_frame`] (frame-LRU →
+/// decoder session → codec → F32 frame). `pub(crate)` because the decode
+/// service runs exactly this on its own thread, with the request's own
+/// `allow_import` flag (the montage requests stage on purpose, §M5 audit).
+///
+/// `allow_import` false forces the CPU staging path: the resolve step's
+/// fallback re-decodes with it off, and the montage compositor never
+/// wants a GPU frame.
+pub(crate) fn render_footage_frame_inner_opts(
+    filename: &str,
+    stream_index: i32,
+    time: Rational,
+    size: (i32, i32),
+    format: PixelFormat,
+    allow_import: bool,
 ) -> Result<Texture> {
    // (file, stream, time) repeatedly (pre-render restarts, repeated
    // scale-up at the same time, graph + montage interleaving); the
@@ -1327,10 +1494,16 @@ pub(crate) fn render_footage_frame_inner(
    {
        let mut cache = decoded_frames();
        let tick = DECODER_TICK.fetch_add(1, std::sync::atomic::Ordering::Relaxed);
-        let frame = cache.get(&cache_key).map(|(f, _)| f.clone());
-        if let Some(frame) = frame {
-            cache.insert(cache_key.clone(), (frame.clone(), tick));
-            return Ok(Texture::wrap_frame(frame));
+        let texture = cache.get(&cache_key).map(|(f, _)| f.clone());
+        if let Some(texture) = texture {
+            // A planar entry cannot satisfy a staging request (the resolve
+            // step failed and asked for CPU pixels): treat it as a miss and
+            // decode through the scaler instead.
+            if allow_import || !texture.is_planar() {
+                cache.insert(cache_key.clone(), (texture.clone(), tick));
+                return Ok(texture);
+            }
+            cache.remove(&cache_key);
        }
    }
    // Past the frame cache: this call really goes to the codec (the
@@ -1355,6 +1528,20 @@ pub(crate) fn render_footage_frame_inner(
            None
        },
    };
+
+    // M5: hardware frames are offered to the zero-copy import first; the
+    // CPU retrieval below is the per-frame staging fallback (the decoder
+    // session cache holds the raw surface, so the fallback never
+    // re-decodes).
+    if allow_import {
+        if let Ok(Some(imported)) = decoder.retrieve_video_frame_gpu(&params) {
+            let mut cache = decoded_frames();
+            let tick = DECODER_TICK.fetch_add(1, std::sync::atomic::Ordering::Relaxed);
+            insert_cached_frame(&mut cache, cache_key, imported.clone(), tick);
+            return Ok(imported);
+        }
+    }
+
    let decoded = decoder
        .retrieve_video_frame(&params)
        .map_err(|e| Error::Failed(format!("footage decode at {time:?}: {e:?}")))?;
@@ -1398,24 +1585,127 @@ pub(crate) fn render_footage_frame_inner(
    // by default; the legacy sRGB working space keeps the pass-through).
    convert_decoded_to_working(&mut dst, &decoded);
    // Memoize the finished working-space frame (LRU-capped).
+    let texture = Texture::wrap_frame(dst);
    {
        let mut cache = decoded_frames();
        let tick = DECODER_TICK.fetch_add(1, std::sync::atomic::Ordering::Relaxed);
-        if !cache.contains_key(&cache_key) {
-            if cache.len() >= MAX_CACHED_FRAMES {
-                if let Some(victim) = cache
-                    .iter()
-                    .filter(|(_, (_, t))| *t > 0)
-                    .min_by_key(|(_, (_, t))| *t)
-                    .map(|(k, _)| k.clone())
-                {
-                    cache.remove(&victim);
-                }
+        insert_cached_frame(&mut cache, cache_key, texture.clone(), tick);
+    }
+    Ok(texture)
+}
+
+/// Resolve an imported planar hardware frame (M5) to working-space RGBA
+/// on the GPU: the YUV→RGB pass (matrix + range from the frame's own
+/// colorimetry) followed by the source→working transform. In the legacy
+/// sRGB working space the transform is a pass-through (matching the CPU
+/// path's `convert_decoded_to_working`); otherwise it is applied as a CPU
+/// reference-baked 3D LUT, exactly like the OCIO color transforms.
+fn resolve_planar_footage(texture: &Texture) -> Result<Texture> {
+    let planar = texture
+        .as_planar()
+        .ok_or_else(|| Error::Failed("resolve_planar_footage: not planar".into()))?;
+    let Some(gpu) = planar
+        .ctx
+        .as_any()
+        .and_then(|a| a.downcast_ref::<oak_core::backend::GpuContext>())
+    else {
+        return Err(Error::Failed(
+            "planar texture belongs to a non-wgpu context".into(),
+        ));
+    };
+    let (w, h) = (planar.width.max(1), planar.height.max(1));
+    let dst = gpu
+        .create_texture(w, h)
+        .map_err(|e| Error::Failed(format!("planar resolve target: {e:?}")))?;
+    if let Err(e) = gpu.run_planar_yuv_to_rgb(planar.y, planar.uv, dst, &planar.transform) {
+        gpu.destroy_texture(dst);
+        return Err(Error::Failed(format!("planar YUV pass: {e:?}")));
+    }
+    let token = match footage_working_lut(planar.color_primaries, planar.color_trc) {
+        Some((key, lut)) => match gpu.apply_color_lut(dst, &key, &lut) {
+            Ok(transformed) => {
+                gpu.destroy_texture(dst);
+                transformed
+            }
+            Err(e) => {
+                eprintln!("footage working-space LUT failed, using source RGB: {e:?}");
+                dst
+            }
+        },
+        None => dst,
+    };
+    Ok(Texture::gpu(planar.ctx.clone(), token, w, h, PixelFormat::F32))
+}
+
+/// The source→working-space 3D LUT for a decoded frame (M5 GPU import
+/// path): baked from the same CPU reference (`colormath::decode_to_acescg`)
+/// the staging path uses, cached per colorimetry. `None` in the legacy
+/// sRGB working space (pass-through) or when the LUT cannot be built.
+fn footage_working_lut(
+    color_primaries: i32,
+    color_trc: i32,
+) -> Option<(String, std::sync::Arc<oak_core::lut::Lut3d>)> {
+    use oak_core::colormath::WorkingColorSpace;
+    if oak_core::color::pipeline_working_space() == WorkingColorSpace::SrgbLegacy {
+        return None;
+    }
+    let key = format!("footage/{color_primaries}/{color_trc}");
+    type Cache = std::sync::Mutex<
+        std::collections::HashMap<String, std::sync::Arc<oak_core::lut::Lut3d>>,
+    >;
+    static CACHE: std::sync::OnceLock<Cache> = std::sync::OnceLock::new();
+    let mut cache = CACHE
+        .get_or_init(|| std::sync::Mutex::new(std::collections::HashMap::new()))
+        .lock()
+        .unwrap_or_else(|e| e.into_inner());
+    if let Some(lut) = cache.get(&key) {
+        return Some((key, lut.clone()));
+    }
+    let lut = build_footage_working_lut(color_primaries, color_trc)?;
+    if cache.len() >= 16 {
+        cache.clear();
+    }
+    cache.insert(key.clone(), lut.clone());
+    Some((key, lut))
+}
+
+/// Bake the source→ACEScg transform into a 3D LUT over the display
+/// domain (the same domain the color-transform LUTs use).
+fn build_footage_working_lut(
+    color_primaries: i32,
+    color_trc: i32,
+) -> Option<std::sync::Arc<oak_core::lut::Lut3d>> {
+    use oak_core::colormath::{decode_to_acescg, source_primaries_from_av, source_transfer_from_av};
+    use oak_core::lut::Lut3d;
+    let edge = Lut3d::DISPLAY_EDGE;
+    let (lo, hi) = (Lut3d::DISPLAY_LO, Lut3d::DISPLAY_HI);
+    let n = (edge as usize).pow(3);
+    let mut samples = vec![0.0f32; n * 4];
+    let step = |i: usize, axis: usize| -> f32 {
+        let t = i as f32 / (edge - 1) as f32;
+        lo[axis] + (hi[axis] - lo[axis]) * t
+    };
+    for b in 0..edge as usize {
+        for g in 0..edge as usize {
+            for r in 0..edge as usize {
+                let idx = ((b * edge as usize + g) * edge as usize + r) * 4;
+                samples[idx] = step(r, 0);
+                samples[idx + 1] = step(g, 1);
+                samples[idx + 2] = step(b, 2);
+                samples[idx + 3] = 1.0;
            }
-            cache.insert(cache_key.clone(), (dst.clone(), tick));
        }
    }
-    Ok(Texture::wrap_frame(dst))
+    decode_to_acescg(
+        &mut samples,
+        source_primaries_from_av(color_primaries),
+        source_transfer_from_av(color_trc),
+    );
+    let mut data = Vec::with_capacity(n * 3);
+    for px in samples.chunks_exact(4) {
+        data.extend_from_slice(&px[..3]);
+    }
+    Some(std::sync::Arc::new(Lut3d { edge, lo, hi, data }))
 }

 /// Convert a decoded footage frame (display-referred RGB in the source's
@@ -1532,6 +1822,11 @@ fn composite_tracks_gpu(
 					scratch.push(t);
 					t
 				}
+				Texture::Planar(_) => {
+					return Err(Error::Failed(
+						"unresolved planar texture in composite".into(),
+					))
+				}
 			};
 			let out = ctx.create_texture(w, h)?;
 			ctx.run_shader_pass(&program, &[], &[current, src_token], out)?;
@@ -2503,7 +2798,8 @@ pub fn render_montage_frame_into(
            continue;
        }
        let media_time = clip.media_in + (time - clip.in_time);
-        let decoded = render_footage_frame(
+        // CPU compositor: stage on purpose (never an imported GPU frame).
+        let decoded = render_footage_frame_staged(
            &clip.filename,
            clip.stream_index,
            media_time,
@@ -2514,9 +2810,24 @@ pub fn render_montage_frame_into(
        // (C++ semantics: the clip texture passes through the chain
        // bottom-up, the chain top feeds the track composite).
        let effected = apply_clip_effects(decoded, clip, time);
+        // The compositor is CPU-side. The decode is staged above, but a
+        // clip effect may still hand back a GPU texture: read it back
+        // (an explicit CPU boundary) instead of silently dropping the
+        // clip (M5 audit: that skip produced all-black sequences).
+        let downloaded: Frame;
        let (src_data, src_stride) = match &effected {
            Texture::Cpu(src) => (&src.data, src.linesize_bytes() as i32),
-            _ => continue,
+            Texture::Gpu { .. } => {
+                downloaded = effected.to_frame().map_err(|e| {
+                    Error::Failed(format!("montage GPU readback failed: {e:?}"))
+                })?;
+                (&downloaded.data, downloaded.linesize_bytes() as i32)
+            }
+            Texture::Planar(_) => {
+                return Err(Error::Failed(
+                    "unresolved planar texture reached the montage compositor".into(),
+                ))
+            }
        };
        composite_over(dst, dst_stride, w, h, src_data, src_stride, clip.gain);
    }
@@ -4420,3 +4731,4 @@ mod tests {

 }

+    static LOCK: Mutex<()> = Mutex::new(());
@@ -153,6 +153,9 @@ fn playback_decode_requests(
 			time: clip.media_in + (time - clip.in_time),
 			size,
 			format: PixelFormat::F32,
+			// The montage compositor is CPU-side: importing would only
+			// be downloaded again.
+			allow_import: false,
 		});
 	}
 	if let Some((filename, stream_index)) = &params.footage {
@@ -165,6 +168,7 @@ fn playback_decode_requests(
 			time,
 			size,
 			format: params.force_format.unwrap_or(PixelFormat::F32),
+			allow_import: true,
 		});
 	}
 	requests
@@ -191,6 +195,11 @@ pub struct DecodeRequest {
 	pub size: (i32, i32),
 	/// Target pixel format.
 	pub format: PixelFormat,
+	/// Whether the M5 zero-copy hardware import may serve this request.
+	/// The CPU montage compositor sets `false` (it needs CPU pixels); the
+	/// flag is part of the cache key, so the service never hands an
+	/// imported texture to a staging consumer (M5 audit).
+	pub allow_import: bool,
 }

 /// Decode-service counters (M1's evidence that the service really decodes
@@ -553,12 +562,13 @@ fn prefetch_into(
 /// decode thread.
 fn decode(request: &DecodeRequest, inner: &DecodeInner) -> Result<Texture> {
 	inner.counters.decodes.fetch_add(1, Ordering::Relaxed);
-	let result = crate::eval::render_footage_frame_inner(
+	let result = crate::eval::render_footage_frame_inner_opts(
 		&request.filename,
 		request.stream_index,
 		request.time,
 		request.size,
 		request.format,
+		request.allow_import,
 	);
 	if result.is_err() {
 		inner.counters.errors.fetch_add(1, Ordering::Relaxed);
@@ -295,6 +295,13 @@ impl ExportTask {
 					"Render frame readback for the encoder failed: {e:?}"
 				))
 			})?,
+			// Unresolved planar frames never reach an encoder; the
+			// decode path resolves them before delivery.
+			Texture::Planar(_) => {
+				return Err(Error::Failed(
+					"unresolved planar texture for the encoder".into(),
+				))
+			}
 		};
 		let frame = &frame;
 		let params = CommonVideoParams::new_basic(
@@ -1181,3 +1181,4 @@ pub fn add_track_command(project: ProjectRef, list: NodeId) -> UndoCommand {
 		kind: TrackType::Video,
 	})
 }
+ 		kind: TrackType::Video,
@@ -1308,6 +1308,16 @@ fn render_f32_into(
 									}
 								}
 							}
+							Ok(oak_core::texture::Texture::Planar(_)) => {
+								// The footage path resolves imported planar
+								// frames before returning; a planar texture
+								// here is a bug, not a frame.
+								warn_graph_fallback(
+									spec.viewer_node,
+									"unresolved planar texture",
+								);
+								None
+							}
 							Err(e) => {
 								warn_graph_fallback(spec.viewer_node, &e.to_string());
 								None
@@ -1356,6 +1366,11 @@ fn render_f32_into(
 			gpu @ oak_core::texture::Texture::Gpu { .. } => gpu
 				.to_frame()
 				.map_err(|e| format!("decode readback: {e}"))?,
+			// Resolved by the footage path; a planar frame cannot cross
+			// into the shm slot without a readback.
+			oak_core::texture::Texture::Planar(_) => {
+				return Err("unresolved planar decode texture".to_string())
+			}
 		};
 		let src_stride = frame.linesize_bytes() as usize;
 		let row_bytes = (w as usize) * 16;