gpu/
web_dev.rs

1// web_dev -- WebGPU in the browser: instance, adapter, device and queue, through web/gpu.js
2//
3//   Instance::new() -> Result<Instance>     Err(Disabled) with jfgui_GPU=0, Err(NoLoader) without
4//        navigator.gpu; starts requestAdapter and requestDevice, which answer later
5//   instance.prepare() -> Result<bool>      true: adapter and device are there (never waits; the
6//        worker has to return to the browser's event loop for it); Err(NoAdapter), Err(Failed)
7//   instance.adapter() / adapters() / adapters_low_power(display)   Err(NotReady) before that; one
8//        adapter (the browser chooses, "low-power" asked); a fallback adapter (software) is Cpu and
9//        not among adapters_low_power
10//   adapter.create_device(&DeviceDesc)      the device that was prepared, level Basis; once per
11//        instance: a new device takes a new Instance (a WebGPU adapter gives one)
12//   device.queue() -> &Queue                queue.submit([cmd, ..]) -> Result<SubmitId>
13//   queue.write_buffer, write_texture, write_texture_rects, write_texture_rects_at   the browser's
14//        queue.writeBuffer / writeTexture: copied at the call, before the next submission, from
15//        the WASM memory itself (no staging ring here)
16//   device.is_done(id), completed(), poll()  never wait
17//
18// What the API promises and the browser cannot do, said as errors and not faked: the worker cannot
19// wait for the GPU (the answer needs its event loop), so wait(id, _) answers at once (false: not done
20// yet), wait_idle and queue.read_buffer are Unsupported, Buffer::read says NotReady until the
21// submission has finished.  SPIR-V is Unsupported (ShaderModule::from_wgsl).  Errors the browser finds
22// (validation, memory) come later than the call that caused them, as its uncaptured errors, and are
23// returned by the next submit, try_acquire or poll.  Bindless, indirect count, multi-draw, timestamps
24// and subgroup sizes are not built here (the level Basis; Features says so).
25//
26// WebGPU's objects belong to this thread: the types of this backend hold Rc, so the compiler keeps
27// them from the compute workers of the threads bundle (docs/GPU.md "Ausgabe im Browser").
28
29use alloc::format;
30use alloc::rc::Rc;
31use alloc::string::String;
32use alloc::vec;
33use alloc::vec::Vec;
34use core::cell::{Cell, RefCell};
35
36use crate::*;
37
38#[link(wasm_import_module = "jfgui")]
39unsafe extern "C" {
40    #[link_name = "gpu_call"]
41    fn host_call(op: u32, args: *const u32, n: usize) -> u32;
42}
43
44// the calls of web/gpu.js
45pub(crate) const CALL_PREPARE: u32 = 0;
46pub(crate) const CALL_INFO: u32 = 1;
47pub(crate) const CALL_DEVICE: u32 = 2;
48pub(crate) const CALL_FREE: u32 = 3;
49pub(crate) const CALL_CREATE: u32 = 4;
50pub(crate) const CALL_SHADER: u32 = 5;
51pub(crate) const CALL_WRITE_BUFFER: u32 = 6;
52pub(crate) const CALL_READ_BUFFER: u32 = 7;
53pub(crate) const CALL_WRITE_TEXTURE: u32 = 8;
54pub(crate) const CALL_WRITE_RECTS: u32 = 9;
55pub(crate) const CALL_CONTEXT: u32 = 10;
56pub(crate) const CALL_CURRENT: u32 = 11;
57pub(crate) const CALL_RESIZE: u32 = 12;
58pub(crate) const CALL_SUBMIT: u32 = 13;
59pub(crate) const CALL_STATE: u32 = 14;
60pub(crate) const CALL_DONE: u32 = 16;
61pub(crate) const CALL_LOSE: u32 = 17;
62pub(crate) const CALL_RESERVE: u32 = 18;
63pub(crate) const CALL_PIPELINE_LAYOUT: u32 = 19;
64pub(crate) const CALL_RGBA: u32 = 20;
65
66/// a call of web/gpu.js with its numbers
67pub(crate) fn call(op: u32, args: &[u32]) -> u32 {
68    // SAFETY: the host reads args during the call; addresses among them are live for it (their
69    // callers hold the slices), it keeps nothing.
70    unsafe { host_call(op, args.as_ptr(), args.len()) }
71}
72
73/// an object of web/gpu.js's table: freed when the last Rc goes; buffers and textures destroyed (what
74/// was submitted before finishes first: WebGPU's rule)
75pub(crate) struct Handle {
76    pub h: u32,
77    pub destroy: bool,
78}
79
80impl Drop for Handle {
81    fn drop(&mut self) {
82        call(CALL_FREE, &[self.h, self.destroy as u32]);
83    }
84}
85
86/// an object from a descriptor (JSON; {"$":n} names object n): `method` of the device, or 100, a view
87/// of texture `parent` -> its number
88pub(crate) fn create(parent: u32, method: u32, json: &str) -> Result<u32> {
89    match call(CALL_CREATE, &[parent, method, json.as_ptr() as u32, json.len() as u32]) {
90        0 => Err(GpuError::Failed(0)),
91        h => Ok(h),
92    }
93}
94
95const STATE_LOST: u32 = 1;
96const STATE_VALIDATION: u32 = 2;
97const STATE_MEMORY: u32 = 4;
98const STATE_INTERNAL: u32 = 8;
99
100pub(crate) struct DevInner {
101    pub h: u32,
102    pub info: AdapterInfo,
103    pub features: Features,
104    pub limits: Limits,
105    /// a test said so (the browser's answer comes later)
106    lost: Cell<bool>,
107    last: Cell<u64>,
108    done: Cell<u64>,
109    /// what the buffers and textures asked for (the browser tells nothing)
110    pub reserved: Cell<u64>,
111    /// the words of finished command buffers are thrown away: this keeps one allocation
112    pub spare: RefCell<Vec<u32>>,
113}
114
115impl DevInner {
116    /// the browser's word on this device since the last look: lost, or an error of an earlier call
117    pub fn check(&self) -> Result<()> {
118        if self.lost.get() {
119            return Err(GpuError::DeviceLost);
120        }
121        let s = call(CALL_STATE, &[self.h]);
122        if s & STATE_LOST != 0 {
123            self.lost.set(true);
124            return Err(GpuError::DeviceLost);
125        }
126        if s & STATE_MEMORY != 0 {
127            return Err(GpuError::OutOfMemory);
128        }
129        if s & STATE_VALIDATION != 0 {
130            return Err(GpuError::Invalid("the browser's WebGPU refused a call (the console says which)"));
131        }
132        if s & STATE_INTERNAL != 0 {
133            return Err(GpuError::Failed(1));
134        }
135        Ok(())
136    }
137}
138
139/// the browser's WebGPU; adapter and device come asynchronously (prepare)
140#[derive(Clone)]
141pub struct Instance {
142    adapter: Rc<RefCell<Option<Adapter>>>,
143}
144
145/// the names of the Features WebGPU has, and what they stand for here
146fn features_of(list: &str) -> Features {
147    let has = |n: &str| list.split(',').any(|f| f == n);
148    Features {
149        draw_indirect_first_instance: has("indirect-first-instance"),
150        subgroups_webgpu: has("subgroups"),
151        shader_f16: has("shader-f16"),
152        texture_bc: has("texture-compression-bc"),
153        anisotropy: true,
154        depth_clamp: has("depth-clip-control"),
155        fragment_stores: true,
156        clip_distances: has("clip-distances"),
157        primitive_index: has("primitive-index"),
158        swapchain: true,
159        ..Features::default()
160    }
161}
162
163fn vendor_name(v: &str) -> &'static str {
164    match v {
165        "nvidia" => "NVIDIA",
166        "amd" | "ati" => "AMD",
167        "intel" => "Intel",
168        "apple" => "Apple",
169        "arm" => "ARM",
170        "qualcomm" => "Qualcomm",
171        "imagination" | "img-tec" => "Imagination",
172        "samsung" => "Samsung",
173        "google" => "Google",
174        "microsoft" => "Microsoft",
175        "mesa" => "Mesa",
176        _ => "",
177    }
178}
179
180fn read_adapter() -> Option<Adapter> {
181    let mut buf = vec![0u8; 8192];
182    let n = call(CALL_INFO, &[buf.as_mut_ptr() as u32, buf.len() as u32]) as usize;
183    let text = core::str::from_utf8(&buf[..n]).ok()?;
184    if text.is_empty() {
185        return None;
186    }
187    let get = |k: &str| text.lines().find_map(|l| l.strip_prefix(k).and_then(|r| r.strip_prefix('='))).unwrap_or("");
188    let num = |k: &str| get(k).parse::<f64>().map_or(0, |v| v.min(u32::MAX as f64) as u32);
189    let (vendor, arch, device, desc) = (get("vendor"), get("arch"), get("device"), get("desc"));
190    let name = if !desc.is_empty() {
191        String::from(desc)
192    } else {
193        let mut s = String::from(vendor);
194        for p in [arch, device] {
195            if !p.is_empty() {
196                s.push(' ');
197                s.push_str(p);
198            }
199        }
200        s
201    };
202    let fallback = get("fallback") == "1";
203    let info = AdapterInfo {
204        name,
205        vendor_id: 0,
206        device_id: 0,
207        vendor: vendor_name(vendor),
208        // (the browser does not say integrated or discrete)
209        kind: if fallback { AdapterKind::Cpu } else { AdapterKind::Other },
210        driver: String::from(vendor),
211        driver_info: String::from(arch),
212        driver_version: String::new(),
213        api_version: (1, 0, 0),
214        backend: "WebGPU",
215        missing: None,
216    };
217    let limits = Limits {
218        max_texture_2d: num("maxTextureDimension2D"),
219        max_texture_3d: num("maxTextureDimension3D"),
220        max_texture_layers: num("maxTextureArrayLayers"),
221        max_storage_binding: num("maxStorageBufferBindingSize"),
222        max_uniform_binding: num("maxUniformBufferBindingSize"),
223        max_push_constants: PUSH_CONSTANTS_MAX,
224        // (with push constants the last of them is theirs: create_pipeline_layout)
225        max_bind_groups: num("maxBindGroups"),
226        max_compute_workgroup_size: [num("maxComputeWorkgroupSizeX"), num("maxComputeWorkgroupSizeY"), num("maxComputeWorkgroupSizeZ")],
227        max_compute_invocations: num("maxComputeInvocationsPerWorkgroup"),
228        max_compute_workgroups: [num("maxComputeWorkgroupsPerDimension"); 3],
229        max_compute_shared_memory: num("maxComputeWorkgroupStorageSize"),
230        min_uniform_offset_alignment: num("minUniformBufferOffsetAlignment"),
231        min_storage_offset_alignment: num("minStorageBufferOffsetAlignment"),
232        max_color_targets: num("maxColorAttachments"),
233        max_vertex_buffers: num("maxVertexBuffers"),
234        max_vertex_attributes: num("maxVertexAttributes"),
235        max_draw_indirect_count: 1,
236        max_anisotropy: 16.0,
237        timestamp_period_ns: 0.0,
238        subgroup_size: 0,
239        subgroup_min_size: num("minSubgroupSize"),
240        subgroup_max_size: num("maxSubgroupSize"),
241        ..Limits::default()
242    };
243    Some(Adapter(Rc::new(AdapterInner { info, features: features_of(get("features")), limits, taken: Cell::new(false) })))
244}
245
246impl Instance {
247    /// the browser's WebGPU, adapter and device asked for (prepare says when they are there)
248    pub fn new() -> Result<Instance> {
249        Instance::with(&InstanceDesc::DEFAULT)
250    }
251
252    /// (app_name and validate mean nothing here: the browser validates every call)
253    pub fn with(_desc: &InstanceDesc) -> Result<Instance> {
254        if matches!(plat::env("jfgui_GPU").as_deref(), Some("0") | Some("off")) {
255            return Err(GpuError::Disabled);
256        }
257        match call(CALL_PREPARE, &[1]) {
258            0 => Err(GpuError::NoLoader),
259            _ => Ok(Instance { adapter: Rc::new(RefCell::new(None)) }),
260        }
261    }
262
263    /// have the browser's adapter and device arrived?  Never waits; a worker gets them by
264    /// returning to the browser's event loop.  (true at once on the other targets.)
265    pub fn prepare(&self) -> Result<bool> {
266        if self.adapter.borrow().is_some() {
267            return Ok(true);
268        }
269        match call(CALL_PREPARE, &[0]) {
270            0 => Err(GpuError::NoLoader),
271            1 => Ok(false),
272            2 => {
273                let a = read_adapter().ok_or(GpuError::NoAdapter)?;
274                *self.adapter.borrow_mut() = Some(a);
275                Ok(true)
276            }
277            3 => Err(GpuError::NoAdapter),
278            _ => Err(GpuError::Failed(0)),
279        }
280    }
281
282    pub fn validating(&self) -> bool {
283        false
284    }
285
286    /// the one adapter the browser offers
287    pub fn adapters(&self) -> Result<Vec<Adapter>> {
288        Ok(vec![self.adapter()?])
289    }
290
291    pub fn adapter(&self) -> Result<Adapter> {
292        self.prepare()?;
293        self.adapter.borrow().clone().ok_or(GpuError::NotReady)
294    }
295
296    /// the adapter for showing windows: not a software one, which is no gain over the CPU path
297    pub fn adapters_low_power(&self, _display: Option<(u32, u32)>) -> Result<Vec<Adapter>> {
298        let a = self.adapter()?;
299        Ok(if a.info().kind == AdapterKind::Cpu { Vec::new() } else { vec![a] })
300    }
301}
302
303/// the order of adapters for showing windows (the browser has one; see adapters_low_power)
304pub fn low_power_rank(kind: AdapterKind, _drives_display: bool) -> Option<u32> {
305    (kind != AdapterKind::Cpu).then_some(0)
306}
307
308/// nothing counts them here: the browser's console shows what its validation says
309pub fn validation_messages() -> (u32, u32) {
310    (0, 0)
311}
312
313struct AdapterInner {
314    info: AdapterInfo,
315    features: Features,
316    limits: Limits,
317    taken: Cell<bool>,
318}
319
320/// the GPU the browser offers
321#[derive(Clone)]
322pub struct Adapter(Rc<AdapterInner>);
323
324impl Adapter {
325    pub fn info(&self) -> &AdapterInfo {
326        &self.0.info
327    }
328    pub fn features(&self) -> &Features {
329        &self.0.features
330    }
331    pub fn limits(&self) -> &Limits {
332        &self.0.limits
333    }
334    /// always Basis in the browser
335    pub fn level(&self) -> Level {
336        Level::Basis
337    }
338    pub fn can_present(&self, _surface: &crate::Surface) -> bool {
339        true
340    }
341
342    /// the device of this adapter, which prepare got (WebGPU gives one per adapter: a second one
343    /// takes a new Instance)
344    pub fn create_device(&self, _desc: &DeviceDesc) -> Result<Device> {
345        if self.0.taken.replace(true) {
346            return Err(GpuError::Unsupported("one device per Instance in the browser"));
347        }
348        let h = call(CALL_DEVICE, &[]);
349        if h == 0 {
350            return Err(GpuError::NotReady);
351        }
352        let d = Rc::new(DevInner {
353            h,
354            info: self.0.info.clone(),
355            features: self.0.features,
356            limits: self.0.limits,
357            lost: Cell::new(false),
358            last: Cell::new(0),
359            done: Cell::new(0),
360            reserved: Cell::new(0),
361            spare: RefCell::new(Vec::new()),
362        });
363        Ok(Device { d: d.clone(), queue: Queue { d } })
364    }
365}
366
367/// a logical GPU device: creates resources, records commands, owns the queue
368#[derive(Clone)]
369pub struct Device {
370    pub(crate) d: Rc<DevInner>,
371    queue: Queue,
372}
373
374/// the device's queue
375#[derive(Clone)]
376pub struct Queue {
377    pub(crate) d: Rc<DevInner>,
378}
379
380impl Device {
381    pub fn queue(&self) -> &Queue {
382        &self.queue
383    }
384    pub fn info(&self) -> &AdapterInfo {
385        &self.d.info
386    }
387    pub fn features(&self) -> &Features {
388        &self.d.features
389    }
390    pub fn limits(&self) -> &Limits {
391        &self.d.limits
392    }
393    pub fn level(&self) -> Level {
394        Level::Basis
395    }
396    /// the device was lost: every call fails with DeviceLost; create a new device
397    pub fn is_lost(&self) -> bool {
398        matches!(self.d.check(), Err(GpuError::DeviceLost))
399    }
400    /// (test aid) the device is destroyed: the browser's own way of losing it
401    pub fn test_lose(&self) {
402        call(CALL_LOSE, &[self.d.h]);
403        self.d.lost.set(true);
404    }
405    pub fn validating(&self) -> bool {
406        false
407    }
408    /// what the browser has said since the last look (nothing to free here)
409    pub fn poll(&self) -> Result<()> {
410        self.d.check()
411    }
412    /// the browser cannot wait on this thread: at once whether the submission is done
413    pub fn wait(&self, id: SubmitId, _timeout_ns: u64) -> Result<bool> {
414        self.d.check()?;
415        Ok(self.is_done(id))
416    }
417    pub fn wait_idle(&self) -> Result<()> {
418        Err(GpuError::Unsupported("waiting for the GPU in the browser: Device::is_done"))
419    }
420    pub fn completed(&self) -> SubmitId {
421        SubmitId(self.d.done.get())
422    }
423    /// the submission id is done on the GPU (the browser's answer comes with its event loop)
424    pub fn is_done(&self, id: SubmitId) -> bool {
425        if self.d.done.get() >= id.0 {
426            return true;
427        }
428        if call(CALL_DONE, &[self.d.h, id.0 as u32]) != 0 {
429            self.d.done.set(id.0);
430            return true;
431        }
432        false
433    }
434    /// what the buffers and textures asked for (the browser says nothing about its own memory)
435    pub fn memory_reserved(&self) -> u64 {
436        self.d.reserved.get()
437    }
438}
439
440fn split(v: u64) -> (u32, u32) {
441    (v as u32, (v >> 32) as u32)
442}
443
444impl Queue {
445    /// run command buffers in order -> their number
446    pub fn submit<'a, I: IntoIterator<Item = crate::CommandBuffer<'a>>>(&self, cmds: I) -> Result<SubmitId> {
447        self.d.check()?;
448        let mut words: Vec<u32> = core::mem::take(&mut *self.d.spare.borrow_mut());
449        words.clear();
450        for c in cmds {
451            words.extend_from_slice(&c.words);
452            words.push(0);
453        }
454        let id = call(CALL_SUBMIT, &[self.d.h, words.as_ptr() as u32, words.len() as u32]);
455        *self.d.spare.borrow_mut() = words;
456        self.d.last.set(id as u64);
457        Ok(SubmitId(id as u64))
458    }
459
460    /// nothing waits for a submission here
461    pub fn flush(&self) -> Result<SubmitId> {
462        self.d.check()?;
463        Ok(SubmitId(self.d.last.get()))
464    }
465
466    /// write data into a buffer at offset, before the next submission
467    pub fn write_buffer<T: Pod>(&self, buf: &crate::Buffer, offset: u64, data: &[T]) -> Result<()> {
468        let bytes = bytes_of(data);
469        if offset % 4 != 0 || bytes.len() % 4 != 0 {
470            return Err(GpuError::Invalid("write_buffer: offset and size must be multiples of 4 in WebGPU"));
471        }
472        if offset.checked_add(bytes.len() as u64).is_none_or(|e| e > buf.size()) {
473            return Err(GpuError::Invalid("write_buffer: outside the buffer"));
474        }
475        let (lo, hi) = split(offset);
476        call(CALL_WRITE_BUFFER, &[self.d.h, buf.handle(), lo, hi, bytes.as_ptr() as u32, bytes.len() as u32]);
477        Ok(())
478    }
479
480    /// write texels into a texture region, before the next submission (no layouts here)
481    pub fn write_texture(&self, dst: &TextureRegion, data: &[u8], layout: DataLayout, _before: Layout, _after: Layout) -> Result<()> {
482        let f = dst.texture.format();
483        let row = layout.bytes_per_row.max(((dst.size[0] / f.block_dim()).max(1)) * f.block_bytes());
484        let rows = if layout.rows_per_image != 0 { layout.rows_per_image } else { dst.size[1] / f.block_dim() };
485        let z = if dst.texture.depth() > 1 { dst.origin[2] } else { dst.base_layer };
486        let d = if dst.texture.depth() > 1 { dst.size[2].max(1) } else { dst.layers.max(1) };
487        if layout.offset > data.len() as u64 {
488            return Err(GpuError::Invalid("write_texture: the offset lies outside the data"));
489        }
490        let base = data.as_ptr() as u32 + layout.offset as u32;
491        call(CALL_WRITE_TEXTURE, &[self.d.h, dst.texture.handle(), dst.mip, dst.origin[0], dst.origin[1], z, dst.size[0], dst.size[1], d, base, data.len() as u32, row, rows]);
492        Ok(())
493    }
494
495    /// write rectangles (x, y, w, h; clipped to the texture, empty ones skipped) of mip 0, layer 0 from
496    /// `data`, an image of the texture's size in rows of `stride` bytes (a framebuffer)
497    pub fn write_texture_rects(&self, t: &crate::Texture, rects: &[[i32; 4]], data: &[u8], stride: u32, before: Layout, after: Layout) -> Result<()> {
498        self.write_texture_rects_at(t, rects, (0, 0), data, stride, before, after)
499    }
500
501    /// write_texture_rects from an image of another size: rectangles of `data` go to (x - origin.0,
502    /// y - origin.1) of the texture; what lies outside the image or the texture is skipped
503    pub fn write_texture_rects_at(&self, t: &crate::Texture, rects: &[[i32; 4]], origin: (i32, i32), data: &[u8], stride: u32, _before: Layout, _after: Layout) -> Result<()> {
504        let (tw, th) = (t.width() as i32, t.height() as i32);
505        let bytes = t.format().block_bytes() as i32;
506        if stride == 0 || bytes != 4 {
507            return Err(GpuError::Invalid("write_texture_rects: a 32-bit format and a stride"));
508        }
509        let (img_w, img_h) = (stride as i32 / 4, (data.len() / stride as usize) as i32);
510        let mut list: Vec<[i32; 4]> = Vec::with_capacity(rects.len());
511        for r in rects {
512            let x0 = r[0].max(origin.0).max(0);
513            let y0 = r[1].max(origin.1).max(0);
514            let x1 = r[0].saturating_add(r[2].max(0)).min(origin.0.saturating_add(tw)).min(img_w);
515            let y1 = r[1].saturating_add(r[3].max(0)).min(origin.1.saturating_add(th)).min(img_h);
516            if x0 < x1 && y0 < y1 {
517                list.push([x0, y0, x1 - x0, y1 - y0]);
518            }
519        }
520        if !list.is_empty() {
521            call(CALL_WRITE_RECTS, &[self.d.h, t.handle(), list.as_ptr() as u32, list.len() as u32, data.as_ptr() as u32, data.len() as u32, stride, origin.0 as u32, origin.1 as u32]);
522        }
523        Ok(())
524    }
525
526    /// the worker cannot wait for the GPU: a Readback buffer and Buffer::read
527    pub fn read_buffer<T: Pod>(&self, _buf: &crate::Buffer, _offset: u64, _out: &mut [T]) -> Result<()> {
528        Err(GpuError::Unsupported("Queue::read_buffer waits for the GPU; in the browser a Readback buffer and Buffer::read"))
529    }
530}
531
532/// a number of a JSON text
533pub(crate) fn json_f64(v: f64) -> String {
534    if v.is_finite() { format!("{v}") } else { String::from("0") }
535}
gpu/web_dev.rs

gpu/
web_dev.rs
