From 312828f12f1a1497bc0387a72a5346ef911acad7 Mon Sep 17 00:00:00 2001 From: Zicklag Date: Thu, 7 Oct 2021 15:18:09 -0500 Subject: [PATCH] Implement WebGL2 Backend (#1686) * Implement WebGL Backend * Add WebGL Fixes by @mrk-its * Update Limits for WASM and Examples * Address Review Points --- .github/workflows/ci.yml | 3 + run-wasm-example.sh | 34 ++++ wasm-resources/README.md | 3 + wasm-resources/index.template.html | 14 ++ wgpu-core/Cargo.toml | 3 + wgpu-core/build.rs | 7 +- wgpu-core/src/device/mod.rs | 4 + wgpu-core/src/hub.rs | 1 + wgpu-core/src/lib.rs | 7 +- wgpu-hal/Cargo.toml | 5 + wgpu-hal/src/gles/adapter.rs | 38 +++- wgpu-hal/src/gles/command.rs | 11 +- wgpu-hal/src/gles/device.rs | 131 +++++++++---- wgpu-hal/src/gles/mod.rs | 37 +++- wgpu-hal/src/gles/queue.rs | 105 +++++++++-- wgpu-hal/src/gles/web.rs | 273 ++++++++++++++++++++++++++++ wgpu-hal/src/gles/web/present.frag | 16 ++ wgpu-hal/src/gles/web/present.vert | 18 ++ wgpu-hal/src/lib.rs | 16 +- wgpu-types/src/lib.rs | 62 +++++-- wgpu/Cargo.toml | 2 + wgpu/examples/boids/main.rs | 11 ++ wgpu/examples/capture/main.rs | 2 +- wgpu/examples/framework.rs | 57 +++++- wgpu/examples/hello-compute/main.rs | 2 +- wgpu/examples/hello-windows/main.rs | 2 +- wgpu/examples/hello/main.rs | 2 +- wgpu/examples/shadow/main.rs | 35 +++- wgpu/examples/shadow/shader.wgsl | 32 ++++ wgpu/src/backend/direct.rs | 22 ++- wgpu/src/lib.rs | 4 +- wgpu/tests/common/mod.rs | 14 +- 32 files changed, 853 insertions(+), 120 deletions(-) create mode 100755 run-wasm-example.sh create mode 100644 wasm-resources/README.md create mode 100644 wasm-resources/index.template.html create mode 100644 wgpu-hal/src/gles/web.rs create mode 100644 wgpu-hal/src/gles/web/present.frag create mode 100644 wgpu-hal/src/gles/web/present.vert diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 5e1b06f133..6980229314 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -130,6 +130,9 @@ jobs: run: | cargo clippy --target ${{ matrix.target }} -p wgpu + # Build for WebGL + cargo clippy --target ${{ matrix.target }} -p wgpu --features webgl -- -D warnings + # build docs cargo doc --target ${{ matrix.target }} -p wgpu --no-deps diff --git a/run-wasm-example.sh b/run-wasm-example.sh new file mode 100755 index 0000000000..4032616353 --- /dev/null +++ b/run-wasm-example.sh @@ -0,0 +1,34 @@ +#!/bin/env bash + +set -e + +echo "Compiling..." +cargo build --example $1 --target wasm32-unknown-unknown --features webgl + +echo "Generating bindings..." +mkdir -p target/wasm-examples/$1 +wasm-bindgen --target web --out-dir target/wasm-examples/$1 target/wasm32-unknown-unknown/debug/examples/$1.wasm +cp wasm-resources/index.template.html target/wasm-examples/$1/index.html +sed -i "s/{{example}}/$1/g" target/wasm-examples/$1/index.html + +# Find a serving tool to host the example +SERVE_CMD="" +SERVE_ARGS="" +if which basic-http-server; then + SERVE_CMD="basic-http-server" + SERVE_ARGS="target/wasm-examples/$1 -a 127.0.0.1:1234" +elif which miniserve && python3 -m http.server --help > /dev/null; then + SERVE_CMD="miniserve" + SERVE_ARGS="target/wasm-examples/$1 -p 1234 --index index.html" +elif python3 -m http.server --help > /dev/null; then + SERVE_CMD="python3" + SERVE_ARGS="-m http.server --directory target/wasm-examples/$1 1234" +fi + +# Exit if we couldn't find a tool to serve the example with +if [ "$SERVE_CMD" = "" ]; then + echo "Couldn't find a utility to use to serve the example web page. You can serve the `target/wasm-examples/$1` folder yourself using any simple static http file server." +fi + +echo "Serving example with $SERVE_CMD at http://localhost:1234" +$SERVE_CMD $SERVE_ARGS \ No newline at end of file diff --git a/wasm-resources/README.md b/wasm-resources/README.md new file mode 100644 index 0000000000..8b2b328fb7 --- /dev/null +++ b/wasm-resources/README.md @@ -0,0 +1,3 @@ +# WASM Resources + +This directory contains resources used when building the WGPU examples for web. diff --git a/wasm-resources/index.template.html b/wasm-resources/index.template.html new file mode 100644 index 0000000000..155072d10e --- /dev/null +++ b/wasm-resources/index.template.html @@ -0,0 +1,14 @@ + + + + + + + + + \ No newline at end of file diff --git a/wgpu-core/Cargo.toml b/wgpu-core/Cargo.toml index 19576346c2..c55fc60f2c 100644 --- a/wgpu-core/Cargo.toml +++ b/wgpu-core/Cargo.toml @@ -50,6 +50,9 @@ path = "../wgpu-hal" package = "wgpu-hal" version = "0.10.1" +[target.'cfg(target_arch = "wasm32")'.dependencies] +hal = { path = "../wgpu-hal", package = "wgpu-hal", version = "0.10", features = ["gles"] } + [target.'cfg(all(not(target_arch = "wasm32"), any(target_os = "ios", target_os = "macos")))'.dependencies] hal = { path = "../wgpu-hal", package = "wgpu-hal", version = "0.10", features = ["metal"] } #Note: could also enable "vulkan" for Vulkan Portability diff --git a/wgpu-core/build.rs b/wgpu-core/build.rs index 38dad2eac7..37cf80b794 100644 --- a/wgpu-core/build.rs +++ b/wgpu-core/build.rs @@ -11,6 +11,11 @@ fn main() { metal: { all(not(wasm), apple) }, dx12: { all(not(wasm), windows) }, dx11: { all(false, not(wasm), windows) }, - gl: { all(not(wasm), unix_wo_apple) }, + gl: { + any( + all(not(wasm), unix_wo_apple), + wasm + ) + }, } } diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs index cd32d8634c..cf3135e52e 100644 --- a/wgpu-core/src/device/mod.rs +++ b/wgpu-core/src/device/mod.rs @@ -4584,6 +4584,10 @@ impl Global { { self.poll_devices::(force_wait, &mut closures)?; } + #[cfg(gl)] + { + self.poll_devices::(force_wait, &mut closures)?; + } unsafe { closures.fire(); diff --git a/wgpu-core/src/hub.rs b/wgpu-core/src/hub.rs index d43c68a7de..41f263e378 100644 --- a/wgpu-core/src/hub.rs +++ b/wgpu-core/src/hub.rs @@ -1022,6 +1022,7 @@ impl HalApi for hal::api::Dx11 { impl HalApi for hal::api::Gles { const VARIANT: Backend = Backend::Gl; fn create_instance_from_hal(name: &str, hal_instance: Self::Instance) -> Instance { + #[allow(clippy::needless_update)] Instance { name: name.to_owned(), gl: Some(hal_instance), diff --git a/wgpu-core/src/lib.rs b/wgpu-core/src/lib.rs index e01d4a2ced..88aaf899f7 100644 --- a/wgpu-core/src/lib.rs +++ b/wgpu-core/src/lib.rs @@ -45,7 +45,7 @@ pub mod resource; mod track; mod validation; -pub use hal::api; +pub use hal::{api, MAX_BIND_GROUPS, MAX_COLOR_TARGETS, MAX_VERTEX_BUFFERS}; use atomic::{AtomicUsize, Ordering}; @@ -211,7 +211,10 @@ macro_rules! gfx_select { wgt::Backend::Dx12 => $global.$method::<$crate::api::Dx12>( $($param),* ), //#[cfg(all(not(target_arch = "wasm32"), windows))] //wgt::Backend::Dx11 => $global.$method::<$crate::api::Dx11>( $($param),* ), - #[cfg(all(not(target_arch = "wasm32"), unix, not(any(target_os = "ios", target_os = "macos"))))] + #[cfg(any( + all(unix, not(target_os = "macos"), not(target_os = "ios")), + target_arch = "wasm32" + ))] wgt::Backend::Gl => $global.$method::<$crate::api::Gles>( $($param),+ ), other => panic!("Unexpected backend {:?}", other), diff --git a/wgpu-hal/Cargo.toml b/wgpu-hal/Cargo.toml index 4e2aef02f7..18b6dc7a88 100644 --- a/wgpu-hal/Cargo.toml +++ b/wgpu-hal/Cargo.toml @@ -68,6 +68,11 @@ mtl = { package = "metal", version = "0.23.1" } objc = "0.2.5" core-graphics-types = "0.1" +[target.'cfg(target_arch = "wasm32")'.dependencies] +wasm-bindgen = { version = "0.2" } +web-sys = { version = "0.3", features = ["Window", "HtmlCanvasElement", "WebGl2RenderingContext"] } +js-sys = { version = "0.3" } + [dependencies.naga] git = "https://github.com/gfx-rs/naga" rev = "2e7d629" diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs index 373866623d..3cbaf4b5ef 100644 --- a/wgpu-hal/src/gles/adapter.rs +++ b/wgpu-hal/src/gles/adapter.rs @@ -192,7 +192,7 @@ impl super::Adapter { let shading_language_version = { let sl_version = gl.get_parameter_string(glow::SHADING_LANGUAGE_VERSION); - log::info!("SL version: {}", sl_version); + log::info!("SL version: {}", &sl_version); let (sl_major, sl_minor) = Self::parse_version(&sl_version).ok()?; let value = sl_major as u16 * 100 + sl_minor as u16 * 10; naga::back::glsl::Version::Embedded(value) @@ -209,9 +209,11 @@ impl super::Adapter { let max_storage_block_size = gl.get_parameter_i32(glow::MAX_SHADER_STORAGE_BLOCK_SIZE) as u32; - // WORKAROUND: - // In order to work around an issue with GL on RPI4 and similar, we ignore a zero vertex ssbo count if there are vertex sstos. (more info: https://github.com/gfx-rs/wgpu/pull/1607#issuecomment-874938961) - // The hardware does not want us to write to these SSBOs, but GLES cannot express that. We detect this case and disable writing to SSBOs. + // WORKAROUND: In order to work around an issue with GL on RPI4 and similar, we ignore a + // zero vertex ssbo count if there are vertex sstos. (more info: + // https://github.com/gfx-rs/wgpu/pull/1607#issuecomment-874938961) The hardware does not + // want us to write to these SSBOs, but GLES cannot express that. We detect this case and + // disable writing to SSBOs. let vertex_ssbo_false_zero = vertex_shader_storage_blocks == 0 && vertex_shader_storage_textures != 0; if vertex_ssbo_false_zero { @@ -254,6 +256,7 @@ impl super::Adapter { && max_storage_block_size != 0 && (vertex_shader_storage_blocks != 0 || vertex_ssbo_false_zero), ); + downlevel_flags.set(wgt::DownlevelFlags::FRAGMENT_STORAGE, ver >= (3, 1)); let mut features = wgt::Features::empty() | wgt::Features::TEXTURE_COMPRESSION_ETC2 @@ -283,6 +286,14 @@ impl super::Adapter { super::PrivateCapabilities::VERTEX_BUFFER_LAYOUT, ver >= (3, 1), ); + private_caps.set( + super::PrivateCapabilities::INDEX_BUFFER_ROLE_CHANGE, + cfg!(not(target_arch = "wasm32")), + ); + private_caps.set( + super::PrivateCapabilities::CAN_DISABLE_DRAW_BUFFER, + cfg!(not(target_arch = "wasm32")), + ); let max_texture_size = gl.get_parameter_i32(glow::MAX_TEXTURE_SIZE) as u32; let max_texture_3d_size = gl.get_parameter_i32(glow::MAX_3D_TEXTURE_SIZE) as u32; @@ -340,6 +351,12 @@ impl super::Adapter { }; let mut workarounds = super::Workarounds::empty(); + + workarounds.set( + super::Workarounds::EMULATE_BUFFER_MAP, + cfg!(target_arch = "wasm32"), + ); + let r = renderer.to_lowercase(); // Check for Mesa sRGB clear bug. See // [`super::PrivateCapabilities::MESA_I915_SRGB_SHADER_CLEAR`]. @@ -358,6 +375,9 @@ impl super::Adapter { let downlevel_defaults = wgt::DownlevelLimits {}; // Drop the GL guard so we can move the context into AdapterShared + // ( on WASM the gl handle is just a ref so we tell clippy to allow + // dropping the ref ) + #[allow(clippy::drop_ref)] drop(gl); Some(crate::ExposedAdapter { @@ -365,6 +385,7 @@ impl super::Adapter { shared: Arc::new(super::AdapterShared { context, private_caps, + downlevel_flags, workarounds, shading_language_version, }), @@ -462,6 +483,7 @@ impl crate::Adapter for super::Adapter { zero_buffer, temp_query_results: Vec::new(), draw_buffer_count: 1, + current_index_buffer: None, }, }) } @@ -561,11 +583,13 @@ impl crate::Adapter for super::Adapter { formats: if surface.enable_srgb { vec![ wgt::TextureFormat::Rgba8UnormSrgb, + #[cfg(not(target_arch = "wasm32"))] wgt::TextureFormat::Bgra8UnormSrgb, ] } else { vec![ wgt::TextureFormat::Rgba8Unorm, + #[cfg(not(target_arch = "wasm32"))] wgt::TextureFormat::Bgra8Unorm, ] }, @@ -590,6 +614,12 @@ impl crate::Adapter for super::Adapter { } } +// SAFE: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for super::Adapter {} +#[cfg(target_arch = "wasm32")] +unsafe impl Send for super::Adapter {} + #[cfg(test)] mod tests { use super::super::Adapter; diff --git a/wgpu-hal/src/gles/command.rs b/wgpu-hal/src/gles/command.rs index bb2b50f9a7..aba27e1361 100644 --- a/wgpu-hal/src/gles/command.rs +++ b/wgpu-hal/src/gles/command.rs @@ -266,14 +266,17 @@ impl crate::CommandEncoder for super::CommandEncoder { ) where T: Iterator, { - //TODO: preserve `src.target` and `dst.target` - // at least for the buffers that require it. + let (src_target, dst_target) = if src.target == dst.target { + (glow::COPY_READ_BUFFER, glow::COPY_WRITE_BUFFER) + } else { + (src.target, dst.target) + }; for copy in regions { self.cmd_buffer.commands.push(C::CopyBufferToBuffer { src: src.raw, - src_target: glow::COPY_READ_BUFFER, + src_target, dst: dst.raw, - dst_target: glow::COPY_WRITE_BUFFER, + dst_target, copy, }) } diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs index b96f4f61ff..476b043e6a 100644 --- a/wgpu-hal/src/gles/device.rs +++ b/wgpu-hal/src/gles/device.rs @@ -1,7 +1,10 @@ use super::conv; use crate::auxil::map_naga_stage; use glow::HasContext; -use std::{convert::TryInto, iter, mem, ptr, sync::Arc}; +use std::{convert::TryInto, iter, ptr, sync::Arc}; + +#[cfg(not(target_arch = "wasm32"))] +use std::mem; type ShaderStage<'a> = ( naga::ShaderStage, @@ -81,7 +84,7 @@ impl super::Device { gl: &glow::Context, shader: &str, naga_stage: naga::ShaderStage, - label: Option<&str>, + #[cfg_attr(target_arch = "wasm32", allow(unused))] label: Option<&str>, ) -> Result { let target = match naga_stage { naga::ShaderStage::Vertex => glow::VERTEX_SHADER, @@ -90,6 +93,7 @@ impl super::Device { }; let raw = gl.create_shader(target).unwrap(); + #[cfg(not(target_arch = "wasm32"))] if gl.supports_debug() { //TODO: remove all transmutes from `object_label` // https://github.com/grovesNL/glow/issues/186 @@ -170,9 +174,10 @@ impl super::Device { gl: &glow::Context, shaders: I, layout: &super::PipelineLayout, - label: crate::Label, + #[cfg_attr(target_arch = "wasm32", allow(unused))] label: Option<&str>, ) -> Result { let program = gl.create_program().unwrap(); + #[cfg(not(target_arch = "wasm32"))] if let Some(label) = label { if gl.supports_debug() { gl.object_label(glow::PROGRAM, mem::transmute(program), Some(label)); @@ -325,26 +330,46 @@ impl crate::Device for super::Device { .contains(crate::MemoryFlags::PREFER_COHERENT); let mut map_flags = 0; - if is_host_visible { - map_flags |= glow::MAP_PERSISTENT_BIT; - if is_coherent { - map_flags |= glow::MAP_COHERENT_BIT; - } - } - if desc.usage.contains(crate::BufferUses::MAP_READ) { - map_flags |= glow::MAP_READ_BIT; - } - if desc.usage.contains(crate::BufferUses::MAP_WRITE) { - map_flags |= glow::MAP_WRITE_BIT; - } - let raw = gl.create_buffer().unwrap(); gl.bind_buffer(target, Some(raw)); let raw_size = desc .size .try_into() .map_err(|_| crate::DeviceError::OutOfMemory)?; - gl.buffer_storage(target, raw_size, None, map_flags); + + if self + .shared + .downlevel_flags + .contains(wgt::DownlevelFlags::VERTEX_STORAGE | wgt::DownlevelFlags::FRAGMENT_STORAGE) + { + if is_host_visible { + map_flags |= glow::MAP_PERSISTENT_BIT; + if is_coherent { + map_flags |= glow::MAP_COHERENT_BIT; + } + } + if desc.usage.contains(crate::BufferUses::MAP_READ) { + map_flags |= glow::MAP_READ_BIT; + } + if desc.usage.contains(crate::BufferUses::MAP_WRITE) { + map_flags |= glow::MAP_WRITE_BIT; + } + + gl.buffer_storage(target, raw_size, None, map_flags); + } else { + assert!(!is_coherent); + let usage = if is_host_visible { + if desc.usage.contains(crate::BufferUses::MAP_READ) { + glow::STREAM_READ + } else { + glow::DYNAMIC_DRAW + } + } else { + glow::STATIC_DRAW + }; + gl.buffer_data_size(target, raw_size, usage); + } + gl.bind_buffer(target, None); if !is_coherent && desc.usage.contains(crate::BufferUses::MAP_WRITE) { @@ -352,6 +377,7 @@ impl crate::Device for super::Device { } //TODO: do we need `glow::MAP_UNSYNCHRONIZED_BIT`? + #[cfg(not(target_arch = "wasm32"))] if let Some(label) = desc.label { if gl.supports_debug() { gl.object_label(glow::BUFFER, mem::transmute(raw), Some(label)); @@ -363,6 +389,7 @@ impl crate::Device for super::Device { target, size: desc.size, map_flags, + emulate_map_allocation: Default::default(), }) } unsafe fn destroy_buffer(&self, buffer: super::Buffer) { @@ -379,14 +406,28 @@ impl crate::Device for super::Device { let is_coherent = buffer.map_flags & glow::MAP_COHERENT_BIT != 0; - gl.bind_buffer(buffer.target, Some(buffer.raw)); - let ptr = gl.map_buffer_range( - buffer.target, - range.start as i32, - (range.end - range.start) as i32, - buffer.map_flags, - ); - gl.bind_buffer(buffer.target, None); + let ptr = if self + .shared + .workarounds + .contains(super::Workarounds::EMULATE_BUFFER_MAP) + { + let mut buf = vec![0; buffer.size as usize]; + let ptr = buf.as_mut_ptr(); + *buffer.emulate_map_allocation.lock().unwrap() = Some(buf); + + ptr + } else { + gl.bind_buffer(buffer.target, Some(buffer.raw)); + let ptr = gl.map_buffer_range( + buffer.target, + range.start as i32, + (range.end - range.start) as i32, + buffer.map_flags, + ); + gl.bind_buffer(buffer.target, None); + + ptr + }; Ok(crate::BufferMapping { ptr: ptr::NonNull::new(ptr).ok_or(crate::DeviceError::Lost)?, @@ -396,7 +437,14 @@ impl crate::Device for super::Device { unsafe fn unmap_buffer(&self, buffer: &super::Buffer) -> Result<(), crate::DeviceError> { let gl = &self.shared.context.lock(); gl.bind_buffer(buffer.target, Some(buffer.raw)); - gl.unmap_buffer(buffer.target); + + if let Some(buf) = buffer.emulate_map_allocation.lock().unwrap().take() { + gl.buffer_sub_data_u8_slice(buffer.target, 0, &buf); + drop(buf); + } else { + gl.unmap_buffer(buffer.target); + } + gl.bind_buffer(buffer.target, None); Ok(()) } @@ -407,11 +455,15 @@ impl crate::Device for super::Device { let gl = &self.shared.context.lock(); gl.bind_buffer(buffer.target, Some(buffer.raw)); for range in ranges { - gl.flush_mapped_buffer_range( - buffer.target, - range.start as i32, - (range.end - range.start) as i32, - ); + if let Some(buf) = buffer.emulate_map_allocation.lock().unwrap().as_ref() { + gl.buffer_sub_data_u8_slice(buffer.target, range.start as i32, buf); + } else { + gl.flush_mapped_buffer_range( + buffer.target, + range.start as i32, + (range.end - range.start) as i32, + ); + } } } unsafe fn invalidate_mapped_ranges(&self, _buffer: &super::Buffer, _ranges: I) { @@ -458,6 +510,7 @@ impl crate::Device for super::Device { ); } + #[cfg(not(target_arch = "wasm32"))] if let Some(label) = desc.label { if gl.supports_debug() { gl.object_label(glow::RENDERBUFFER, mem::transmute(raw), Some(label)); @@ -537,6 +590,7 @@ impl crate::Device for super::Device { } }; + #[cfg(not(target_arch = "wasm32"))] if let Some(label) = desc.label { if gl.supports_debug() { gl.object_label(glow::TEXTURE, mem::transmute(raw), Some(label)); @@ -672,6 +726,7 @@ impl crate::Device for super::Device { ); } + #[cfg(not(target_arch = "wasm32"))] if let Some(label) = desc.label { if gl.supports_debug() { gl.object_label(glow::SAMPLER, mem::transmute(raw), Some(label)); @@ -959,11 +1014,11 @@ impl crate::Device for super::Device { gl.delete_program(pipeline.inner.program); } + #[cfg_attr(target_arch = "wasm32", allow(unused))] unsafe fn create_query_set( &self, desc: &wgt::QuerySetDescriptor, ) -> Result { - use std::fmt::Write; let gl = &self.shared.context.lock(); let mut temp_string = String::new(); @@ -972,7 +1027,10 @@ impl crate::Device for super::Device { let query = gl .create_query() .map_err(|_| crate::DeviceError::OutOfMemory)?; + #[cfg(not(target_arch = "wasm32"))] if gl.supports_debug() { + use std::fmt::Write; + if let Some(label) = desc.label { temp_string.clear(); let _ = write!(temp_string, "{}[{}]", label, i); @@ -1012,6 +1070,7 @@ impl crate::Device for super::Device { &self, fence: &super::Fence, ) -> Result { + #[cfg_attr(target_arch = "wasm32", allow(clippy::needless_borrow))] Ok(fence.get_latest(&self.shared.context.lock())) } unsafe fn wait( @@ -1020,7 +1079,7 @@ impl crate::Device for super::Device { wait_value: crate::FenceValue, timeout_ms: u32, ) -> Result { - if fence.last_completed < wait_value { + if cfg!(not(target_arch = "wasm32")) && fence.last_completed < wait_value { let gl = &self.shared.context.lock(); let timeout_ns = (timeout_ms as u64 * 1_000_000).min(!0u32 as u64); let &(_, sync) = fence @@ -1053,3 +1112,9 @@ impl crate::Device for super::Device { .end_frame_capture(ptr::null_mut(), ptr::null_mut()) } } + +// SAFE: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for super::Device {} +#[cfg(target_arch = "wasm32")] +unsafe impl Send for super::Device {} diff --git a/wgpu-hal/src/gles/mod.rs b/wgpu-hal/src/gles/mod.rs index 631d86eb4d..8b78c345c8 100644 --- a/wgpu-hal/src/gles/mod.rs +++ b/wgpu-hal/src/gles/mod.rs @@ -58,6 +58,8 @@ To address this, we invalidate the vertex buffers based on: #[cfg(not(target_arch = "wasm32"))] mod egl; +#[cfg(target_arch = "wasm32")] +mod web; mod adapter; mod command; @@ -68,6 +70,9 @@ mod queue; #[cfg(not(target_arch = "wasm32"))] use self::egl::{AdapterContext, Instance, Surface}; +#[cfg(target_arch = "wasm32")] +use self::web::{AdapterContext, Instance, Surface}; + use arrayvec::ArrayVec; use glow::HasContext; @@ -122,6 +127,11 @@ bitflags::bitflags! { const MEMORY_BARRIERS = 1 << 2; /// Vertex buffer layouts separate from the data. const VERTEX_BUFFER_LAYOUT = 1 << 3; + /// Indicates that buffers used as ELEMENT_ARRAY_BUFFER may be created / initialized / used + /// as other targets, if not present they must not be mixed with other targets. + const INDEX_BUFFER_ROLE_CHANGE = 1 << 4; + /// Indicates that the device supports disabling draw buffers + const CAN_DISABLE_DRAW_BUFFER = 1 << 5; } } @@ -135,6 +145,8 @@ bitflags::bitflags! { // (https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4972/diffs?diff_id=75888#22f5d1004713c9bbf857988c7efb81631ab88f99_323_327) // seems to indicate all skylake models are effected. const MESA_I915_SRGB_SHADER_CLEAR = 1 << 0; + /// Buffer map must emulated becuase it is not supported natively + const EMULATE_BUFFER_MAP = 1 << 1; } } @@ -163,6 +175,7 @@ struct TextureFormatDesc { struct AdapterShared { context: AdapterContext, private_caps: PrivateCapabilities, + downlevel_flags: wgt::DownlevelFlags, workarounds: Workarounds, shading_language_version: naga::back::glsl::Version, } @@ -193,6 +206,7 @@ pub struct Queue { zero_buffer: glow::Buffer, temp_query_results: Vec, draw_buffer_count: u8, + current_index_buffer: Option, } #[derive(Debug)] @@ -201,8 +215,15 @@ pub struct Buffer { target: BindTarget, size: wgt::BufferAddress, map_flags: u32, + emulate_map_allocation: std::sync::Mutex>>, } +// Safe: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for Buffer {} +#[cfg(target_arch = "wasm32")] +unsafe impl Send for Buffer {} + #[derive(Clone, Debug)] enum TextureInner { Renderbuffer { @@ -217,7 +238,9 @@ enum TextureInner { impl TextureInner { fn as_native(&self) -> (glow::Texture, BindTarget) { match *self { - Self::Renderbuffer { raw, .. } => panic!("Unexpected renderbuffer {:?}", raw), + Self::Renderbuffer { .. } => { + panic!("Unexpected renderbuffer"); + } Self::Texture { raw, target } => (raw, target), } } @@ -400,10 +423,22 @@ pub struct RenderPipeline { stencil: Option, } +// SAFE: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Send for RenderPipeline {} +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for RenderPipeline {} + pub struct ComputePipeline { inner: PipelineInner, } +// SAFE: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Send for ComputePipeline {} +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for ComputePipeline {} + #[derive(Debug)] pub struct QuerySet { queries: Box<[glow::Query]>, diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs index b39e022917..57a4c6b59f 100644 --- a/wgpu-hal/src/gles/queue.rs +++ b/wgpu-hal/src/gles/queue.rs @@ -1,8 +1,9 @@ use super::Command as C; use arrayvec::ArrayVec; use glow::HasContext; -use std::{mem, ops::Range, slice, sync::Arc}; +use std::{mem, slice, sync::Arc}; +#[cfg(not(target_arch = "wasm32"))] const DEBUG_ID: u32 = 0; const CUBEMAP_FACES: [u32; 6] = [ @@ -14,7 +15,8 @@ const CUBEMAP_FACES: [u32; 6] = [ glow::TEXTURE_CUBE_MAP_NEGATIVE_Z, ]; -fn extract_marker<'a>(data: &'a [u8], range: &Range) -> &'a str { +#[cfg(not(target_arch = "wasm32"))] +fn extract_marker<'a>(data: &'a [u8], range: &std::ops::Range) -> &'a str { std::str::from_utf8(&data[range.start as usize..range.end as usize]).unwrap() } @@ -49,6 +51,7 @@ impl super::Queue { .map(|i| glow::COLOR_ATTACHMENT0 + i) .collect::>(); gl.draw_buffers(&indices); + #[cfg(not(target_arch = "wasm32"))] for draw_buffer in 0..self.draw_buffer_count as u32 { gl.disable_draw_buffer(glow::BLEND, draw_buffer); } @@ -105,7 +108,7 @@ impl super::Queue { &mut self, gl: &glow::Context, command: &C, - data_bytes: &[u8], + #[cfg_attr(target_arch = "wasm32", allow(unused))] data_bytes: &[u8], queries: &[glow::Query], ) { match *command { @@ -231,16 +234,54 @@ impl super::Queue { dst_target, copy, } => { - gl.bind_buffer(src_target, Some(src)); - gl.bind_buffer(dst_target, Some(dst)); + let is_index_buffer_only_element_dst = !self + .shared + .private_caps + .contains(super::PrivateCapabilities::INDEX_BUFFER_ROLE_CHANGE) + && dst_target == glow::ELEMENT_ARRAY_BUFFER + || src_target == glow::ELEMENT_ARRAY_BUFFER; - gl.copy_buffer_sub_data( - src_target, - dst_target, - copy.src_offset as i32, - copy.dst_offset as i32, - copy.size.get() as i32, - ); + let copy_src_target = glow::COPY_READ_BUFFER; + + // WebGL not allowed to copy data from other targets to element buffer and can't copy element data to other buffers + let copy_dst_target = if is_index_buffer_only_element_dst { + glow::ELEMENT_ARRAY_BUFFER + } else { + glow::COPY_WRITE_BUFFER + }; + + gl.bind_buffer(copy_src_target, Some(src)); + gl.bind_buffer(copy_dst_target, Some(dst)); + + if is_index_buffer_only_element_dst { + let mut buffer_data = vec![0; copy.size.get() as usize]; + gl.get_buffer_sub_data( + copy_src_target, + copy.src_offset as i32, + &mut buffer_data, + ); + gl.buffer_sub_data_u8_slice( + copy_dst_target, + copy.dst_offset as i32, + &buffer_data, + ); + } else { + gl.copy_buffer_sub_data( + copy_src_target, + copy_dst_target, + copy.src_offset as _, + copy.dst_offset as _, + copy.size.get() as _, + ); + } + + gl.bind_buffer(copy_src_target, None); + + if is_index_buffer_only_element_dst { + gl.bind_buffer(glow::ELEMENT_ARRAY_BUFFER, self.current_index_buffer); + } else { + gl.bind_buffer(copy_dst_target, None); + } } C::CopyTextureToTexture { src, @@ -513,6 +554,7 @@ impl super::Queue { } C::SetIndexBuffer(buffer) => { gl.bind_buffer(glow::ELEMENT_ARRAY_BUFFER, Some(buffer)); + self.current_index_buffer = Some(buffer); } C::BeginQuery(query, target) => { gl.begin_query(target, query); @@ -603,8 +645,15 @@ impl super::Queue { .map(|i| glow::COLOR_ATTACHMENT0 + i) .collect::>(); gl.draw_buffers(&indices); - for draw_buffer in 0..count as u32 { - gl.disable_draw_buffer(glow::BLEND, draw_buffer); + + if self + .shared + .private_caps + .contains(super::PrivateCapabilities::CAN_DISABLE_DRAW_BUFFER) + { + for draw_buffer in 0..count as u32 { + gl.disable_draw_buffer(glow::BLEND, draw_buffer); + } } } C::ClearColorF { @@ -863,7 +912,11 @@ impl super::Queue { gl.blend_equation_draw_buffer(index, blend.color.equation); gl.blend_func_draw_buffer(index, blend.color.src, blend.color.dst); } - } else { + } else if self + .shared + .private_caps + .contains(super::PrivateCapabilities::CAN_DISABLE_DRAW_BUFFER) + { gl.disable_draw_buffer(index, glow::BLEND); } } else { @@ -923,6 +976,7 @@ impl super::Queue { binding.format, ); } + #[cfg(not(target_arch = "wasm32"))] C::InsertDebugMarker(ref range) => { let marker = extract_marker(data_bytes, range); gl.debug_message_insert( @@ -933,11 +987,17 @@ impl super::Queue { marker, ); } + #[cfg(target_arch = "wasm32")] + C::InsertDebugMarker(_) => (), + #[cfg_attr(target_arch = "wasm32", allow(unused))] C::PushDebugGroup(ref range) => { + #[cfg(not(target_arch = "wasm32"))] let marker = extract_marker(data_bytes, range); + #[cfg(not(target_arch = "wasm32"))] gl.push_debug_group(glow::DEBUG_SOURCE_APPLICATION, DEBUG_ID, marker); } C::PopDebugGroup => { + #[cfg(not(target_arch = "wasm32"))] gl.pop_debug_group(); } } @@ -954,12 +1014,16 @@ impl crate::Queue for super::Queue { let gl = &shared.context.lock(); self.reset_state(gl); for cmd_buf in command_buffers.iter() { + #[cfg(not(target_arch = "wasm32"))] if let Some(ref label) = cmd_buf.label { gl.push_debug_group(glow::DEBUG_SOURCE_APPLICATION, DEBUG_ID, label); } + for command in cmd_buf.commands.iter() { self.process(gl, command, &cmd_buf.data_bytes, &cmd_buf.queries); } + + #[cfg(not(target_arch = "wasm32"))] if cmd_buf.label.is_some() { gl.pop_debug_group(); } @@ -981,7 +1045,12 @@ impl crate::Queue for super::Queue { surface: &mut super::Surface, texture: super::Texture, ) -> Result<(), crate::SurfaceError> { + #[cfg(not(target_arch = "wasm32"))] let gl = &self.shared.context.get_without_egl_lock(); + + #[cfg(target_arch = "wasm32")] + let gl = &self.shared.context.glow_context; + surface.present(texture, gl) } @@ -989,3 +1058,9 @@ impl crate::Queue for super::Queue { 1.0 } } + +// SAFE: WASM doesn't have threads +#[cfg(target_arch = "wasm32")] +unsafe impl Sync for super::Queue {} +#[cfg(target_arch = "wasm32")] +unsafe impl Send for super::Queue {} diff --git a/wgpu-hal/src/gles/web.rs b/wgpu-hal/src/gles/web.rs new file mode 100644 index 0000000000..5f26029e68 --- /dev/null +++ b/wgpu-hal/src/gles/web.rs @@ -0,0 +1,273 @@ +use glow::HasContext; +use parking_lot::Mutex; +use wasm_bindgen::JsCast; + +use super::TextureFormatDesc; + +/// A wrapper around a [`glow::Context`] to provide a fake `lock()` api that makes it compatible +/// with the `AdapterContext` API fromt the EGL implementation. +pub struct AdapterContext { + pub glow_context: glow::Context, +} + +impl AdapterContext { + /// Obtain a lock to the EGL context and get handle to the [`glow::Context`] that can be used to + /// do rendering. + #[track_caller] + pub fn lock(&self) -> &glow::Context { + &self.glow_context + } +} + +#[derive(Debug)] +pub struct Instance { + canvas: Mutex>, +} + +// SAFE: WASM doesn't have threads +unsafe impl Sync for Instance {} +unsafe impl Send for Instance {} + +impl crate::Instance for Instance { + unsafe fn init(_desc: &crate::InstanceDescriptor) -> Result { + Ok(Instance { + canvas: Mutex::new(None), + }) + } + + unsafe fn enumerate_adapters(&self) -> Vec> { + let canvas_guard = self.canvas.lock(); + let gl = match *canvas_guard { + Some(ref canvas) => { + let context_options = js_sys::Object::new(); + js_sys::Reflect::set( + &context_options, + &"antialias".into(), + &wasm_bindgen::JsValue::FALSE, + ) + .expect("Cannot create context options"); + let webgl2_context = canvas + .get_context_with_context_options("webgl2", &context_options) + .expect("Cannot create WebGL2 context") + .and_then(|context| context.dyn_into::().ok()) + .expect("Cannot convert into WebGL2 context"); + glow::Context::from_webgl2_context(webgl2_context) + } + None => return Vec::new(), + }; + + super::Adapter::expose(AdapterContext { glow_context: gl }) + .into_iter() + .collect() + } + + unsafe fn create_surface( + &self, + has_handle: &impl raw_window_handle::HasRawWindowHandle, + ) -> Result { + if let raw_window_handle::RawWindowHandle::Web(handle) = has_handle.raw_window_handle() { + let canvas: web_sys::HtmlCanvasElement = web_sys::window() + .and_then(|win| win.document()) + .expect("Cannot get document") + .query_selector(&format!("canvas[data-raw-handle=\"{}\"]", handle.id)) + .expect("Cannot query for canvas") + .expect("Canvas is not found") + .dyn_into() + .expect("Failed to downcast to canvas type"); + + *self.canvas.lock() = Some(canvas.clone()); + + Ok(Surface { + canvas, + present_program: None, + swapchain: None, + texture: None, + presentable: true, + enable_srgb: true, // WebGL only supports sRGB + }) + } else { + unreachable!() + } + } + + unsafe fn destroy_surface(&self, surface: Surface) { + let mut canvas_option_ref = self.canvas.lock(); + + if let Some(canvas) = canvas_option_ref.as_ref() { + if canvas == &surface.canvas { + *canvas_option_ref = None; + } + } + } +} + +#[derive(Clone, Debug)] +pub struct Surface { + canvas: web_sys::HtmlCanvasElement, + pub(super) swapchain: Option, + texture: Option, + pub(super) presentable: bool, + pub(super) enable_srgb: bool, + present_program: Option, +} + +// SAFE: Because web doesn't have threads ( yet ) +unsafe impl Sync for Surface {} +unsafe impl Send for Surface {} + +#[derive(Clone, Debug)] +pub struct Swapchain { + pub(crate) extent: wgt::Extent3d, + // pub(crate) channel: f::ChannelType, + pub(super) format: wgt::TextureFormat, + pub(super) framebuffer: glow::Framebuffer, + pub(super) format_desc: TextureFormatDesc, +} + +impl Surface { + pub(super) unsafe fn present( + &mut self, + _suf_texture: super::Texture, + gl: &glow::Context, + ) -> Result<(), crate::SurfaceError> { + gl.bind_framebuffer(glow::DRAW_FRAMEBUFFER, None); + gl.bind_sampler(0, None); + gl.active_texture(glow::TEXTURE0); + gl.bind_texture(glow::TEXTURE_2D, self.texture); + gl.use_program(self.present_program); + gl.disable(glow::DEPTH_TEST); + gl.disable(glow::STENCIL_TEST); + gl.disable(glow::SCISSOR_TEST); + gl.disable(glow::BLEND); + gl.disable(glow::CULL_FACE); + gl.draw_buffers(&[glow::BACK]); + gl.draw_arrays(glow::TRIANGLES, 0, 3); + + Ok(()) + } + + unsafe fn create_present_program(gl: &glow::Context) -> glow::Program { + let program = gl + .create_program() + .expect("Could not create shader program"); + let vertex = gl + .create_shader(glow::VERTEX_SHADER) + .expect("Could not create shader"); + gl.shader_source(vertex, include_str!("./web/present.vert")); + gl.compile_shader(vertex); + let fragment = gl + .create_shader(glow::FRAGMENT_SHADER) + .expect("Could not create shader"); + gl.shader_source(fragment, include_str!("./web/present.frag")); + gl.compile_shader(fragment); + gl.attach_shader(program, vertex); + gl.attach_shader(program, fragment); + gl.link_program(program); + gl.delete_shader(vertex); + gl.delete_shader(fragment); + gl.bind_texture(glow::TEXTURE_2D, None); + + program + } +} + +impl crate::Surface for Surface { + unsafe fn configure( + &mut self, + device: &super::Device, + config: &crate::SurfaceConfiguration, + ) -> Result<(), crate::SurfaceError> { + let gl = &device.shared.context.lock(); + + if let Some(swapchain) = self.swapchain.take() { + // delete all frame buffers already allocated + gl.delete_framebuffer(swapchain.framebuffer); + } + + if self.present_program.is_none() { + self.present_program = Some(Self::create_present_program(gl)); + } + + if self.texture.is_none() { + self.texture = Some(gl.create_texture().unwrap()); + } + + let desc = device.shared.describe_texture_format(config.format); + gl.bind_texture(glow::TEXTURE_2D, self.texture); + gl.tex_parameter_i32( + glow::TEXTURE_2D, + glow::TEXTURE_MIN_FILTER, + glow::NEAREST as _, + ); + gl.tex_parameter_i32( + glow::TEXTURE_2D, + glow::TEXTURE_MAG_FILTER, + glow::NEAREST as _, + ); + gl.tex_storage_2d( + glow::TEXTURE_2D, + 1, + desc.internal, + config.extent.width as i32, + config.extent.height as i32, + ); + + let framebuffer = gl.create_framebuffer().unwrap(); + gl.bind_framebuffer(glow::READ_FRAMEBUFFER, Some(framebuffer)); + gl.framebuffer_texture_2d( + glow::READ_FRAMEBUFFER, + glow::COLOR_ATTACHMENT0, + glow::TEXTURE_2D, + self.texture, + 0, + ); + gl.bind_texture(glow::TEXTURE_2D, None); + + self.swapchain = Some(Swapchain { + extent: config.extent, + // channel: config.format.base_format().1, + format: config.format, + format_desc: desc, + framebuffer, + }); + Ok(()) + } + + unsafe fn unconfigure(&mut self, device: &super::Device) { + let gl = device.shared.context.lock(); + if let Some(swapchain) = self.swapchain.take() { + gl.delete_framebuffer(swapchain.framebuffer); + } + if let Some(renderbuffer) = self.texture.take() { + gl.delete_texture(renderbuffer); + } + } + + unsafe fn acquire_texture( + &mut self, + _timeout_ms: u32, + ) -> Result>, crate::SurfaceError> { + let sc = self.swapchain.as_ref().unwrap(); + let texture = super::Texture { + inner: super::TextureInner::Texture { + raw: self.texture.unwrap(), + target: glow::TEXTURE_2D, + }, + array_layer_count: 1, + mip_level_count: 1, + format: sc.format, + format_desc: sc.format_desc.clone(), + copy_size: crate::CopyExtent { + width: sc.extent.width, + height: sc.extent.height, + depth: 1, + }, + }; + Ok(Some(crate::AcquiredSurfaceTexture { + texture, + suboptimal: false, + })) + } + + unsafe fn discard_texture(&mut self, _texture: super::Texture) {} +} diff --git a/wgpu-hal/src/gles/web/present.frag b/wgpu-hal/src/gles/web/present.frag new file mode 100644 index 0000000000..853f82a6ae --- /dev/null +++ b/wgpu-hal/src/gles/web/present.frag @@ -0,0 +1,16 @@ +#version 300 es +precision mediump float; +in vec2 uv; +uniform sampler2D present_texture; +out vec4 frag; +vec4 linear_to_srgb(vec4 linear) { + vec3 color_linear = linear.rgb; + vec3 selector = ceil(color_linear - 0.0031308); // 0 if under value, 1 if over + vec3 under = 12.92 * color_linear; + vec3 over = 1.055 * pow(color_linear, vec3(0.41666)) - 0.055; + vec3 result = mix(under, over, selector); + return vec4(result, linear.a); +} +void main() { + frag = linear_to_srgb(texture(present_texture, uv)); +} \ No newline at end of file diff --git a/wgpu-hal/src/gles/web/present.vert b/wgpu-hal/src/gles/web/present.vert new file mode 100644 index 0000000000..922f2a1848 --- /dev/null +++ b/wgpu-hal/src/gles/web/present.vert @@ -0,0 +1,18 @@ +#version 300 es +precision mediump float; +// A triangle that fills the whole screen +const vec2[3] TRIANGLE_POS = vec2[]( + vec2( 0.0, -3.0), + vec2(-3.0, 1.0), + vec2( 3.0, 1.0) +); +const vec2[3] TRIANGLE_UV = vec2[]( + vec2( 0.5, 1.), + vec2( -1.0, -1.0), + vec2( 2.0, -1.0) +); +out vec2 uv; +void main() { + uv = TRIANGLE_UV[gl_VertexID]; + gl_Position = vec4(TRIANGLE_POS[gl_VertexID], 0.0, 1.0); +} \ No newline at end of file diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs index 5eda0973e0..a90dae8002 100644 --- a/wgpu-hal/src/lib.rs +++ b/wgpu-hal/src/lib.rs @@ -53,7 +53,13 @@ compile_error!("DX12 API enabled on non-Windows OS. If your project is not using #[cfg(all(feature = "dx12", windows))] mod dx12; mod empty; -#[cfg(feature = "gles")] +#[cfg(all( + feature = "gles", + any( + target_arch = "wasm32", + all(unix, not(target_os = "ios"), not(target_os = "macos")) + ) +))] mod gles; #[cfg(all(feature = "metal", any(target_os = "macos", target_os = "ios")))] mod metal; @@ -65,7 +71,13 @@ pub mod api { #[cfg(feature = "dx12")] pub use super::dx12::Api as Dx12; pub use super::empty::Api as Empty; - #[cfg(feature = "gles")] + #[cfg(all( + feature = "gles", + any( + target_arch = "wasm32", + all(unix, not(target_os = "ios"), not(target_os = "macos")) + ) + ))] pub use super::gles::Api as Gles; #[cfg(feature = "metal")] pub use super::metal::Api as Metal; diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index 1d0ba1b9b7..6f061b666c 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -546,28 +546,31 @@ impl Features { /// Represents the sets of limits an adapter/device supports. /// -/// We provide two different defaults. -/// - [`Limits::downlevel_defaults()]. This is a set of limits that is guaranteed to -/// work on all backends, including "downlevel" backends such -/// as OpenGL and D3D11. For most applications we recommend using these -/// limits, assuming they are high enough for your application. -/// - [`Limits::default()`]. This is the set of limits that is guaranteed to -/// work on all modern backends and is guaranteed to be supported by WebGPU. -/// Applications needing more modern features can use this as a reasonable set of -/// limits if they are targeting only desktop and modern mobile devices. +/// We provide three different defaults. +/// - [`Limits::downlevel_defaults()`]. This is a set of limits that is guarenteed to work on almost +/// all backends, including "downlevel" backends such as OpenGL and D3D11, other than WebGL. For +/// most applications we recommend using these limits, assuming they are high enough for your +/// application, and you do not intent to support WebGL. +/// - [`Limits::downlevel_webgl2_defaults()`] This is a set of limits that is lower even than the +/// [`downlevel_defaults()`], configured to be low enough to support running in the browser using +/// WebGL2. +/// - [`Limits::default()`]. This is the set of limits that is guarenteed to work on all modern +/// backends and is guarenteed to be supported by WebGPU. Applications needing more modern +/// features can use this as a reasonable set of limits if they are targetting only desktop and +/// modern mobile devices. /// -/// We recommend starting with the most restrictive limits you can and manually -/// increasing the limits you need boosted. This will let you stay running on -/// all hardware that supports the limits you need. +/// We recommend starting with the most restrictive limits you can and manually increasing the +/// limits you need boosted. This will let you stay running on all hardware that supports the limits +/// you need. /// /// Limits "better" than the default must be supported by the adapter and requested when requesting -/// a device. If limits "better" than the adapter supports are requested, requesting a device will panic. -/// Once a device is requested, you may only use resources up to the limits requested _even_ if the -/// adapter supports "better" limits. +/// a device. If limits "better" than the adapter supports are requested, requesting a device will +/// panic. Once a device is requested, you may only use resources up to the limits requested _even_ +/// if the adapter supports "better" limits. /// /// Requesting limits that are "better" than you need may cause performance to decrease because the -/// implementation needs to support more than is needed. You should ideally only request exactly what -/// you need. +/// implementation needs to support more than is needed. You should ideally only request exactly +/// what you need. /// /// See also: #[repr(C)] @@ -668,7 +671,7 @@ impl Default for Limits { } impl Limits { - /// These default limits are guaranteed to be compatible with GLES3, WebGL, and D3D11 + /// These default limits are guarenteed to be compatible with GLES3, and D3D11 pub fn downlevel_defaults() -> Self { Self { max_texture_dimension_1d: 2096, @@ -694,6 +697,26 @@ impl Limits { } } + /// These default limits are guarenteed to be compatible with GLES3, and D3D11, and WebGL2 + pub fn downlevel_webgl2_defaults() -> Self { + #[cfg(target_arch = "wasm32")] + let defaults = Self { + max_storage_buffers_per_shader_stage: 0, + max_storage_textures_per_shader_stage: 0, + max_dynamic_storage_buffers_per_pipeline_layout: 0, + max_storage_buffer_binding_size: 0, + max_vertex_buffer_array_stride: 255, + + // Most of the values should be the same as the downlevel defaults + ..Self::downlevel_defaults() + }; + + #[cfg(not(target_arch = "wasm32"))] + let defaults = Self::downlevel_defaults(); + + defaults + } + /// Modify the current limits to use the resolution limits of the other. /// /// This is useful because the swapchain might need to be larger than any other image in the application. @@ -809,6 +832,9 @@ bitflags::bitflags! { /// WebGPU, the implementation is allowed to completely ignore aniso clamp. This flag is /// here for native backends so they can comunicate to the user of aniso is enabled. const ANISOTROPIC_FILTERING = 1 << 11; + + /// Supports storage buffers in fragment shaders. + const FRAGMENT_STORAGE = 1 << 12; } } diff --git a/wgpu/Cargo.toml b/wgpu/Cargo.toml index 16a4c38428..743d068790 100644 --- a/wgpu/Cargo.toml +++ b/wgpu/Cargo.toml @@ -280,3 +280,5 @@ parking_lot = { version = "0.11", features = ["wasm-bindgen"] } [target.'cfg(target_arch = "wasm32")'.dev-dependencies] console_error_panic_hook = "0.1.6" console_log = "0.1.2" +# We need the Location feature in the framework examples +web-sys = { version = "0.3.53", features = ["Location"] } diff --git a/wgpu/examples/boids/main.rs b/wgpu/examples/boids/main.rs index 331e63c819..54c1c47638 100644 --- a/wgpu/examples/boids/main.rs +++ b/wgpu/examples/boids/main.rs @@ -31,6 +31,17 @@ struct Example { } impl framework::Example for Example { + fn required_limits() -> wgpu::Limits { + wgpu::Limits::downlevel_defaults() + } + + fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities { + wgpu::DownlevelCapabilities { + flags: wgpu::DownlevelFlags::COMPUTE_SHADERS, + ..Default::default() + } + } + /// constructs initial instance of Example struct fn init( config: &wgpu::SurfaceConfiguration, diff --git a/wgpu/examples/capture/main.rs b/wgpu/examples/capture/main.rs index 6df1fe70f7..108f2d2360 100644 --- a/wgpu/examples/capture/main.rs +++ b/wgpu/examples/capture/main.rs @@ -28,7 +28,7 @@ async fn create_red_image_with_dimensions( width: usize, height: usize, ) -> (Device, Buffer, BufferDimensions) { - let adapter = wgpu::Instance::new(wgpu::Backends::PRIMARY) + let adapter = wgpu::Instance::new(wgpu::Backends::all()) .request_adapter(&wgpu::RequestAdapterOptions::default()) .await .unwrap(); diff --git a/wgpu/examples/framework.rs b/wgpu/examples/framework.rs index ce64049da5..c19e2066ec 100644 --- a/wgpu/examples/framework.rs +++ b/wgpu/examples/framework.rs @@ -40,8 +40,15 @@ pub trait Example: 'static + Sized { fn required_features() -> wgpu::Features { wgpu::Features::empty() } + fn required_downlevel_capabilities() -> wgpu::DownlevelCapabilities { + wgpu::DownlevelCapabilities { + flags: wgpu::DownlevelFlags::empty(), + shader_model: wgpu::ShaderModel::Sm5, + ..wgpu::DownlevelCapabilities::default() + } + } fn required_limits() -> wgpu::Limits { - wgpu::Limits::downlevel_defaults() // These downlevel limits will allow the code to run on all possible hardware + wgpu::Limits::downlevel_webgl2_defaults() // These downlevel limits will allow the code to run on all possible hardware } fn init( config: &wgpu::SurfaceConfiguration, @@ -95,7 +102,12 @@ async fn setup(title: &str) -> Setup { #[cfg(target_arch = "wasm32")] { use winit::platform::web::WindowExtWebSys; - console_log::init().expect("could not initialize logger"); + let query_string = web_sys::window().unwrap().location().search().unwrap(); + let level: log::Level = parse_url_query_string(&query_string, "RUST_LOG") + .map(|x| x.parse().ok()) + .flatten() + .unwrap_or(log::Level::Error); + console_log::init_with_level(level).expect("could not initialize logger"); std::panic::set_hook(Box::new(console_error_panic_hook::hook)); // On wasm, append the canvas to the document body web_sys::window() @@ -110,7 +122,7 @@ async fn setup(title: &str) -> Setup { log::info!("Initializing the surface..."); - let backend = wgpu::util::backend_bits_from_env().unwrap_or(wgpu::Backends::PRIMARY); + let backend = wgpu::util::backend_bits_from_env().unwrap_or_else(wgpu::Backends::all); let instance = wgpu::Instance::new(backend); let (size, surface) = unsafe { @@ -138,6 +150,21 @@ async fn setup(title: &str) -> Setup { required_features - adapter_features ); + let required_downlevel_capabilities = E::required_downlevel_capabilities(); + let downlevel_capabilities = adapter.get_downlevel_properties(); + assert!( + downlevel_capabilities.shader_model >= required_downlevel_capabilities.shader_model, + "Adapter does not support the minimum shader model required to run this example: {:?}", + required_downlevel_capabilities.shader_model + ); + assert!( + downlevel_capabilities + .flags + .contains(required_downlevel_capabilities.flags), + "Adapter does not support the downlevel capabilities required to run this example: {:?}", + required_downlevel_capabilities.flags - downlevel_capabilities.flags + ); + // Make sure we use the texture resolution limits from the adapter, so we can support images the size of the surface. let needed_limits = E::required_limits().using_resolution(adapter.limits()); @@ -389,6 +416,25 @@ pub fn run(title: &str) { }); } +#[cfg(target_arch = "wasm32")] +/// Parse the query string as returned by `web_sys::window()?.location().search()?` and get a +/// specific key out of it. +pub fn parse_url_query_string<'a>(query: &'a str, search_key: &str) -> Option<&'a str> { + let query_string = query.strip_prefix('?')?; + + for pair in query_string.split('&') { + let mut pair = pair.split('='); + let key = pair.next()?; + let value = pair.next()?; + + if key == search_key { + return Some(value); + } + } + + None +} + #[cfg(test)] pub struct FrameworkRefTest { pub image_path: &'static str, @@ -408,12 +454,9 @@ pub fn test(mut params: FrameworkRefTest) { assert_eq!(params.width % 64, 0, "width needs to be aligned 64"); let features = E::required_features() | params.optional_features; - let limits = E::required_limits(); test_common::initialize_test( - mem::take(&mut params.base_test_parameters) - .features(features) - .limits(limits), + mem::take(&mut params.base_test_parameters).features(features), |ctx| { let spawner = Spawner::new(); diff --git a/wgpu/examples/hello-compute/main.rs b/wgpu/examples/hello-compute/main.rs index 8bcb9d3102..fb45db79ea 100644 --- a/wgpu/examples/hello-compute/main.rs +++ b/wgpu/examples/hello-compute/main.rs @@ -33,7 +33,7 @@ async fn run() { async fn execute_gpu(numbers: &[u32]) -> Option> { // Instantiates instance of WebGPU - let instance = wgpu::Instance::new(wgpu::Backends::PRIMARY); + let instance = wgpu::Instance::new(wgpu::Backends::all()); // `request_adapter` instantiates the general connection to the GPU let adapter = instance diff --git a/wgpu/examples/hello-windows/main.rs b/wgpu/examples/hello-windows/main.rs index 807a019553..6cef470275 100644 --- a/wgpu/examples/hello-windows/main.rs +++ b/wgpu/examples/hello-windows/main.rs @@ -58,7 +58,7 @@ impl Viewport { } async fn run(event_loop: EventLoop<()>, viewports: Vec<(Window, wgpu::Color)>) { - let instance = wgpu::Instance::new(wgpu::Backends::PRIMARY); + let instance = wgpu::Instance::new(wgpu::Backends::all()); let viewports: Vec<_> = viewports .into_iter() .map(|(window, color)| ViewportDesc::new(window, color, &instance)) diff --git a/wgpu/examples/hello/main.rs b/wgpu/examples/hello/main.rs index a4c5d3f87f..1a9928d21b 100644 --- a/wgpu/examples/hello/main.rs +++ b/wgpu/examples/hello/main.rs @@ -1,7 +1,7 @@ /// This example shows how to describe the adapter in use. async fn run() { #[cfg_attr(target_arch = "wasm32", allow(unused_variables))] - let adapter = wgpu::Instance::new(wgpu::Backends::PRIMARY) + let adapter = wgpu::Instance::new(wgpu::Backends::all()) .request_adapter(&wgpu::RequestAdapterOptions::default()) .await .unwrap(); diff --git a/wgpu/examples/shadow/main.rs b/wgpu/examples/shadow/main.rs index be6ae25ec4..78b88854a8 100644 --- a/wgpu/examples/shadow/main.rs +++ b/wgpu/examples/shadow/main.rs @@ -214,11 +214,17 @@ impl framework::Example for Example { } fn init( - config: &wgpu::SurfaceConfiguration, - _adapter: &wgpu::Adapter, + sc_desc: &wgpu::SurfaceConfiguration, + adapter: &wgpu::Adapter, device: &wgpu::Device, _queue: &wgpu::Queue, ) -> Self { + let supports_storage_resources = adapter + .get_downlevel_properties() + .flags + .contains(wgpu::DownlevelFlags::VERTEX_STORAGE) + && device.limits().max_storage_buffers_per_shader_stage > 0; + // Create the vertex and index buffers let vertex_size = mem::size_of::(); let (cube_vertex_data, cube_index_data) = create_cube(); @@ -429,8 +435,11 @@ impl framework::Example for Example { let light_storage_buf = device.create_buffer(&wgpu::BufferDescriptor { label: None, size: light_uniform_size, - usage: wgpu::BufferUsages::STORAGE - | wgpu::BufferUsages::COPY_SRC + usage: if supports_storage_resources { + wgpu::BufferUsages::STORAGE + } else { + wgpu::BufferUsages::UNIFORM + } | wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false, }); @@ -546,7 +555,11 @@ impl framework::Example for Example { binding: 1, // lights visibility: wgpu::ShaderStages::FRAGMENT, ty: wgpu::BindingType::Buffer { - ty: wgpu::BufferBindingType::Storage { read_only: true }, + ty: if supports_storage_resources { + wgpu::BufferBindingType::Storage { read_only: true } + } else { + wgpu::BufferBindingType::Uniform + }, has_dynamic_offset: false, min_binding_size: wgpu::BufferSize::new(light_uniform_size), }, @@ -580,7 +593,7 @@ impl framework::Example for Example { push_constant_ranges: &[], }); - let mx_total = Self::generate_matrix(config.width as f32 / config.height as f32); + let mx_total = Self::generate_matrix(sc_desc.width as f32 / sc_desc.height as f32); let forward_uniforms = GlobalUniforms { proj: *mx_total.as_ref(), num_lights: [lights.len() as u32, 0, 0, 0], @@ -626,8 +639,12 @@ impl framework::Example for Example { }, fragment: Some(wgpu::FragmentState { module: &shader, - entry_point: "fs_main", - targets: &[config.format.into()], + entry_point: if supports_storage_resources { + "fs_main" + } else { + "fs_main_without_storage" + }, + targets: &[sc_desc.format.into()], }), primitive: wgpu::PrimitiveState { front_face: wgpu::FrontFace::Ccw, @@ -651,7 +668,7 @@ impl framework::Example for Example { } }; - let forward_depth = Self::create_depth_texture(config, device); + let forward_depth = Self::create_depth_texture(sc_desc, device); Example { entities, diff --git a/wgpu/examples/shadow/shader.wgsl b/wgpu/examples/shadow/shader.wgsl index 8fbd5b5b95..a4984bf33b 100644 --- a/wgpu/examples/shadow/shader.wgsl +++ b/wgpu/examples/shadow/shader.wgsl @@ -54,8 +54,16 @@ struct Lights { data: [[stride(96)]] array; }; +// Used when storage types are not supported +[[block]] +struct LightsWithoutStorage { + data: array; +}; + [[group(0), binding(1)]] var s_lights: Lights; +[[group(0), binding(1)]] +var u_lights: LightsWithoutStorage; [[group(0), binding(2)]] var t_shadow: texture_depth_2d_array; [[group(0), binding(3)]] @@ -102,3 +110,27 @@ fn fs_main(in: VertexOutput) -> [[location(0)]] vec4 { // multiply the light by material color return vec4(color, 1.0) * u_entity.color; } + +// The fragment entrypoint used when storage buffers are not available for the lights +[[stage(fragment)]] +fn fs_main_without_storage(in: VertexOutput) -> [[location(0)]] vec4 { + let normal = normalize(in.world_normal); + var color: vec3 = c_ambient; + var i: u32 = 0u; + loop { + if (i >= min(u_globals.num_lights.x, c_max_lights)) { + break; + } + // This line is the only difference from the entrypoint above. It uses the lights + // uniform instead of the lights storage buffer + let light = u_lights.data[i]; + let shadow = fetch_shadow(i, light.proj * in.world_position); + let light_dir = normalize(light.pos.xyz - in.world_position.xyz); + let diffuse = max(0.0, dot(normal, light_dir)); + color = color + shadow * diffuse * light.color.xyz; + continuing { + i = i + 1u; + } + } + return vec4(color, 1.0) * u_entity.color; +} \ No newline at end of file diff --git a/wgpu/src/backend/direct.rs b/wgpu/src/backend/direct.rs index 3b671ce25f..22796c2f5a 100644 --- a/wgpu/src/backend/direct.rs +++ b/wgpu/src/backend/direct.rs @@ -38,6 +38,7 @@ impl fmt::Debug for Context { } impl Context { + #[cfg(not(target_arch = "wasm32"))] pub unsafe fn from_hal_instance(hal_instance: A::Instance) -> Self { Self(wgc::hub::Global::from_hal_instance::( "wgpu", @@ -50,6 +51,7 @@ impl Context { &self.0 } + #[cfg(not(target_arch = "wasm32"))] pub fn enumerate_adapters(&self, backends: wgt::Backends) -> Vec { self.0 .enumerate_adapters(wgc::instance::AdapterInputs::Mask(backends, |_| { @@ -57,6 +59,7 @@ impl Context { })) } + #[cfg(not(target_arch = "wasm32"))] pub unsafe fn create_adapter_from_hal( &self, hal_adapter: hal::ExposedAdapter, @@ -64,6 +67,7 @@ impl Context { self.0.create_adapter_from_hal(hal_adapter, PhantomData) } + #[cfg(not(target_arch = "wasm32"))] pub unsafe fn create_device_from_hal( &self, adapter: &wgc::id::AdapterId, @@ -90,6 +94,7 @@ impl Context { Ok((device, device_id)) } + #[cfg(not(target_arch = "wasm32"))] pub unsafe fn create_texture_from_hal( &self, hal_texture: A::Texture, @@ -118,6 +123,7 @@ impl Context { } } + #[cfg(not(target_arch = "wasm32"))] pub unsafe fn texture_as_hal)>( &self, texture: &Texture, @@ -127,6 +133,7 @@ impl Context { .texture_as_hal::(texture.id, hal_texture_callback) } + #[cfg(not(target_arch = "wasm32"))] pub fn generate_report(&self) -> wgc::hub::GlobalReport { self.0.generate_report() } @@ -1172,17 +1179,17 @@ impl crate::Context for Context { // Limit is always less or equal to hal::MAX_BIND_GROUPS, so this is always right // Guards following ArrayVec assert!( - desc.bind_group_layouts.len() <= hal::MAX_BIND_GROUPS, + desc.bind_group_layouts.len() <= wgc::MAX_BIND_GROUPS, "Bind group layout count {} exceeds device bind group limit {}", desc.bind_group_layouts.len(), - hal::MAX_BIND_GROUPS + wgc::MAX_BIND_GROUPS ); let temp_layouts = desc .bind_group_layouts .iter() .map(|bgl| bgl.id) - .collect::>(); + .collect::>(); let descriptor = wgc::binding_model::PipelineLayoutDescriptor { label: desc.label.map(Borrowed), bind_group_layouts: Borrowed(&temp_layouts), @@ -1214,7 +1221,7 @@ impl crate::Context for Context { ) -> Self::RenderPipelineId { use wgc::pipeline as pipe; - let vertex_buffers: ArrayVec<_, { hal::MAX_VERTEX_BUFFERS }> = desc + let vertex_buffers: ArrayVec<_, { wgc::MAX_VERTEX_BUFFERS }> = desc .vertex .buffers .iter() @@ -1229,7 +1236,7 @@ impl crate::Context for Context { Some(_) => None, None => Some(wgc::device::ImplicitPipelineIds { root_id: PhantomData, - group_ids: &[PhantomData; hal::MAX_BIND_GROUPS], + group_ids: &[PhantomData; wgc::MAX_BIND_GROUPS], }), }; let descriptor = pipe::RenderPipelineDescriptor { @@ -1288,7 +1295,7 @@ impl crate::Context for Context { Some(_) => None, None => Some(wgc::device::ImplicitPipelineIds { root_id: PhantomData, - group_ids: &[PhantomData; hal::MAX_BIND_GROUPS], + group_ids: &[PhantomData; wgc::MAX_BIND_GROUPS], }), }; let descriptor = pipe::ComputePipelineDescriptor { @@ -1480,6 +1487,7 @@ impl crate::Context for Context { } } + #[cfg_attr(target_arch = "wasm32", allow(unused))] fn device_drop(&self, device: &Self::DeviceId) { #[cfg(not(target_arch = "wasm32"))] { @@ -1916,7 +1924,7 @@ impl crate::Context for Context { resolve_target: ca.resolve_target.map(|rt| rt.id), channel: map_pass_channel(Some(&ca.ops)), }) - .collect::>(); + .collect::>(); let depth_stencil = desc.depth_stencil_attachment.as_ref().map(|dsa| { wgc::command::RenderPassDepthStencilAttachment { diff --git a/wgpu/src/lib.rs b/wgpu/src/lib.rs index 56910ec81d..ffadabf293 100644 --- a/wgpu/src/lib.rs +++ b/wgpu/src/lib.rs @@ -1483,7 +1483,7 @@ impl Instance { /// # Safety /// /// - canvas must be a valid element to create a surface upon. - #[cfg(target_arch = "wasm32")] + #[cfg(all(target_arch = "wasm32", not(feature = "webgl")))] pub unsafe fn create_surface_from_canvas( &self, canvas: &web_sys::HtmlCanvasElement, @@ -1499,7 +1499,7 @@ impl Instance { /// # Safety /// /// - canvas must be a valid OffscreenCanvas to create a surface upon. - #[cfg(target_arch = "wasm32")] + #[cfg(all(target_arch = "wasm32", not(feature = "webgl")))] pub unsafe fn create_surface_from_offscreen_canvas( &self, canvas: &web_sys::OffscreenCanvas, diff --git a/wgpu/tests/common/mod.rs b/wgpu/tests/common/mod.rs index 20ece7f843..ee2ce39e71 100644 --- a/wgpu/tests/common/mod.rs +++ b/wgpu/tests/common/mod.rs @@ -83,7 +83,6 @@ pub struct FailureCase { // This information determines if a test should run. pub struct TestParameters { pub required_features: Features, - pub required_limits: Limits, pub required_downlevel_properties: DownlevelCapabilities, // Backends where test should fail. pub failures: Vec, @@ -93,7 +92,6 @@ impl Default for TestParameters { fn default() -> Self { Self { required_features: Features::empty(), - required_limits: Limits::downlevel_defaults(), required_downlevel_properties: lowest_downlevel_properties(), failures: Vec::new(), } @@ -122,12 +120,6 @@ impl TestParameters { self } - /// Set the list - pub fn limits(mut self, limits: Limits) -> Self { - self.required_limits = limits; - self - } - pub fn downlevel_flags(mut self, downlevel_flags: DownlevelFlags) -> Self { self.required_downlevel_properties.flags |= downlevel_flags; self @@ -178,7 +170,6 @@ impl TestParameters { self } } - pub fn initialize_test(parameters: TestParameters, test_function: impl FnOnce(TestingContext)) { // We don't actually care if it fails let _ = env_logger::try_init(); @@ -192,6 +183,7 @@ pub fn initialize_test(parameters: TestParameters, test_function: impl FnOnce(Te )) .expect("could not find sutable adapter on the system"); + let required_limits = Limits::downlevel_defaults(); let adapter_info = adapter.get_info(); let adapter_lowercase_name = adapter_info.name.to_lowercase(); let adapter_features = adapter.features(); @@ -204,7 +196,7 @@ pub fn initialize_test(parameters: TestParameters, test_function: impl FnOnce(Te return; } - if adapter_limits < parameters.required_limits { + if adapter_limits < required_limits { println!("TEST SKIPPED: LIMIT TOO LOW"); return; } @@ -232,7 +224,7 @@ pub fn initialize_test(parameters: TestParameters, test_function: impl FnOnce(Te let (device, queue) = pollster::block_on(initialize_device( &adapter, parameters.required_features, - parameters.required_limits, + required_limits, )); let context = TestingContext {