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    /// (G5d) as create_device: WebGPU has one queue, so a compute queue asked for stays None
343    pub fn create_device_with(&self, desc: &DeviceDesc, _opts: &DeviceOptions) -> Result<Device> {
344        self.create_device(desc)
345    }
346
347    /// the device of this adapter, which prepare got (WebGPU gives one per adapter: a second one
348    /// takes a new Instance)
349    pub fn create_device(&self, _desc: &DeviceDesc) -> Result<Device> {
350        if self.0.taken.replace(true) {
351            return Err(GpuError::Unsupported("one device per Instance in the browser"));
352        }
353        let h = call(CALL_DEVICE, &[]);
354        if h == 0 {
355            return Err(GpuError::NotReady);
356        }
357        let d = Rc::new(DevInner {
358            h,
359            info: self.0.info.clone(),
360            features: self.0.features,
361            limits: self.0.limits,
362            lost: Cell::new(false),
363            last: Cell::new(0),
364            done: Cell::new(0),
365            reserved: Cell::new(0),
366            spare: RefCell::new(Vec::new()),
367        });
368        Ok(Device { d: d.clone(), queue: Queue { d } })
369    }
370}
371
372/// a logical GPU device: creates resources, records commands, owns the queue
373#[derive(Clone)]
374pub struct Device {
375    pub(crate) d: Rc<DevInner>,
376    queue: Queue,
377}
378
379/// the device's queue
380#[derive(Clone)]
381pub struct Queue {
382    pub(crate) d: Rc<DevInner>,
383}
384
385/// (G5d) a queue of its own for compute: WebGPU has one queue, so Device::compute_queue is None here and there is no
386/// value of this type
387#[derive(Clone)]
388pub struct ComputeQueue {
389    _none: (),
390}
391
392impl ComputeQueue {
393    pub fn encoder(&self) -> Result<crate::CommandEncoder<'_>> {
394        Err(GpuError::Unsupported("a compute queue in the browser"))
395    }
396    pub fn submit<'a, I: IntoIterator<Item = crate::CommandBuffer<'a>>>(&self, _cmds: I) -> Result<SubmitId> {
397        Err(GpuError::Unsupported("a compute queue in the browser"))
398    }
399    pub fn submit_after<'a, I: IntoIterator<Item = crate::CommandBuffer<'a>>>(&self, _wait: &[SubmitId], _cmds: I) -> Result<SubmitId> {
400        Err(GpuError::Unsupported("a compute queue in the browser"))
401    }
402}
403
404impl Device {
405    pub fn queue(&self) -> &Queue {
406        &self.queue
407    }
408    /// (G5d) WebGPU has one queue: None
409    pub fn compute_queue(&self) -> Option<&ComputeQueue> {
410        None
411    }
412    pub fn info(&self) -> &AdapterInfo {
413        &self.d.info
414    }
415    pub fn features(&self) -> &Features {
416        &self.d.features
417    }
418    pub fn limits(&self) -> &Limits {
419        &self.d.limits
420    }
421    pub fn level(&self) -> Level {
422        Level::Basis
423    }
424    /// the device was lost: every call fails with DeviceLost; create a new device
425    pub fn is_lost(&self) -> bool {
426        matches!(self.d.check(), Err(GpuError::DeviceLost))
427    }
428    /// (test aid) the device is destroyed: the browser's own way of losing it
429    pub fn test_lose(&self) {
430        call(CALL_LOSE, &[self.d.h]);
431        self.d.lost.set(true);
432    }
433    pub fn validating(&self) -> bool {
434        false
435    }
436    /// what the browser has said since the last look (nothing to free here)
437    pub fn poll(&self) -> Result<()> {
438        self.d.check()
439    }
440    /// the browser cannot wait on this thread: at once whether the submission is done
441    pub fn wait(&self, id: SubmitId, _timeout_ns: u64) -> Result<bool> {
442        self.d.check()?;
443        Ok(self.is_done(id))
444    }
445    pub fn wait_idle(&self) -> Result<()> {
446        Err(GpuError::Unsupported("waiting for the GPU in the browser: Device::is_done"))
447    }
448    pub fn completed(&self) -> SubmitId {
449        SubmitId(self.d.done.get())
450    }
451    /// the submission id is done on the GPU (the browser's answer comes with its event loop)
452    pub fn is_done(&self, id: SubmitId) -> bool {
453        if self.d.done.get() >= id.0 {
454            return true;
455        }
456        if call(CALL_DONE, &[self.d.h, id.0 as u32]) != 0 {
457            self.d.done.set(id.0);
458            return true;
459        }
460        false
461    }
462    /// what the buffers and textures asked for (the browser says nothing about its own memory)
463    pub fn memory_reserved(&self) -> u64 {
464        self.d.reserved.get()
465    }
466}
467
468fn split(v: u64) -> (u32, u32) {
469    (v as u32, (v >> 32) as u32)
470}
471
472impl Queue {
473    /// run command buffers in order -> their number
474    pub fn submit<'a, I: IntoIterator<Item = crate::CommandBuffer<'a>>>(&self, cmds: I) -> Result<SubmitId> {
475        self.d.check()?;
476        let mut words: Vec<u32> = core::mem::take(&mut *self.d.spare.borrow_mut());
477        words.clear();
478        for c in cmds {
479            words.extend_from_slice(&c.words);
480            words.push(0);
481        }
482        let id = call(CALL_SUBMIT, &[self.d.h, words.as_ptr() as u32, words.len() as u32]);
483        *self.d.spare.borrow_mut() = words;
484        self.d.last.set(id as u64);
485        Ok(SubmitId(id as u64))
486    }
487
488    /// (G5d) one queue: its submissions run in order, so this is submit
489    pub fn submit_after<'a, I: IntoIterator<Item = crate::CommandBuffer<'a>>>(&self, _wait: &[SubmitId], cmds: I) -> Result<SubmitId> {
490        self.submit(cmds)
491    }
492
493    /// nothing waits for a submission here
494    pub fn flush(&self) -> Result<SubmitId> {
495        self.d.check()?;
496        Ok(SubmitId(self.d.last.get()))
497    }
498
499    /// write data into a buffer at offset, before the next submission
500    pub fn write_buffer<T: Pod>(&self, buf: &crate::Buffer, offset: u64, data: &[T]) -> Result<()> {
501        let bytes = bytes_of(data);
502        if offset % 4 != 0 || bytes.len() % 4 != 0 {
503            return Err(GpuError::Invalid("write_buffer: offset and size must be multiples of 4 in WebGPU"));
504        }
505        if offset.checked_add(bytes.len() as u64).is_none_or(|e| e > buf.size()) {
506            return Err(GpuError::Invalid("write_buffer: outside the buffer"));
507        }
508        let (lo, hi) = split(offset);
509        call(CALL_WRITE_BUFFER, &[self.d.h, buf.handle(), lo, hi, bytes.as_ptr() as u32, bytes.len() as u32]);
510        Ok(())
511    }
512
513    /// write texels into a texture region, before the next submission (no layouts here)
514    pub fn write_texture(&self, dst: &TextureRegion, data: &[u8], layout: DataLayout, _before: Layout, _after: Layout) -> Result<()> {
515        let f = dst.texture.format();
516        let row = layout.bytes_per_row.max(((dst.size[0] / f.block_dim()).max(1)) * f.block_bytes());
517        let rows = if layout.rows_per_image != 0 { layout.rows_per_image } else { dst.size[1] / f.block_dim() };
518        let z = if dst.texture.depth() > 1 { dst.origin[2] } else { dst.base_layer };
519        let d = if dst.texture.depth() > 1 { dst.size[2].max(1) } else { dst.layers.max(1) };
520        if layout.offset > data.len() as u64 {
521            return Err(GpuError::Invalid("write_texture: the offset lies outside the data"));
522        }
523        let base = data.as_ptr() as u32 + layout.offset as u32;
524        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]);
525        Ok(())
526    }
527
528    /// write rectangles (x, y, w, h; clipped to the texture, empty ones skipped) of mip 0, layer 0 from
529    /// `data`, an image of the texture's size in rows of `stride` bytes (a framebuffer)
530    pub fn write_texture_rects(&self, t: &crate::Texture, rects: &[[i32; 4]], data: &[u8], stride: u32, before: Layout, after: Layout) -> Result<()> {
531        self.write_texture_rects_at(t, rects, (0, 0), data, stride, before, after)
532    }
533
534    /// write_texture_rects from an image of another size: rectangles of `data` go to (x - origin.0,
535    /// y - origin.1) of the texture; what lies outside the image or the texture is skipped
536    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<()> {
537        let (tw, th) = (t.width() as i32, t.height() as i32);
538        let bytes = t.format().block_bytes() as i32;
539        if stride == 0 || bytes != 4 {
540            return Err(GpuError::Invalid("write_texture_rects: a 32-bit format and a stride"));
541        }
542        let (img_w, img_h) = (stride as i32 / 4, (data.len() / stride as usize) as i32);
543        let mut list: Vec<[i32; 4]> = Vec::with_capacity(rects.len());
544        for r in rects {
545            let x0 = r[0].max(origin.0).max(0);
546            let y0 = r[1].max(origin.1).max(0);
547            let x1 = r[0].saturating_add(r[2].max(0)).min(origin.0.saturating_add(tw)).min(img_w);
548            let y1 = r[1].saturating_add(r[3].max(0)).min(origin.1.saturating_add(th)).min(img_h);
549            if x0 < x1 && y0 < y1 {
550                list.push([x0, y0, x1 - x0, y1 - y0]);
551            }
552        }
553        if !list.is_empty() {
554            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]);
555        }
556        Ok(())
557    }
558
559    /// the worker cannot wait for the GPU: a Readback buffer and Buffer::read
560    pub fn read_buffer<T: Pod>(&self, _buf: &crate::Buffer, _offset: u64, _out: &mut [T]) -> Result<()> {
561        Err(GpuError::Unsupported("Queue::read_buffer waits for the GPU; in the browser a Readback buffer and Buffer::read"))
562    }
563}
564
565/// a number of a JSON text
566pub(crate) fn json_f64(v: f64) -> String {
567    if v.is_finite() { format!("{v}") } else { String::from("0") }
568}
gpu/web_dev.rs

gpu/
web_dev.rs
