Unchecked shaders cause VK_ERROR_DEVICE_LOST on Pixel 9 Pro
## Summary
Rendering **any** scene containing at least one path loses the Vulkan device on
the first frame on a Pixel 9 Pro (Tensor G4, ARM Immortalis-G715 MC7, stock
Android 16). Bisection shows the cause is an out-of-bounds access in one of the
GPU geometry stages, which desktop drivers tolerate but Mali's MMU faults on —
and it only faults because vello disables shader bounds checks:
`create_compute_pipeline` builds every pipeline with
`create_shader_module_trusted(..., ShaderRuntimeChecks::unchecked())`
(`vello/src/wgpu_engine.rs`).
**Changing that one call to `ShaderRuntimeChecks::checked()` makes every case
pass on this device.** This looks like the same class as #1465 (Mali-TKRX page
fault, also fixed by enabling bounds checks), reproduced here on current Google
flagship hardware with a stock OS.
## Environment
- Device: Google Pixel 9 Pro (caiman), Tensor G4, ARM Immortalis-G715 MC7
- Driver (via wgpu): `Mali-G715 v1.r54p2-00eac0.7001668a7a05889fc53765c682b36a3f`
- OS: stock Android 16, `google/caiman/caiman:16/CP1A.260305.018` (release keys)
- vello 0.9.0 (crates.io), wgpu 29.0.4, Vulkan backend
- rustc 1.99.0-nightly (9f36de775 2026-07-19), `aarch64-linux-android`, NDK r29
- Headless standalone binary run from `adb shell` (also reproduces inside an
app rendering through a SurfaceView)
## Reproduction
```rust
use std::sync::atomic::{AtomicBool, Ordering};
use std::sync::Arc;
use vello::kurbo::{Affine, Circle, Line, Stroke};
use vello::peniko::Color;
static DEVICE_LOST: AtomicBool = AtomicBool::new(false);
fn main() {
let mut descriptor = wgpu::InstanceDescriptor::new_without_display_handle();
descriptor.backends = wgpu::Backends::VULKAN;
let instance = wgpu::Instance::new(descriptor);
let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
power_preference: wgpu::PowerPreference::HighPerformance,
compatible_surface: None,
force_fallback_adapter: false,
}))
.expect("no Vulkan adapter");
let info = adapter.get_info();
println!(
"adapter: {} ({:?}, driver {} {})",
info.name, info.device_type, info.driver, info.driver_info
);
let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
label: Some("repro device"),
required_features: wgpu::Features::empty(),
required_limits: wgpu::Limits::default(),
memory_hints: wgpu::MemoryHints::Performance,
experimental_features: wgpu::ExperimentalFeatures::default(),
trace: wgpu::Trace::default(),
}))
.expect("request_device failed");
let device = Arc::new(device);
device.set_device_lost_callback(|reason, message| {
DEVICE_LOST.store(true, Ordering::SeqCst);
eprintln!("DEVICE LOST: {reason:?}: {message}");
});
// Control: a trivial compute dispatch, to show plain compute is healthy on
// this device before vello runs.
run_trivial_compute(&device, &queue);
if DEVICE_LOST.load(Ordering::SeqCst) {
eprintln!("device lost during the trivial compute control");
std::process::exit(2);
}
println!("trivial compute dispatch: OK");
let use_cpu = std::env::var("REPRO_CPU").is_ok();
println!("use_cpu: {use_cpu}");
let mut renderer = vello::Renderer::new(
&device,
vello::RendererOptions {
use_cpu,
antialiasing_support: vello::AaSupport::all(),
num_init_threads: Some(std::num::NonZeroUsize::new(1).unwrap()),
pipeline_cache: None,
},
)
.expect("vello renderer creation failed");
let width = 54u32;
let height = 54u32;
let texture = device.create_texture(&wgpu::TextureDescriptor {
label: Some("repro target"),
size: wgpu::Extent3d {
width,
height,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba8Unorm,
usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
view_formats: &[],
});
let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
// Scene content selected by REPRO_SCENE: empty | fill | line | icon (default).
let variant = std::env::var("REPRO_SCENE").unwrap_or_else(|_| "icon".into());
println!("scene variant: {variant}");
let mut scene = vello::Scene::new();
let stroke = Stroke::new(2.0);
let fg = Color::from_rgb8(0x20, 0x20, 0x20);
match variant.as_str() {
"empty" => {}
"fill" => {
scene.fill(
vello::peniko::Fill::NonZero,
Affine::IDENTITY,
fg,
None,
&Circle::new((27.0, 27.0), 18.0),
);
}
"line" => {
scene.stroke(
&stroke,
Affine::IDENTITY,
fg,
None,
&Line::new((27.0, 4.0), (27.0, 14.0)),
);
}
_ => {
scene.stroke(
&stroke,
Affine::IDENTITY,
fg,
None,
&Circle::new((27.0, 27.0), 18.0),
);
scene.stroke(
&stroke,
Affine::IDENTITY,
fg,
None,
&Line::new((27.0, 4.0), (27.0, 14.0)),
);
scene.stroke(
&stroke,
Affine::IDENTITY,
fg,
None,
&Line::new((4.0, 27.0), (14.0, 27.0)),
);
}
}
for frame in 0..10 {
renderer
.render_to_texture(
&device,
&queue,
&scene,
&view,
&vello::RenderParams {
base_color: Color::from_rgba8(0, 0, 0, 0),
width,
height,
antialiasing_method: match std::env::var("REPRO_AA").as_deref() {
Ok("msaa8") => vello::AaConfig::Msaa8,
Ok("msaa16") => vello::AaConfig::Msaa16,
_ => vello::AaConfig::Area,
},
},
)
.expect("vello render_to_texture failed");
device
.poll(wgpu::PollType::Wait {
submission_index: None,
timeout: None,
})
.expect("device poll failed");
if DEVICE_LOST.load(Ordering::SeqCst) {
eprintln!("device lost after frame {frame}");
std::process::exit(1);
}
println!("frame {frame}: OK");
}
println!("all frames rendered, device healthy");
}
fn run_trivial_compute(device: &wgpu::Device, queue: &wgpu::Queue) {
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("control"),
source: wgpu::ShaderSource::Wgsl(
"@group(0) @binding(0) var<storage, read_write> data: array<u32>;\n\
@compute @workgroup_size(64)\n\
fn main(@builtin(global_invocation_id) id: vec3<u32>) {\n\
data[id.x] = id.x * 2u;\n\
}"
.into(),
),
});
let buffer = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("control buffer"),
size: 256 * 4,
usage: wgpu::BufferUsages::STORAGE,
mapped_at_creation: false,
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("control pipeline"),
layout: None,
module: &shader,
entry_point: Some("main"),
compilation_options: wgpu::PipelineCompilationOptions::default(),
cache: None,
});
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("control bind group"),
layout: &pipeline.get_bind_group_layout(0),
entries: &[wgpu::BindGroupEntry {
binding: 0,
resource: buffer.as_entire_binding(),
}],
});
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
{
let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
pass.set_pipeline(&pipeline);
pass.set_bind_group(0, &bind_group, &[]);
pass.dispatch_workgroups(4, 1, 1);
}
queue.submit([encoder.finish()]);
device
.poll(wgpu::PollType::Wait {
submission_index: None,
timeout: None,
})
.expect("control compute poll failed");
}
```
关闭于 13 天前 0 条评论