diff --git a/examples/with_winit/src/main.rs b/examples/with_winit/src/main.rs
index 4046538..9a28f01 100644
--- a/examples/with_winit/src/main.rs
+++ b/examples/with_winit/src/main.rs
@@ -225,6 +225,9 @@ enum UserEvent {
 
 fn main() {
     let args = Args::parse();
+    // TODO: initializing both env_logger and console_logger fails on wasm.
+    // Figure out a more principled approach.
+    #[cfg(not(target_arch = "wasm32"))]
     env_logger::init();
     #[cfg(not(target_arch = "wasm32"))]
     {
@@ -261,6 +264,6 @@ fn main() {
             .and_then(|doc| doc.body())
             .and_then(|body| body.append_child(&web_sys::Element::from(canvas)).ok())
             .expect("couldn't append canvas to document body");
-        wasm_bindgen_futures::spawn_local(run(event_loop, window));
+        wasm_bindgen_futures::spawn_local(run(event_loop, window, args));
     }
 }
diff --git a/examples/with_winit/src/test_scene.rs b/examples/with_winit/src/test_scene.rs
index cd74f15..96fe0c2 100644
--- a/examples/with_winit/src/test_scene.rs
+++ b/examples/with_winit/src/test_scene.rs
@@ -82,9 +82,11 @@ pub fn render_svg_scene(
 ) {
     let scene_frag = scene.get_or_insert_with(|| {
         use super::pico_svg::*;
+        #[cfg(not(target_arch = "wasm32"))]
         let start = Instant::now();
         eprintln!("Starting to parse svg");
         let svg = PicoSvg::load(svg, scale).unwrap();
+        #[cfg(not(target_arch = "wasm32"))]
         eprintln!("Parsing svg took {:?}", start.elapsed());
         let mut new_scene = SceneFragment::new();
         let mut builder = SceneBuilder::for_fragment(&mut new_scene);
diff --git a/shader/coarse.wgsl b/shader/coarse.wgsl
index cea3637..df09de9 100644
--- a/shader/coarse.wgsl
+++ b/shader/coarse.wgsl
@@ -148,7 +148,17 @@ fn main(
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
     // we still want to know this workgroup's memory requirement.   
-    if (atomicLoad(&bump.failed) & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u {
+    if local_id.x == 0u {
+        // Reuse sh_part_count to hold failed flag, shmem is tight
+        sh_part_count[0] = atomicLoad(&bump.failed);
+    }
+#ifdef have_uniform
+    let failed = workgroupUniformLoad(&sh_part_count[0]);
+#else
+    workgroupBarrier();
+    let failed = sh_part_count[0];
+#endif
+    if (failed & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u {
         return;
     }
     let width_in_bins = (config.width_in_tiles + N_TILE_X - 1u) / N_TILE_X;
@@ -207,8 +217,12 @@ fn main(
                     workgroupBarrier();
                 }
                 sh_part_count[local_id.x] = part_start_ix + count;
+#ifdef have_uniform
+                ready_ix = workgroupUniformLoad(&sh_part_count[WG_SIZE - 1u]);
+#else
                 workgroupBarrier();
                 ready_ix = sh_part_count[WG_SIZE - 1u];
+#endif
                 partition_ix += WG_SIZE;
             }
             // use binary search to find draw object to read
diff --git a/shader/tile_alloc.wgsl b/shader/tile_alloc.wgsl
index b28166e..8d39e7c 100644
--- a/shader/tile_alloc.wgsl
+++ b/shader/tile_alloc.wgsl
@@ -29,6 +29,7 @@ let WG_SIZE = 256u;
 
 var<workgroup> sh_tile_count: array<u32, WG_SIZE>;
 var<workgroup> sh_tile_offset: u32;
+var<workgroup> sh_atomic_failed: u32;
 
 @compute @workgroup_size(256)
 fn main(
@@ -37,8 +38,17 @@ fn main(
 ) {
     // Exit early if prior stages failed, as we can't run this stage.
     // We need to check only prior stages, as if this stage has failed in another workgroup, 
-    // we still want to know this workgroup's memory requirement.   
-    if (atomicLoad(&bump.failed) & STAGE_BINNING) != 0u {
+    // we still want to know this workgroup's memory requirement.
+    if local_id.x == 0u {
+        sh_atomic_failed = atomicLoad(&bump.failed);
+    }
+#ifdef have_uniform
+    let failed = workgroupUniformLoad(&sh_atomic_failed);
+#else
+    workgroupBarrier();
+    let failed = sh_atomic_failed;
+#endif
+    if (failed & STAGE_BINNING) != 0u {
         return;
     }    
     // scale factors useful for converting coordinates to tiles
diff --git a/src/engine.rs b/src/engine.rs
index 2039d9d..d3750a6 100644
--- a/src/engine.rs
+++ b/src/engine.rs
@@ -352,7 +352,19 @@ impl Engine {
                 }
                 Command::Clear(proxy, offset, size) => {
                     let buffer = bind_map.get_or_create(*proxy, device, &mut self.pool)?;
+                    #[cfg(not(target_arch = "wasm32"))]
                     encoder.clear_buffer(buffer, *offset, *size);
+                    #[cfg(target_arch = "wasm32")]
+                    {
+                        // TODO: remove this workaround when wgpu implements clear_buffer
+                        // Also note: semantics are wrong, it's queue order rather than encoder.
+                        let size = match size {
+                            Some(size) => size.get(),
+                            None => proxy.size,
+                        };
+                        let zeros = vec![0; size as usize];
+                        queue.write_buffer(buffer, *offset, &zeros);
+                    }
                 }
             }
         }
diff --git a/src/shaders.rs b/src/shaders.rs
index 6e09a97..4c7b3da 100644
--- a/src/shaders.rs
+++ b/src/shaders.rs
@@ -160,6 +160,11 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let mut small_config = HashSet::new();
     small_config.insert("full".into());
     small_config.insert("small".into());
+    // TODO: remove this workaround when workgroupUniformLoad lands in naga
+    #[allow(unused_mut)]
+    let mut uniform = HashSet::new();
+    #[cfg(target_arch = "wasm32")]
+    uniform.insert("have_uniform".into());
     let pathtag_reduce = engine.add_shader(
         device,
         "pathtag_reduce",
@@ -286,7 +291,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let tile_alloc = engine.add_shader(
         device,
         "tile_alloc",
-        preprocess::preprocess(shader!("tile_alloc"), &empty, &imports).into(),
+        preprocess::preprocess(shader!("tile_alloc"), &uniform, &imports).into(),
         &[
             BindType::Uniform,
             BindType::BufReadOnly,
@@ -321,7 +326,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
     let coarse = engine.add_shader(
         device,
         "coarse",
-        preprocess::preprocess(shader!("coarse"), &empty, &imports).into(),
+        preprocess::preprocess(shader!("coarse"), &uniform, &imports).into(),
         &[
             BindType::Uniform,
             BindType::BufReadOnly,