diff --git a/Cargo.lock b/Cargo.lock
index 43aaa61367..f50f4e6780 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -1039,7 +1039,7 @@ dependencies = [
 [[package]]
 name = "naga"
 version = "0.8.0"
-source = "git+https://github.com/gfx-rs/naga?rev=a1840be#a1840beb1a96c9a49980fe3efa8e2f3dcd88abe6"
+source = "git+https://github.com/gfx-rs/naga?rev=81dc674#81dc67402a743b4dc6b2d24c0d306cd18799238b"
 dependencies = [
  "bit-set",
  "bitflags",
diff --git a/cts_runner/examples/hello-compute.js b/cts_runner/examples/hello-compute.js
index d2101fe66b..2f30c376e8 100644
--- a/cts_runner/examples/hello-compute.js
+++ b/cts_runner/examples/hello-compute.js
@@ -4,11 +4,12 @@ const numbers = [1, 4, 3, 295];
 
 const device = await adapter.requestDevice();
 
-const shaderCode = `[[block]]
+const shaderCode = `@block
 struct PrimeIndices {
-    data: [[stride(4)]] array<u32>;
+    data: @stride(4) array<u32>;
 }; // this is used as both input and output for convenience
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<storage, read_write> v_indices: PrimeIndices;
 // The Collatz Conjecture states that for any integer n:
 // If n is even, n = n/2
@@ -37,8 +38,9 @@ fn collatz_iterations(n_base: u32) -> u32{
     }
     return i;
 }
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
 }`;
 
diff --git a/player/tests/data/empty.wgsl b/player/tests/data/empty.wgsl
index 2d6cffb18e..b328193f5d 100644
--- a/player/tests/data/empty.wgsl
+++ b/player/tests/data/empty.wgsl
@@ -1,3 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute)
+@workgroup_size(1)
 fn main() {
 }
diff --git a/player/tests/data/quad.wgsl b/player/tests/data/quad.wgsl
index 8998ce3cfe..a735389886 100644
--- a/player/tests/data/quad.wgsl
+++ b/player/tests/data/quad.wgsl
@@ -1,5 +1,5 @@
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> @builtin(position) vec4<f32> {
     // hacky way to draw a large triangle
     let tmp1 = i32(vertex_index) / 2;
     let tmp2 = i32(vertex_index) & 1;
@@ -10,7 +10,7 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> [[builtin(position)]]
     return vec4<f32>(pos, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main() -> @location(0) vec4<f32> {
     return vec4<f32>(1.0, 1.0, 1.0, 1.0);
 }
diff --git a/player/tests/data/zero-init-buffer-for-binding.wgsl b/player/tests/data/zero-init-buffer-for-binding.wgsl
index 41d33705a0..aee562cf9c 100644
--- a/player/tests/data/zero-init-buffer-for-binding.wgsl
+++ b/player/tests/data/zero-init-buffer-for-binding.wgsl
@@ -1,11 +1,13 @@
 struct InOutBuffer {
-    data: [[stride(4)]] array<u32>;
+    data: @stride(4) array<u32>;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<storage, read_write> buffer: InOutBuffer;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     buffer.data[global_id.x] = buffer.data[global_id.x] + global_id.x;
 }
diff --git a/player/tests/data/zero-init-texture-binding.wgsl b/player/tests/data/zero-init-texture-binding.wgsl
index a7ff6b7b6c..5ddd6473cf 100644
--- a/player/tests/data/zero-init-texture-binding.wgsl
+++ b/player/tests/data/zero-init-texture-binding.wgsl
@@ -1,6 +1,7 @@
-[[group(0), binding(0)]] var tex: texture_2d<f32>;
-[[group(0), binding(1)]] var tex_storage: texture_storage_2d<rgba8uint, write>;
+@group(0) @binding(0) var tex: texture_2d<f32>;
+@group(0) @binding(1) var tex_storage: texture_storage_2d<rgba8uint, write>;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
 }
diff --git a/wgpu-core/Cargo.toml b/wgpu-core/Cargo.toml
index 45977a517b..fecdbec031 100644
--- a/wgpu-core/Cargo.toml
+++ b/wgpu-core/Cargo.toml
@@ -38,7 +38,7 @@ thiserror = "1"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["span", "validate", "wgsl-in"]
 
diff --git a/wgpu-hal/Cargo.toml b/wgpu-hal/Cargo.toml
index 9f69af5098..79b1a39175 100644
--- a/wgpu-hal/Cargo.toml
+++ b/wgpu-hal/Cargo.toml
@@ -82,14 +82,14 @@ js-sys = { version = "0.3" }
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 
 # DEV dependencies
 
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["wgsl-in"]
 
diff --git a/wgpu-hal/examples/halmark/shader.wgsl b/wgpu-hal/examples/halmark/shader.wgsl
index a72128f762..00341fde0f 100644
--- a/wgpu-hal/examples/halmark/shader.wgsl
+++ b/wgpu-hal/examples/halmark/shader.wgsl
@@ -9,20 +9,22 @@ struct Locals {
     color: u32;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> globals: Globals;
 
-[[group(1), binding(0)]]
+@group(1)
+@binding(0)
 var<uniform> locals: Locals;
 
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] tex_coords: vec2<f32>;
-    [[location(1)]] color: vec4<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) tex_coords: vec2<f32>;
+    @location(1) color: vec4<f32>;
 };
 
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vi: u32) -> VertexOutput {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vi: u32) -> VertexOutput {
     let tc = vec2<f32>(f32(vi & 1u), 0.5 * f32(vi & 2u));
     let offset = vec2<f32>(tc.x * globals.size.x, tc.y * globals.size.y);
     let pos = globals.mvp * vec4<f32>(locals.position + offset, 0.0, 1.0);
@@ -30,12 +32,14 @@ fn vs_main([[builtin(vertex_index)]] vi: u32) -> VertexOutput {
     return VertexOutput(pos, tc, color);
 }
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var texture: texture_2d<f32>;
-[[group(0), binding(2)]]
+@group(0)
+@binding(2)
 var sam: sampler;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     return in.color * textureSampleLevel(texture, sam, in.tex_coords, 0.0);
 }
diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs
index 42a8b5bccd..bcac8bf318 100644
--- a/wgpu-hal/src/gles/adapter.rs
+++ b/wgpu-hal/src/gles/adapter.rs
@@ -286,7 +286,8 @@ impl super::Adapter {
 
         let mut features = wgt::Features::empty()
             | wgt::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES
-            | wgt::Features::CLEAR_TEXTURE;
+            | wgt::Features::CLEAR_TEXTURE
+            | wgt::Features::PUSH_CONSTANTS;
         features.set(
             wgt::Features::ADDRESS_MODE_CLAMP_TO_BORDER | wgt::Features::ADDRESS_MODE_CLAMP_TO_ZERO,
             extensions.contains("GL_EXT_texture_border_clamp"),
@@ -399,7 +400,7 @@ impl super::Adapter {
             } else {
                 !0
             },
-            max_push_constant_size: 0,
+            max_push_constant_size: super::MAX_PUSH_CONSTANTS as u32 * 4,
             min_uniform_buffer_offset_alignment,
             min_storage_buffer_offset_alignment,
             max_inter_stage_shader_components: gl.get_parameter_i32(glow::MAX_VARYING_COMPONENTS)
diff --git a/wgpu-hal/src/gles/command.rs b/wgpu-hal/src/gles/command.rs
index a2e80e0f6c..b0ef0d6de6 100644
--- a/wgpu-hal/src/gles/command.rs
+++ b/wgpu-hal/src/gles/command.rs
@@ -28,6 +28,7 @@ pub(super) struct State {
     has_pass_label: bool,
     instance_vbuf_mask: usize,
     dirty_vbuf_mask: usize,
+    push_offset_to_uniform: ArrayVec<super::UniformDesc, { super::MAX_PUSH_CONSTANTS }>,
 }
 
 impl super::CommandBuffer {
@@ -43,6 +44,21 @@ impl super::CommandBuffer {
         self.data_bytes.extend(marker.as_bytes());
         start..self.data_bytes.len() as u32
     }
+
+    fn add_push_constant_data(&mut self, data: &[u32]) -> Range<u32> {
+        let data_raw = unsafe {
+            std::slice::from_raw_parts(
+                data.as_ptr() as *const _,
+                data.len() * mem::size_of::<u32>(),
+            )
+        };
+        let start = self.data_bytes.len();
+        assert!(start < u32::MAX as usize);
+        self.data_bytes.extend_from_slice(data_raw);
+        let end = self.data_bytes.len();
+        assert!(end < u32::MAX as usize);
+        (start as u32)..(end as u32)
+    }
 }
 
 impl super::CommandEncoder {
@@ -148,8 +164,10 @@ impl super::CommandEncoder {
     fn set_pipeline_inner(&mut self, inner: &super::PipelineInner) {
         self.cmd_buffer.commands.push(C::SetProgram(inner.program));
 
-        //TODO: push constants
-        let _ = &inner.uniforms;
+        self.state.push_offset_to_uniform.clear();
+        self.state
+            .push_offset_to_uniform
+            .extend(inner.uniforms.iter().cloned());
 
         // rebind textures, if needed
         let mut dirty_textures = 0u32;
@@ -603,10 +621,25 @@ impl crate::CommandEncoder<super::Api> for super::CommandEncoder {
         &mut self,
         _layout: &super::PipelineLayout,
         _stages: wgt::ShaderStages,
-        _offset: u32,
-        _data: &[u32],
+        start_offset: u32,
+        data: &[u32],
     ) {
-        unimplemented!()
+        let range = self.cmd_buffer.add_push_constant_data(data);
+
+        let end = start_offset + data.len() as u32 * 4;
+        let mut offset = start_offset;
+        while offset < end {
+            let uniform = self.state.push_offset_to_uniform[offset as usize / 4].clone();
+            let size = uniform.size;
+            if uniform.location.is_none() {
+                panic!("No uniform for push constant");
+            }
+            self.cmd_buffer.commands.push(C::SetPushConstants {
+                uniform,
+                offset: range.start + offset,
+            });
+            offset += size;
+        }
     }
 
     unsafe fn insert_debug_marker(&mut self, label: &str) {
diff --git a/wgpu-hal/src/gles/conv.rs b/wgpu-hal/src/gles/conv.rs
index 09c269143a..daf637b5e4 100644
--- a/wgpu-hal/src/gles/conv.rs
+++ b/wgpu-hal/src/gles/conv.rs
@@ -375,3 +375,59 @@ pub(super) fn map_storage_access(access: wgt::StorageTextureAccess) -> u32 {
         wgt::StorageTextureAccess::ReadWrite => glow::READ_WRITE,
     }
 }
+
+pub(super) fn is_sampler(glsl_uniform_type: u32) -> bool {
+    match glsl_uniform_type {
+        glow::INT_SAMPLER_1D
+        | glow::INT_SAMPLER_1D_ARRAY
+        | glow::INT_SAMPLER_2D
+        | glow::INT_SAMPLER_2D_ARRAY
+        | glow::INT_SAMPLER_2D_MULTISAMPLE
+        | glow::INT_SAMPLER_2D_MULTISAMPLE_ARRAY
+        | glow::INT_SAMPLER_2D_RECT
+        | glow::INT_SAMPLER_3D
+        | glow::INT_SAMPLER_CUBE
+        | glow::INT_SAMPLER_CUBE_MAP_ARRAY
+        | glow::UNSIGNED_INT_SAMPLER_1D
+        | glow::UNSIGNED_INT_SAMPLER_1D_ARRAY
+        | glow::UNSIGNED_INT_SAMPLER_2D
+        | glow::UNSIGNED_INT_SAMPLER_2D_ARRAY
+        | glow::UNSIGNED_INT_SAMPLER_2D_MULTISAMPLE
+        | glow::UNSIGNED_INT_SAMPLER_2D_MULTISAMPLE_ARRAY
+        | glow::UNSIGNED_INT_SAMPLER_2D_RECT
+        | glow::UNSIGNED_INT_SAMPLER_3D
+        | glow::UNSIGNED_INT_SAMPLER_CUBE
+        | glow::UNSIGNED_INT_SAMPLER_CUBE_MAP_ARRAY
+        | glow::SAMPLER_1D
+        | glow::SAMPLER_1D_SHADOW
+        | glow::SAMPLER_1D_ARRAY
+        | glow::SAMPLER_1D_ARRAY_SHADOW
+        | glow::SAMPLER_2D
+        | glow::SAMPLER_2D_SHADOW
+        | glow::SAMPLER_2D_ARRAY
+        | glow::SAMPLER_2D_ARRAY_SHADOW
+        | glow::SAMPLER_2D_MULTISAMPLE
+        | glow::SAMPLER_2D_MULTISAMPLE_ARRAY
+        | glow::SAMPLER_2D_RECT
+        | glow::SAMPLER_2D_RECT_SHADOW
+        | glow::SAMPLER_3D
+        | glow::SAMPLER_CUBE
+        | glow::SAMPLER_CUBE_MAP_ARRAY
+        | glow::SAMPLER_CUBE_MAP_ARRAY_SHADOW
+        | glow::SAMPLER_CUBE_SHADOW => true,
+        _ => false,
+    }
+}
+
+pub(super) fn uniform_byte_size(glsl_uniform_type: u32) -> u32 {
+    match glsl_uniform_type {
+        glow::FLOAT | glow::INT => 4,
+        glow::FLOAT_VEC2 | glow::INT_VEC2 => 8,
+        glow::FLOAT_VEC3 | glow::INT_VEC3 => 12,
+        glow::FLOAT_VEC4 | glow::INT_VEC4 => 16,
+        glow::FLOAT_MAT2 => 16,
+        glow::FLOAT_MAT3 => 36,
+        glow::FLOAT_MAT4 => 64,
+        _ => panic!("Unsupported uniform datatype!"),
+    }
+}
diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs
index 29a40df31c..47998364a7 100644
--- a/wgpu-hal/src/gles/device.rs
+++ b/wgpu-hal/src/gles/device.rs
@@ -272,30 +272,35 @@ impl super::Device {
             }
         }
 
-        let uniforms = {
-            let count = gl.get_active_uniforms(program);
-            let mut offset = 0;
-            let mut uniforms = Vec::new();
-
-            for uniform in 0..count {
-                let glow::ActiveUniform { size, utype, name } =
-                    gl.get_active_uniform(program, uniform).unwrap();
-
-                if let Some(location) = gl.get_uniform_location(program, &name) {
-                    // Sampler2D won't show up in UniformLocation and the only other uniforms
-                    // should be push constants
-                    uniforms.push(super::UniformDesc {
-                        location,
-                        offset,
-                        utype,
-                    });
+        let mut uniforms: [super::UniformDesc; super::MAX_PUSH_CONSTANTS] = Default::default();
+        let count = gl.get_active_uniforms(program);
+        let mut offset = 0;
 
-                    offset += size as u32;
-                }
+        for uniform in 0..count {
+            let glow::ActiveUniform { utype, name, .. } =
+                gl.get_active_uniform(program, uniform).unwrap();
+
+            if conv::is_sampler(utype) {
+                continue;
             }
 
-            uniforms.into_boxed_slice()
-        };
+            if let Some(location) = gl.get_uniform_location(program, &name) {
+                if uniforms[offset / 4].location.is_some() {
+                    panic!("Offset already occupied")
+                }
+
+                // `size` will always be 1 so we need to guess the real size from the type
+                let uniform_size = conv::uniform_byte_size(utype);
+
+                uniforms[offset / 4] = super::UniformDesc {
+                    location: Some(location),
+                    size: uniform_size,
+                    utype,
+                };
+
+                offset += uniform_size as usize;
+            }
+        }
 
         Ok(super::PipelineInner {
             program,
@@ -858,7 +863,6 @@ impl crate::Device<super::Api> for super::Device {
                 version: self.shared.shading_language_version,
                 writer_flags,
                 binding_map,
-                push_constant_binding: 0, //TODO?
             },
         })
     }
diff --git a/wgpu-hal/src/gles/mod.rs b/wgpu-hal/src/gles/mod.rs
index d5e1980b21..7a89b2bbed 100644
--- a/wgpu-hal/src/gles/mod.rs
+++ b/wgpu-hal/src/gles/mod.rs
@@ -88,6 +88,7 @@ const MAX_TEXTURE_SLOTS: usize = 16;
 const MAX_SAMPLERS: usize = 16;
 const MAX_VERTEX_ATTRIBUTES: usize = 16;
 const ZERO_BUFFER_SIZE: usize = 256 << 10;
+const MAX_PUSH_CONSTANTS: usize = 16;
 
 impl crate::Api for Api {
     type Instance = Instance;
@@ -392,14 +393,19 @@ struct VertexBufferDesc {
     stride: u32,
 }
 
-#[allow(unused)]
-#[derive(Clone)]
+#[derive(Clone, Debug, Default)]
 struct UniformDesc {
-    location: glow::UniformLocation,
-    offset: u32,
+    location: Option<glow::UniformLocation>,
+    size: u32,
     utype: u32,
 }
 
+// Safe: WASM doesn't have threads
+#[cfg(target_arch = "wasm32")]
+unsafe impl Sync for UniformDesc {}
+#[cfg(target_arch = "wasm32")]
+unsafe impl Send for UniformDesc {}
+
 /// For each texture in the pipeline layout, store the index of the only
 /// sampler (in this layout) that the texture is used with.
 type SamplerBindMap = [Option<u8>; MAX_TEXTURE_SLOTS];
@@ -407,7 +413,7 @@ type SamplerBindMap = [Option<u8>; MAX_TEXTURE_SLOTS];
 struct PipelineInner {
     program: glow::Program,
     sampler_map: SamplerBindMap,
-    uniforms: Box<[UniformDesc]>,
+    uniforms: [UniformDesc; MAX_PUSH_CONSTANTS],
 }
 
 #[derive(Clone, Debug)]
@@ -716,6 +722,11 @@ enum Command {
     InsertDebugMarker(Range<u32>),
     PushDebugGroup(Range<u32>),
     PopDebugGroup,
+    SetPushConstants {
+        uniform: UniformDesc,
+        /// Offset from the start of the `data_bytes`
+        offset: u32,
+    },
 }
 
 #[derive(Default)]
diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs
index b3acddc9f8..731ff99185 100644
--- a/wgpu-hal/src/gles/queue.rs
+++ b/wgpu-hal/src/gles/queue.rs
@@ -1067,6 +1067,70 @@ impl super::Queue {
                 #[cfg(not(target_arch = "wasm32"))]
                 gl.pop_debug_group();
             }
+            C::SetPushConstants {
+                ref uniform,
+                offset,
+            } => {
+                fn get_data<T>(data: &[u8], offset: u32) -> &[T] {
+                    let raw = &data[(offset as usize)..];
+                    unsafe {
+                        slice::from_raw_parts(
+                            raw.as_ptr() as *const _,
+                            raw.len() / mem::size_of::<T>(),
+                        )
+                    }
+                }
+
+                let location = uniform.location.as_ref();
+
+                match uniform.utype {
+                    glow::FLOAT => {
+                        let data = get_data::<f32>(data_bytes, offset)[0];
+                        gl.uniform_1_f32(location, data);
+                    }
+                    glow::FLOAT_VEC2 => {
+                        let data = get_data::<[f32; 2]>(data_bytes, offset)[0];
+                        gl.uniform_2_f32_slice(location, &data);
+                    }
+                    glow::FLOAT_VEC3 => {
+                        let data = get_data::<[f32; 3]>(data_bytes, offset)[0];
+                        gl.uniform_3_f32_slice(location, &data);
+                    }
+                    glow::FLOAT_VEC4 => {
+                        let data = get_data::<[f32; 4]>(data_bytes, offset)[0];
+                        gl.uniform_4_f32_slice(location, &data);
+                    }
+                    glow::INT => {
+                        let data = get_data::<i32>(data_bytes, offset)[0];
+                        gl.uniform_1_i32(location, data);
+                    }
+                    glow::INT_VEC2 => {
+                        let data = get_data::<[i32; 2]>(data_bytes, offset)[0];
+                        gl.uniform_2_i32_slice(location, &data);
+                    }
+                    glow::INT_VEC3 => {
+                        let data = get_data::<[i32; 3]>(data_bytes, offset)[0];
+                        gl.uniform_3_i32_slice(location, &data);
+                    }
+                    glow::INT_VEC4 => {
+                        let data = get_data::<[i32; 4]>(data_bytes, offset)[0];
+                        gl.uniform_4_i32_slice(location, &data);
+                    }
+                    glow::FLOAT_MAT2 => {
+                        let data = get_data::<[f32; 4]>(data_bytes, offset)[0];
+                        gl.uniform_matrix_2_f32_slice(location, false, &data);
+                    }
+                    glow::FLOAT_MAT3 => {
+                        let data = get_data::<[f32; 9]>(data_bytes, offset)[0];
+                        gl.uniform_matrix_3_f32_slice(location, false, &data);
+                    }
+                    glow::FLOAT_MAT4 => {
+                        let data = get_data::<[f32; 16]>(data_bytes, offset)[0];
+                        gl.uniform_matrix_4_f32_slice(location, false, &data);
+                    }
+                    _ => panic!("Unsupported uniform datatype!"),
+                }
+            }
         }
     }
 }
diff --git a/wgpu/Cargo.toml b/wgpu/Cargo.toml
index dc6f4eb455..fe69c54961 100644
--- a/wgpu/Cargo.toml
+++ b/wgpu/Cargo.toml
@@ -136,20 +136,20 @@ env_logger = "0.8"
 
 [dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 optional = true
 
 # used to test all the example shaders
 [dev-dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["wgsl-in"]
 
 [target.'cfg(target_arch = "wasm32")'.dependencies.naga]
 git = "https://github.com/gfx-rs/naga"
-rev = "a1840be"
+rev = "81dc674"
 #version = "0.8"
 features = ["wgsl-out"]
 
diff --git a/wgpu/examples/boids/compute.wgsl b/wgpu/examples/boids/compute.wgsl
index d69750d68e..4277b5aefa 100644
--- a/wgpu/examples/boids/compute.wgsl
+++ b/wgpu/examples/boids/compute.wgsl
@@ -14,16 +14,17 @@ struct SimParams {
 };
 
 struct Particles {
-  particles : [[stride(16)]] array<Particle>;
+  particles : @stride(16) array<Particle>;
 };
 
-[[group(0), binding(0)]] var<uniform> params : SimParams;
-[[group(0), binding(1)]] var<storage, read> particlesSrc : Particles;
-[[group(0), binding(2)]] var<storage, read_write> particlesDst : Particles;
+@group(0) @binding(0) var<uniform> params : SimParams;
+@group(0) @binding(1) var<storage, read> particlesSrc : Particles;
+@group(0) @binding(2) var<storage, read_write> particlesDst : Particles;
 
 // https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
-[[stage(compute), workgroup_size(64)]]
-fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(64)
+fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
   let total = arrayLength(&particlesSrc.particles);
   let index = global_invocation_id.x;
   if (index >= total) {
diff --git a/wgpu/examples/boids/draw.wgsl b/wgpu/examples/boids/draw.wgsl
index 6822c57bbf..72056df838 100644
--- a/wgpu/examples/boids/draw.wgsl
+++ b/wgpu/examples/boids/draw.wgsl
@@ -1,9 +1,9 @@
-[[stage(vertex)]]
+@stage(vertex)
 fn main_vs(
-    [[location(0)]] particle_pos: vec2<f32>,
-    [[location(1)]] particle_vel: vec2<f32>,
-    [[location(2)]] position: vec2<f32>,
-) -> [[builtin(position)]] vec4<f32> {
+    @location(0) particle_pos: vec2<f32>,
+    @location(1) particle_vel: vec2<f32>,
+    @location(2) position: vec2<f32>,
+) -> @builtin(position) vec4<f32> {
     let angle = -atan2(particle_vel.x, particle_vel.y);
     let pos = vec2<f32>(
         position.x * cos(angle) - position.y * sin(angle),
@@ -12,7 +12,7 @@ fn main_vs(
     return vec4<f32>(pos + particle_pos, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn main_fs() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn main_fs() -> @location(0) vec4<f32> {
     return vec4<f32>(1.0, 1.0, 1.0, 1.0);
 }
diff --git a/wgpu/examples/conservative-raster/triangle_and_lines.wgsl b/wgpu/examples/conservative-raster/triangle_and_lines.wgsl
index a553ff4f69..5593cee430 100644
--- a/wgpu/examples/conservative-raster/triangle_and_lines.wgsl
+++ b/wgpu/examples/conservative-raster/triangle_and_lines.wgsl
@@ -1,22 +1,22 @@
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> @builtin(position) vec4<f32> {
     let i: i32 = i32(vertex_index % 3u);
     let x: f32 = f32(i - 1) * 0.75;
     let y: f32 = f32((i & 1) * 2 - 1) * 0.75 + x * 0.2 + 0.1;
     return vec4<f32>(x, y, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main_red() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main_red() -> @location(0) vec4<f32> {
     return vec4<f32>(1.0, 0.0, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main_blue() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main_blue() -> @location(0) vec4<f32> {
     return vec4<f32>(0.13, 0.31, 0.85, 1.0); // cornflower blue in linear space
 }
 
-[[stage(fragment)]]
-fn fs_main_white() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main_white() -> @location(0) vec4<f32> {
     return vec4<f32>(1.0, 1.0, 1.0, 1.0);
 }
\ No newline at end of file
diff --git a/wgpu/examples/conservative-raster/upscale.wgsl b/wgpu/examples/conservative-raster/upscale.wgsl
index edb87df173..addfbe2701 100644
--- a/wgpu/examples/conservative-raster/upscale.wgsl
+++ b/wgpu/examples/conservative-raster/upscale.wgsl
@@ -1,10 +1,10 @@
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] tex_coords: vec2<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) tex_coords: vec2<f32>;
 };
 
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
     let x: f32 = f32(i32(vertex_index & 1u) << 2u) - 1.0;
     let y: f32 = f32(i32(vertex_index & 2u) << 1u) - 1.0;
     var output: VertexOutput;
@@ -13,12 +13,14 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
     return output;
 }
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var r_color: texture_2d<f32>;
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var r_sampler: sampler;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     return textureSample(r_color, r_sampler, in.tex_coords);
-}
\ No newline at end of file
+}
diff --git a/wgpu/examples/cube/shader.wgsl b/wgpu/examples/cube/shader.wgsl
index 4bd4305abf..0efe215f08 100644
--- a/wgpu/examples/cube/shader.wgsl
+++ b/wgpu/examples/cube/shader.wgsl
@@ -1,18 +1,19 @@
 struct VertexOutput {
-    [[location(0)]] tex_coord: vec2<f32>;
-    [[builtin(position)]] position: vec4<f32>;
+    @location(0) tex_coord: vec2<f32>;
+    @builtin(position) position: vec4<f32>;
 };
 
 struct Locals {
     transform: mat4x4<f32>;
 };
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> r_locals: Locals;
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(
-    [[location(0)]] position: vec4<f32>,
-    [[location(1)]] tex_coord: vec2<f32>,
+    @location(0) position: vec4<f32>,
+    @location(1) tex_coord: vec2<f32>,
 ) -> VertexOutput {
     var out: VertexOutput;
     out.tex_coord = tex_coord;
@@ -20,17 +21,18 @@ fn vs_main(
     return out;
 }
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var r_color: texture_2d<u32>;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     let tex = textureLoad(r_color, vec2<i32>(in.tex_coord * 256.0), 0);
     let v = f32(tex.x) / 255.0;
     return vec4<f32>(1.0 - (v * 5.0), 1.0 - (v * 15.0), 1.0 - (v * 50.0), 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_wire() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_wire() -> @location(0) vec4<f32> {
     return vec4<f32>(0.0, 0.5, 0.0, 0.5);
 }
diff --git a/wgpu/examples/hello-compute/shader.wgsl b/wgpu/examples/hello-compute/shader.wgsl
index df541aff5e..4bcaf40b8f 100644
--- a/wgpu/examples/hello-compute/shader.wgsl
+++ b/wgpu/examples/hello-compute/shader.wgsl
@@ -1,8 +1,9 @@
 struct PrimeIndices {
-    data: [[stride(4)]] array<u32>;
+    data: @stride(4) array<u32>;
 }; // this is used as both input and output for convenience
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<storage, read_write> v_indices: PrimeIndices;
 
 // The Collatz Conjecture states that for any integer n:
@@ -34,7 +35,8 @@ fn collatz_iterations(n_base: u32) -> u32{
     return i;
 }
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+@stage(compute)
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
 }
diff --git a/wgpu/examples/hello-triangle/shader.wgsl b/wgpu/examples/hello-triangle/shader.wgsl
index a1ef447d6c..54d05914a5 100644
--- a/wgpu/examples/hello-triangle/shader.wgsl
+++ b/wgpu/examples/hello-triangle/shader.wgsl
@@ -1,11 +1,11 @@
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] in_vertex_index: u32) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> @builtin(position) vec4<f32> {
     let x = f32(i32(in_vertex_index) - 1);
     let y = f32(i32(in_vertex_index & 1u) * 2 - 1);
     return vec4<f32>(x, y, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main() -> @location(0) vec4<f32> {
     return vec4<f32>(1.0, 0.0, 0.0, 1.0);
 }
diff --git a/wgpu/examples/mipmap/blit.wgsl b/wgpu/examples/mipmap/blit.wgsl
index f8a91dfc59..80f76b1f6e 100644
--- a/wgpu/examples/mipmap/blit.wgsl
+++ b/wgpu/examples/mipmap/blit.wgsl
@@ -1,10 +1,10 @@
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] tex_coords: vec2<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) tex_coords: vec2<f32>;
 };
 
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
     var out: VertexOutput;
     let x = i32(vertex_index) / 2;
     let y = i32(vertex_index) & 1;
@@ -21,12 +21,14 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
     return out;
 }
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var r_color: texture_2d<f32>;
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var r_sampler: sampler;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     return textureSample(r_color, r_sampler, in.tex_coords);
 }
diff --git a/wgpu/examples/mipmap/draw.wgsl b/wgpu/examples/mipmap/draw.wgsl
index 5d8bb2df1a..aa569f1459 100644
--- a/wgpu/examples/mipmap/draw.wgsl
+++ b/wgpu/examples/mipmap/draw.wgsl
@@ -1,16 +1,17 @@
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] tex_coords: vec2<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) tex_coords: vec2<f32>;
 };
 
 struct Locals {
     transform: mat4x4<f32>;
 };
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> r_data: Locals;
 
-[[stage(vertex)]]
-fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
+@stage(vertex)
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
     let pos = vec2<f32>(
         100.0 * (1.0 - f32(vertex_index & 2u)),
         1000.0 * f32(vertex_index & 1u)
@@ -21,12 +22,14 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput {
     return out;
 }
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var r_color: texture_2d<f32>;
-[[group(0), binding(2)]]
+@group(0)
+@binding(2)
 var r_sampler: sampler;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     return textureSample(r_color, r_sampler, in.tex_coords);
 }
diff --git a/wgpu/examples/msaa-line/shader.wgsl b/wgpu/examples/msaa-line/shader.wgsl
index 8dbce73b35..30f29c905b 100644
--- a/wgpu/examples/msaa-line/shader.wgsl
+++ b/wgpu/examples/msaa-line/shader.wgsl
@@ -1,12 +1,12 @@
 struct VertexOutput {
-    [[location(0)]] color: vec4<f32>;
-    [[builtin(position)]] position: vec4<f32>;
+    @location(0) color: vec4<f32>;
+    @builtin(position) position: vec4<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(
-    [[location(0)]] position: vec2<f32>,
-    [[location(1)]] color: vec4<f32>,
+    @location(0) position: vec2<f32>,
+    @location(1) color: vec4<f32>,
 ) -> VertexOutput {
     var out: VertexOutput;
     out.position = vec4<f32>(position, 0.0, 1.0);
@@ -14,7 +14,7 @@ fn vs_main(
     return out;
 }
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     return in.color;
 }
diff --git a/wgpu/examples/shadow/shader.wgsl b/wgpu/examples/shadow/shader.wgsl
index d4dc0d04e5..bc70a93a22 100644
--- a/wgpu/examples/shadow/shader.wgsl
+++ b/wgpu/examples/shadow/shader.wgsl
@@ -3,7 +3,8 @@ struct Globals {
     num_lights: vec4<u32>;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> u_globals: Globals;
 
 struct Entity {
@@ -11,24 +12,25 @@ struct Entity {
     color: vec4<f32>;
 };
 
-[[group(1), binding(0)]]
+@group(1)
+@binding(0)
 var<uniform> u_entity: Entity;
 
-[[stage(vertex)]]
-fn vs_bake([[location(0)]] position: vec4<i32>) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_bake(@location(0) position: vec4<i32>) -> @builtin(position) vec4<f32> {
     return u_globals.view_proj * u_entity.world * vec4<f32>(position);
 }
 
 struct VertexOutput {
-    [[builtin(position)]] proj_position: vec4<f32>;
-    [[location(0)]] world_normal: vec3<f32>;
-    [[location(1)]] world_position: vec4<f32>;
+    @builtin(position) proj_position: vec4<f32>;
+    @location(0) world_normal: vec3<f32>;
+    @location(1) world_position: vec4<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(
-    [[location(0)]] position: vec4<i32>,
-    [[location(1)]] normal: vec4<i32>,
+    @location(0) position: vec4<i32>,
+    @location(1) normal: vec4<i32>,
 ) -> VertexOutput {
     let w = u_entity.world;
     let world_pos = u_entity.world * vec4<f32>(position);
@@ -48,7 +50,7 @@ struct Light {
 };
 
 struct Lights {
-    data: [[stride(96)]] array<Light>;
+    data: @stride(96) array<Light>;
 };
 
 // Used when storage types are not supported
@@ -56,13 +58,17 @@ struct LightsWithoutStorage {
     data: array<Light, 10>;
 };
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var<storage, read> s_lights: Lights;
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var<uniform> u_lights: LightsWithoutStorage;
-[[group(0), binding(2)]]
+@group(0)
+@binding(2)
 var t_shadow: texture_depth_2d_array;
-[[group(0), binding(3)]]
+@group(0)
+@binding(3)
 var sampler_shadow: sampler_comparison;
 
 fn fetch_shadow(light_id: u32, homogeneous_coords: vec4<f32>) -> f32 {
@@ -81,8 +87,8 @@ fn fetch_shadow(light_id: u32, homogeneous_coords: vec4<f32>) -> f32 {
 let c_ambient: vec3<f32> = vec3<f32>(0.05, 0.05, 0.05);
 let c_max_lights: u32 = 10u;
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     let normal = normalize(in.world_normal);
     // accumulate color
     var color: vec3<f32> = c_ambient;
@@ -101,8 +107,8 @@ fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
 }
 
 // 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> {
+@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;
     for(var i = 0u; i < min(u_globals.num_lights.x, c_max_lights); i += 1u) {
@@ -115,4 +121,4 @@ fn fs_main_without_storage(in: VertexOutput) -> [[location(0)]] vec4<f32> {
         color += shadow * diffuse * light.color.xyz;
     }
     return vec4<f32>(color, 1.0) * u_entity.color;
-}
\ No newline at end of file
+}
diff --git a/wgpu/examples/skybox/shader.wgsl b/wgpu/examples/skybox/shader.wgsl
index d10f293e70..6aff11102b 100644
--- a/wgpu/examples/skybox/shader.wgsl
+++ b/wgpu/examples/skybox/shader.wgsl
@@ -1,6 +1,6 @@
 struct SkyOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] uv: vec3<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) uv: vec3<f32>;
 };
 
 struct Data {
@@ -13,11 +13,12 @@ struct Data {
     // camera position
     cam_pos: vec4<f32>;
 };
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> r_data: Data;
 
-[[stage(vertex)]]
-fn vs_sky([[builtin(vertex_index)]] vertex_index: u32) -> SkyOutput {
+@stage(vertex)
+fn vs_sky(@builtin(vertex_index) vertex_index: u32) -> SkyOutput {
     // hacky way to draw a large triangle
     let tmp1 = i32(vertex_index) / 2;
     let tmp2 = i32(vertex_index) & 1;
@@ -39,15 +40,15 @@ fn vs_sky([[builtin(vertex_index)]] vertex_index: u32) -> SkyOutput {
 }
 
 struct EntityOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(1)]] normal: vec3<f32>;
-    [[location(3)]] view: vec3<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(1) normal: vec3<f32>;
+    @location(3) view: vec3<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_entity(
-    [[location(0)]] pos: vec3<f32>,
-    [[location(1)]] normal: vec3<f32>,
+    @location(0) pos: vec3<f32>,
+    @location(1) normal: vec3<f32>,
 ) -> EntityOutput {
     var out: EntityOutput;
     out.normal = normal;
@@ -56,18 +57,20 @@ fn vs_entity(
     return out;
 }
 
-[[group(0), binding(1)]]
+@group(0)
+@binding(1)
 var r_texture: texture_cube<f32>;
-[[group(0), binding(2)]]
+@group(0)
+@binding(2)
 var r_sampler: sampler;
 
-[[stage(fragment)]]
-fn fs_sky(in: SkyOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_sky(in: SkyOutput) -> @location(0) vec4<f32> {
     return textureSample(r_texture, r_sampler, in.uv);
 }
 
-[[stage(fragment)]]
-fn fs_entity(in: EntityOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_entity(in: EntityOutput) -> @location(0) vec4<f32> {
     let incident = normalize(in.view);
     let normal = normalize(in.normal);
     let reflected = incident - 2.0 * dot(normal, incident) * normal;
diff --git a/wgpu/examples/water/terrain.wgsl b/wgpu/examples/water/terrain.wgsl
index ce752cea97..c21685b194 100644
--- a/wgpu/examples/water/terrain.wgsl
+++ b/wgpu/examples/water/terrain.wgsl
@@ -3,7 +3,8 @@ struct Uniforms {
     clipping_plane: vec4<f32>;
 };
 
-[[group(0), binding(0)]]
+@group(0)
+@binding(0)
 var<uniform> uniforms: Uniforms;
 
 let light = vec3<f32>(150.0, 70.0, 0.0);
@@ -11,17 +12,17 @@ let light_colour = vec3<f32>(1.0, 0.98, 0.82);
 let ambient = 0.2;
 
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] colour: vec4<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) colour: vec4<f32>;
     // Comment this out if using user-clipping planes:
-    [[location(1)]] clip_dist: f32;
+    @location(1) clip_dist: f32;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(
-    [[location(0)]] position: vec3<f32>,
-    [[location(1)]] normal: vec3<f32>,
-    [[location(2)]] colour: vec4<f32>,
+    @location(0) position: vec3<f32>,
+    @location(1) normal: vec3<f32>,
+    @location(2) colour: vec4<f32>,
 ) -> VertexOutput {
     var out: VertexOutput;
     out.position = uniforms.projection_view * vec4<f32>(position, 1.0);
@@ -35,10 +36,11 @@ fn vs_main(
     return out;
 }
 
-[[stage(fragment), early_depth_test]]
+@stage(fragment)
+@early_depth_test
 fn fs_main(
     in: VertexOutput,
-) -> [[location(0)]] vec4<f32> {
+) -> @location(0) vec4<f32> {
     // Comment this out if using user-clipping planes:
     if(in.clip_dist < 0.0) {
         discard;
diff --git a/wgpu/examples/water/water.wgsl b/wgpu/examples/water/water.wgsl
index 859eb38fb9..80daad7d69 100644
--- a/wgpu/examples/water/water.wgsl
+++ b/wgpu/examples/water/water.wgsl
@@ -4,7 +4,7 @@ struct Uniforms {
     time_size_width: vec4<f32>;
     viewport_height: f32;
 };
-[[group(0), binding(0)]] var<uniform> uniforms: Uniforms;
+@group(0) @binding(0) var<uniform> uniforms: Uniforms;
 
 let light_point = vec3<f32>(150.0, 70.0, 0.0);
 let light_colour = vec3<f32>(1.0, 0.98, 0.82);
@@ -181,16 +181,16 @@ fn calc_specular(eye: vec3<f32>, normal: vec3<f32>, light: vec3<f32>) -> f32 {
 }
 
 struct VertexOutput {
-    [[builtin(position)]] position: vec4<f32>;
-    [[location(0)]] f_WaterScreenPos: vec2<f32>;
-    [[location(1)]] f_Fresnel: f32;
-    [[location(2)]] f_Light: vec3<f32>;
+    @builtin(position) position: vec4<f32>;
+    @location(0) f_WaterScreenPos: vec2<f32>;
+    @location(1) f_Fresnel: f32;
+    @location(2) f_Light: vec3<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(
-    [[location(0)]] position: vec2<i32>,
-    [[location(1)]] offsets: vec4<i32>,
+    @location(0) position: vec2<i32>,
+    @location(1) offsets: vec4<i32>,
 ) -> VertexOutput {
     let p_pos = vec2<f32>(position);
     let b_pos = make_position(p_pos + vec2<f32>(offsets.xy));
@@ -222,9 +222,9 @@ let water_colour = vec3<f32>(0.0, 0.46, 0.95);
 let zNear = 10.0;
 let zFar = 400.0;
 
-[[group(0), binding(1)]] var reflection: texture_2d<f32>;
-[[group(0), binding(2)]] var terrain_depth_tex: texture_2d<f32>;
-[[group(0), binding(3)]] var colour_sampler: sampler;
+@group(0) @binding(1) var reflection: texture_2d<f32>;
+@group(0) @binding(2) var terrain_depth_tex: texture_2d<f32>;
+@group(0) @binding(3) var colour_sampler: sampler;
 
 fn to_linear_depth(depth: f32) -> f32 {
     let z_n = 2.0 * depth - 1.0;
@@ -232,8 +232,8 @@ fn to_linear_depth(depth: f32) -> f32 {
     return z_e;
 }
 
-[[stage(fragment)]]
-fn fs_main(in: VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in: VertexOutput) -> @location(0) vec4<f32> {
     let reflection_colour = textureSample(reflection, colour_sampler, in.f_WaterScreenPos.xy).xyz;
 
     let pixel_depth = to_linear_depth(in.position.z);
diff --git a/wgpu/tests/vertex_indices/draw.vert.wgsl b/wgpu/tests/vertex_indices/draw.vert.wgsl
index a4babee980..197c4dbb51 100644
--- a/wgpu/tests/vertex_indices/draw.vert.wgsl
+++ b/wgpu/tests/vertex_indices/draw.vert.wgsl
@@ -2,17 +2,17 @@ struct Indices {
     arr: array<u32>;
 }; // this is used as both input and output for convenience
 
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read_write> indices: Indices;
 
-[[stage(vertex)]]
-fn vs_main([[builtin(instance_index)]] instance: u32, [[builtin(vertex_index)]] index: u32) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn vs_main(@builtin(instance_index) instance: u32, @builtin(vertex_index) index: u32) -> @builtin(position) vec4<f32> {
     let idx = instance * 3u + index;
     indices.arr[idx] = idx;
     return vec4<f32>(0.0, 0.0, 0.0, 1.0);
 }
 
-[[stage(fragment)]]
-fn fs_main() -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main() -> @location(0) vec4<f32> {
     return vec4<f32>(0.0);
 }