Implement WebGL2 Backend (#1686)

* Implement WebGL Backend * Add WebGL Fixes by @mrk-its * Update Limits for WASM and Examples * Address Review Points
2024-11-25 08:13:27 +00:00 · 2021-10-07 15:18:09 -05:00 · 2021-10-07 15:18:09 -05:00 · 312828f12f
commit 312828f12f
parent c36e08073f
32 changed files with 853 additions and 120 deletions
--- 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

--- a/run-wasm-example.sh
+++ 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
--- a/wasm-resources/README.md
+++ b/wasm-resources/README.md
@ -0,0 +1,3 @@
+# WASM Resources
+
+This directory contains resources used when building the WGPU examples for web.
--- a/wasm-resources/index.template.html
+++ b/wasm-resources/index.template.html
@ -0,0 +1,14 @@
+<html>
+  <head>
+    <meta charset="UTF-8" />
+    <meta name="viewport" content="width=device-width, initial-scale=1.0" />
+  </head>
+  <body>
+    <script type="module">
+      import init from "./{{example}}.js";
+      window.addEventListener("load", () => {
+        init();
+      });
+    </script>
+  </body>
+</html>
--- 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
--- 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
+            )
+        },
    }
 }
--- a/wgpu-core/src/device/mod.rs
+++ b/wgpu-core/src/device/mod.rs
@ -4584,6 +4584,10 @@ impl<G: GlobalIdentityHandlerFactory> Global<G> {
        {
            self.poll_devices::<hal::api::Dx11>(force_wait, &mut closures)?;
        }
+        #[cfg(gl)]
+        {
+            self.poll_devices::<hal::api::Gles>(force_wait, &mut closures)?;
+        }

        unsafe {
            closures.fire();
--- 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),
--- 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),

--- 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"
--- 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<super::Api> 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<super::Api> 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<super::Api> 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;
--- a/wgpu-hal/src/gles/command.rs
+++ b/wgpu-hal/src/gles/command.rs
@ -266,14 +266,17 @@ impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
    ) where
        T: Iterator<Item = crate::BufferCopy>,
    {
-        //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,
            })
        }
--- 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<glow::Shader, crate::PipelineError> {
        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<super::PipelineInner, crate::PipelineError> {
        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,6 +330,18 @@ impl crate::Device<super::Api> for super::Device {
            .contains(crate::MemoryFlags::PREFER_COHERENT);
        let mut map_flags = 0;

+        let raw = gl.create_buffer().unwrap();
+        gl.bind_buffer(target, Some(raw));
+        let raw_size = desc
+            .size
+            .try_into()
+            .map_err(|_| crate::DeviceError::OutOfMemory)?;
+
+        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 {
@ -338,13 +355,21 @@ impl crate::Device<super::Api> for super::Device {
                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);
+        } 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<super::Api> 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<super::Api> for super::Device {
            target,
            size: desc.size,
            map_flags,
+            emulate_map_allocation: Default::default(),
        })
    }
    unsafe fn destroy_buffer(&self, buffer: super::Buffer) {
@ -379,6 +406,17 @@ impl crate::Device<super::Api> for super::Device {

        let is_coherent = buffer.map_flags & glow::MAP_COHERENT_BIT != 0;

+        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,
@ -388,6 +426,9 @@ impl crate::Device<super::Api> for super::Device {
            );
            gl.bind_buffer(buffer.target, None);

+            ptr
+        };
+
        Ok(crate::BufferMapping {
            ptr: ptr::NonNull::new(ptr).ok_or(crate::DeviceError::Lost)?,
            is_coherent,
@ -396,7 +437,14 @@ impl crate::Device<super::Api> 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));
+
+        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,6 +455,9 @@ impl crate::Device<super::Api> for super::Device {
        let gl = &self.shared.context.lock();
        gl.bind_buffer(buffer.target, Some(buffer.raw));
        for range in ranges {
+            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,
@ -414,6 +465,7 @@ impl crate::Device<super::Api> for super::Device {
                );
            }
        }
+    }
    unsafe fn invalidate_mapped_ranges<I>(&self, _buffer: &super::Buffer, _ranges: I) {
        //TODO: do we need to do anything?
    }
@ -458,6 +510,7 @@ impl crate::Device<super::Api> 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<super::Api> 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<super::Api> 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<super::Api> 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<crate::Label>,
    ) -> Result<super::QuerySet, crate::DeviceError> {
-        use std::fmt::Write;
        let gl = &self.shared.context.lock();
        let mut temp_string = String::new();

@ -972,7 +1027,10 @@ impl crate::Device<super::Api> 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<super::Api> for super::Device {
        &self,
        fence: &super::Fence,
    ) -> Result<crate::FenceValue, crate::DeviceError> {
+        #[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<super::Api> for super::Device {
        wait_value: crate::FenceValue,
        timeout_ms: u32,
    ) -> Result<bool, crate::DeviceError> {
-        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<super::Api> 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 {}
--- 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<u64>,
    draw_buffer_count: u8,
+    current_index_buffer: Option<glow::Buffer>,
 }

 #[derive(Debug)]
@ -201,8 +215,15 @@ pub struct Buffer {
    target: BindTarget,
    size: wgt::BufferAddress,
    map_flags: u32,
+    emulate_map_allocation: std::sync::Mutex<Option<Vec<u8>>>,
 }

+// 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<StencilState>,
 }

+// 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]>,
--- 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<u32>) -> &'a str {
+#[cfg(not(target_arch = "wasm32"))]
+fn extract_marker<'a>(data: &'a [u8], range: &std::ops::Range<u32>) -> &'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::<ArrayVec<_, { crate::MAX_COLOR_TARGETS }>>();
        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,
+                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,
-                    copy.dst_offset as i32,
-                    copy.size.get() 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,10 +645,17 @@ impl super::Queue {
                    .map(|i| glow::COLOR_ATTACHMENT0 + i)
                    .collect::<ArrayVec<_, { crate::MAX_COLOR_TARGETS }>>();
                gl.draw_buffers(&indices);
+
+                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 {
                draw_buffer,
                ref color,
@ -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<super::Api> 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<super::Api> 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<super::Api> 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 {}
--- a/wgpu-hal/src/gles/web.rs
+++ 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<Option<web_sys::HtmlCanvasElement>>,
+}
+
+// SAFE: WASM doesn't have threads
+unsafe impl Sync for Instance {}
+unsafe impl Send for Instance {}
+
+impl crate::Instance<super::Api> for Instance {
+    unsafe fn init(_desc: &crate::InstanceDescriptor) -> Result<Self, crate::InstanceError> {
+        Ok(Instance {
+            canvas: Mutex::new(None),
+        })
+    }
+
+    unsafe fn enumerate_adapters(&self) -> Vec<crate::ExposedAdapter<super::Api>> {
+        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::<web_sys::WebGl2RenderingContext>().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<Surface, crate::InstanceError> {
+        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<Swapchain>,
+    texture: Option<glow::Texture>,
+    pub(super) presentable: bool,
+    pub(super) enable_srgb: bool,
+    present_program: Option<glow::Program>,
+}
+
+// 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<super::Api> 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<Option<crate::AcquiredSurfaceTexture<super::Api>>, 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) {}
+}
--- a/wgpu-hal/src/gles/web/present.frag
+++ 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));
+}
--- a/wgpu-hal/src/gles/web/present.vert
+++ 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);
+}
--- 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;
--- 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: <https://gpuweb.github.io/gpuweb/#dictdef-gpulimits>
 #[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;
    }
 }

--- 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"] }
--- 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,
--- 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();
--- 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<E: Example>(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<E: Example>(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<E: Example>(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<E: Example>(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<E: Example>(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();

--- 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<Vec<u32>> {
    // 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
--- 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))
--- 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();
--- 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::<Vertex>();
        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,
--- a/wgpu/examples/shadow/shader.wgsl
+++ b/wgpu/examples/shadow/shader.wgsl
@ -54,8 +54,16 @@ struct Lights {
    data: [[stride(96)]] array<Light>;
 };

+// Used when storage types are not supported
+[[block]]
+struct LightsWithoutStorage {
+    data: array<Light, 10>;
+};
+
 [[group(0), binding(1)]]
 var<storage, read> s_lights: Lights;
+[[group(0), binding(1)]]
+var<uniform> 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<f32> {
    // multiply the light by material color
    return vec4<f32>(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<f32> {
+    let normal = normalize(in.world_normal);
+    var color: vec3<f32> = 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<f32>(color, 1.0) * u_entity.color;
+}
--- 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<A: wgc::hub::HalApi>(hal_instance: A::Instance) -> Self {
        Self(wgc::hub::Global::from_hal_instance::<A>(
            "wgpu",
@ -50,6 +51,7 @@ impl Context {
        &self.0
    }

+    #[cfg(not(target_arch = "wasm32"))]
    pub fn enumerate_adapters(&self, backends: wgt::Backends) -> Vec<wgc::id::AdapterId> {
        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<A: wgc::hub::HalApi>(
        &self,
        hal_adapter: hal::ExposedAdapter<A>,
@ -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<A: wgc::hub::HalApi>(
        &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<A: wgc::hub::HalApi>(
        &self,
        hal_texture: A::Texture,
@ -118,6 +123,7 @@ impl Context {
        }
    }

+    #[cfg(not(target_arch = "wasm32"))]
    pub unsafe fn texture_as_hal<A: wgc::hub::HalApi, F: FnOnce(Option<&A::Texture>)>(
        &self,
        texture: &Texture,
@ -127,6 +133,7 @@ impl Context {
            .texture_as_hal::<A, F>(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::<ArrayVec<_, { hal::MAX_BIND_GROUPS }>>();
+            .collect::<ArrayVec<_, { wgc::MAX_BIND_GROUPS }>>();
        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::<ArrayVec<_, { hal::MAX_COLOR_TARGETS }>>();
+            .collect::<ArrayVec<_, { wgc::MAX_COLOR_TARGETS }>>();

        let depth_stencil = desc.depth_stencil_attachment.as_ref().map(|dsa| {
            wgc::command::RenderPassDepthStencilAttachment {
--- a/wgpu/src/lib.rs
+++ b/wgpu/src/lib.rs
@ -1483,7 +1483,7 @@ impl Instance {
    /// # Safety
    ///
    /// - canvas must be a valid <canvas> 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,
--- 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<FailureCase>,
@ -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 {