From 94dbccf69f68fb4fae693162841848a3a52fe742 Mon Sep 17 00:00:00 2001
From: Dexter <37185622+dexterlabs1@users.noreply.github.com>
Date: Tue, 6 Oct 2026 16:40:12 -0500
Subject: [PATCH] fix: contain GPU initialization and retire failed readbacks

---
 Cargo.lock                        |   3 +
 apps/effectcraft-web/src/app.rs   |  17 +
 apps/effectcraft/src/main.rs      |  18 +
 crates/gpu/Cargo.toml             |   5 +
 crates/gpu/src/adv3d.rs           | 371 +++++++++++++++------
 crates/gpu/src/context.rs         | 433 ++++++++++++++++--------
 crates/gpu/src/deferred.rs        | 120 +++++--
 crates/gpu/src/lib.rs             |  44 ++-
 crates/gpu/src/particles.rs       |  23 +-
 crates/gpu/src/readback.rs        | 529 ++++++++++++++++++++++++++++++
 crates/gpu/src/tests.rs           |   2 +
 crates/gpu/src/tests_deferred.rs  |  10 +-
 crates/gpu/src/tests_init.rs      | 161 +++++++++
 crates/ui-egui/src/frames.rs      |  73 ++++-
 crates/ui-egui/src/gpu_failure.rs | 158 +++++++++
 crates/ui-egui/src/lib.rs         | 109 ++++++
 docs/architecture.md              |  16 +-
 17 files changed, 1776 insertions(+), 316 deletions(-)
 create mode 100644 crates/gpu/src/readback.rs
 create mode 100644 crates/gpu/src/tests_init.rs
 create mode 100644 crates/ui-egui/src/gpu_failure.rs

diff --git a/Cargo.lock b/Cargo.lock
index 564eaacd..90cb4e02 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -1314,8 +1314,11 @@ dependencies = [
  "effectcraft-render",
  "effectcraft-text",
  "effectcraft-time",
+ "js-sys",
  "log",
  "pollster 0.4.0",
+ "wasm-bindgen",
+ "web-time",
  "wgpu",
 ]
 
diff --git a/apps/effectcraft-web/src/app.rs b/apps/effectcraft-web/src/app.rs
index c4a5dbac..ded2c56c 100644
--- a/apps/effectcraft-web/src/app.rs
+++ b/apps/effectcraft-web/src/app.rs
@@ -204,6 +204,22 @@ pub async fn start(canvas_id: String) -> Result<(), JsValue> {
             eframe::WebOptions::default(),
             Box::new(move |cc| {
                 CTX.with(|c| *c.borrow_mut() = Some(cc.egui_ctx.clone()));
+                // The page owns this shared presentation device; worker devices retain their
+                // separate owned-device handlers. No browser error-scope future is blocked on.
+                let gpu_failures = effectcraft_ui_egui::gpu_failure::GpuFailureBridge::new(&cc.egui_ctx);
+                if let Some(rs) = &cc.wgpu_render_state {
+                    let errors = gpu_failures.clone();
+                    rs.device.on_uncaptured_error(Arc::new(move |error| {
+                        let message = format!("uncaptured GPU error: {error}");
+                        if errors.report(&message, false) {
+                            log::error!("{message}");
+                        }
+                    }));
+                    let lost = gpu_failures.clone();
+                    rs.device.set_device_lost_callback(move |reason, message| {
+                        lost.report(&format!("GPU device lost ({reason:?}): {message}"), true);
+                    });
+                }
                 let mut session = session();
                 session.load_settings();
                 let recovery = session.begin_recovery();
@@ -215,6 +231,7 @@ pub async fn start(canvas_id: String) -> Result<(), JsValue> {
                 }
                 api::set_info("restored", json!(restored));
                 let mut app = EffectcraftApp::new(session);
+                app.set_gpu_failure_bridge(gpu_failures);
                 app.ui.start_screen = home;
                 if let Some(r) = recovery.filter(|r| !restored && r.autosave.is_some()) {
                     app.offer_recovery(r);
diff --git a/apps/effectcraft/src/main.rs b/apps/effectcraft/src/main.rs
index 5afb3c79..057e6f12 100644
--- a/apps/effectcraft/src/main.rs
+++ b/apps/effectcraft/src/main.rs
@@ -77,6 +77,23 @@ fn main() -> eframe::Result {
         "EffectCraft",
         options,
         Box::new(move |cc| {
+            // This executable owns the shared device's handlers. Install them before
+            // EffectCraft's compositor pipelines are built, without changing handlers from
+            // inside Gpu::new. eframe has already constructed its presentation renderer here.
+            let gpu_failures = effectcraft_ui_egui::gpu_failure::GpuFailureBridge::new(&cc.egui_ctx);
+            if let Some(rs) = &cc.wgpu_render_state {
+                let errors = gpu_failures.clone();
+                rs.device.on_uncaptured_error(std::sync::Arc::new(move |error| {
+                    let message = format!("uncaptured GPU error: {error}");
+                    if errors.report(&message, false) {
+                        log::error!("{message}");
+                    }
+                }));
+                let lost = gpu_failures.clone();
+                rs.device.set_device_lost_callback(move |reason, message| {
+                    lost.report(&format!("GPU device lost ({reason:?}): {message}"), true);
+                });
+            }
             let mut session = effectcraft_host::session();
             // Lazy open and non-blocking auto-save: footage is checked and auto-saves are
             // written on background threads (M13.14).
@@ -110,6 +127,7 @@ fn main() -> eframe::Result {
             }
             let show_home = home.unwrap_or(session.prefs.startup.show_home_on_launch && files.is_empty() && control_port.is_none());
             let mut app = EffectcraftApp::new(session);
+            app.set_gpu_failure_bridge(gpu_failures);
             app.ui.start_screen = show_home;
             if let Some(r) = recovery {
                 app.offer_recovery(r);
diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml
index 36c751df..49fe51c1 100644
--- a/crates/gpu/Cargo.toml
+++ b/crates/gpu/Cargo.toml
@@ -21,9 +21,14 @@ effectcraft-keyframe = { workspace = true }
 wgpu = { workspace = true }
 bytemuck = { workspace = true }
 log = { workspace = true }
+web-time = { workspace = true }
 
 [target.'cfg(not(target_arch = "wasm32"))'.dependencies]
 pollster = { workspace = true }
 
 [dev-dependencies]
 effectcraft-text = { workspace = true }
+
+[target.'cfg(target_arch = "wasm32")'.dependencies]
+wasm-bindgen = "=0.2.129"
+js-sys = "0.3"
diff --git a/crates/gpu/src/adv3d.rs b/crates/gpu/src/adv3d.rs
index 1b059f47..f3b9fee2 100644
--- a/crates/gpu/src/adv3d.rs
+++ b/crates/gpu/src/adv3d.rs
@@ -21,10 +21,9 @@ use crate::context::{Enc, GpuContext, GpuImage};
 
 /// Pipelines (built on first use).
 pub(crate) struct Pipes {
-    bgl: wgpu::BindGroupLayout,
     /// (opaque, transparent); `None` when the adapter can't rasterise ([`raster_unsupported`]):
     /// scenes then render on the CPU.
-    raster: Option<(wgpu::RenderPipeline, wgpu::RenderPipeline)>,
+    raster: Option<(wgpu::BindGroupLayout, wgpu::RenderPipeline, wgpu::RenderPipeline)>,
     post: Post,
 }
 
@@ -45,21 +44,171 @@ const VERTEX_SIZE: u64 = 52;
 /// would be a validation error, which release builds only log (and then draw nothing), so the
 /// caller keeps the scenes on the CPU instead.
 pub(crate) fn raster_unsupported(adapter: &wgpu::Adapter) -> Option<String> {
+    raster_capability_reason(adapter.get_downlevel_capabilities().flags, |format| adapter.get_texture_format_features(format))
+}
+
+fn raster_capability_reason(flags: wgpu::DownlevelFlags, features: impl Fn(wgpu::TextureFormat) -> wgpu::TextureFormatFeatures) -> Option<String> {
     use wgpu::TextureFormatFeatureFlags as F;
-    let needs = [(COLOR, F::BLENDABLE), (DEPTH_OUT, F::empty()), (DEPTH, F::empty())];
-    for (format, flags) in needs {
-        let f = adapter.get_texture_format_features(format);
-        if !f.allowed_usages.contains(wgpu::TextureUsages::RENDER_ATTACHMENT) || !f.flags.contains(flags) {
-            return Some(format!("{format:?} is not a {}render target", if flags.is_empty() { "" } else { "blendable " }));
+    use wgpu::TextureUsages as U;
+    if !flags.contains(wgpu::DownlevelFlags::FRAGMENT_STORAGE) {
+        return Some("no fragment storage buffers".into());
+    }
+    let targets = U::RENDER_ATTACHMENT | U::TEXTURE_BINDING | U::COPY_SRC;
+    let needs = [(COLOR, targets, F::BLENDABLE), (DEPTH_OUT, targets, F::empty()), (DEPTH, U::RENDER_ATTACHMENT, F::empty())];
+    for (format, usages, required_flags) in needs {
+        let f = features(format);
+        if !f.allowed_usages.contains(usages) || !f.flags.contains(required_flags) {
+            return Some(format!("{format:?} lacks raster usages {usages:?} or flags {required_flags:?}"));
         }
     }
     // The transparent pipeline blends colour and masks the depth target off.
-    if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::INDEPENDENT_BLEND) {
+    if !flags.contains(wgpu::DownlevelFlags::INDEPENDENT_BLEND) {
         return Some("no independent blending".into());
     }
     None
 }
 
+/// Bound one allocation by its actual byte stride, without imposing a global image limit.
+fn buffer_bytes(count: u64, stride: u64, limit: u64) -> Option<u64> {
+    let size = count.checked_mul(stride)?;
+    (size > 0 && size <= limit && usize::try_from(size).is_ok()).then_some(size)
+}
+
+/// Sizes of exactly the arrays serialized by pack, including its empty storage padding.
+/// Checking these before serialization bounds host packing as well as the GPU uploads.
+fn packed_bytes(counts: [u64; 6], uniform_limit: u64, storage_limit: u64, buffer_limit: u64) -> Option<[u64; 7]> {
+    let mut sizes = [0; 7];
+    sizes[0] = buffer_bytes(40, 4, uniform_limit.min(buffer_limit))?;
+    for (index, (count, stride)) in counts.into_iter().zip([96, 16, 16, 80, 80, 4]).enumerate() {
+        let size = count.checked_mul(stride)?.max(16);
+        sizes[index + 1] = buffer_bytes(size, 1, storage_limit.min(buffer_limit))?;
+    }
+    Some(sizes)
+}
+
+#[cfg(test)]
+mod capability_tests {
+    use super::*;
+
+    fn supported(format: wgpu::TextureFormat) -> wgpu::TextureFormatFeatures {
+        let rt = wgpu::TextureUsages::RENDER_ATTACHMENT;
+        wgpu::TextureFormatFeatures {
+            allowed_usages: if format == DEPTH { rt } else { rt | wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_SRC },
+            flags: if format == COLOR { wgpu::TextureFormatFeatureFlags::BLENDABLE } else { wgpu::TextureFormatFeatureFlags::empty() },
+        }
+    }
+
+    fn downlevel() -> wgpu::DownlevelFlags {
+        wgpu::DownlevelFlags::FRAGMENT_STORAGE | wgpu::DownlevelFlags::INDEPENDENT_BLEND
+    }
+
+    #[test]
+    fn matching_raster_capabilities_are_accepted_without_optional_features() {
+        assert!(raster_capability_reason(downlevel(), supported).is_none());
+    }
+
+    #[test]
+    fn fragment_storage_and_independent_blending_are_required() {
+        for flag in [wgpu::DownlevelFlags::FRAGMENT_STORAGE, wgpu::DownlevelFlags::INDEPENDENT_BLEND] {
+            assert!(raster_capability_reason(downlevel() - flag, supported).is_some());
+        }
+    }
+
+    #[test]
+    fn color_and_integer_depth_require_their_actual_texture_usages() {
+        for target in [COLOR, DEPTH_OUT] {
+            for usage in [wgpu::TextureUsages::RENDER_ATTACHMENT, wgpu::TextureUsages::TEXTURE_BINDING, wgpu::TextureUsages::COPY_SRC] {
+                assert!(
+                    raster_capability_reason(downlevel(), |format| {
+                        let mut features = supported(format);
+                        if format == target {
+                            features.allowed_usages.remove(usage);
+                        }
+                        features
+                    })
+                    .is_some()
+                );
+            }
+        }
+    }
+
+    #[test]
+    fn color_blending_and_depth_attachment_are_required() {
+        for target in [COLOR, DEPTH] {
+            assert!(
+                raster_capability_reason(downlevel(), |format| {
+                    let mut features = supported(format);
+                    if format == target {
+                        if target == COLOR {
+                            features.flags.remove(wgpu::TextureFormatFeatureFlags::BLENDABLE);
+                        } else {
+                            features.allowed_usages.remove(wgpu::TextureUsages::RENDER_ATTACHMENT);
+                        }
+                    }
+                    features
+                })
+                .is_some()
+            );
+        }
+    }
+
+    #[test]
+    fn byte_guards_use_each_buffer_stride_and_reject_overflow() {
+        assert_eq!(buffer_bytes(2, VERTEX_SIZE, 104), Some(104));
+        assert_eq!(buffer_bytes(2, VERTEX_SIZE, 103), None);
+        assert_eq!(buffer_bytes(26, 4, 104), Some(104));
+        assert_eq!(buffer_bytes(27, 4, 104), None);
+        assert_eq!(buffer_bytes(0, 4, 104), None);
+        assert_eq!(buffer_bytes(u64::MAX, 4, u64::MAX), None);
+        assert_eq!(buffer_bytes(1, 16, 15), None);
+    }
+
+    #[test]
+    fn packed_arrays_preflight_exact_strides_and_empty_padding() {
+        assert_eq!(packed_bytes([0; 6], 160, 16, 160), Some([160, 16, 16, 16, 16, 16, 16]));
+        assert_eq!(packed_bytes([2; 6], 160, 192, 192), Some([160, 192, 32, 32, 160, 160, 16]));
+        assert!(packed_bytes([0; 6], 159, 16, 160).is_none());
+        assert!(packed_bytes([0; 6], 160, 15, 160).is_none());
+        assert!(packed_bytes([2; 6], 160, 192, 191).is_none());
+        for index in 0..6 {
+            let mut counts = [0; 6];
+            counts[index] = u64::MAX;
+            assert!(packed_bytes(counts, u64::MAX, u64::MAX, u64::MAX).is_none());
+        }
+    }
+
+    #[test]
+    fn predicted_sizes_match_the_real_scene_serializer() {
+        use effectcraft_render::three_d::adv::{Light, Material, ShadowMap, TexInfo};
+        let scene = Scene {
+            materials: vec![Material::default()],
+            textures: vec![TexInfo { offset: 0, width: 1, height: 1, wrap_u: 0, wrap_v: 0 }],
+            texels: vec![[0.0; 4]],
+            lights: vec![Light {
+                kind: 0,
+                pos: [0.0; 3],
+                dir: [0.0; 3],
+                color: [0.0; 3],
+                cos_inner: 0.0,
+                cos_outer: 0.0,
+                falloff: 0,
+                radius: 0.0,
+                falloff_distance: 0.0,
+                shadow: -1,
+                shadow_darkness: 0.0,
+                shadow_diffusion: 0.0,
+            }],
+            shadows: vec![ShadowMap { view: [[0.0; 4]; 3], focal: 0.0, ortho: false, size: 1, offset: 0, bias: 0.0, radius: 0.0 }],
+            shadow_texels: vec![0.0],
+            ..Scene::default()
+        };
+        for (scene, counts) in [(Scene::default(), [0; 6]), (scene, [1; 6])] {
+            let actual = pack(&scene).map(|bytes| bytes.len() as u64);
+            assert_eq!(Some(actual), packed_bytes(counts, 160, 160, 160));
+        }
+    }
+}
+
 fn storage(binding: u32) -> wgpu::BindGroupLayoutEntry {
     wgpu::BindGroupLayoutEntry {
         binding,
@@ -71,70 +220,78 @@ fn storage(binding: u32) -> wgpu::BindGroupLayoutEntry {
 
 impl Pipes {
     fn new(device: &wgpu::Device, raster: bool) -> Pipes {
-        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("effectcraft advanced 3d"),
-            source: wgpu::ShaderSource::Wgsl(include_str!("shaders/advanced3d.wgsl").into()),
+        let raster = raster.then(|| {
+            let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: Some("effectcraft advanced 3d"),
+                source: wgpu::ShaderSource::Wgsl(include_str!("shaders/advanced3d.wgsl").into()),
+            });
+            let mut entries = vec![wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
+                ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
+                count: None,
+            }];
+            entries.extend((1..=6).map(storage));
+            let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("advanced 3d"), entries: &entries });
+            let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("advanced 3d"),
+                bind_group_layouts: &[Some(&bgl)],
+                immediate_size: 0,
+            });
+            let attrs = [
+                wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: 0, shader_location: 0 },
+                wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: 12, shader_location: 1 },
+                wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x2, offset: 24, shader_location: 2 },
+                wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x4, offset: 32, shader_location: 3 },
+                wgpu::VertexAttribute { format: wgpu::VertexFormat::Uint32, offset: 48, shader_location: 4 },
+            ];
+            let make = |transparent: bool| {
+                let blend = transparent.then_some(wgpu::BlendState::PREMULTIPLIED_ALPHA_BLENDING);
+                device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+                    label: Some(if transparent { "advanced 3d transparent" } else { "advanced 3d opaque" }),
+                    layout: Some(&layout),
+                    vertex: wgpu::VertexState {
+                        module: &module,
+                        entry_point: Some("vs"),
+                        compilation_options: Default::default(),
+                        buffers: &[Some(wgpu::VertexBufferLayout { array_stride: VERTEX_SIZE, step_mode: wgpu::VertexStepMode::Vertex, attributes: &attrs })],
+                    },
+                    primitive: wgpu::PrimitiveState {
+                        topology: wgpu::PrimitiveTopology::TriangleList,
+                        front_face: wgpu::FrontFace::Ccw,
+                        cull_mode: None,
+                        ..Default::default()
+                    },
+                    depth_stencil: Some(wgpu::DepthStencilState {
+                        format: DEPTH,
+                        depth_write_enabled: Some(!transparent),
+                        depth_compare: Some(wgpu::CompareFunction::Greater),
+                        stencil: Default::default(),
+                        bias: Default::default(),
+                    }),
+                    multisample: Default::default(),
+                    fragment: Some(wgpu::FragmentState {
+                        module: &module,
+                        entry_point: Some("fs"),
+                        compilation_options: Default::default(),
+                        targets: &[
+                            Some(wgpu::ColorTargetState { format: COLOR, blend, write_mask: wgpu::ColorWrites::ALL }),
+                            Some(wgpu::ColorTargetState {
+                                format: DEPTH_OUT,
+                                blend: None,
+                                write_mask: if transparent { wgpu::ColorWrites::empty() } else { wgpu::ColorWrites::ALL },
+                            }),
+                        ],
+                    }),
+                    multiview_mask: None,
+                    cache: None,
+                })
+            };
+            let opaque = make(false);
+            let transparent = make(true);
+            (bgl, opaque, transparent)
         });
-        let mut entries = vec![wgpu::BindGroupLayoutEntry {
-            binding: 0,
-            visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
-            ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
-            count: None,
-        }];
-        entries.extend((1..=6).map(storage));
-        let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { label: Some("advanced 3d"), entries: &entries });
-        let layout =
-            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("advanced 3d"), bind_group_layouts: &[Some(&bgl)], immediate_size: 0 });
-        let attrs = [
-            wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: 0, shader_location: 0 },
-            wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x3, offset: 12, shader_location: 1 },
-            wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x2, offset: 24, shader_location: 2 },
-            wgpu::VertexAttribute { format: wgpu::VertexFormat::Float32x4, offset: 32, shader_location: 3 },
-            wgpu::VertexAttribute { format: wgpu::VertexFormat::Uint32, offset: 48, shader_location: 4 },
-        ];
-        let make = |transparent: bool| {
-            let blend = transparent.then_some(wgpu::BlendState::PREMULTIPLIED_ALPHA_BLENDING);
-            device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
-                label: Some(if transparent { "advanced 3d transparent" } else { "advanced 3d opaque" }),
-                layout: Some(&layout),
-                vertex: wgpu::VertexState {
-                    module: &module,
-                    entry_point: Some("vs"),
-                    compilation_options: Default::default(),
-                    buffers: &[Some(wgpu::VertexBufferLayout { array_stride: VERTEX_SIZE, step_mode: wgpu::VertexStepMode::Vertex, attributes: &attrs })],
-                },
-                primitive: wgpu::PrimitiveState {
-                    topology: wgpu::PrimitiveTopology::TriangleList,
-                    front_face: wgpu::FrontFace::Ccw,
-                    cull_mode: None,
-                    ..Default::default()
-                },
-                depth_stencil: Some(wgpu::DepthStencilState {
-                    format: DEPTH,
-                    depth_write_enabled: Some(!transparent),
-                    depth_compare: Some(wgpu::CompareFunction::Greater),
-                    stencil: Default::default(),
-                    bias: Default::default(),
-                }),
-                multisample: Default::default(),
-                fragment: Some(wgpu::FragmentState {
-                    module: &module,
-                    entry_point: Some("fs"),
-                    compilation_options: Default::default(),
-                    targets: &[
-                        Some(wgpu::ColorTargetState { format: COLOR, blend, write_mask: wgpu::ColorWrites::ALL }),
-                        Some(wgpu::ColorTargetState {
-                            format: DEPTH_OUT,
-                            blend: None,
-                            write_mask: if transparent { wgpu::ColorWrites::empty() } else { wgpu::ColorWrites::ALL },
-                        }),
-                    ],
-                }),
-                multiview_mask: None,
-                cache: None,
-            })
-        };
-        Pipes { raster: raster.then(|| (make(false), make(true))), bgl, post: Post::new(device) }
+        Pipes { raster, post: Post::new(device) }
     }
 }
 
@@ -251,16 +408,28 @@ fn half_to_f32(h: u16) -> f32 {
 /// (render targets, size or buffer limits).
 fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)> {
     let g = e.g;
-    if !g.fits(s.width, s.height) || s.indices.is_empty() {
+    if !g.adv3d_raster || !g.fits(s.width, s.height) || s.indices.is_empty() {
         return None;
     }
     let limits = g.device.limits();
+    buffer_bytes(u64::try_from(s.vertices.len()).ok()?, VERTEX_SIZE, limits.max_buffer_size)?;
+    buffer_bytes(u64::try_from(s.indices.len()).ok()?, 4, limits.max_buffer_size)?;
+    let n = u32::try_from(s.indices.len()).ok()?;
+    if s.opaque_count > n {
+        return None;
+    }
+    let counts = [s.materials.len(), s.textures.len(), s.texels.len(), s.lights.len(), s.shadows.len(), s.shadow_texels.len()];
+    let mut checked_counts = [0; 6];
+    for (out, count) in checked_counts.iter_mut().zip(counts) {
+        *out = u64::try_from(count).ok()?;
+    }
+    let expected_sizes = packed_bytes(checked_counts, limits.max_uniform_buffer_binding_size, limits.max_storage_buffer_binding_size, limits.max_buffer_size)?;
     let bufs = pack(s);
-    if bufs.iter().any(|b| b.len() as u64 > limits.max_storage_buffer_binding_size) {
+    if bufs.iter().zip(expected_sizes).any(|(bytes, expected)| bytes.len() as u64 != expected) {
         return None;
     }
     let p = pipes(g);
-    let Some((opaque, transparent)) = &p.raster else { return None };
+    let Some((bgl, opaque, transparent)) = &p.raster else { return None };
     let dev = &g.device;
     let vb: Vec<u8> = s
         .vertices
@@ -286,7 +455,7 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
     for (i, b) in sbs.iter().enumerate() {
         entries.push(wgpu::BindGroupEntry { binding: i as u32 + 1, resource: b.as_entire_binding() });
     }
-    let bg = dev.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("adv3d"), layout: &p.bgl, entries: &entries });
+    let bg = dev.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("adv3d"), layout: bgl, entries: &entries });
     let tex = |format: wgpu::TextureFormat, usage: wgpu::TextureUsages| {
         dev.create_texture(&wgpu::TextureDescriptor {
             label: Some("adv3d target"),
@@ -334,7 +503,6 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
         pass.set_bind_group(0, &bg, &[]);
         pass.set_vertex_buffer(0, vbuf.slice(..));
         pass.set_index_buffer(ibuf.slice(..), wgpu::IndexFormat::Uint32);
-        let n = s.indices.len() as u32;
         if s.opaque_count > 0 {
             pass.set_pipeline(opaque);
             pass.draw_indexed(0..s.opaque_count, 0, 0..1);
@@ -604,6 +772,11 @@ fn read_buffer(e: &mut Enc, buf: &wgpu::Buffer, size: u64) -> Option<Vec<u8>> {
     if !e.g.can_wait() {
         return None;
     }
+    buffer_bytes(size, 1, e.g.device.limits().max_buffer_size)?;
+    let row = usize::try_from(size).ok()?;
+    if size > buf.size() || !size.is_multiple_of(wgpu::COPY_BUFFER_ALIGNMENT) {
+        return None;
+    }
     let rb = e.g.device.create_buffer(&wgpu::BufferDescriptor {
         label: Some("adv3d readback"),
         size,
@@ -612,14 +785,7 @@ fn read_buffer(e: &mut Enc, buf: &wgpu::Buffer, size: u64) -> Option<Vec<u8>> {
     });
     e.encoder().copy_buffer_to_buffer(buf, 0, &rb, 0, size);
     e.submit();
-    let (tx, rx) = std::sync::mpsc::channel();
-    rb.map_async(wgpu::MapMode::Read, .., move |r| {
-        let _ = tx.send(r);
-    });
-    e.g.device.poll(wgpu::PollType::wait_indefinitely()).ok()?;
-    rx.recv().ok()?.ok()?;
-    let out = rb.get_mapped_range(..).ok()?.to_vec();
-    rb.unmap();
+    let out = e.g.read_buffer(&rb, row, row, 1).map_err(|e| log::error!("gpu 3d readback: {e}")).ok()?;
     Some(out)
 }
 
@@ -645,12 +811,12 @@ pub(crate) fn resolve_run(e: &mut Enc, prep: &Prepared) -> Option<Option<Resolve
         return None;
     }
     let px = w as u64 * h as u64;
-    let max = g.device.limits().max_storage_buffer_binding_size;
-    if px * 16 > max {
-        return None;
-    }
-    let acc = storage_buf(g, px * 16, "adv3d acc");
-    let depth = storage_buf(g, px * 4, "adv3d depth");
+    let limits = g.device.limits();
+    let max = limits.max_storage_buffer_binding_size.min(limits.max_buffer_size);
+    let acc_bytes = buffer_bytes(px, 16, max)?;
+    let depth_bytes = buffer_bytes(px, 4, max)?;
+    let acc = storage_buf(g, acc_bytes, "adv3d acc");
+    let depth = storage_buf(g, depth_bytes, "adv3d depth");
     let n = prep.samples.max(1);
     let weight = if n <= 1 { 1.0 } else { 1.0 / n as f32 };
     let mut first = true;
@@ -659,7 +825,7 @@ pub(crate) fn resolve_run(e: &mut Enc, prep: &Prepared) -> Option<Option<Resolve
         if s.indices.is_empty() {
             continue;
         }
-        if s.width != w * s.ssaa.max(1) || s.height != h * s.ssaa.max(1) {
+        if Some(s.width) != w.checked_mul(s.ssaa.max(1)) || Some(s.height) != h.checked_mul(s.ssaa.max(1)) {
             return None;
         }
         let (color, zout) = raster_into(e, &s)?;
@@ -684,9 +850,10 @@ pub(crate) fn resolve_run(e: &mut Enc, prep: &Prepared) -> Option<Option<Resolve
 fn depth_of_field(e: &mut Enc, acc: &wgpu::Buffer, depth: &wgpu::Buffer, (w, h): (u32, u32), dof: &Dof, scale: f64) -> Option<()> {
     let g = e.g;
     let px = w as u64 * h as u64;
-    if (w as u64 + 1) * h as u64 * 16 > g.device.limits().max_storage_buffer_binding_size {
-        return None;
-    }
+    let limits = g.device.limits();
+    let max = limits.max_storage_buffer_binding_size.min(limits.max_buffer_size);
+    let prefix_bytes = buffer_bytes((u64::from(w) + 1).checked_mul(u64::from(h))?, 16, max)?;
+    buffer_bytes(1, 16, max)?;
     let radius = storage_buf(g, px * 4, "adv3d radius");
     let mut init = Vec::with_capacity(16);
     for v in [f32::INFINITY.to_bits(), 0, 0, 0] {
@@ -724,12 +891,12 @@ fn depth_of_field(e: &mut Enc, acc: &wgpu::Buffer, depth: &wgpu::Buffer, (w, h):
         }
     }
     let table: Vec<u8> = heads.iter().chain(&spans).flatten().flat_map(|x| x.to_le_bytes()).collect();
-    if table.len() as u64 > g.device.limits().max_storage_buffer_binding_size {
+    if table.len() as u64 > max {
         return None;
     }
     let table = e.bytes(table);
     let src = storage_buf(g, px * 16, "adv3d dof source");
-    let prefix = storage_buf(g, (w as u64 + 1) * h as u64 * 16, "adv3d prefix");
+    let prefix = storage_buf(g, prefix_bytes, "adv3d prefix");
     let hl = dof.highlight;
     let b = Binds { bufs: [Some(acc), None, Some(&src), None, Some(&radius), None, None], ..Default::default() };
     dispatch(e, "boost", [w, h, 0, 0], [hl.gain as f32, hl.threshold as f32, hl.saturation as f32, 0.0], b, pixel_groups(w, h));
@@ -788,6 +955,8 @@ fn render_prepared_deferred(g: &GpuContext, d: &std::sync::Arc<crate::deferred::
         return None;
     }
     let (w, h) = prep.out;
+    let size = buffer_bytes(u64::from(w).checked_mul(u64::from(h))?, 4, g.device.limits().max_buffer_size)?;
+    let row = usize::try_from(size).ok()?;
     let mut k = crate::deferred::Key::new(5);
     k.u64(w as u64);
     k.u64(h as u64);
@@ -836,7 +1005,6 @@ fn render_prepared_deferred(g: &GpuContext, d: &std::sync::Arc<crate::deferred::
     let Some(r) = resolve_run(&mut e, prep)? else { return Some(None) };
     let img = finish(&mut e, &r, !prep.linear, None);
     let _ = e.read_keyed(&img, ki, [0.0; 2], 1.0);
-    let size = w as u64 * h as u64 * 4;
     let rb = g.device.create_buffer(&wgpu::BufferDescriptor {
         label: Some("adv3d depth readback"),
         size,
@@ -846,14 +1014,9 @@ fn render_prepared_deferred(g: &GpuContext, d: &std::sync::Arc<crate::deferred::
     e.encoder().copy_buffer_to_buffer(&r.depth, 0, &rb, 0, size);
     e.submit();
     let done = d.start(kz, w, h, [0.0; 2], 1.0);
-    let b = rb.clone();
-    rb.map_async(wgpu::MapMode::Read, .., move |res| {
-        let out = res.ok().and_then(|_| b.get_mapped_range(..).ok().map(|v| v.to_vec()));
-        b.unmap();
-        done(out);
-    });
+    g.map_readback(&rb, row, row, 1, move |r| done(r.map_err(|e| log::error!("gpu depth readback: {e}")).ok()));
     #[cfg(not(target_arch = "wasm32"))]
-    let _ = g.device.poll(wgpu::PollType::Poll);
+    let _ = g.poll_readbacks(wgpu::PollType::Poll);
     d.miss();
     Some(None)
 }
diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs
index 2b2a978f..119648cf 100644
--- a/crates/gpu/src/context.rs
+++ b/crates/gpu/src/context.rs
@@ -267,6 +267,7 @@ pub struct GpuContext {
     pub(crate) particle_states: crate::particles::StatesCell,
     /// Deferred readbacks (a browser worker's WebGPU device), see [`crate::deferred`].
     pub(crate) deferred: Option<Arc<crate::deferred::Deferred>>,
+    readbacks: Arc<crate::readback::Readbacks>,
 }
 
 fn tex_entry(binding: u32) -> wgpu::BindGroupLayoutEntry {
@@ -291,24 +292,44 @@ fn storage_tex_entry(binding: u32, format: wgpu::TextureFormat) -> wgpu::BindGro
     }
 }
 
-/// Log the device's uncaptured errors instead of wgpu's default panic.
-fn log_uncaptured_errors(device: &wgpu::Device) {
-    device.on_uncaptured_error(Arc::new(|e| {
-        // Under test a validation error (a WGSL typo invalidates every kernel) fails loudly.
-        #[cfg(test)]
-        #[allow(clippy::panic)] // test builds only
-        {
-            panic!("wgpu: {e}")
-        }
-        #[cfg(not(test))]
-        log::error!("wgpu: {e}")
-    }));
+/// Contain synchronous native initialization errors without replacing the device owner's
+/// handlers. Scopes are thread-local with wgpu's enabled `std` feature. The pinned native
+/// wgpu-core backend returns an immediately-ready future from pop_error_scope, so these
+/// drains do not poll or wait for unrelated submitted GPU work.
+#[cfg(not(target_arch = "wasm32"))]
+pub(crate) fn init_resource<T>(device: &wgpu::Device, label: &str, create: impl FnOnce() -> T) -> Result<T, String> {
+    let oom = device.push_error_scope(wgpu::ErrorFilter::OutOfMemory);
+    let internal = device.push_error_scope(wgpu::ErrorFilter::Internal);
+    let validation = device.push_error_scope(wgpu::ErrorFilter::Validation);
+    // Some native drivers/wgpu paths can still unwind instead of reporting a scoped error.
+    // Keep the last-resort guard inside the scopes, then drain every scope even on failure.
+    let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(create));
+    let validation_error = pollster::block_on(validation.pop());
+    let internal_error = pollster::block_on(internal.pop());
+    let oom_error = pollster::block_on(oom.pop());
+    if let Some(error) = validation_error.or(internal_error).or(oom_error) {
+        return Err(format!("GPU initialization ({label}): {error}"));
+    }
+    result.map_err(|payload| {
+        let reason =
+            payload.downcast_ref::<String>().map(String::as_str).or_else(|| payload.downcast_ref::<&str>().copied()).unwrap_or("native backend panicked");
+        format!("GPU initialization ({label}): {reason}")
+    })
+}
+
+/// Browser scope completion is asynchronous. Preserve the existing browser constructor;
+/// never block its event loop or introduce a guard across an asynchronous boundary here.
+#[cfg(target_arch = "wasm32")]
+fn init_resource<T>(_device: &wgpu::Device, _label: &str, create: impl FnOnce() -> T) -> Result<T, String> {
+    Ok(create())
 }
 
 impl GpuContext {
     /// Build on an existing device (the desktop app shares egui-wgpu's). `Err` when the
     /// adapter cannot run the compositor (no compute shaders, e.g. WebGL2; no float storage
     /// textures) or the device was created with limits below what its pipelines bind.
+    /// Preserves the host's device handlers. The host must forward failures using
+    /// [`Self::device_failure_notifier`] for immediate readback cancellation.
     pub fn new(adapter: &wgpu::Adapter, device: wgpu::Device, queue: wgpu::Queue) -> Result<GpuContext, String> {
         let info = adapter.get_info();
         if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
@@ -334,6 +355,7 @@ impl GpuContext {
                 return Err(format!("{} ({:?}): {f:?} storage textures unsupported", info.name, info.backend));
             }
         }
+        let readbacks = crate::readback::Readbacks::new()?;
         let name = format!("{} ({:?})", info.name, info.backend);
         let adv3d_unsupported = crate::adv3d::raster_unsupported(adapter);
         if let Some(why) = &adv3d_unsupported {
@@ -366,52 +388,61 @@ impl GpuContext {
             include_str!("shaders/sky.wgsl"),
         ]
         .concat();
-        let module =
-            device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) });
-        let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("effectcraft kernels"),
-            entries: &[
-                wgpu::BindGroupLayoutEntry {
-                    binding: 0,
-                    visibility: wgpu::ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
-                    count: None,
-                },
-                tex_entry(1),
-                tex_entry(2),
-                storage_tex_entry(3, FORMAT),
-                wgpu::BindGroupLayoutEntry {
-                    binding: 4,
-                    visibility: wgpu::ShaderStages::COMPUTE,
-                    ty: wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Storage { read_only: true },
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
+        let module = init_resource(&device, "kernels shader module", || {
+            device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) })
+        })?;
+        let bgl = init_resource(&device, "kernels bind-group layout", || {
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("effectcraft kernels"),
+                entries: &[
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 0,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
+                        count: None,
                     },
-                    count: None,
-                },
-            ],
-        });
-        let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("effectcraft kernels"),
-            bind_group_layouts: &[Some(&bgl)],
-            immediate_size: 0,
-        });
+                    tex_entry(1),
+                    tex_entry(2),
+                    storage_tex_entry(3, FORMAT),
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 4,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: true },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    },
+                ],
+            })
+        })?;
+        let layout = init_resource(&device, "kernels pipeline layout", || {
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("effectcraft kernels"),
+                bind_group_layouts: &[Some(&bgl)],
+                immediate_size: 0,
+            })
+        })?;
         let storage = |binding: u32| wgpu::BindGroupLayoutEntry {
             binding,
             visibility: wgpu::ShaderStages::COMPUTE,
             ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: true }, has_dynamic_offset: false, min_binding_size: None },
             count: None,
         };
-        let bgl_ext = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("effectcraft kernels (ext)"),
-            entries: &[storage(0), storage(1), storage(2), storage(3)],
-        });
-        let layout_ext = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("effectcraft kernels (ext)"),
-            bind_group_layouts: &[Some(&bgl), Some(&bgl_ext)],
-            immediate_size: 0,
-        });
+        let bgl_ext = init_resource(&device, "extended kernels bind-group layout", || {
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("effectcraft kernels (ext)"),
+                entries: &[storage(0), storage(1), storage(2), storage(3)],
+            })
+        })?;
+        let layout_ext = init_resource(&device, "extended kernels pipeline layout", || {
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("effectcraft kernels (ext)"),
+                bind_group_layouts: &[Some(&bgl), Some(&bgl_ext)],
+                immediate_size: 0,
+            })
+        })?;
         let pipelines = ENTRIES
             .iter()
             .chain(crate::fx_color::KERNELS)
@@ -438,57 +469,65 @@ impl GpuContext {
             .map(|e| (e, &layout))
             .chain(EXT_ENTRIES.iter().map(|e| (e, &layout_ext)))
             .map(|(e, layout)| {
-                let p = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-                    label: Some(e),
-                    layout: Some(layout),
-                    module: &module,
-                    entry_point: Some(e),
-                    compilation_options: Default::default(),
-                    cache: None,
-                });
-                (*e, p)
+                let p = init_resource(&device, e, || {
+                    device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                        label: Some(e),
+                        layout: Some(layout),
+                        module: &module,
+                        entry_point: Some(e),
+                        compilation_options: Default::default(),
+                        cache: None,
+                    })
+                })?;
+                Ok((*e, p))
             })
-            .collect();
-        let dmodule = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("effectcraft display"),
-            source: wgpu::ShaderSource::Wgsl(include_str!("shaders/display.wgsl").into()),
-        });
-        let display_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("effectcraft display"),
-            entries: &[tex_entry(1), storage_tex_entry(3, wgpu::TextureFormat::Rgba8Unorm)],
-        });
-        let dlayout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
-            label: Some("effectcraft display"),
-            bind_group_layouts: &[Some(&display_bgl)],
-            immediate_size: 0,
-        });
-        let display = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
-            label: Some("display"),
-            layout: Some(&dlayout),
-            module: &dmodule,
-            entry_point: Some("display"),
-            compilation_options: Default::default(),
-            cache: None,
-        });
-        let dummy = device.create_texture(&wgpu::TextureDescriptor {
-            label: Some("dummy"),
-            size: wgpu::Extent3d { width: 1, height: 1, depth_or_array_layers: 1 },
-            mip_level_count: 1,
-            sample_count: 1,
-            dimension: wgpu::TextureDimension::D2,
-            format: FORMAT,
-            usage: wgpu::TextureUsages::TEXTURE_BINDING,
-            view_formats: &[],
-        });
-        let dummy_buf =
-            device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("dummy"), contents: &[0u8; 16], usage: wgpu::BufferUsages::STORAGE });
+            .collect::<Result<HashMap<_, _>, String>>()?;
+        let dmodule = init_resource(&device, "display shader module", || {
+            device.create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: Some("effectcraft display"),
+                source: wgpu::ShaderSource::Wgsl(include_str!("shaders/display.wgsl").into()),
+            })
+        })?;
+        let display_bgl = init_resource(&device, "display bind-group layout", || {
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some("effectcraft display"),
+                entries: &[tex_entry(1), storage_tex_entry(3, wgpu::TextureFormat::Rgba8Unorm)],
+            })
+        })?;
+        let dlayout = init_resource(&device, "display pipeline layout", || {
+            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: Some("effectcraft display"),
+                bind_group_layouts: &[Some(&display_bgl)],
+                immediate_size: 0,
+            })
+        })?;
+        let display = init_resource(&device, "display compute pipeline", || {
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some("display"),
+                layout: Some(&dlayout),
+                module: &dmodule,
+                entry_point: Some("display"),
+                compilation_options: Default::default(),
+                cache: None,
+            })
+        })?;
+        let dummy = init_resource(&device, "dummy texture", || {
+            device.create_texture(&wgpu::TextureDescriptor {
+                label: Some("dummy"),
+                size: wgpu::Extent3d { width: 1, height: 1, depth_or_array_layers: 1 },
+                mip_level_count: 1,
+                sample_count: 1,
+                dimension: wgpu::TextureDimension::D2,
+                format: FORMAT,
+                usage: wgpu::TextureUsages::TEXTURE_BINDING,
+                view_formats: &[],
+            })
+        })?;
+        let dummy_buf = init_resource(&device, "dummy storage buffer", || {
+            device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("dummy"), contents: &[0u8; 16], usage: wgpu::BufferUsages::STORAGE })
+        })?;
+        let dummy_tex = init_resource(&device, "dummy texture view", || dummy.create_view(&Default::default()))?;
         let max_dim = device.limits().max_texture_dimension_2d;
-        // From here on errors are logged, not fatal: wgpu's default handler panics, and egui-wgpu
-        // installs none on the device the desktop app shares. Running out of video memory on a
-        // frame thread must not crash the app (#106); frames catch it with
-        // [`crate::Gpu::within_memory`]. (Errors while the pipelines are built above still panic
-        // on that device, which the app catches to fall back to the CPU.)
-        log_uncaptured_errors(&device);
         Ok(GpuContext {
             device,
             queue,
@@ -499,7 +538,7 @@ impl GpuContext {
             pipelines,
             display_bgl,
             display,
-            dummy_tex: dummy.create_view(&Default::default()),
+            dummy_tex,
             dummy,
             dummy_buf,
             uploads: Mutex::new(Uploads::default()),
@@ -512,11 +551,109 @@ impl GpuContext {
             particles: Default::default(),
             particle_states: Default::default(),
             deferred: None,
+            readbacks,
         })
     }
 
+    /// Whether this context has observed a readback timeout or device failure.
+    /// A successful check is not a guarantee against faults after it returns.
+    /// Shared-device hosts must forward their device failures through the notifier.
+    pub fn check_health(&self) -> Result<(), String> {
+        self.readbacks.check()
+    }
+
+    /// A notification hook for the owner of a shared device's error/loss handlers.
+    /// Invoke it with diagnostic context when the host observes a device failure.
+    /// It retires this context and settles pending readbacks without replacing any
+    /// host handler or keeping this context/device alive. It is harmless after drop.
+    ///
+    /// Shared-device construction does not automatically subscribe to device faults:
+    /// without host forwarding, readbacks rely on poll errors and their deadlines.
+    pub fn device_failure_notifier(&self) -> impl Fn(String) + Send + Sync + 'static {
+        crate::readback::failure_notifier(&self.readbacks)
+    }
+
+    /// All device polling routes through this failure bridge.
+    pub(crate) fn poll_readbacks(&self, kind: wgpu::PollType) -> Result<wgpu::PollStatus, String> {
+        self.readbacks.check()?;
+        self.device.poll(kind).map_err(|e| {
+            let error = format!("GPU readback poll: {e}");
+            self.readbacks.retire(error.clone());
+            log::error!("{error}");
+            error
+        })
+    }
+
+    /// Start a bounded buffer mapping. Retired contexts reject new requests immediately.
+    pub(crate) fn map_readback(
+        &self,
+        buffer: &wgpu::Buffer,
+        row: usize,
+        stride: usize,
+        height: usize,
+        done: impl FnOnce(Result<Vec<u8>, String>) + wgpu::WasmNotSend + 'static,
+    ) {
+        let ticket = match self.readbacks.start(done) {
+            Ok(t) => t,
+            Err(e) => {
+                log::error!("gpu: {e}");
+                return;
+            }
+        };
+        if !ticket.active() {
+            return;
+        }
+        let b = buffer.clone();
+        buffer.map_async(wgpu::MapMode::Read, .., move |r| {
+            if !ticket.active() {
+                b.unmap();
+                return;
+            }
+            let result = r.map_err(|e| format!("GPU buffer mapping: {e}")).and_then(|_| {
+                let view = b.get_mapped_range(..).map_err(|e| format!("GPU mapped range: {e}"))?;
+                crate::readback::packed(&view, row, stride, height)
+            });
+            b.unmap();
+            ticket.complete(result);
+        });
+    }
+
+    /// A bounded native readback: polling and receiving share one deadline.
+    #[cfg(not(target_arch = "wasm32"))]
+    pub(crate) fn read_buffer(&self, buffer: &wgpu::Buffer, row: usize, stride: usize, height: usize) -> Result<Vec<u8>, String> {
+        let deadline = std::time::Instant::now() + crate::readback::TIMEOUT;
+        let (tx, rx) = std::sync::mpsc::channel();
+        self.map_readback(buffer, row, stride, height, move |r| {
+            let _ = tx.send(r);
+        });
+        loop {
+            self.readbacks.check()?;
+            self.poll_readbacks(wgpu::PollType::Poll)?;
+            let remaining = deadline.saturating_duration_since(std::time::Instant::now());
+            if remaining.is_zero() {
+                let error = "GPU readback receive timed out; context retired".to_string();
+                self.readbacks.retire(error.clone());
+                return Err(error);
+            }
+            match rx.recv_timeout(remaining.min(std::time::Duration::from_millis(10))) {
+                Ok(result) => return result,
+                Err(std::sync::mpsc::RecvTimeoutError::Timeout) => {}
+                Err(e) => return Err(format!("GPU readback completion channel: {e}")),
+            }
+        }
+    }
+
+    #[cfg(target_arch = "wasm32")]
+    pub(crate) fn read_buffer(&self, _buffer: &wgpu::Buffer, _row: usize, _stride: usize, _height: usize) -> Result<Vec<u8>, String> {
+        Err("synchronous GPU readback is unavailable in the browser".into())
+    }
+
     /// Switch to deferred readbacks (see [`crate::deferred`]): readbacks are never waited for.
     pub fn set_deferred(&mut self, on: bool) {
+        self.readbacks.cancel_pending("GPU readback mode changed".into());
+        if let Some(d) = self.deferred.take() {
+            d.cancel();
+        }
         self.deferred = on.then(Default::default);
     }
 
@@ -554,8 +691,10 @@ impl GpuContext {
             .request_device(&wgpu::DeviceDescriptor { label: Some("effectcraft gpu"), required_limits, ..Default::default() })
             .await
             .map_err(|e| format!("no device: {e}"))?;
-        log_uncaptured_errors(&device);
-        GpuContext::new(&adapter, device, queue)
+        let context = GpuContext::new(&adapter, device, queue)?;
+        // This path owns the device; install handlers only after construction succeeds.
+        crate::readback::install(&context.device, &context.readbacks);
+        Ok(context)
     }
 
     /// A device of its own on the best adapter (CLI, tests, benchmarks). `None` when no
@@ -568,13 +707,13 @@ impl GpuContext {
     /// Readbacks (GPU → CPU) are possible: waited for natively, or deferred (a browser
     /// worker, [`crate::deferred`]). Impossible on the browser's main thread.
     pub fn can_readback(&self) -> bool {
-        self.can_wait() || self.deferred.is_some()
+        self.readbacks.check().is_ok() && (self.can_wait() || self.deferred.is_some())
     }
 
     /// Readbacks can be waited for (natively, unless deferred). Steps that read back without a
     /// key ([`crate::deferred`]) need this.
     pub fn can_wait(&self) -> bool {
-        cfg!(not(target_arch = "wasm32")) && self.deferred.is_none()
+        self.readbacks.check().is_ok() && cfg!(not(target_arch = "wasm32")) && self.deferred.is_none()
     }
 
     pub(crate) fn fits(&self, w: u32, h: u32) -> bool {
@@ -876,7 +1015,13 @@ impl<'g> Enc<'g> {
 
     /// Read an image back to the CPU (submits and waits). `None` where readback is impossible.
     pub fn download(&mut self, img: &GpuImage) -> Option<Image> {
-        let mut out = Image::new(img.width, img.height);
+        let layout = crate::readback::Layout::new(img.width, img.height, 16, self.g.device.limits().max_buffer_size)
+            .map_err(|e| log::error!("gpu image readback: {e}"))
+            .ok()?;
+        let mut data = Vec::new();
+        data.try_reserve_exact(layout.len / 16).map_err(|e| log::error!("gpu image allocation: {e}")).ok()?;
+        data.resize(layout.len / 16, [0.0; 4]);
+        let mut out = Image { width: img.width, height: img.height, data };
         let dst: &mut [u8] = bytemuck::cast_slice_mut(&mut out.data);
         self.read_texture_into(&img.texture, img.width, img.height, 16, dst)?;
         Some(out)
@@ -884,7 +1029,11 @@ impl<'g> Enc<'g> {
 
     /// Tightly packed texel bytes of a texture (`bpp` bytes per pixel).
     pub fn read_texture(&mut self, texture: &wgpu::Texture, w: u32, h: u32, bpp: u32) -> Option<Vec<u8>> {
-        let mut out = vec![0u8; (w * bpp) as usize * h as usize];
+        let layout =
+            crate::readback::Layout::new(w, h, bpp, self.g.device.limits().max_buffer_size).map_err(|e| log::error!("gpu texture readback: {e}")).ok()?;
+        let mut out = Vec::new();
+        out.try_reserve_exact(layout.len).map_err(|e| log::error!("gpu readback allocation: {e}")).ok()?;
+        out.resize(layout.len, 0);
         self.read_texture_into(texture, w, h, bpp, &mut out)?;
         Some(out)
     }
@@ -895,9 +1044,15 @@ impl<'g> Enc<'g> {
             return None;
         }
         self.g.transfers.count(false, out.len());
-        let row = (w * bpp) as usize;
-        let padded = (w * bpp).div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
-        let size = padded as u64 * h as u64;
+        let layout =
+            crate::readback::Layout::new(w, h, bpp, self.g.device.limits().max_buffer_size).map_err(|e| log::error!("gpu texture readback: {e}")).ok()?;
+        if out.len() != layout.len {
+            log::error!("gpu readback destination size mismatch");
+            return None;
+        }
+        let row = layout.row as usize;
+        let padded = layout.stride;
+        let size = layout.size;
         let buf = self.g.staging.lock().ok().and_then(|mut s| s.get_mut(&size).and_then(Vec::pop)).unwrap_or_else(|| {
             self.g.device.create_buffer(&wgpu::BufferDescriptor {
                 label: Some("readback"),
@@ -912,27 +1067,12 @@ impl<'g> Enc<'g> {
             wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
         );
         self.submit();
-        let (tx, rx) = std::sync::mpsc::channel();
-        buf.map_async(wgpu::MapMode::Read, .., move |r| {
-            let _ = tx.send(r);
-        });
-        if let Err(e) = self.g.device.poll(wgpu::PollType::wait_indefinitely()) {
-            log::error!("gpu: poll: {e}");
+        let bytes = self.g.read_buffer(&buf, row, padded as usize, h as usize).map_err(|e| log::error!("gpu texture readback: {e}")).ok()?;
+        if bytes.len() != out.len() {
+            log::error!("gpu texture readback size mismatch");
             return None;
         }
-        rx.recv().ok()?.ok()?;
-        {
-            let view = buf.get_mapped_range(..).ok()?;
-            if row == padded as usize {
-                out.copy_from_slice(&view[..out.len()]);
-            } else {
-                for (y, dst) in out.chunks_exact_mut(row).enumerate() {
-                    let s = y * padded as usize;
-                    dst.copy_from_slice(&view[s..s + row]);
-                }
-            }
-        }
-        buf.unmap();
+        out.copy_from_slice(&bytes);
         // Keep a few staging buffers per size (frames read back at the same size every time).
         if let Ok(mut s) = self.g.staging.lock() {
             let v = s.entry(size).or_default();
@@ -947,11 +1087,24 @@ impl<'g> Enc<'g> {
     /// gets the bytes once the copy is mapped (from the browser's event loop on the web; from
     /// a later device poll natively).
     pub fn read_texture_async(&mut self, texture: &wgpu::Texture, w: u32, h: u32, bpp: u32, done: impl FnOnce(Option<Vec<u8>>) + wgpu::WasmNotSend + 'static) {
-        let row = w * bpp;
-        let padded = row.div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
+        if let Err(e) = self.g.readbacks.check() {
+            log::error!("gpu readback: {e}");
+            crate::readback::reject(done, e);
+            return;
+        }
+        let layout = match crate::readback::Layout::new(w, h, bpp, self.g.device.limits().max_buffer_size) {
+            Ok(layout) => layout,
+            Err(e) => {
+                log::error!("gpu texture readback: {e}");
+                crate::readback::reject(done, e);
+                return;
+            }
+        };
+        let row = layout.row;
+        let padded = layout.stride;
         let buf = self.g.device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("readback-async"),
-            size: padded as u64 * h as u64,
+            size: layout.size,
             usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
             mapped_at_creation: false,
         });
@@ -961,21 +1114,11 @@ impl<'g> Enc<'g> {
             wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
         );
         self.submit();
-        let b = buf.clone();
-        buf.map_async(wgpu::MapMode::Read, .., move |r| {
-            let out = r.ok().and_then(|_| {
-                let view = b.get_mapped_range(..).ok()?;
-                let mut out = Vec::with_capacity((row * h) as usize);
-                for y in 0..h as usize {
-                    out.extend_from_slice(&view[y * padded as usize..y * padded as usize + row as usize]);
-                }
-                Some(out)
-            });
-            b.unmap();
-            done(out);
+        self.g.map_readback(&buf, row as usize, padded as usize, h as usize, move |r| {
+            done(r.map_err(|e| log::error!("gpu texture readback: {e}")).ok());
         });
         #[cfg(not(target_arch = "wasm32"))]
-        let _ = self.g.device.poll(wgpu::PollType::Poll);
+        let _ = self.g.poll_readbacks(wgpu::PollType::Poll);
     }
 
     /// A deferred readback of `img` (RGBA f32) under `key` (see [`crate::deferred`]): the bytes
diff --git a/crates/gpu/src/deferred.rs b/crates/gpu/src/deferred.rs
index 1dab0654..3ac2c67e 100644
--- a/crates/gpu/src/deferred.rs
+++ b/crates/gpu/src/deferred.rs
@@ -14,7 +14,7 @@
 //! of its key, and a chain whose kernels read other data through the effect host (other layers,
 //! other times) is not started when that data missed while it ran (the miss counter moved).
 
-use std::collections::{HashMap, HashSet};
+use std::collections::HashMap;
 use std::future::Future;
 use std::hash::{Hash, Hasher};
 use std::pin::Pin;
@@ -48,11 +48,9 @@ pub(crate) enum Lookup {
 
 #[derive(Default)]
 struct State {
-    /// Frame generation: results of an older frame's readbacks are dropped.
-    generation: u64,
     ready: HashMap<u128, Option<Readback>>,
-    waiting: HashSet<u128>,
-    in_flight: usize,
+    /// Each request owns a unique token, even when a key is reused after cancellation.
+    waiting: HashMap<u128, Arc<()>>,
     /// Misses so far in this frame (a counter, so a chain sees misses during its own work).
     misses: u64,
     wakers: Vec<Waker>,
@@ -71,11 +69,15 @@ pub struct Deferred {
 
 impl Deferred {
     pub fn frame_begin(&self) {
-        if let Ok(mut s) = self.st.lock() {
-            s.generation += 1;
+        let wakers = {
+            let mut s = self.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
             s.ready.clear();
             s.waiting.clear();
             s.misses = 0;
+            std::mem::take(&mut s.wakers)
+        };
+        for w in wakers {
+            w.wake();
         }
         self.gate.store(false, Ordering::Relaxed);
     }
@@ -111,32 +113,25 @@ impl Deferred {
         if let Some(r) = s.ready.get(&key) {
             return Lookup::Ready(r.clone());
         }
-        if s.waiting.contains(&key) { Lookup::InFlight } else { Lookup::Absent }
+        if s.waiting.contains_key(&key) { Lookup::InFlight } else { Lookup::Absent }
     }
 
     /// A readback for `key` is being started; returns the callback that stores its bytes.
     pub(crate) fn start(self: &Arc<Self>, key: u128, width: u32, height: u32, offset: [f64; 2], scale: f64) -> impl FnOnce(Option<Vec<u8>>) + 'static {
-        let generation = match self.st.lock() {
-            Ok(mut s) => {
-                s.waiting.insert(key);
-                s.in_flight += 1;
-                s.generation
-            }
-            Err(_) => 0,
-        };
+        let token = Arc::new(());
+        self.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner).waiting.insert(key, token.clone());
         self.readbacks.fetch_add(1, Ordering::Relaxed);
         let me = self.clone();
         move |bytes: Option<Vec<u8>>| {
-            let wakers = match me.st.lock() {
-                Ok(mut s) => {
-                    s.in_flight = s.in_flight.saturating_sub(1);
-                    if s.generation == generation {
-                        s.waiting.remove(&key);
-                        s.ready.insert(key, bytes.map(|bytes| Readback { bytes, width, height, offset, scale }));
-                    }
-                    if s.in_flight == 0 { std::mem::take(&mut s.wakers) } else { vec![] }
+            let wakers = {
+                let mut s = me.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+                // Reject stale completion before touching either results or accounting.
+                if !s.waiting.get(&key).is_some_and(|current| Arc::ptr_eq(current, &token)) {
+                    return;
                 }
-                Err(_) => vec![],
+                s.waiting.remove(&key);
+                s.ready.insert(key, bytes.map(|bytes| Readback { bytes, width, height, offset, scale }));
+                if s.waiting.is_empty() { std::mem::take(&mut s.wakers) } else { vec![] }
             };
             for w in wakers {
                 w.wake();
@@ -144,9 +139,24 @@ impl Deferred {
         }
     }
 
+    /// Cancel pending requests without exposing placeholder pixels as successful results.
+    /// Existing waiters wake; late callbacks cannot change a newer request's accounting.
+    pub(crate) fn cancel(&self) {
+        let wakers = {
+            let mut s = self.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            for (key, _) in std::mem::take(&mut s.waiting) {
+                s.ready.insert(key, None);
+            }
+            std::mem::take(&mut s.wakers)
+        };
+        for w in wakers {
+            w.wake();
+        }
+    }
+
     /// Readbacks still in flight.
     pub fn in_flight(&self) -> usize {
-        self.st.lock().map(|s| s.in_flight).unwrap_or(0)
+        self.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner).waiting.len()
     }
 
     /// Resolves once no readback is in flight (on the web the browser's event loop delivers
@@ -162,11 +172,13 @@ pub struct Settled(Arc<Deferred>);
 impl Future for Settled {
     type Output = ();
     fn poll(self: Pin<&mut Self>, cx: &mut Context<'_>) -> Poll<()> {
-        let Ok(mut s) = self.0.st.lock() else { return Poll::Ready(()) };
-        if s.in_flight == 0 {
+        let mut s = self.0.st.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+        if s.waiting.is_empty() {
             return Poll::Ready(());
         }
-        s.wakers.push(cx.waker().clone());
+        if !s.wakers.iter().any(|w| w.will_wake(cx.waker())) {
+            s.wakers.push(cx.waker().clone());
+        }
         Poll::Pending
     }
 }
@@ -249,6 +261,56 @@ pub(crate) fn chain_key(chain: &[FxStep], buf: &Buf, levels: Option<f32>) -> u12
 mod tests {
     use super::*;
 
+    struct WakeCount(std::sync::atomic::AtomicUsize);
+    impl std::task::Wake for WakeCount {
+        fn wake(self: Arc<Self>) {
+            self.0.fetch_add(1, Ordering::SeqCst);
+        }
+    }
+
+    #[test]
+    fn cancellation_wakes_waiters_and_stale_same_key_does_not_drain_new_work() {
+        let d = Arc::new(Deferred::default());
+        let old = d.start(7, 1, 1, [0.0; 2], 1.0);
+        let wake = Arc::new(WakeCount(std::sync::atomic::AtomicUsize::new(0)));
+        let waker = Waker::from(wake.clone());
+        let mut settled = std::pin::pin!(d.settled());
+        for _ in 0..4 {
+            assert!(settled.as_mut().poll(&mut Context::from_waker(&waker)).is_pending());
+        }
+        d.cancel();
+        assert_eq!(wake.0.load(Ordering::SeqCst), 1);
+        assert_eq!(d.in_flight(), 0);
+        assert!(settled.as_mut().poll(&mut Context::from_waker(&waker)).is_ready());
+        assert!(matches!(d.lookup(7), Lookup::Ready(None)));
+        d.frame_begin();
+        let new = d.start(7, 1, 1, [0.0; 2], 1.0);
+        old(Some(vec![99; 4]));
+        assert_eq!(d.in_flight(), 1);
+        assert!(matches!(d.lookup(7), Lookup::InFlight));
+        new(None);
+        assert_eq!(d.in_flight(), 0);
+        assert!(matches!(d.lookup(7), Lookup::Ready(None)));
+    }
+
+    #[test]
+    fn device_failure_without_map_callback_drains_and_wakes_deferred() {
+        let d = Arc::new(Deferred::default());
+        let registry = crate::readback::Readbacks::new().unwrap();
+        let done = d.start(11, 1, 1, [0.0; 2], 1.0);
+        let ticket = registry.start(move |r| done(r.ok())).unwrap();
+        let wake = Arc::new(WakeCount(std::sync::atomic::AtomicUsize::new(0)));
+        let waker = Waker::from(wake.clone());
+        let mut settled = std::pin::pin!(d.settled());
+        assert!(settled.as_mut().poll(&mut Context::from_waker(&waker)).is_pending());
+        registry.retire("synthetic device loss".into());
+        assert_eq!(wake.0.load(Ordering::SeqCst), 1);
+        assert!(settled.as_mut().poll(&mut Context::from_waker(&waker)).is_ready());
+        assert_eq!(d.in_flight(), 0);
+        ticket.complete(Ok(vec![99; 4]));
+        assert!(matches!(d.lookup(11), Lookup::Ready(None)));
+    }
+
     #[test]
     fn keys_are_stable_and_distinct() {
         let mut a = Key::new(1);
diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs
index 36b724a7..ece74bb1 100644
--- a/crates/gpu/src/lib.rs
+++ b/crates/gpu/src/lib.rs
@@ -81,6 +81,7 @@ mod fx_vr;
 mod fx_warp;
 mod ops;
 mod particles;
+mod readback;
 mod walk;
 
 use std::sync::Arc;
@@ -157,7 +158,7 @@ impl Gpu {
     /// Deliver finished readbacks without blocking (native; the browser delivers them itself).
     pub fn poll(&self) {
         #[cfg(not(target_arch = "wasm32"))]
-        let _ = self.ctx.device.poll(wgpu::PollType::Poll);
+        let _ = self.ctx.poll_readbacks(wgpu::PollType::Poll);
     }
 
     /// `{passes, readbacks, inFlight}` of deferred rendering (diagnostics).
@@ -177,10 +178,14 @@ impl Gpu {
     /// Render a top-level comp frame and leave it on the GPU as a display texture (no
     /// readback unless a CPU fallback step needs one). `None` = render on the CPU.
     pub fn render_display(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option<DisplayFrame> {
+        self.ctx.check_health().map_err(|e| log::error!("gpu display: {e}")).ok()?;
         let mut e = Enc::new(&self.ctx);
         let img = walk::render(&mut e, r, comp, t)?;
         let texture = e.display(&img);
         e.submit();
+        // Reject failures observed during rendering/submission before publishing the frame.
+        // A later device fault still requires the presentation host's lifecycle handling.
+        self.ctx.check_health().map_err(|e| log::error!("gpu display: {e}")).ok()?;
         Some(DisplayFrame { texture, width: img.width, height: img.height })
     }
 
@@ -192,11 +197,17 @@ impl Gpu {
     /// [`Gpu::read_display`] without blocking (the browser's main thread): `done` gets the bytes
     /// once the GPU has them.
     pub fn read_display_async(&self, f: &DisplayFrame, done: impl FnOnce(Option<Vec<u8>>) + wgpu::WasmNotSend + 'static) {
+        if let Err(e) = self.ctx.check_health() {
+            log::error!("gpu display readback: {e}");
+            crate::readback::reject(done, e);
+            return;
+        }
         Enc::new(&self.ctx).read_texture_async(&f.texture, f.width, f.height, 4, done);
     }
 
     /// Render a frame on the GPU and read it back (`None` = not handled).
     pub fn render(&self, r: &Renderer, comp: ItemId, t: Tick) -> Option<Image> {
+        self.ctx.check_health().ok()?;
         if !self.ctx.can_readback() {
             return None;
         }
@@ -244,11 +255,12 @@ impl Gpu {
         }
     }
 
-    /// Run GPU work `f` on this thread, catching the device running out of video memory
-    /// meanwhile (#106). `Err` (why) means the result is incomplete: a texture that failed to
-    /// allocate holds nothing. The uploaded layers and pooled textures are then freed, so the
-    /// next render has room. The browser reports it to the device's error handler instead.
+    /// Run work in a native thread-local out-of-memory scope. A failed result is discarded
+    /// and cached allocations are released; callers can retry with a fresh CPU renderer.
+    /// Browser callers use the host/owned-device failure bridge and bounded readbacks instead
+    /// of blocking their event loop on an error-scope future.
     pub fn within_memory<T>(&self, f: impl FnOnce() -> T) -> Result<T, String> {
+        self.ctx.check_health()?;
         #[cfg(not(target_arch = "wasm32"))]
         {
             let scope = self.ctx.device.push_error_scope(wgpu::ErrorFilter::OutOfMemory);
@@ -258,16 +270,19 @@ impl Gpu {
                     self.ctx.clear_uploads();
                     Err(e.to_string())
                 }
-                None => Ok(out),
+                None => self.ctx.check_health().map(|()| out),
             }
         }
         #[cfg(target_arch = "wasm32")]
-        Ok(f())
+        {
+            let out = f();
+            self.ctx.check_health().map(|()| out)
+        }
     }
 
-    /// Wait for all submitted GPU work (benchmarks).
+    /// Wait for submitted GPU work with the readback deadline (benchmarks).
     pub fn wait(&self) {
-        let _ = self.ctx.device.poll(wgpu::PollType::wait_indefinitely());
+        let _ = self.ctx.poll_readbacks(wgpu::PollType::Wait { submission_index: None, timeout: Some(crate::readback::TIMEOUT) });
     }
 }
 
@@ -281,22 +296,26 @@ impl Accelerator for Gpu {
     }
 
     fn supports_effect(&self, id: &str) -> bool {
-        self.ctx.can_readback() && effects::supports(id)
+        self.ctx.check_health().is_ok() && self.ctx.can_readback() && effects::supports(id)
     }
 
     fn effects(&self, chain: &[FxStep], buf: &Buf, levels: Option<f32>) -> Option<Buf> {
+        self.ctx.check_health().ok()?;
         effects::run_chain(&mut Enc::new(&self.ctx), chain, buf, levels)
     }
 
     fn raster_3d(&self, scene: &effectcraft_render::three_d::adv::Scene) -> Option<effectcraft_render::three_d::adv::Target> {
+        self.ctx.check_health().ok()?;
         adv3d::render(&self.ctx, scene)
     }
 
     fn render_3d(&self, run: &effectcraft_render::three_d::adv::Prepared) -> Option<effectcraft_render::three_d::adv::Rendered> {
+        self.ctx.check_health().ok()?;
         adv3d::render_prepared(&self.ctx, run)
     }
 
     fn particles(&self) -> Option<&dyn effectcraft_effects::psim::ParticleSim> {
+        self.ctx.check_health().ok()?;
         self.ctx.can_readback().then_some(self as &dyn effectcraft_effects::psim::ParticleSim)
     }
 
@@ -329,7 +348,7 @@ impl Accelerator for Gpu {
         // Natively the device delivers the readbacks when polled: wait for them here.
         #[cfg(not(target_arch = "wasm32"))]
         if d.in_flight() > 0 {
-            let _ = self.ctx.device.poll(wgpu::PollType::wait_indefinitely());
+            let _ = self.ctx.poll_readbacks(wgpu::PollType::Wait { submission_index: None, timeout: Some(crate::readback::TIMEOUT) });
         }
         Some(Box::pin(async move { d.settled().await }))
     }
@@ -337,6 +356,7 @@ impl Accelerator for Gpu {
 
 impl effectcraft_effects::psim::ParticleSim for Gpu {
     fn simulate(&self, req: &effectcraft_effects::psim::SimRequest) -> Option<Vec<effectcraft_effects::psim::SimParticle>> {
+        self.ctx.check_health().ok()?;
         particles::simulate(&self.ctx, req)
     }
 }
@@ -393,5 +413,7 @@ mod tests_fx_transition;
 mod tests_fx_vr;
 #[cfg(test)]
 mod tests_fx_warp;
+#[cfg(all(test, not(target_arch = "wasm32")))]
+mod tests_init;
 #[cfg(test)]
 mod tests_particles;
diff --git a/crates/gpu/src/particles.rs b/crates/gpu/src/particles.rs
index 70aaee9f..015edc82 100644
--- a/crates/gpu/src/particles.rs
+++ b/crates/gpu/src/particles.rs
@@ -193,29 +193,18 @@ pub(crate) fn simulate(g: &GpuContext, req: &SimRequest) -> Option<Vec<SimPartic
     enc.copy_buffer_to_buffer(&next, 0, &read, 0, size);
     g.queue.submit([enc.finish()]);
     if let (Some(d), Some(k)) = (&deferred, key) {
+        let row = usize::try_from(size).ok()?;
         let done = d.start(k, n, 1, [0.0; 2], 1.0);
-        let b = read.clone();
-        read.map_async(wgpu::MapMode::Read, .., move |r| {
-            let out = r.ok().and_then(|_| b.get_mapped_range(..).ok().map(|v| v.to_vec()));
-            b.unmap();
-            done(out);
-        });
+        g.map_readback(&read, row, row, 1, move |r| done(r.map_err(|e| log::error!("gpu particle readback: {e}")).ok()));
         #[cfg(not(target_arch = "wasm32"))]
-        let _ = g.device.poll(wgpu::PollType::Poll);
+        let _ = g.poll_readbacks(wgpu::PollType::Poll);
         keep_checkpoint(g, req, n, next);
         d.miss();
         return Some(vec![]);
     }
-    let (tx, rx) = std::sync::mpsc::channel();
-    read.map_async(wgpu::MapMode::Read, .., move |r| {
-        let _ = tx.send(r);
-    });
-    g.device.poll(wgpu::PollType::wait_indefinitely()).ok()?;
-    rx.recv().ok()?.ok()?;
-    let view = read.get_mapped_range(..).ok()?;
-    let parts = parse(&view);
-    drop(view);
-    read.unmap();
+    let row = usize::try_from(size).ok()?;
+    let bytes = g.read_buffer(&read, row, row, 1).map_err(|e| log::error!("gpu particle readback: {e}")).ok()?;
+    let parts = parse(&bytes);
     keep_checkpoint(g, req, n, next);
     Some(parts)
 }
diff --git a/crates/gpu/src/readback.rs b/crates/gpu/src/readback.rs
new file mode 100644
index 00000000..0ca01135
--- /dev/null
+++ b/crates/gpu/src/readback.rs
@@ -0,0 +1,529 @@
+//! Bounded, exactly-once readback completion, independent of GPU callbacks.
+// Browser callbacks stay on their worker thread. The workspace's wgpu feature still
+// requires Send callbacks on non-atomic wasm, and ParticleSim requires Send + Sync.
+use std::collections::BTreeMap;
+use std::sync::{Arc, Mutex, Weak};
+use std::time::Duration;
+
+pub(crate) const TIMEOUT: Duration = Duration::from_secs(10);
+pub(crate) type Error = String;
+type Done = Box<dyn FnOnce(Result<Vec<u8>, Error>) + Send>;
+
+struct Entry {
+    deadline: web_time::Instant,
+    done: Done,
+}
+#[derive(Default)]
+struct State {
+    next: u64,
+    pending: BTreeMap<u64, Entry>,
+    retired: Option<Error>,
+    #[cfg(target_arch = "wasm32")]
+    timer_armed: bool,
+}
+#[derive(Default)]
+pub(crate) struct Readbacks(Mutex<State>);
+pub(crate) struct Ticket {
+    owner: Weak<Readbacks>,
+    id: u64,
+}
+
+/// User callbacks run outside locks and are isolated one at a time. On panic=abort
+/// (including the usual wasm build), Rust panics cannot be caught: the host must
+/// avoid panicking callbacks. This branch does not claim abort-mode recovery.
+fn notify(done: Done, result: Result<Vec<u8>, Error>) {
+    #[cfg(panic = "unwind")]
+    if std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| done(result))).is_err() {
+        log::error!("GPU readback callback panicked; remaining callbacks will still settle");
+    }
+    #[cfg(not(panic = "unwind"))]
+    done(result);
+}
+
+/// Reject an optional-pixel callback through the same isolation boundary as mapped results.
+pub(crate) fn reject(done: impl FnOnce(Option<Vec<u8>>) + wgpu::WasmNotSend + 'static, error: Error) {
+    notify(Box::new(move |result| done(result.ok())), Err(error));
+}
+
+impl Readbacks {
+    pub(crate) fn new() -> Result<Arc<Self>, Error> {
+        let me = Arc::new(Self::default());
+        #[cfg(not(target_arch = "wasm32"))]
+        {
+            let weak = Arc::downgrade(&me);
+            std::thread::Builder::new()
+                .name("gpu-readback-deadlines".into())
+                .spawn(move || {
+                    loop {
+                        std::thread::sleep(Duration::from_millis(100));
+                        let Some(me) = weak.upgrade() else { break };
+                        me.expire(web_time::Instant::now());
+                    }
+                })
+                .map_err(|e| format!("starting GPU readback deadline worker: {e}"))?;
+        }
+        Ok(me)
+    }
+
+    pub(crate) fn check(&self) -> Result<(), Error> {
+        let s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+        match &s.retired {
+            Some(e) => Err(e.clone()),
+            None => Ok(()),
+        }
+    }
+
+    pub(crate) fn start(self: &Arc<Self>, done: impl FnOnce(Result<Vec<u8>, Error>) + wgpu::WasmNotSend + 'static) -> Result<Ticket, Error> {
+        let mut done: Option<Done> = Some(Box::new(done));
+        let result = {
+            let mut s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            if let Some(e) = &s.retired {
+                Err(e.clone())
+            } else if s.pending.len() >= 256 {
+                Err("too many pending GPU readbacks".into())
+            } else if let Some(id) = s.next.checked_add(1) {
+                s.next = id;
+                if let Some(done) = done.take() {
+                    s.pending.insert(id, Entry { deadline: web_time::Instant::now() + TIMEOUT, done });
+                }
+                Ok(Ticket { owner: Arc::downgrade(self), id })
+            } else {
+                Err("GPU readback request id exhausted".into())
+            }
+        };
+        if let Err(e) = &result
+            && let Some(done) = done
+        {
+            notify(done, Err(e.clone()));
+        }
+        #[cfg(target_arch = "wasm32")]
+        if result.is_ok() {
+            self.arm_timer();
+        }
+        result
+    }
+
+    #[cfg(target_arch = "wasm32")]
+    fn arm_timer(self: &Arc<Self>) {
+        let delay = {
+            let mut s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            if s.timer_armed || s.retired.is_some() {
+                return;
+            }
+            let Some(deadline) = s.pending.values().map(|e| e.deadline).min() else {
+                return;
+            };
+            s.timer_armed = true;
+            deadline.saturating_duration_since(web_time::Instant::now())
+        };
+        if let Err(e) = schedule(Arc::downgrade(self), delay) {
+            self.retire(e);
+        }
+    }
+
+    pub(crate) fn retire(&self, error: Error) {
+        let pending = {
+            let mut s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            if s.retired.is_some() {
+                return;
+            }
+            s.retired = Some(error.clone());
+            std::mem::take(&mut s.pending)
+        };
+        for (_, entry) in pending {
+            notify(entry.done, Err(error.clone()));
+        }
+    }
+
+    pub(crate) fn cancel_pending(&self, error: Error) {
+        let pending = {
+            let mut s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            std::mem::take(&mut s.pending)
+        };
+        for (_, entry) in pending {
+            notify(entry.done, Err(error.clone()));
+        }
+    }
+
+    fn expire(&self, now: web_time::Instant) {
+        let error = "GPU readback timed out; context retired".to_string();
+        let pending = {
+            let mut s = self.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            // Decide and drain atomically: a cancelled old deadline must not retire
+            // a newer request inserted between the deadline check and retirement.
+            if s.retired.is_some() || !s.pending.values().any(|e| now >= e.deadline) {
+                return;
+            }
+            s.retired = Some(error.clone());
+            std::mem::take(&mut s.pending)
+        };
+        for (_, entry) in pending {
+            notify(entry.done, Err(error.clone()));
+        }
+    }
+}
+impl Drop for Readbacks {
+    fn drop(&mut self) {
+        self.retire("GPU readback context dropped".into());
+    }
+}
+
+pub(crate) struct Layout {
+    pub row: u32,
+    pub stride: u32,
+    pub size: u64,
+    pub len: usize,
+}
+impl Layout {
+    pub(crate) fn new(w: u32, h: u32, bpp: u32, limit: u64) -> Result<Self, Error> {
+        if w == 0 || h == 0 || bpp == 0 {
+            return Err("empty GPU texture readback".into());
+        }
+        let row = w.checked_mul(bpp).ok_or("GPU readback row size overflow")?;
+        let alignment = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
+        let stride = row.checked_add(alignment - 1).map(|n| n / alignment * alignment).ok_or("GPU readback row alignment overflow")?;
+        let size = u64::from(stride) * u64::from(h);
+        // Bound host allocations even when the adapter advertises much larger buffers.
+        if size > limit.min(1 << 30) {
+            return Err("GPU readback exceeds buffer budget".into());
+        }
+        let len = usize::try_from(u64::from(row) * u64::from(h)).map_err(|e| format!("GPU readback host size: {e}"))?;
+        Ok(Self { row, stride, size, len })
+    }
+}
+
+impl Ticket {
+    pub(crate) fn active(&self) -> bool {
+        self.owner.upgrade().is_some_and(|me| me.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner).pending.contains_key(&self.id))
+    }
+    pub(crate) fn complete(self, result: Result<Vec<u8>, Error>) {
+        let Some(me) = self.owner.upgrade() else { return };
+        let entry = me.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner).pending.remove(&self.id);
+        if let Some(entry) = entry {
+            notify(entry.done, result);
+        }
+    }
+}
+
+#[cfg(target_arch = "wasm32")]
+fn schedule(me: Weak<Readbacks>, delay: Duration) -> Result<(), Error> {
+    use wasm_bindgen::{JsCast, closure::Closure};
+    let global = js_sys::global();
+    let function = js_sys::Reflect::get(&global, &"setTimeout".into())
+        .map_err(|e| format!("GPU deadline timer lookup: {e:?}"))?
+        .dyn_into::<js_sys::Function>()
+        .map_err(|_| "GPU deadline timer unavailable".to_string())?;
+    let callback = Closure::once(move || {
+        if let Some(me) = me.upgrade() {
+            me.0.lock().unwrap_or_else(std::sync::PoisonError::into_inner).timer_armed = false;
+            me.expire(web_time::Instant::now());
+            // If the earliest request completed, arm for the next live deadline.
+            me.arm_timer();
+        }
+    });
+    function
+        .call2(&global, callback.as_ref(), &wasm_bindgen::JsValue::from_f64(delay.as_secs_f64() * 1000.0 + 1.0))
+        .map_err(|e| format!("scheduling GPU readback deadline: {e:?}"))?;
+    // Transfer the wrapper to JS GC rather than leaking a wrapper per readback.
+    let _ = callback.into_js_value();
+    Ok(())
+}
+
+/// A weak notification bridge. Shared-device hosts keep ownership of their handlers.
+pub(crate) fn failure_notifier(me: &Arc<Readbacks>) -> impl Fn(String) + Send + Sync + 'static {
+    #[cfg(not(target_arch = "wasm32"))]
+    let fail = {
+        let me = Arc::downgrade(me);
+        move |error: String| {
+            if let Some(me) = me.upgrade() {
+                me.retire(error);
+            }
+        }
+    };
+    #[cfg(target_arch = "wasm32")]
+    let fail = {
+        // GPU callbacks require Send even on wasm. Capture only the index; the registry
+        // (which can contain JS callbacks) stays on this worker's event-loop thread.
+        let index = REGISTRIES.with(|r| {
+            let mut r = r.borrow_mut();
+            r.push(Arc::downgrade(me));
+            r.len() - 1
+        });
+        move |error: String| {
+            let me = REGISTRIES.with(|r| r.borrow().get(index).and_then(Weak::upgrade));
+            if let Some(me) = me {
+                me.retire(error);
+            }
+        }
+    };
+    fail
+}
+
+/// Replacement handlers are installed only by the successful owned-device request path.
+pub(crate) fn install(device: &wgpu::Device, me: &Arc<Readbacks>) {
+    let fail = Arc::new(failure_notifier(me));
+    let lost = fail.clone();
+    device.set_device_lost_callback(move |reason, message| lost(format!("GPU device lost ({reason:?}): {message}")));
+    device.on_uncaptured_error(Arc::new(move |e| {
+        let message = format!("uncaptured GPU error: {e}");
+        fail(message.clone());
+        #[cfg(test)]
+        #[allow(clippy::panic)]
+        if matches!(e, wgpu::Error::Validation { .. }) {
+            panic!("{message}");
+        }
+        log::error!("{message}");
+    }));
+}
+#[cfg(target_arch = "wasm32")]
+thread_local! {
+    static REGISTRIES: std::cell::RefCell<Vec<Weak<Readbacks>>> = const { std::cell::RefCell::new(Vec::new()) };
+}
+
+/// Validate every row before publishing tightly packed bytes.
+pub(crate) fn packed(view: &[u8], row: usize, stride: usize, height: usize) -> Result<Vec<u8>, Error> {
+    if row == 0 || height == 0 || stride < row {
+        return Err("invalid GPU readback layout".into());
+    }
+    let size = row.checked_mul(height).ok_or("GPU readback output size overflow")?;
+    let end = (height - 1).checked_mul(stride).and_then(|s| s.checked_add(row)).ok_or("GPU readback mapped size overflow")?;
+    if view.len() < end {
+        return Err(format!("short GPU readback: {} bytes, needs {end}", view.len()));
+    }
+    let mut out = Vec::new();
+    out.try_reserve_exact(size).map_err(|e| format!("allocating GPU readback output: {e}"))?;
+    for y in 0..height {
+        let start = y.checked_mul(stride).ok_or("GPU readback row overflow")?;
+        let bytes = view.get(start..start + row).ok_or("short GPU readback row")?;
+        out.extend_from_slice(bytes);
+    }
+    Ok(out)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use std::sync::atomic::{AtomicUsize, Ordering};
+
+    #[test]
+    fn early_rejection_reports_no_pixels_exactly_once() {
+        let calls = Arc::new(AtomicUsize::new(0));
+        let observed = calls.clone();
+        reject(
+            move |pixels| {
+                assert!(pixels.is_none());
+                observed.fetch_add(1, Ordering::SeqCst);
+            },
+            "synthetic early rejection".into(),
+        );
+        assert_eq!(calls.load(Ordering::SeqCst), 1);
+    }
+
+    #[test]
+    #[cfg(all(not(target_arch = "wasm32"), panic = "unwind"))]
+    fn early_rejection_contains_callback_panic_and_allows_later_completion() {
+        let calls = Arc::new(AtomicUsize::new(0));
+        let first = calls.clone();
+        reject(
+            move |pixels| {
+                assert!(pixels.is_none());
+                first.fetch_add(1, Ordering::SeqCst);
+                panic!("synthetic rejected callback panic");
+            },
+            "synthetic retired context".into(),
+        );
+        let second = calls.clone();
+        reject(
+            move |pixels| {
+                assert!(pixels.is_none());
+                second.fetch_add(1, Ordering::SeqCst);
+            },
+            "synthetic invalid layout".into(),
+        );
+        assert_eq!(calls.load(Ordering::SeqCst), 2);
+    }
+
+    #[test]
+    #[cfg(all(not(target_arch = "wasm32"), panic = "unwind"))]
+    fn panicking_callback_does_not_strand_waiters_or_kill_deadline_worker() {
+        use std::future::Future;
+        use std::task::{Context, Poll, Wake, Waker};
+        struct WakeCount(AtomicUsize);
+        impl Wake for WakeCount {
+            fn wake(self: Arc<Self>) {
+                self.0.fetch_add(1, Ordering::SeqCst);
+            }
+        }
+        let me = Readbacks::new().unwrap();
+        let deferred = Arc::new(crate::deferred::Deferred::default());
+        let (tx, rx) = std::sync::mpsc::channel();
+        let _first = me.start(|_| panic!("synthetic public callback panic")).unwrap();
+        for key in 1..=2 {
+            let done = deferred.start(key, 1, 1, [0.0; 2], 1.0);
+            let tx = tx.clone();
+            let _ticket = me
+                .start(move |r| {
+                    assert!(r.is_err());
+                    done(None);
+                    tx.send(()).unwrap();
+                })
+                .unwrap();
+        }
+        let wakes = Arc::new(WakeCount(AtomicUsize::new(0)));
+        let waker = Waker::from(wakes.clone());
+        let mut settled = std::pin::pin!(deferred.settled());
+        assert!(settled.as_mut().poll(&mut Context::from_waker(&waker)).is_pending());
+        me.cancel_pending("synthetic cancellation".into());
+        assert_eq!(deferred.in_flight(), 0);
+        assert_eq!(wakes.0.load(Ordering::SeqCst), 1);
+        assert!(matches!(settled.as_mut().poll(&mut Context::from_waker(&waker)), Poll::Ready(())));
+        assert!(me.0.lock().unwrap().pending.is_empty());
+        rx.recv_timeout(Duration::from_secs(2)).unwrap();
+        rx.recv_timeout(Duration::from_secs(2)).unwrap();
+
+        // Drive an actual later watchdog expiry without waiting for the production 10s.
+        let _first = me.start(|_| panic!("synthetic watchdog callback panic")).unwrap();
+        let tx2 = tx.clone();
+        let _second = me
+            .start(move |r| {
+                assert!(r.is_err());
+                tx2.send(()).unwrap();
+            })
+            .unwrap();
+        {
+            let mut state = me.0.lock().unwrap();
+            for entry in state.pending.values_mut() {
+                entry.deadline = web_time::Instant::now();
+            }
+        }
+        rx.recv_timeout(Duration::from_secs(2)).unwrap();
+        assert!(me.0.lock().unwrap().pending.is_empty());
+        assert!(me.check().is_err());
+    }
+
+    #[test]
+    fn missing_callback_times_out_once_and_rejects_late_success() {
+        let me = Arc::new(Readbacks::default());
+        let calls = Arc::new(AtomicUsize::new(0));
+        let c = calls.clone();
+        let ticket = me
+            .start(move |r| {
+                assert!(r.is_err());
+                c.fetch_add(1, Ordering::SeqCst);
+            })
+            .unwrap();
+        me.expire(web_time::Instant::now() + TIMEOUT + Duration::from_secs(1));
+        assert!(!ticket.active());
+        ticket.complete(Ok(vec![42]));
+        assert_eq!(calls.load(Ordering::SeqCst), 1);
+        assert!(me.check().is_err());
+    }
+    #[test]
+    fn cancellation_drains_all_and_new_context_is_independent() {
+        let old = Arc::new(Readbacks::default());
+        let calls = Arc::new(AtomicUsize::new(0));
+        let mut tickets = Vec::new();
+        for _ in 0..3 {
+            let c = calls.clone();
+            tickets.push(
+                old.start(move |r| {
+                    assert!(r.is_err());
+                    c.fetch_add(1, Ordering::SeqCst);
+                })
+                .unwrap(),
+            );
+        }
+        old.retire("synthetic device failure".into());
+        assert_eq!(calls.load(Ordering::SeqCst), 3);
+        let new = Arc::new(Readbacks::default());
+        let c = calls.clone();
+        let fresh = new
+            .start(move |r| {
+                assert_eq!(r.unwrap(), vec![7]);
+                c.fetch_add(1, Ordering::SeqCst);
+            })
+            .unwrap();
+        for t in tickets {
+            t.complete(Ok(vec![99]));
+        }
+        assert!(fresh.active());
+        fresh.complete(Ok(vec![7]));
+        assert_eq!(calls.load(Ordering::SeqCst), 4);
+    }
+    #[test]
+    fn cancelled_ticket_cannot_complete_a_new_request_on_same_context() {
+        let me = Arc::new(Readbacks::default());
+        let calls = Arc::new(AtomicUsize::new(0));
+        let c = calls.clone();
+        let old = me
+            .start(move |r| {
+                assert!(r.is_err());
+                c.fetch_add(1, Ordering::SeqCst);
+            })
+            .unwrap();
+        me.cancel_pending("cancelled".into());
+        let c = calls.clone();
+        let new = me
+            .start(move |r| {
+                assert_eq!(r.unwrap(), vec![7]);
+                c.fetch_add(1, Ordering::SeqCst);
+            })
+            .unwrap();
+        old.complete(Ok(vec![99]));
+        assert!(new.active());
+        new.complete(Ok(vec![7]));
+        assert_eq!(calls.load(Ordering::SeqCst), 2);
+    }
+
+    #[test]
+    #[cfg(not(target_arch = "wasm32"))]
+    fn completion_and_failure_race_invokes_callback_once() {
+        for _ in 0..32 {
+            let me = Arc::new(Readbacks::default());
+            let calls = Arc::new(AtomicUsize::new(0));
+            let c = calls.clone();
+            let ticket = me
+                .start(move |_| {
+                    c.fetch_add(1, Ordering::SeqCst);
+                })
+                .unwrap();
+            let barrier = Arc::new(std::sync::Barrier::new(2));
+            let b = barrier.clone();
+            let t = std::thread::spawn(move || {
+                b.wait();
+                ticket.complete(Ok(vec![1]));
+            });
+            barrier.wait();
+            me.retire("concurrent device failure".into());
+            t.join().unwrap();
+            assert_eq!(calls.load(Ordering::SeqCst), 1);
+        }
+    }
+
+    #[test]
+    fn malformed_layouts_fail_without_allocating() {
+        assert!(Layout::new(u32::MAX, 2, 16, u64::MAX).is_err());
+        assert!(Layout::new(1, 0, 4, u64::MAX).is_err());
+        assert!(Layout::new(8192, 8192, 16, 128).is_err());
+    }
+    #[test]
+    fn completed_request_survives_later_failure_without_second_callback() {
+        let me = Arc::new(Readbacks::default());
+        let calls = Arc::new(AtomicUsize::new(0));
+        let c = calls.clone();
+        let ticket = me
+            .start(move |r| {
+                assert_eq!(r.unwrap(), vec![1]);
+                c.fetch_add(1, Ordering::SeqCst);
+            })
+            .unwrap();
+        ticket.complete(Ok(vec![1]));
+        me.retire("later failure".into());
+        assert_eq!(calls.load(Ordering::SeqCst), 1);
+    }
+    #[test]
+    fn no_partial_rows() {
+        assert!(packed(&[0; 7], 4, 4, 2).is_err());
+        assert_eq!(packed(&[1, 2, 0, 0, 3, 4], 2, 4, 2).unwrap(), vec![1, 2, 3, 4]);
+    }
+}
diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs
index e77d6c18..f2a2d155 100644
--- a/crates/gpu/src/tests.rs
+++ b/crates/gpu/src/tests.rs
@@ -55,6 +55,8 @@ pub(crate) fn gpu() -> Option<&'static Gpu> {
         let g = Gpu::headless();
         if g.is_none() {
             eprintln!("effectcraft-gpu tests: no GPU adapter, skipping GPU comparisons");
+        } else if let Some(g) = &g {
+            eprintln!("effectcraft-gpu tests: adapter {}", g.ctx.name);
         }
         g
     })
diff --git a/crates/gpu/src/tests_deferred.rs b/crates/gpu/src/tests_deferred.rs
index 5f80affe..e798bd8f 100644
--- a/crates/gpu/src/tests_deferred.rs
+++ b/crates/gpu/src/tests_deferred.rs
@@ -21,7 +21,15 @@ pub(crate) fn deferred_gpu() -> Option<(&'static Gpu, std::sync::MutexGuard<'sta
     static LOCK: std::sync::Mutex<()> = std::sync::Mutex::new(());
     // Also no other GPU test while it runs (see `hold_gpu_lock`).
     crate::tests::hold_gpu_lock();
-    let g = G.get_or_init(|| pollster::block_on(Gpu::request_deferred()).map_err(|e| eprintln!("deferred gpu tests skipped: {e}")).ok()).as_ref()?;
+    let g = G
+        .get_or_init(|| {
+            let g = pollster::block_on(Gpu::request_deferred()).map_err(|e| eprintln!("deferred gpu tests skipped: {e}")).ok();
+            if let Some(g) = &g {
+                eprintln!("deferred gpu tests: adapter {}", g.ctx.name);
+            }
+            g
+        })
+        .as_ref()?;
     Some((g, LOCK.lock().unwrap_or_else(|e| e.into_inner())))
 }
 
diff --git a/crates/gpu/src/tests_init.rs b/crates/gpu/src/tests_init.rs
new file mode 100644
index 00000000..38c278c2
--- /dev/null
+++ b/crates/gpu/src/tests_init.rs
@@ -0,0 +1,161 @@
+//! Native constructor containment. Tiny original shaders; no frames or large allocations.
+//! Set EFFECTCRAFT_REQUIRE_GPU=1 for acceptance: unavailable adapters/devices then fail
+//! instead of taking the explicitly logged headless-CI skip path.
+#![cfg(not(target_arch = "wasm32"))]
+#![allow(clippy::expect_used)] // Test helpers require successful setup to exercise rejection.
+
+use std::sync::Arc;
+use std::sync::atomic::{AtomicUsize, Ordering};
+
+#[cfg(target_os = "windows")]
+use crate::context::GpuContext;
+use crate::context::init_resource;
+
+fn brief(message: impl std::fmt::Display) -> String {
+    message.to_string().chars().take(1024).collect()
+}
+
+fn native_device(desc: wgpu::InstanceDescriptor, required: bool) -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
+    crate::tests::hold_gpu_lock();
+    let required = required || std::env::var("EFFECTCRAFT_REQUIRE_GPU").is_ok_and(|value| value == "1");
+    let instance = wgpu::Instance::new(desc);
+    let adapter = match pollster::block_on(
+        instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }),
+    ) {
+        Ok(adapter) => adapter,
+        Err(error) => {
+            assert!(!required, "required initialization-test adapter unavailable: {}", brief(error));
+            eprintln!("initialization tests: no adapter, explicitly skipping");
+            return None;
+        }
+    };
+    let info = adapter.get_info();
+    eprintln!("initialization tests: adapter {} ({:?})", brief(&info.name), info.backend);
+    let limits = adapter.limits();
+    let desc = wgpu::DeviceDescriptor {
+        label: Some("initialization containment test"),
+        required_limits: wgpu::Limits {
+            max_texture_dimension_2d: limits.max_texture_dimension_2d.min(16384),
+            max_buffer_size: limits.max_buffer_size,
+            ..wgpu::Limits::default()
+        },
+        ..Default::default()
+    };
+    match pollster::block_on(adapter.request_device(&desc)) {
+        Ok((device, queue)) => Some((adapter, device, queue)),
+        Err(error) => {
+            assert!(!required, "required initialization-test device unavailable: {}", brief(error));
+            eprintln!("initialization tests: no device, explicitly skipping");
+            None
+        }
+    }
+}
+
+fn count_host_errors(device: &wgpu::Device) -> Arc<AtomicUsize> {
+    let calls = Arc::new(AtomicUsize::new(0));
+    let count = calls.clone();
+    device.on_uncaptured_error(Arc::new(move |_| {
+        count.fetch_add(1, Ordering::SeqCst);
+    }));
+    calls
+}
+
+fn invalid_module(device: &wgpu::Device) -> wgpu::ShaderModule {
+    device.create_shader_module(wgpu::ShaderModuleDescriptor {
+        label: Some("original deliberately invalid initialization-test module"),
+        source: wgpu::ShaderSource::Wgsl("this is deliberately not WGSL".into()),
+    })
+}
+
+fn valid_module(device: &wgpu::Device) -> wgpu::ShaderModule {
+    init_resource(device, "valid tiny module", || {
+        device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("original empty initialization-test kernel"),
+            source: wgpu::ShaderSource::Wgsl("@compute @workgroup_size(1) fn main() {}".into()),
+        })
+    })
+    .expect("a valid tiny module still compiles after rejection")
+}
+
+fn pipeline(device: &wgpu::Device, module: &wgpu::ShaderModule, entry: &str) -> wgpu::ComputePipeline {
+    device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+        label: Some("initialization-test pipeline"),
+        layout: None,
+        module,
+        entry_point: Some(entry),
+        compilation_options: Default::default(),
+        cache: None,
+    })
+}
+
+fn valid_pipeline(device: &wgpu::Device) {
+    let module = valid_module(device);
+    init_resource(device, "valid tiny pipeline", || pipeline(device, &module, "main")).expect("a valid pipeline still compiles after rejection");
+}
+
+#[test]
+fn scoped_module_and_pipeline_errors_preserve_outer_scope_and_host_handler() {
+    let Some((_, device, _queue)) = native_device(wgpu::InstanceDescriptor::new_without_display_handle().with_env(), false) else { return };
+    let host = count_host_errors(&device);
+    let outer = device.push_error_scope(wgpu::ErrorFilter::Validation);
+
+    let error = init_resource(&device, "invalid tiny module", || invalid_module(&device)).expect_err("invalid module must return Err");
+    assert!(error.contains("GPU initialization (invalid tiny module)"), "{}", brief(error));
+    let module = valid_module(&device);
+    let error = init_resource(&device, "missing tiny entry", || pipeline(&device, &module, "missing_entry")).expect_err("missing entry point must return Err");
+    assert!(error.contains("GPU initialization (missing tiny entry)"), "{}", brief(error));
+    assert_eq!(host.load(Ordering::SeqCst), 0, "contained errors must not reach the host");
+    valid_pipeline(&device);
+
+    // This belongs to our pre-existing caller scope. Its successful pop also proves that
+    // none of the three nested helper scopes remain on the native thread-local stack.
+    let _invalid = invalid_module(&device);
+    assert!(pollster::block_on(outer.pop()).is_some(), "the caller's own scope survives");
+    assert_eq!(host.load(Ordering::SeqCst), 0);
+    let _invalid = invalid_module(&device);
+    assert_eq!(host.load(Ordering::SeqCst), 1, "the original host handler still receives unscoped errors");
+}
+
+#[test]
+fn initialization_unwind_drains_all_scopes_before_the_next_operation() {
+    let Some((_, device, _queue)) = native_device(wgpu::InstanceDescriptor::new_without_display_handle().with_env(), false) else { return };
+    let host = count_host_errors(&device);
+    let outer = device.push_error_scope(wgpu::ErrorFilter::Validation);
+    let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| {
+        init_resource::<()>(&device, "synthetic native unwind", || panic!("original synthetic initialization failure"))
+    }));
+    let error = result.expect("the initialization unwind is contained").expect_err("unwind becomes an error");
+    assert!(error.contains("synthetic native unwind") && error.contains("original synthetic initialization failure"), "{}", brief(error));
+    valid_pipeline(&device);
+    let _invalid = invalid_module(&device);
+    assert!(pollster::block_on(outer.pop()).is_some(), "all helper scopes must be drained after unwind");
+    assert_eq!(host.load(Ordering::SeqCst), 0);
+    let _invalid = invalid_module(&device);
+    assert_eq!(host.load(Ordering::SeqCst), 1, "unwind must neither replace the host handler nor leak a scope");
+}
+
+/// Acceptance for the captured RTX 5090 / FXC failure. Future compiler/shader fixes can
+/// legitimately make FXC succeed; this diagnostic is deliberately outside portable CI.
+#[test]
+#[cfg(target_os = "windows")]
+#[ignore = "requires the reproduced Windows DX12/FXC pointwise compiler failure; no adapter skips"]
+fn reproduced_fxc_pointwise_failure_returns_error_without_escaping_or_notifying_host() {
+    let mut desc = wgpu::InstanceDescriptor::new_without_display_handle();
+    desc.backends = wgpu::Backends::DX12;
+    desc.backend_options.dx12.shader_compiler = wgpu::Dx12Compiler::Fxc;
+    let (adapter, device, queue) = native_device(desc, true).expect("required DX12 device");
+    assert_eq!(adapter.get_info().backend, wgpu::Backend::Dx12);
+    let host = count_host_errors(&device);
+    let outer = device.push_error_scope(wgpu::ErrorFilter::Validation);
+    let result = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| GpuContext::new(&adapter, device.clone(), queue)));
+    let error = match result.expect("FXC compiler failure must not escape the constructor") {
+        Err(error) => error,
+        Ok(_) => panic!("the platform-specific FXC failure no longer reproduces; reassess this diagnostic"),
+    };
+    assert!(error.contains("GPU initialization (pointwise)") && error.contains("FXC D3DCompile"), "{}", brief(error));
+    assert!(pollster::block_on(outer.pop()).is_none(), "constructor must drain its own scopes and leave the caller's scope clean");
+    assert_eq!(host.load(Ordering::SeqCst), 0);
+    valid_pipeline(&device);
+    let _invalid = invalid_module(&device);
+    assert_eq!(host.load(Ordering::SeqCst), 1);
+}
diff --git a/crates/ui-egui/src/frames.rs b/crates/ui-egui/src/frames.rs
index 6d14f66f..9fa79594 100644
--- a/crates/ui-egui/src/frames.rs
+++ b/crates/ui-egui/src/frames.rs
@@ -24,6 +24,7 @@
 //! them (Info panel, eyedroppers, histograms).
 
 use std::collections::{BTreeSet, HashMap, HashSet, VecDeque};
+use std::sync::atomic::{AtomicBool, Ordering};
 use std::sync::{Arc, Mutex};
 
 use effectcraft_engine::project::{Comp, ItemId, ItemKind, LayerSource, Node, Project, PropGroup};
@@ -365,6 +366,7 @@ pub struct Frames {
     /// Render time of the last viewer (urgent) frame in ms, for the Info panel / perf readout.
     pub last_ms: Arc<Mutex<f64>>,
     content_keys: ContentKeys,
+    gpu_retired: Arc<AtomicBool>,
     /// [`comp_content`] identities by (revision, comp), with the project that keeps an identity
     /// found by [`Frames::identity`] in another project.
     identities: Mutex<HashMap<(u64, u64), Identity>>,
@@ -398,6 +400,7 @@ impl Default for Frames {
             ctx: None,
             last_ms: Arc::new(Mutex::new(0.0)),
             content_keys: Arc::default(),
+            gpu_retired: Arc::default(),
             identities: Mutex::default(),
             bases: Mutex::default(),
             keepers: Mutex::default(),
@@ -423,6 +426,24 @@ pub fn to_color_image(img: &effectcraft_engine::render::Image) -> egui::ColorIma
 }
 
 impl Frames {
+    /// Retire only GPU preview state; documents, CPU frames and running job leases survive.
+    /// Publication checks the same latch while holding the cache lock, so a late GPU frame
+    /// cannot restore an invalid texture after this removal.
+    pub fn retire_gpu(&self) {
+        self.gpu_retired.store(true, Ordering::Release);
+        if let Ok(mut cache) = self.cache.lock() {
+            let keys: Vec<_> = cache.map.iter().filter_map(|(key, image)| image.is_gpu().then_some(*key)).collect();
+            for key in keys {
+                cache.remove(&key);
+            }
+            let Cache { map, order, .. } = &mut *cache;
+            order.retain(|key| map.contains_key(key));
+        }
+        if let Some(ctx) = &self.ctx {
+            ctx.request_repaint();
+        }
+    }
+
     /// Actual native frame-worker capacity, including the OS fallback pool.
     #[cfg(not(target_arch = "wasm32"))]
     pub fn parallelism(&self) -> usize {
@@ -649,6 +670,7 @@ impl Frames {
             ctx: self.ctx.clone(),
             last: self.last_ms.clone(),
             content_keys: self.content_keys.clone(),
+            gpu_retired: self.gpu_retired.clone(),
         }
     }
 
@@ -777,6 +799,7 @@ struct Worker {
     ctx: Option<egui::Context>,
     last: Arc<Mutex<f64>>,
     content_keys: ContentKeys,
+    gpu_retired: Arc<AtomicBool>,
 }
 
 impl Worker {
@@ -830,7 +853,9 @@ impl Worker {
         if urgent && let Ok(mut l) = self.last.lock() {
             *l = t0.elapsed().as_secs_f64() * 1000.0;
         }
-        if let Ok(mut c) = self.cache.lock() {
+        if let Ok(mut c) = self.cache.lock()
+            && (!img.is_gpu() || !self.gpu_retired.load(Ordering::Acquire))
+        {
             c.insert(key, img);
         }
         self.release(key, urgent);
@@ -862,8 +887,13 @@ impl Worker {
         })
     }
 
-    /// The GPU ran out of video memory rendering a frame (#106): fewer GPU frames are kept.
-    fn out_of_memory(&self, e: &str) {
+    /// A failed scoped attempt retries on CPU. Retired contexts cannot recover by reducing
+    /// their budget; only a healthy device's allocation failure uses the upstream backoff.
+    fn gpu_attempt_failed(&self, gpu: &Gpu, e: &str) {
+        if gpu.context().check_health().is_err() {
+            log::warn!("the GPU context is retired ({e}); the frame renders on the CPU");
+            return;
+        }
         log::warn!("the GPU ran out of memory ({e}); the frame renders on the CPU and fewer GPU frames are kept");
         if let Ok(mut c) = self.cache.lock() {
             c.gpu_out_of_memory();
@@ -872,6 +902,9 @@ impl Worker {
 
     /// The frame on the GPU, kept there for the viewer, if the compositor can do it alone.
     fn render_gpu(&self, job: &Job) -> Option<FrameImage> {
+        if self.gpu_retired.load(Ordering::Acquire) {
+            return None;
+        }
         let g = job.src.gpu.as_ref()?;
         let mut r = Renderer::new(&job.src.project, job.src.footage.as_ref(), job.opts);
         r.expr = job.src.expr.as_deref();
@@ -881,7 +914,7 @@ impl Worker {
         match g.within_memory(|| g.render_display(&r, job.comp, job.t)) {
             Ok(f) => f.map(|f| FrameImage::Gpu(Arc::new(f))),
             Err(e) => {
-                self.out_of_memory(&e);
+                self.gpu_attempt_failed(g, &e);
                 None
             }
         }
@@ -891,8 +924,11 @@ impl Worker {
     fn render_job(&self, job: &Job, gpu: bool) -> FrameImage {
         let mut r = Renderer::new(&job.src.project, job.src.footage.as_ref(), job.opts);
         r.expr = job.src.expr.as_deref();
-        r.cache = Some(&job.src.layer_cache);
-        let gpu = job.src.gpu.as_ref().filter(|_| gpu);
+        let gpu_allowed = gpu && !self.gpu_retired.load(Ordering::Acquire);
+        // A failed GPU attempt may have populated intermediate entries before its scope
+        // reported failure. The CPU retry must compute its own pixels from the snapshot.
+        r.cache = gpu_allowed.then_some(job.src.layer_cache.as_ref());
+        let gpu = job.src.gpu.as_ref().filter(|_| gpu_allowed);
         r.accel = gpu.map(|g| g as &dyn effectcraft_engine::render::Accelerator);
         // The GPU leaves the frame in a texture for the viewer; otherwise (Software Only, no
         // adapter, or a frame the GPU cannot finish here) the CPU renders it.
@@ -911,6 +947,9 @@ impl Worker {
                         if let Some((p, key)) = auto {
                             // Time the GPU's work, not just its recording.
                             g.wait();
+                            if g.context().check_health().is_err() {
+                                return self.render_job(job, false);
+                            }
                             p.record(key, true, t0.elapsed().as_secs_f64() * 1000.0);
                         }
                         return FrameImage::Gpu(Arc::new(f));
@@ -921,7 +960,7 @@ impl Worker {
                         }
                     }
                     Err(e) => {
-                        self.out_of_memory(&e);
+                        self.gpu_attempt_failed(g, &e);
                         return self.render_job(job, false);
                     }
                 }
@@ -938,7 +977,7 @@ impl Worker {
             return match g.within_memory(|| r.comp_frame(job.comp, job.t)) {
                 Ok(img) => FrameImage::Cpu(Arc::new(to_color_image(&img))),
                 Err(e) => {
-                    self.out_of_memory(&e);
+                    self.gpu_attempt_failed(g, &e);
                     self.render_job(job, false)
                 }
             };
@@ -976,6 +1015,24 @@ impl Queue {
 mod tests {
     use super::*;
 
+    #[test]
+    fn retirement_reaches_existing_workers_and_keeps_cpu_frames_and_completion_live() {
+        let frames = Frames::default();
+        let worker = frames.worker();
+        let first = key(1, 0, 0);
+        insert(&frames, first);
+        frames.retire_gpu();
+        assert!(worker.gpu_retired.load(Ordering::Acquire), "already queued workers share retirement");
+        assert!(frames.is_cached(&first), "healthy materialized CPU frames survive");
+        let next = key(1, 1, 0);
+        frames.inflight.lock().unwrap().insert(next);
+        frames.queue.lock().unwrap().running.insert(next);
+        worker.finish_key(next, false, img(), web_time::Instant::now());
+        assert!(frames.is_cached(&next), "CPU retry can still publish");
+        assert_eq!(frames.inflight(), 0);
+        assert!(frames.queue.lock().unwrap().running.is_empty());
+    }
+
     fn key(content: u64, frame: i64, opts: u64) -> FrameKey {
         FrameKey { revision: content, content, comp: 1, frame, scale: 1000, view: 0, opts }
     }
diff --git a/crates/ui-egui/src/gpu_failure.rs b/crates/ui-egui/src/gpu_failure.rs
new file mode 100644
index 00000000..27511291
--- /dev/null
+++ b/crates/ui-egui/src/gpu_failure.rs
@@ -0,0 +1,158 @@
+//! Host-owned shared-device failure forwarding. Device handlers belong to the executable or
+//! embedding host, never to a library borrowing its device. This bridge does no device work,
+//! document mutation, file I/O or renderer locking inside a notification.
+
+use std::sync::{Arc, Mutex};
+
+type Notifier = Arc<dyn Fn(String) + Send + Sync>;
+
+#[derive(Clone, Debug, PartialEq, Eq)]
+pub struct GpuFailure {
+    pub reason: String,
+    /// Device loss requires restarting/recreating the presentation backend too. CPU
+    /// compositing alone cannot make a lost eframe presentation device usable again.
+    pub presentation_lost: bool,
+}
+
+#[derive(Default)]
+struct State {
+    failure: Option<GpuFailure>,
+    notifier: Option<Notifier>,
+}
+
+/// One bridge per host device. Failure is latched; duplicate reports cannot grow a queue.
+#[derive(Clone)]
+pub struct GpuFailureBridge {
+    state: Arc<Mutex<State>>,
+    ctx: egui::Context,
+}
+
+impl GpuFailureBridge {
+    pub fn new(ctx: &egui::Context) -> Self {
+        Self { state: Arc::default(), ctx: ctx.clone() }
+    }
+
+    pub fn failure(&self) -> Option<GpuFailure> {
+        self.state.lock().unwrap_or_else(std::sync::PoisonError::into_inner).failure.clone()
+    }
+
+    /// Forward to a compositor's weak notifier once construction succeeds. A fault that
+    /// preceded registration is replayed; the compositor must not then be attached to the UI.
+    pub fn forward_to(&self, notifier: impl Fn(String) + Send + Sync + 'static) {
+        let notifier: Notifier = Arc::new(notifier);
+        let failure = {
+            let mut state = self.state.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            state.notifier = Some(notifier.clone());
+            state.failure.clone()
+        };
+        if let Some(failure) = failure {
+            notify(&notifier, failure.reason);
+        }
+    }
+
+    /// Called by the device owner for an uncaptured error or device loss. A later device-loss
+    /// notification escalates the recovery message without completing pending work twice.
+    pub fn report(&self, reason: &str, presentation_lost: bool) -> bool {
+        let (notification, changed) = {
+            let mut state = self.state.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
+            if let Some(failure) = &mut state.failure {
+                let changed = presentation_lost && !failure.presentation_lost;
+                failure.presentation_lost |= presentation_lost;
+                (None, changed)
+            } else {
+                let reason: String = reason.chars().take(2048).collect();
+                state.failure = Some(GpuFailure { reason: reason.clone(), presentation_lost });
+                (state.notifier.clone().map(|notifier| (notifier, reason)), true)
+            }
+        };
+        if let Some((notifier, reason)) = notification {
+            notify(&notifier, reason);
+        }
+        if changed {
+            self.ctx.request_repaint();
+        }
+        changed
+    }
+}
+
+fn notify(notifier: &Notifier, reason: String) {
+    // A host callback must never propagate a consumer panic into wgpu's error handling.
+    #[cfg(not(target_arch = "wasm32"))]
+    if std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| notifier(reason))).is_err() {
+        log::error!("GPU failure consumer panicked");
+    }
+    #[cfg(target_arch = "wasm32")]
+    notifier(reason);
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use std::sync::atomic::{AtomicUsize, Ordering};
+
+    #[test]
+    fn a_construction_failure_is_replayed_after_notifier_registration() {
+        let bridge = GpuFailureBridge::new(&egui::Context::default());
+        bridge.report("pipeline allocation failed", false);
+        let observed = Arc::new(Mutex::new(Vec::new()));
+        let copy = observed.clone();
+        let probe = Arc::downgrade(&bridge.state);
+        bridge.forward_to(move |reason| {
+            assert!(probe.upgrade().unwrap().try_lock().is_ok(), "consumer runs outside bridge lock");
+            copy.lock().unwrap().push(reason);
+        });
+        assert_eq!(*observed.lock().unwrap(), ["pipeline allocation failed"]);
+        bridge.report("duplicate", false);
+        assert_eq!(observed.lock().unwrap().len(), 1);
+    }
+
+    #[test]
+    fn a_late_failure_notifies_once_and_later_device_loss_escalates() {
+        let bridge = GpuFailureBridge::new(&egui::Context::default());
+        let calls = Arc::new(AtomicUsize::new(0));
+        let copy = calls.clone();
+        bridge.forward_to(move |_| {
+            copy.fetch_add(1, Ordering::SeqCst);
+        });
+        assert!(bridge.failure().is_none());
+        bridge.report("uncaptured allocation failure", false);
+        bridge.report("device lost", true);
+        bridge.report("another error", false);
+        assert_eq!(calls.load(Ordering::SeqCst), 1);
+        let failure = bridge.failure().unwrap();
+        assert!(failure.presentation_lost);
+        assert_eq!(failure.reason, "uncaptured allocation failure");
+    }
+
+    #[test]
+    fn diagnostics_are_bounded_and_native_consumer_panics_are_isolated() {
+        let bridge = GpuFailureBridge::new(&egui::Context::default());
+        #[cfg(not(target_arch = "wasm32"))]
+        bridge.forward_to(|_| panic!("synthetic consumer failure"));
+        bridge.report(&"é".repeat(3000), false);
+        assert_eq!(bridge.failure().unwrap().reason.chars().count(), 2048);
+        bridge.report("device lost after consumer failure", true);
+        assert!(bridge.failure().unwrap().presentation_lost);
+    }
+
+    #[test]
+    #[cfg(not(target_arch = "wasm32"))]
+    fn concurrent_registration_and_failure_deliver_exactly_once() {
+        let bridge = GpuFailureBridge::new(&egui::Context::default());
+        let barrier = Arc::new(std::sync::Barrier::new(2));
+        let calls = Arc::new(AtomicUsize::new(0));
+        let reporter = bridge.clone();
+        let ready = barrier.clone();
+        let thread = std::thread::spawn(move || {
+            ready.wait();
+            reporter.report("device failed during subscription", false);
+        });
+        let count = calls.clone();
+        barrier.wait();
+        bridge.forward_to(move |_| {
+            count.fetch_add(1, Ordering::SeqCst);
+        });
+        thread.join().unwrap();
+        assert_eq!(calls.load(Ordering::SeqCst), 1);
+    }
+}
diff --git a/crates/ui-egui/src/lib.rs b/crates/ui-egui/src/lib.rs
index 2f25561f..e3cbf26c 100644
--- a/crates/ui-egui/src/lib.rs
+++ b/crates/ui-egui/src/lib.rs
@@ -14,6 +14,7 @@ pub mod control;
 pub mod dock;
 pub mod dock_ui;
 pub mod frames;
+pub mod gpu_failure;
 pub mod header;
 pub mod i18n;
 pub mod icons;
@@ -215,6 +216,8 @@ pub struct EffectcraftApp {
     /// The GPU compositor on that device (None: no usable adapter → CPU only).
     pub(crate) gpu: Option<effectcraft_gpu::Gpu>,
     gpu_checked: bool,
+    gpu_failures: Option<gpu_failure::GpuFailureBridge>,
+    gpu_failure_applied: Option<bool>,
     /// Commands from the native menu bar.
     pub command_inbox: Option<Receiver<String>>,
     /// Pointer position in comp pixels (Info panel).
@@ -286,6 +289,8 @@ impl EffectcraftApp {
             wgpu: None,
             gpu: None,
             gpu_checked: false,
+            gpu_failures: None,
+            gpu_failure_applied: None,
             command_inbox: None,
             pointer_comp: None,
             dialog_state: Default::default(),
@@ -455,6 +460,60 @@ impl EffectcraftApp {
             && self.ui.viewer.extended.is_none()
     }
 
+    /// Provide the executable/embedding host's device notifications. The host installs the
+    /// handlers once; this frontend never replaces handlers on a borrowed device.
+    pub fn set_gpu_failure_bridge(&mut self, bridge: gpu_failure::GpuFailureBridge) {
+        if let Some(gpu) = &self.gpu {
+            bridge.forward_to(gpu.context().device_failure_notifier());
+        }
+        self.gpu_failures = Some(bridge);
+    }
+
+    fn apply_gpu_failure(&mut self, ctx: &egui::Context) {
+        let Some(bridge) = &self.gpu_failures else { return };
+        // Readback deadlines can retire the compositor without a host device notification.
+        if let Some(gpu) = &self.gpu
+            && let Err(reason) = gpu.context().check_health()
+        {
+            bridge.report(&reason, false);
+        }
+        let Some(failure) = bridge.failure() else { return };
+        if self.gpu_failure_applied == Some(failure.presentation_lost) {
+            return;
+        }
+        self.gpu_failure_applied = Some(failure.presentation_lost);
+        self.gpu_checked = true;
+        self.frames.retire_gpu();
+        self.session.accel = None;
+        self.session.layer_cache.clear();
+        self.gpu = None;
+        if let Some(rs) = &self.wgpu {
+            let mut renderer = rs.renderer.write();
+            if let Some((id, _, _)) = &self.viewer_native {
+                renderer.free_texture(id);
+            }
+            for (_, texture) in self.passive_tex.values() {
+                if let panels::viewers::PassiveTexture::Gpu(id, _) = texture {
+                    renderer.free_texture(id);
+                }
+            }
+        }
+        self.viewer_native = None;
+        self.passive_tex.retain(|_, (_, texture)| matches!(texture, panels::viewers::PassiveTexture::Cpu(_)));
+        self.viewer_shown = self.viewer_tex.as_ref().map(|(texture, key)| (texture.id(), *key));
+        self.viewer_image = None;
+        // Old completion closures retain only their detached slot, not the replacement.
+        self.viewer_readback = Default::default();
+        self.wgpu = None;
+        self.ui.status = if failure.presentation_lost {
+            "Graphics device lost. The project is still open; presentation requires restarting the application.".into()
+        } else {
+            "Shared GPU preview disabled; CPU compositing remains available.".into()
+        };
+        log::warn!("GPU preview retired: {}", failure.reason);
+        ctx.request_repaint();
+    }
+
     /// Set up the GPU compositor on egui-wgpu's device (once). Without an adapter that runs
     /// compute shaders (WebGL2, software GL) everything stays on the CPU.
     pub(crate) fn init_gpu(&mut self, frame: &eframe::Frame) {
@@ -462,12 +521,21 @@ impl EffectcraftApp {
             return;
         }
         self.gpu_checked = true;
+        if self.gpu_failures.as_ref().is_some_and(|bridge| bridge.failure().is_some()) {
+            return;
+        }
         let Some(rs) = frame.wgpu_render_state() else { return };
         // Backend validation can panic on shader compilation even when device creation
         // succeeded. Keep the document and the egui renderer, and use CPU compositing.
         let gpu = std::panic::catch_unwind(std::panic::AssertUnwindSafe(|| effectcraft_gpu::Gpu::new(&rs.adapter, rs.device.clone(), rs.queue.clone())));
         match gpu {
             Ok(Ok(g)) => {
+                if let Some(bridge) = &self.gpu_failures {
+                    bridge.forward_to(g.context().device_failure_notifier());
+                    if bridge.failure().is_some() {
+                        return;
+                    }
+                }
                 log::info!("GPU compositor: {}", effectcraft_engine::render::Accelerator::name(&g));
                 self.session.accel = Some(Arc::new(g.clone()));
                 self.gpu = Some(g);
@@ -1263,6 +1331,7 @@ impl EffectcraftApp {
 
     fn frame(&mut self, ui: &mut egui::Ui) {
         let ctx = ui.ctx().clone();
+        self.apply_gpu_failure(&ctx);
         self.auto.begin_frame();
         self.frames.set_context(&ctx);
         if self.session.render_job.is_some() {
@@ -1439,6 +1508,7 @@ impl eframe::App for EffectcraftApp {
     }
 
     fn logic(&mut self, ctx: &egui::Context, _frame: &mut eframe::Frame) {
+        self.apply_gpu_failure(ctx);
         if !self.styled {
             theme::install(ctx, &self.tokens);
             self.styled = true;
@@ -1495,3 +1565,42 @@ impl eframe::App for EffectcraftApp {
 pub(crate) fn tick_playback(app: &mut EffectcraftApp, ctx: &egui::Context, scale: f64) {
     app.advance_playback(ctx, scale);
 }
+
+#[cfg(test)]
+mod gpu_failure_tests {
+    use super::*;
+
+    #[test]
+    fn host_failure_preserves_document_undo_and_cpu_pixels_and_detaches_late_readbacks() {
+        let mut session = Session::default();
+        let comp = ItemId(session.execute("comp.new", json!({"name": "Recovery", "width": 4, "height": 4, "duration": 1})).unwrap()["comp"].as_u64().unwrap());
+        session.execute("layer.newSolid", json!({"comp": comp.0, "color": "#ff0000"})).unwrap();
+        let before = session.render(comp, Tick::ZERO, RenderOpts::default());
+        let project = session.project.clone();
+        let revision = session.revision;
+        let prefs = serde_json::to_value(&session.prefs).unwrap();
+        let mut app = EffectcraftApp::new(session);
+        let ctx = egui::Context::default();
+        let bridge = gpu_failure::GpuFailureBridge::new(&ctx);
+        app.set_gpu_failure_bridge(bridge.clone());
+        let late_slot = app.viewer_readback.1.clone();
+        let key = FrameKey { revision, content: 1, comp: comp.0, frame: 0, scale: 1000, view: 0, opts: 0 };
+        app.viewer_readback.0 = Some(key);
+        bridge.report("synthetic host device failure", true);
+        app.apply_gpu_failure(&ctx);
+        assert!(Arc::ptr_eq(&project, &app.session.project));
+        assert_eq!(app.session.revision, revision);
+        assert_eq!(serde_json::to_value(&app.session.prefs).unwrap(), prefs);
+        assert!(app.gpu_checked, "failed device cannot be initialized again next frame");
+        assert!(app.session.accel.is_none());
+        assert!(app.ui.status.contains("presentation requires restarting"));
+        assert_eq!(app.session.render(comp, Tick::ZERO, RenderOpts::default()).data, before.data);
+        *late_slot.lock().unwrap() = Some((key, Arc::new(egui::ColorImage::new([1, 1], vec![egui::Color32::RED]))));
+        assert!(app.viewer_readback.0.is_none());
+        assert!(app.viewer_readback.1.lock().unwrap().is_none(), "a late completion cannot refill the replacement slot");
+        app.session.execute("edit.undo", json!({})).unwrap();
+        assert!(app.session.project.comp(comp).unwrap().layers.is_empty());
+        app.session.execute("edit.redo", json!({})).unwrap();
+        assert_eq!(app.session.render(comp, Tick::ZERO, RenderOpts::default()).data, before.data);
+    }
+}
diff --git a/docs/architecture.md b/docs/architecture.md
index 441715de..49f9da4a 100644
--- a/docs/architecture.md
+++ b/docs/architecture.md
@@ -272,7 +272,21 @@ as NaNs (a kernel that skips pixels fails the oracle) and validation errors pani
 tests, device errors are logged rather than fatal. A viewer frame that runs out of video memory
 (`Gpu::within_memory`) renders on the CPU, the uploaded layers and pooled textures are freed, and
 the RAM preview keeps half as many GPU frames from then on. A frame whose render panics is
-rendered again on the CPU or released, never left in progress. On the web
+rendered again on the CPU or released, never left in progress. Native readbacks and timing waits
+share a finite deadline; failure retires pending readbacks exactly once, and late results are
+discarded. The desktop executable and browser page own their shared device's error/loss handlers and forward
+failures to the compositor through a weak notifier. Libraries borrowing a device preserve its
+host's handlers. A host fault disables shared-device preview acceleration without editing the project or its
+renderer preference; materialized CPU frames remain usable and late GPU textures cannot return
+to the preview cache. Scoped allocation failures retain the adaptive budget/CPU retry behavior.
+
+**Recovery boundary:** the host installs handlers through eframe's CreationContext, after egui's
+presentation pipelines but before EffectCraft's pipelines. This does not cover that earlier egui
+startup window. eframe 0.36.2 supports lost-surface recreation, not replacing its live device and
+all presentation resources. Device loss therefore requires restarting the presentation backend;
+CPU compositing alone cannot restore that display. No automatic restart or persistent safe-mode
+sentinel is implemented here. Browser scope futures are not synchronously awaited; the browser page
+forwards shared-device errors, while owned worker devices install their own handlers. On the web
 (WebGPU) the GPU composites viewer frames; steps that need a readback fall back to the CPU (the
 Info panel's pixel readout reads GPU frames back asynchronously). **GPU particles**
 (`effects::psim`, `gpu::particles`): CC Particle World, CC Particle Systems II and Particle