wgsl: Replace [[decoration]] with @decoration

Deprecate the old syntax. Migrate everything to the new syntax.

Bug: tint:1382
Change-Id: Ide12b2e927b17dc93b9714c7049090864cc568d3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/77260
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>

test/access/let/matrix.spvasm.expected.wgsl

@@ -3,7 +3,7 @@ fn main_1() {
   return;
 }
 
-[[stage(compute), workgroup_size(1, 1, 1)]]
+@stage(compute) @workgroup_size(1, 1, 1)
 fn main() {
   main_1();
 }

test/access/let/matrix.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(1., 2., 3.), vec3<f32>(4., 5., 6.), vec3<f32>(7., 8., 9.));
   let v : vec3<f32> = m[1];

test/access/let/matrix.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let m : mat3x3<f32> = mat3x3<f32>(vec3<f32>(1.0, 2.0, 3.0), vec3<f32>(4.0, 5.0, 6.0), vec3<f32>(7.0, 8.0, 9.0));
   let v : vec3<f32> = m[1];

test/access/let/vector.spvasm.expected.wgsl

@@ -5,7 +5,7 @@ fn main_1() {
   return;
 }
 
-[[stage(compute), workgroup_size(1, 1, 1)]]
+@stage(compute) @workgroup_size(1, 1, 1)
 fn main() {
   main_1();
 }

test/access/let/vector.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let v : vec3<f32> = vec3<f32>(1., 2., 3.);
   let scalar : f32 = v.y;

test/access/let/vector.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let v : vec3<f32> = vec3<f32>(1.0, 2.0, 3.0);
   let scalar : f32 = v.y;

test/access/var/matrix.spvasm.expected.wgsl

@@ -5,7 +5,7 @@ fn main_1() {
   return;
 }
 
-[[stage(compute), workgroup_size(1, 1, 1)]]
+@stage(compute) @workgroup_size(1, 1, 1)
 fn main() {
   main_1();
 }

test/access/var/matrix.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var m : mat3x3<f32>;
   let v : vec3<f32> = m[1];

test/access/var/matrix.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var m : mat3x3<f32>;
   let v : vec3<f32> = m[1];

test/access/var/vector.spvasm.expected.wgsl

@@ -8,7 +8,7 @@ fn main_1() {
   return;
 }
 
-[[stage(compute), workgroup_size(1, 1, 1)]]
+@stage(compute) @workgroup_size(1, 1, 1)
 fn main() {
   main_1();
 }

test/access/var/vector.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var v : vec3<f32>;
   let scalar : f32 = v.y;

test/access/var/vector.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   var v : vec3<f32>;
   let scalar : f32 = v.y;

test/array/assign_to_function_var.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -6,8 +6,8 @@ struct S {
 
 var<private> src_private : ArrayType;
 var<workgroup> src_workgroup : ArrayType;
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 fn ret_arr() -> ArrayType {
   return ArrayType();

test/array/assign_to_function_var.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -8,9 +8,9 @@ var<private> src_private : ArrayType;
 
 var<workgroup> src_workgroup : ArrayType;
 
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
 
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 fn ret_arr() -> ArrayType {
   return ArrayType();

test/array/assign_to_private_var.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -6,8 +6,8 @@ struct S {
 
 var<private> src_private : ArrayType;
 var<workgroup> src_workgroup : ArrayType;
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 var<private> dst : ArrayType;
 var<private> dst_nested : array<array<array<i32, 2>, 3>, 4>;

test/array/assign_to_private_var.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -8,9 +8,9 @@ var<private> src_private : ArrayType;
 
 var<workgroup> src_workgroup : ArrayType;
 
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
 
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 var<private> dst : ArrayType;
 

test/array/assign_to_storage_var.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -10,11 +10,11 @@ struct S_nested {
 
 var<private> src_private : ArrayType;
 var<workgroup> src_workgroup : ArrayType;
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
-[[group(0), binding(2)]] var<storage, read_write> dst : S;
-[[group(0), binding(3)]] var<storage, read_write> dst_nested : S_nested;
+@group(0) @binding(2) var<storage, read_write> dst : S;
+@group(0) @binding(3) var<storage, read_write> dst_nested : S_nested;
 
 fn ret_arr() -> ArrayType {
   return ArrayType();

test/array/assign_to_storage_var.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -12,13 +12,13 @@ var<private> src_private : ArrayType;
 
 var<workgroup> src_workgroup : ArrayType;
 
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
 
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
-[[group(0), binding(2)]] var<storage, read_write> dst : S;
+@group(0) @binding(2) var<storage, read_write> dst : S;
 
-[[group(0), binding(3)]] var<storage, read_write> dst_nested : S_nested;
+@group(0) @binding(3) var<storage, read_write> dst_nested : S_nested;
 
 fn ret_arr() -> ArrayType {
   return ArrayType();

test/array/assign_to_workgroup_var.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -6,8 +6,8 @@ struct S {
 
 var<private> src_private : ArrayType;
 var<workgroup> src_workgroup : ArrayType;
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 var<workgroup> dst : ArrayType;
 var<workgroup> dst_nested : array<array<array<i32, 2>, 3>, 4>;

test/array/assign_to_workgroup_var.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-type ArrayType = [[stride(16)]] array<i32, 4>;
+type ArrayType = @stride(16) array<i32, 4>;
 
 struct S {
   arr : ArrayType;
@@ -8,9 +8,9 @@ var<private> src_private : ArrayType;
 
 var<workgroup> src_workgroup : ArrayType;
 
-[[group(0), binding(0)]] var<uniform> src_uniform : S;
+@group(0) @binding(0) var<uniform> src_uniform : S;
 
-[[group(0), binding(1)]] var<storage, read_write> src_storage : S;
+@group(0) @binding(1) var<storage, read_write> src_storage : S;
 
 var<workgroup> dst : ArrayType;
 

test/array/function_parameter.wgsl

@@ -10,7 +10,7 @@ fn f3(a : array<array<array<f32, 4>, 3>, 2>) -> f32 {
   return a[1][2][3];
 }
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a1 : array<f32, 4> = array<f32, 4>();
   let a2 : array<array<f32, 4>, 3> = array<array<f32, 4>, 3>();

test/array/function_parameter.wgsl.expected.wgsl

@@ -10,7 +10,7 @@ fn f3(a : array<array<array<f32, 4>, 3>, 2>) -> f32 {
   return a[1][2][3];
 }
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a1 : array<f32, 4> = array<f32, 4>();
   let a2 : array<array<f32, 4>, 3> = array<array<f32, 4>, 3>();

test/array/function_return_type.wgsl

@@ -10,7 +10,7 @@ fn f3() -> array<array<array<f32, 4>, 3>, 2> {
   return array<array<array<f32, 4>, 3>, 2>(f2(), f2());
 }
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a1 : array<f32, 4> = f1();
   let a2 : array<array<f32, 4>, 3> = f2();

test/array/function_return_type.wgsl.expected.wgsl

@@ -10,7 +10,7 @@ fn f3() -> array<array<array<f32, 4>, 3>, 2> {
   return array<array<array<f32, 4>, 3>, 2>(f2(), f2());
 }
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a1 : array<f32, 4> = f1();
   let a2 : array<array<f32, 4>, 3> = f2();

test/array/size.wgsl

@@ -1,7 +1,7 @@
 let slen = 4;
 let ulen = 4u;
 
-[[stage(fragment)]]
+@stage(fragment)
 fn main() {
   var signed_literal : array<f32, 4>;
   var unsigned_literal : array<f32, 4u>;

test/array/size.wgsl.expected.wgsl

@@ -2,7 +2,7 @@ let slen = 4;
 
 let ulen = 4u;
 
-[[stage(fragment)]]
+@stage(fragment)
 fn main() {
   var signed_literal : array<f32, 4>;
   var unsigned_literal : array<f32, 4u>;

test/array/type_constructor.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x : i32 = 42;
 

test/array/type_constructor.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x : i32 = 42;
   let empty : array<i32, 4> = array<i32, 4>();

test/benchmark/particles.wgsl

@@ -17,21 +17,21 @@ struct RenderParams {
   right : vec3<f32>;
   up    : vec3<f32>;
 };
-[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+@binding(0) @group(0) var<uniform> render_params : RenderParams;
 
 struct VertexInput {
-  [[location(0)]] position : vec3<f32>;
-  [[location(1)]] color    : vec4<f32>;
-  [[location(2)]] quad_pos : vec2<f32>; // -1..+1
+  @location(0) position : vec3<f32>;
+  @location(1) color    : vec4<f32>;
+  @location(2) quad_pos : vec2<f32>; // -1..+1
 };
 
 struct VertexOutput {
-  [[builtin(position)]] position : vec4<f32>;
-  [[location(0)]]       color    : vec4<f32>;
-  [[location(1)]]       quad_pos : vec2<f32>; // -1..+1
+  @builtin(position) position : vec4<f32>;
+  @location(0)       color    : vec4<f32>;
+  @location(1)       quad_pos : vec2<f32>; // -1..+1
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(in : VertexInput) -> VertexOutput {
   var quad_pos = mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos;
   var position = in.position + quad_pos * 0.01;
@@ -45,8 +45,8 @@ fn vs_main(in : VertexInput) -> VertexOutput {
 ////////////////////////////////////////////////////////////////////////////////
 // Fragment shader
 ////////////////////////////////////////////////////////////////////////////////
-[[stage(fragment)]]
-fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in : VertexOutput) -> @location(0) vec4<f32> {
   var color = in.color;
   // Apply a circular particle alpha mask
   color.a = color.a * max(1.0 - length(in.quad_pos), 0.0);
@@ -72,12 +72,12 @@ struct Particles {
   particles : array<Particle>;
 };
 
-[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
-[[binding(1), group(0)]] var<storage, read_write> data : Particles;
-[[binding(2), group(0)]] var texture : texture_2d<f32>;
+@binding(0) @group(0) var<uniform> sim_params : SimulationParams;
+@binding(1) @group(0) var<storage, read_write> data : Particles;
+@binding(2) @group(0) var texture : texture_2d<f32>;
 
-[[stage(compute), workgroup_size(64)]]
-fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
   rand_seed = (sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw;
 
   let idx = GlobalInvocationID.x;
@@ -136,11 +136,11 @@ struct Buffer {
   weights : array<f32>;
 };
 
-[[binding(3), group(0)]] var<uniform> ubo : UBO;
-[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
-[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
-[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
-[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+@binding(3) @group(0) var<uniform> ubo : UBO;
+@binding(4) @group(0) var<storage, read> buf_in : Buffer;
+@binding(5) @group(0) var<storage, read_write> buf_out : Buffer;
+@binding(6) @group(0) var tex_in : texture_2d<f32>;
+@binding(7) @group(0) var tex_out : texture_storage_2d<rgba8unorm, write>;
 
 ////////////////////////////////////////////////////////////////////////////////
 // import_level
@@ -148,8 +148,8 @@ struct Buffer {
 // Loads the alpha channel from a texel of the source image, and writes it to
 // the buf_out.weights.
 ////////////////////////////////////////////////////////////////////////////////
-[[stage(compute), workgroup_size(64)]]
-fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn import_level(@builtin(global_invocation_id) coord : vec3<u32>) {
   _ = &buf_in;
   let offset = coord.x + coord.y * ubo.width;
   buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
@@ -163,8 +163,8 @@ fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
 // mip level of tex_out. See simulate() in particle.wgsl to understand the
 // probability logic.
 ////////////////////////////////////////////////////////////////////////////////
-[[stage(compute), workgroup_size(64)]]
-fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn export_level(@builtin(global_invocation_id) coord : vec3<u32>) {
   if (all(coord.xy < vec2<u32>(textureDimensions(tex_out)))) {
     let dst_offset = coord.x    + coord.y    * ubo.width;
     let src_offset = coord.x*2u + coord.y*2u * ubo.width;

test/benchmark/particles.wgsl.expected.wgsl

@@ -12,27 +12,27 @@ struct RenderParams {
   up : vec3<f32>;
 }
 
-[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+@binding(0) @group(0) var<uniform> render_params : RenderParams;
 
 struct VertexInput {
-  [[location(0)]]
+  @location(0)
   position : vec3<f32>;
-  [[location(1)]]
+  @location(1)
   color : vec4<f32>;
-  [[location(2)]]
+  @location(2)
   quad_pos : vec2<f32>;
 }
 
 struct VertexOutput {
-  [[builtin(position)]]
+  @builtin(position)
   position : vec4<f32>;
-  [[location(0)]]
+  @location(0)
   color : vec4<f32>;
-  [[location(1)]]
+  @location(1)
   quad_pos : vec2<f32>;
 }
 
-[[stage(vertex)]]
+@stage(vertex)
 fn vs_main(in : VertexInput) -> VertexOutput {
   var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
   var position = (in.position + (quad_pos * 0.01));
@@ -43,8 +43,8 @@ fn vs_main(in : VertexInput) -> VertexOutput {
   return out;
 }
 
-[[stage(fragment)]]
-fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+@stage(fragment)
+fn fs_main(in : VertexOutput) -> @location(0) vec4<f32> {
   var color = in.color;
   color.a = (color.a * max((1.0 - length(in.quad_pos)), 0.0));
   return color;
@@ -66,14 +66,14 @@ struct Particles {
   particles : array<Particle>;
 }
 
-[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+@binding(0) @group(0) var<uniform> sim_params : SimulationParams;
 
-[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+@binding(1) @group(0) var<storage, read_write> data : Particles;
 
-[[binding(2), group(0)]] var texture : texture_2d<f32>;
+@binding(2) @group(0) var texture : texture_2d<f32>;
 
-[[stage(compute), workgroup_size(64)]]
-fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn simulate(@builtin(global_invocation_id) GlobalInvocationID : vec3<u32>) {
   rand_seed = ((sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
   let idx = GlobalInvocationID.x;
   var particle = data.particles[idx];
@@ -110,25 +110,25 @@ struct Buffer {
   weights : array<f32>;
 }
 
-[[binding(3), group(0)]] var<uniform> ubo : UBO;
+@binding(3) @group(0) var<uniform> ubo : UBO;
 
-[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+@binding(4) @group(0) var<storage, read> buf_in : Buffer;
 
-[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+@binding(5) @group(0) var<storage, read_write> buf_out : Buffer;
 
-[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+@binding(6) @group(0) var tex_in : texture_2d<f32>;
 
-[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+@binding(7) @group(0) var tex_out : texture_storage_2d<rgba8unorm, write>;
 
-[[stage(compute), workgroup_size(64)]]
-fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn import_level(@builtin(global_invocation_id) coord : vec3<u32>) {
   _ = &(buf_in);
   let offset = (coord.x + (coord.y * ubo.width));
   buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
 }
 
-[[stage(compute), workgroup_size(64)]]
-fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+@stage(compute) @workgroup_size(64)
+fn export_level(@builtin(global_invocation_id) coord : vec3<u32>) {
   if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
     let dst_offset = (coord.x + (coord.y * ubo.width));
     let src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo.width));

test/benchmark/simple_compute.wgsl

@@ -2,9 +2,9 @@ struct SB {
   data : array<i32>;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+@group(0) @binding(0) var<storage, read_write> buffer : SB;
 
-[[stage(compute), workgroup_size(1, 2, 3)]]
-fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+@stage(compute) @workgroup_size(1, 2, 3)
+fn main(@builtin(global_invocation_id) id : vec3<u32>) {
   buffer.data[id.x] = buffer.data[id.x] + 1;
 }

test/benchmark/simple_compute.wgsl.expected.wgsl

@@ -2,9 +2,9 @@ struct SB {
   data : array<i32>;
 }
 
-[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+@group(0) @binding(0) var<storage, read_write> buffer : SB;
 
-[[stage(compute), workgroup_size(1, 2, 3)]]
-fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+@stage(compute) @workgroup_size(1, 2, 3)
+fn main(@builtin(global_invocation_id) id : vec3<u32>) {
   buffer.data[id.x] = (buffer.data[id.x] + 1);
 }

test/benchmark/simple_fragment.wgsl

@@ -1,12 +1,12 @@
 struct Input {
-  [[location(0)]] color: vec4<f32>;
+  @location(0) color: vec4<f32>;
 };
 
 struct Output {
-  [[location(0)]] color: vec4<f32>;
+  @location(0) color: vec4<f32>;
 };
 
-[[stage(fragment)]]
+@stage(fragment)
 fn main(in : Input) -> Output {
   return Output(in.color);
 }

test/benchmark/simple_fragment.wgsl.expected.wgsl

@@ -1,14 +1,14 @@
 struct Input {
-  [[location(0)]]
+  @location(0)
   color : vec4<f32>;
 }
 
 struct Output {
-  [[location(0)]]
+  @location(0)
   color : vec4<f32>;
 }
 
-[[stage(fragment)]]
+@stage(fragment)
 fn main(in : Input) -> Output {
   return Output(in.color);
 }

test/benchmark/simple_vertex.wgsl

@@ -1,12 +1,12 @@
 struct Input {
-  [[location(0)]] position: vec4<f32>;
+  @location(0) position: vec4<f32>;
 };
 
 struct Output {
-  [[builtin(position)]] position : vec4<f32>;
+  @builtin(position) position : vec4<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn main(in : Input) -> Output {
   return Output(in.position);
 }

test/benchmark/simple_vertex.wgsl.expected.wgsl

@@ -1,14 +1,14 @@
 struct Input {
-  [[location(0)]]
+  @location(0)
   position : vec4<f32>;
 }
 
 struct Output {
-  [[builtin(position)]]
+  @builtin(position)
   position : vec4<f32>;
 }
 
-[[stage(vertex)]]
+@stage(vertex)
 fn main(in : Input) -> Output {
   return Output(in.position);
 }

test/buffer/storage/dynamic_index/read.wgsl

@@ -7,17 +7,17 @@ struct Inner {
     f : f32;
     g : mat2x3<f32>;
     h : mat3x2<f32>;
-    i : [[stride(16)]] array<vec4<i32>, 4>;
+    i : @stride(16) array<vec4<i32>, 4>;
 };
 
 struct S {
     arr : array<Inner>;
 };
 
-[[binding(0), group(0)]] var<storage, read> s : S;
+@binding(0) @group(0) var<storage, read> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
     let a = s.arr[idx].a;
     let b = s.arr[idx].b;
     let c = s.arr[idx].c;

test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl

@@ -7,17 +7,17 @@ struct Inner {
   f : f32;
   g : mat2x3<f32>;
   h : mat3x2<f32>;
-  i : [[stride(16)]] array<vec4<i32>, 4>;
+  i : @stride(16) array<vec4<i32>, 4>;
 }
 
 struct S {
   arr : array<Inner>;
 }
 
-[[binding(0), group(0)]] var<storage, read> s : S;
+@binding(0) @group(0) var<storage, read> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
   let a = s.arr[idx].a;
   let b = s.arr[idx].b;
   let c = s.arr[idx].c;

test/buffer/storage/dynamic_index/write.wgsl

@@ -7,17 +7,17 @@ struct Inner {
     f : f32;
     g : mat2x3<f32>;
     h : mat3x2<f32>;
-    i : [[stride(16)]] array<vec4<i32>, 4>;
+    i : @stride(16) array<vec4<i32>, 4>;
 };
 
 struct S {
     arr : array<Inner>;
 };
 
-[[binding(0), group(0)]] var<storage, read_write> s : S;
+@binding(0) @group(0) var<storage, read_write> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
     s.arr[idx].a = vec3<i32>();
     s.arr[idx].b = i32();
     s.arr[idx].c = vec3<u32>();
@@ -26,5 +26,5 @@ fn main([[builtin(local_invocation_index)]] idx : u32) {
     s.arr[idx].f = f32();
     s.arr[idx].g = mat2x3<f32>();
     s.arr[idx].h = mat3x2<f32>();
-    s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+    s.arr[idx].i = @stride(16) array<vec4<i32>, 4>();
 }

test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl

@@ -7,17 +7,17 @@ struct Inner {
   f : f32;
   g : mat2x3<f32>;
   h : mat3x2<f32>;
-  i : [[stride(16)]] array<vec4<i32>, 4>;
+  i : @stride(16) array<vec4<i32>, 4>;
 }
 
 struct S {
   arr : array<Inner>;
 }
 
-[[binding(0), group(0)]] var<storage, read_write> s : S;
+@binding(0) @group(0) var<storage, read_write> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
   s.arr[idx].a = vec3<i32>();
   s.arr[idx].b = i32();
   s.arr[idx].c = vec3<u32>();
@@ -26,5 +26,5 @@ fn main([[builtin(local_invocation_index)]] idx : u32) {
   s.arr[idx].f = f32();
   s.arr[idx].g = mat2x3<f32>();
   s.arr[idx].h = mat3x2<f32>();
-  s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+  s.arr[idx].i = @stride(16) array<vec4<i32>, 4>();
 }

test/buffer/storage/static_index/read.wgsl

@@ -12,12 +12,12 @@ struct S {
     g : mat2x3<f32>;
     h : mat3x2<f32>;
     i : Inner;
-    j : [[stride(16)]] array<Inner, 4>;
+    j : @stride(16) array<Inner, 4>;
 };
 
-[[binding(0), group(0)]] var<storage, read> s : S;
+@binding(0) @group(0) var<storage, read> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
     let a = s.a;
     let b = s.b;

test/buffer/storage/static_index/read.wgsl.expected.wgsl

@@ -12,12 +12,12 @@ struct S {
   g : mat2x3<f32>;
   h : mat3x2<f32>;
   i : Inner;
-  j : [[stride(16)]] array<Inner, 4>;
+  j : @stride(16) array<Inner, 4>;
 }
 
-[[binding(0), group(0)]] var<storage, read> s : S;
+@binding(0) @group(0) var<storage, read> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a = s.a;
   let b = s.b;

test/buffer/storage/static_index/write.wgsl

@@ -12,12 +12,12 @@ struct S {
     g : mat2x3<f32>;
     h : mat3x2<f32>;
     i : Inner;
-    j : [[stride(16)]] array<Inner, 4>;
+    j : @stride(16) array<Inner, 4>;
 };
 
-[[binding(0), group(0)]] var<storage, write> s : S;
+@binding(0) @group(0) var<storage, write> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
     s.a = vec3<i32>();
     s.b = i32();
@@ -28,5 +28,5 @@ fn main() {
     s.g = mat2x3<f32>();
     s.h = mat3x2<f32>();
     s.i = Inner();
-    s.j = [[stride(16)]] array<Inner, 4>();
+    s.j = @stride(16) array<Inner, 4>();
 }

test/buffer/storage/static_index/write.wgsl.expected.wgsl

@@ -12,12 +12,12 @@ struct S {
   g : mat2x3<f32>;
   h : mat3x2<f32>;
   i : Inner;
-  j : [[stride(16)]] array<Inner, 4>;
+  j : @stride(16) array<Inner, 4>;
 }
 
-[[binding(0), group(0)]] var<storage, write> s : S;
+@binding(0) @group(0) var<storage, write> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   s.a = vec3<i32>();
   s.b = i32();
@@ -28,5 +28,5 @@ fn main() {
   s.g = mat2x3<f32>();
   s.h = mat3x2<f32>();
   s.i = Inner();
-  s.j = [[stride(16)]] array<Inner, 4>();
+  s.j = @stride(16) array<Inner, 4>();
 }

test/buffer/storage/types/array.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : array<f32, 4>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : array<f32, 4>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/array.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : array<f32, 4>;
+@group(0) @binding(0) var<storage, read> in : array<f32, 4>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : array<f32, 4>;
+@group(0) @binding(1) var<storage, read_write> out : array<f32, 4>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/f32.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : f32;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : f32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/f32.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : f32;
+@group(0) @binding(0) var<storage, read> in : f32;
 
-[[group(0), binding(1)]] var<storage, read_write> out : f32;
+@group(0) @binding(1) var<storage, read_write> out : f32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/i32.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : i32;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : i32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/i32.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : i32;
+@group(0) @binding(0) var<storage, read> in : i32;
 
-[[group(0), binding(1)]] var<storage, read_write> out : i32;
+@group(0) @binding(1) var<storage, read_write> out : i32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat2x2.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : mat2x2<f32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : mat2x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat2x2.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : mat2x2<f32>;
+@group(0) @binding(0) var<storage, read> in : mat2x2<f32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : mat2x2<f32>;
+@group(0) @binding(1) var<storage, read_write> out : mat2x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat2x3.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : mat2x3<f32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : mat2x3<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat2x3.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : mat2x3<f32>;
+@group(0) @binding(0) var<storage, read> in : mat2x3<f32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : mat2x3<f32>;
+@group(0) @binding(1) var<storage, read_write> out : mat2x3<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat3x2.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : mat3x2<f32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : mat3x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat3x2.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : mat3x2<f32>;
+@group(0) @binding(0) var<storage, read> in : mat3x2<f32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : mat3x2<f32>;
+@group(0) @binding(1) var<storage, read_write> out : mat3x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat4x4.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : mat4x4<f32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : mat4x4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/mat4x4.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : mat4x4<f32>;
+@group(0) @binding(0) var<storage, read> in : mat4x4<f32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : mat4x4<f32>;
+@group(0) @binding(1) var<storage, read_write> out : mat4x4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/runtime_array.wgsl

@@ -2,13 +2,13 @@ struct S {
   f : f32;
 };
 
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : array<S>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : array<S>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out[0] = in[0];
 }

test/buffer/storage/types/runtime_array.wgsl.expected.wgsl

@@ -2,11 +2,11 @@ struct S {
   f : f32;
 }
 
-[[group(0), binding(0)]] var<storage, read> in : array<S>;
+@group(0) @binding(0) var<storage, read> in : array<S>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : array<S>;
+@group(0) @binding(1) var<storage, read_write> out : array<S>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out[0] = in[0];
 }

test/buffer/storage/types/struct.wgsl

@@ -5,13 +5,13 @@ struct S {
   inner : Inner;
 };
 
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : S;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/struct.wgsl.expected.wgsl

@@ -6,11 +6,11 @@ struct S {
   inner : Inner;
 }
 
-[[group(0), binding(0)]] var<storage, read> in : S;
+@group(0) @binding(0) var<storage, read> in : S;
 
-[[group(0), binding(1)]] var<storage, read_write> out : S;
+@group(0) @binding(1) var<storage, read_write> out : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/u32.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : u32;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : u32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/u32.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : u32;
+@group(0) @binding(0) var<storage, read> in : u32;
 
-[[group(0), binding(1)]] var<storage, read_write> out : u32;
+@group(0) @binding(1) var<storage, read_write> out : u32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec2.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : vec2<i32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : vec2<i32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec2.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : vec2<i32>;
+@group(0) @binding(0) var<storage, read> in : vec2<i32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : vec2<i32>;
+@group(0) @binding(1) var<storage, read_write> out : vec2<i32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec3.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : vec3<u32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : vec3<u32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec3.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : vec3<u32>;
+@group(0) @binding(0) var<storage, read> in : vec3<u32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : vec3<u32>;
+@group(0) @binding(1) var<storage, read_write> out : vec3<u32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec4.wgsl

@@ -1,10 +1,10 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<storage, read> in : vec4<f32>;
 
-[[group(0), binding(1)]]
+@group(0) @binding(1)
 var<storage, read_write> out : vec4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/storage/types/vec4.wgsl.expected.wgsl

@@ -1,8 +1,8 @@
-[[group(0), binding(0)]] var<storage, read> in : vec4<f32>;
+@group(0) @binding(0) var<storage, read> in : vec4<f32>;
 
-[[group(0), binding(1)]] var<storage, read_write> out : vec4<f32>;
+@group(0) @binding(1) var<storage, read_write> out : vec4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   out = in;
 }

test/buffer/uniform/dynamic_index/read.wgsl

@@ -8,18 +8,18 @@ struct Inner {
     g : vec2<i32>;
     h : vec2<i32>;
     i : mat2x3<f32>;
-    [[align(16)]] j : mat3x2<f32>;
-    [[align(16)]] k : [[stride(16)]] array<vec4<i32>, 4>;
+    @align(16) j : mat3x2<f32>;
+    @align(16) k : @stride(16) array<vec4<i32>, 4>;
 };
 
 struct S {
     arr : array<Inner, 8>;
 };
 
-[[binding(0), group(0)]] var<uniform> s : S;
+@binding(0) @group(0) var<uniform> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
     let a = s.arr[idx].a;
     let b = s.arr[idx].b;
     let c = s.arr[idx].c;

test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl

@@ -8,20 +8,20 @@ struct Inner {
   g : vec2<i32>;
   h : vec2<i32>;
   i : mat2x3<f32>;
-  [[align(16)]]
+  @align(16)
   j : mat3x2<f32>;
-  [[align(16)]]
-  k : [[stride(16)]] array<vec4<i32>, 4>;
+  @align(16)
+  k : @stride(16) array<vec4<i32>, 4>;
 }
 
 struct S {
   arr : array<Inner, 8>;
 }
 
-[[binding(0), group(0)]] var<uniform> s : S;
+@binding(0) @group(0) var<uniform> s : S;
 
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] idx : u32) {
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
   let a = s.arr[idx].a;
   let b = s.arr[idx].b;
   let c = s.arr[idx].c;

test/buffer/uniform/static_index/read.wgsl

@@ -13,13 +13,13 @@ struct S {
     h : vec2<i32>;
     i : mat2x3<f32>;
     j : mat3x2<f32>;
-    [[align(16)]] k : Inner;
-    [[align(16)]] l : [[stride(16)]] array<Inner, 4>;
+    @align(16) k : Inner;
+    @align(16) l : @stride(16) array<Inner, 4>;
 };
 
-[[binding(0), group(0)]] var<uniform> s : S;
+@binding(0) @group(0) var<uniform> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
     let a = s.a;
     let b = s.b;

test/buffer/uniform/static_index/read.wgsl.expected.wgsl

@@ -13,15 +13,15 @@ struct S {
   h : vec2<i32>;
   i : mat2x3<f32>;
   j : mat3x2<f32>;
-  [[align(16)]]
+  @align(16)
   k : Inner;
-  [[align(16)]]
-  l : [[stride(16)]] array<Inner, 4>;
+  @align(16)
+  l : @stride(16) array<Inner, 4>;
 }
 
-[[binding(0), group(0)]] var<uniform> s : S;
+@binding(0) @group(0) var<uniform> s : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let a = s.a;
   let b = s.b;

test/buffer/uniform/types/array.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : array<vec4<f32>, 4>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/array.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : array<vec4<f32>, 4>;
+@group(0) @binding(0) var<uniform> u : array<vec4<f32>, 4>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/f32.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : f32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/f32.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : f32;
+@group(0) @binding(0) var<uniform> u : f32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/i32.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : i32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/i32.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : i32;
+@group(0) @binding(0) var<uniform> u : i32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat2x2.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : mat2x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : mat2x2<f32>;
+@group(0) @binding(0) var<uniform> u : mat2x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat2x3.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : mat2x3<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : mat2x3<f32>;
+@group(0) @binding(0) var<uniform> u : mat2x3<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat3x2.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : mat3x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : mat3x2<f32>;
+@group(0) @binding(0) var<uniform> u : mat3x2<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat4x4.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : mat4x4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : mat4x4<f32>;
+@group(0) @binding(0) var<uniform> u : mat4x4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/struct.wgsl

@@ -5,10 +5,10 @@ struct S {
   inner : Inner;
 };
 
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/struct.wgsl.expected.wgsl

@@ -6,9 +6,9 @@ struct S {
   inner : Inner;
 }
 
-[[group(0), binding(0)]] var<uniform> u : S;
+@group(0) @binding(0) var<uniform> u : S;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/u32.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : u32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/u32.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : u32;
+@group(0) @binding(0) var<uniform> u : u32;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec2.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : vec2<i32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec2.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : vec2<i32>;
+@group(0) @binding(0) var<uniform> u : vec2<i32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec3.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : vec3<u32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec3.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : vec3<u32>;
+@group(0) @binding(0) var<uniform> u : vec3<u32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec4.wgsl

@@ -1,7 +1,7 @@
-[[group(0), binding(0)]]
+@group(0) @binding(0)
 var<uniform> u : vec4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/buffer/uniform/types/vec4.wgsl.expected.wgsl

@@ -1,6 +1,6 @@
-[[group(0), binding(0)]] var<uniform> u : vec4<f32>;
+@group(0) @binding(0) var<uniform> u : vec4<f32>;
 
-[[stage(compute), workgroup_size(1)]]
+@stage(compute) @workgroup_size(1)
 fn main() {
   let x = u;
 }

test/bug/chromium/1251009.wgsl

@@ -1,19 +1,19 @@
 struct VertexInputs0 {
-  [[builtin(vertex_index)]] vertex_index : u32;
-  [[location(0)]] loc0 : i32;
+  @builtin(vertex_index) vertex_index : u32;
+  @location(0) loc0 : i32;
 };
 struct VertexInputs1 {
-  [[location(2)]] loc1 : u32;
-  [[location(3)]] loc3 : vec4<f32>;
+  @location(2) loc1 : u32;
+  @location(3) loc3 : vec4<f32>;
 };
 
-[[stage(vertex)]]
+@stage(vertex)
 fn main(
   inputs0 : VertexInputs0,
-  [[location(1)]] loc1 : u32,
-  [[builtin(instance_index)]] instance_index : u32,
+  @location(1) loc1 : u32,
+  @builtin(instance_index) instance_index : u32,
   inputs1 : VertexInputs1,
-) -> [[builtin(position)]] vec4<f32> {
+) -> @builtin(position) vec4<f32> {
   let foo : u32 = inputs0.vertex_index + instance_index;
   return vec4<f32>();
 }

test/bug/chromium/1251009.wgsl.expected.wgsl

@@ -1,19 +1,19 @@
 struct VertexInputs0 {
-  [[builtin(vertex_index)]]
+  @builtin(vertex_index)
   vertex_index : u32;
-  [[location(0)]]
+  @location(0)
   loc0 : i32;
 }
 
 struct VertexInputs1 {
-  [[location(2)]]
+  @location(2)
   loc1 : u32;
-  [[location(3)]]
+  @location(3)
   loc3 : vec4<f32>;
 }
 
-[[stage(vertex)]]
-fn main(inputs0 : VertexInputs0, [[location(1)]] loc1 : u32, [[builtin(instance_index)]] instance_index : u32, inputs1 : VertexInputs1) -> [[builtin(position)]] vec4<f32> {
+@stage(vertex)
+fn main(inputs0 : VertexInputs0, @location(1) loc1 : u32, @builtin(instance_index) instance_index : u32, inputs1 : VertexInputs1) -> @builtin(position) vec4<f32> {
   let foo : u32 = (inputs0.vertex_index + instance_index);
   return vec4<f32>();
 }