diff --git a/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl b/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl
index 4677d18..a512d50 100644
--- a/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_05ce15();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_05ce15();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_05ce15();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl.expected.wgsl
index 2627d0e..caefa60 100644
--- a/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/05ce15.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_05ce15();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_05ce15();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_05ce15();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl b/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl
index e959086..3843b42 100644
--- a/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_064c7f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_064c7f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_064c7f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl.expected.wgsl
index 042b3af..9d55259 100644
--- a/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/064c7f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_064c7f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_064c7f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_064c7f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/068641.wgsl b/test/tint/builtins/gen/literal/textureStore/068641.wgsl
index 86edaac..0bad849 100644
--- a/test/tint/builtins/gen/literal/textureStore/068641.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/068641.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_068641();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_068641();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_068641();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/068641.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/068641.wgsl.expected.wgsl
index 5b3279a..1e3ae92 100644
--- a/test/tint/builtins/gen/literal/textureStore/068641.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/068641.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_068641();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_068641();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_068641();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl b/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl
index 53f97f4..1ffa433 100644
--- a/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_0af6b5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_0af6b5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_0af6b5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl.expected.wgsl
index 1e67b08..d6371ca 100644
--- a/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/0af6b5.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_0af6b5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_0af6b5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_0af6b5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl b/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl
index f0100ff..1034365 100644
--- a/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_0c3dff();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_0c3dff();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_0c3dff();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl.expected.wgsl
index b115e07..14d2872 100644
--- a/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/0c3dff.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_0c3dff();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_0c3dff();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_0c3dff();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/102722.wgsl b/test/tint/builtins/gen/literal/textureStore/102722.wgsl
index 4c9a7dc..6470bd7 100644
--- a/test/tint/builtins/gen/literal/textureStore/102722.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/102722.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_102722();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_102722();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_102722();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/102722.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/102722.wgsl.expected.wgsl
index 593c342..2d6fc38 100644
--- a/test/tint/builtins/gen/literal/textureStore/102722.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/102722.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_102722();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_102722();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_102722();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl b/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl
index b52633a..9609af7 100644
--- a/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_1bbd08();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_1bbd08();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_1bbd08();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl.expected.wgsl
index a775be7..b6aa3ff 100644
--- a/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/1bbd08.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_1bbd08();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_1bbd08();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_1bbd08();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl b/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl
index a91041c..aebd170 100644
--- a/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_1c02e7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_1c02e7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_1c02e7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl.expected.wgsl
index 732759f..556621d 100644
--- a/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/1c02e7.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_1c02e7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_1c02e7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_1c02e7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/22d955.wgsl b/test/tint/builtins/gen/literal/textureStore/22d955.wgsl
index 5e6f372..f9ce1a3 100644
--- a/test/tint/builtins/gen/literal/textureStore/22d955.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/22d955.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_22d955();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_22d955();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_22d955();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/22d955.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/22d955.wgsl.expected.wgsl
index b9fe255..cfb291c 100644
--- a/test/tint/builtins/gen/literal/textureStore/22d955.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/22d955.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_22d955();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_22d955();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_22d955();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl b/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl
index 8cb6f1e..b253a23 100644
--- a/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_26bf70();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_26bf70();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_26bf70();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl.expected.wgsl
index 896b4f7..10aa757 100644
--- a/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/26bf70.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_26bf70();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_26bf70();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_26bf70();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl b/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl
index 82ce365..2afae6f 100644
--- a/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2796b4();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2796b4();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2796b4();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl.expected.wgsl
index 3fcb394..60ac36f 100644
--- a/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2796b4.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2796b4();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2796b4();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2796b4();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl b/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl
index 39eebf3..7fb5d9f 100644
--- a/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2ac6c7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2ac6c7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2ac6c7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl.expected.wgsl
index 2786ca4..ca7cc4f 100644
--- a/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2ac6c7.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2ac6c7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2ac6c7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2ac6c7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl b/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl
index 05df4c8..dbbc9a0 100644
--- a/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2eb2a4();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2eb2a4();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2eb2a4();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl.expected.wgsl
index 51c5cd9..5399668 100644
--- a/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2eb2a4.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2eb2a4();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2eb2a4();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2eb2a4();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl b/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl
index 9520984..b37f8e4 100644
--- a/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2ed2a3();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2ed2a3();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2ed2a3();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl.expected.wgsl
index c415d58..daad961 100644
--- a/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/2ed2a3.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_2ed2a3();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_2ed2a3();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_2ed2a3();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/31745b.wgsl b/test/tint/builtins/gen/literal/textureStore/31745b.wgsl
index f26b550..115890b 100644
--- a/test/tint/builtins/gen/literal/textureStore/31745b.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/31745b.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_31745b();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_31745b();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_31745b();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/31745b.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/31745b.wgsl.expected.wgsl
index 4288434..0280681 100644
--- a/test/tint/builtins/gen/literal/textureStore/31745b.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/31745b.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_31745b();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_31745b();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_31745b();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/32f368.wgsl b/test/tint/builtins/gen/literal/textureStore/32f368.wgsl
index 33b9e04..83cfdbc 100644
--- a/test/tint/builtins/gen/literal/textureStore/32f368.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/32f368.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_32f368();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_32f368();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_32f368();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/32f368.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/32f368.wgsl.expected.wgsl
index a747822..d0d3f8b 100644
--- a/test/tint/builtins/gen/literal/textureStore/32f368.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/32f368.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_32f368();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_32f368();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_32f368();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/331aee.wgsl b/test/tint/builtins/gen/literal/textureStore/331aee.wgsl
index 1099995..9fc7d04 100644
--- a/test/tint/builtins/gen/literal/textureStore/331aee.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/331aee.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_331aee();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_331aee();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_331aee();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/331aee.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/331aee.wgsl.expected.wgsl
index 0363260..c56aa25 100644
--- a/test/tint/builtins/gen/literal/textureStore/331aee.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/331aee.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_331aee();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_331aee();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_331aee();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl b/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl
index 68b8504..a70f91a 100644
--- a/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_38e8d7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_38e8d7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_38e8d7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl.expected.wgsl
index 28d1791..e586ff1 100644
--- a/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/38e8d7.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_38e8d7();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_38e8d7();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_38e8d7();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl b/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl
index bce15d4..f8a91fc 100644
--- a/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3a52ac();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3a52ac();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3a52ac();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl.expected.wgsl
index 6fbbab0..a3aa110 100644
--- a/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3a52ac.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3a52ac();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3a52ac();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3a52ac();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl b/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl
index 973de1b..d1a8819 100644
--- a/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3bb7a1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3bb7a1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3bb7a1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl.expected.wgsl
index ccb304d..b1b1b8d 100644
--- a/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3bb7a1.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3bb7a1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3bb7a1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3bb7a1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl b/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl
index 1c7fe6a..687108a 100644
--- a/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3bec15();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3bec15();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3bec15();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl.expected.wgsl
index c594be0..96e2996 100644
--- a/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/3bec15.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_3bec15();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_3bec15();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_3bec15();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl b/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl
index 9669e51..ff6c199 100644
--- a/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_441ba8();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_441ba8();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_441ba8();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl.expected.wgsl
index b59255a..6014c40 100644
--- a/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/441ba8.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_441ba8();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_441ba8();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_441ba8();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl b/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl
index 23c5849..32988ef 100644
--- a/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_4fc057();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_4fc057();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_4fc057();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl.expected.wgsl
index 26a0018..a599ff7 100644
--- a/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/4fc057.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_4fc057();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_4fc057();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_4fc057();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl b/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl
index f0ceb8f..2f1873f 100644
--- a/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_5a2f8f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_5a2f8f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_5a2f8f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl.expected.wgsl
index 13d97f4..9e4b11f 100644
--- a/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/5a2f8f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_5a2f8f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_5a2f8f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_5a2f8f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/60975f.wgsl b/test/tint/builtins/gen/literal/textureStore/60975f.wgsl
index 8ad1e46..879c1f4 100644
--- a/test/tint/builtins/gen/literal/textureStore/60975f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/60975f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_60975f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_60975f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_60975f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/60975f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/60975f.wgsl.expected.wgsl
index 47cdd34..c723d57 100644
--- a/test/tint/builtins/gen/literal/textureStore/60975f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/60975f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_60975f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_60975f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_60975f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl b/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl
index fefa207..551f5d7 100644
--- a/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_682fd6();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_682fd6();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_682fd6();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl.expected.wgsl
index 20603cc..15cae1d 100644
--- a/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/682fd6.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_682fd6();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_682fd6();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_682fd6();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl b/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl
index ebceb55..3d991ae 100644
--- a/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6b75c3();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6b75c3();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6b75c3();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl.expected.wgsl
index cd827a3..da091c3 100644
--- a/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6b75c3.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6b75c3();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6b75c3();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6b75c3();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl b/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl
index ab22858..ca34d14 100644
--- a/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6b80d2();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6b80d2();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6b80d2();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl.expected.wgsl
index e7efff2..85e0231 100644
--- a/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6b80d2.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6b80d2();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6b80d2();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6b80d2();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl b/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl
index f00cc00..3f6d4ad 100644
--- a/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6cff2e();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6cff2e();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6cff2e();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl.expected.wgsl
index 4f43b98..bc09df8 100644
--- a/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6cff2e.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6cff2e();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6cff2e();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6cff2e();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6da692.wgsl b/test/tint/builtins/gen/literal/textureStore/6da692.wgsl
index ca7410c..d1561ce 100644
--- a/test/tint/builtins/gen/literal/textureStore/6da692.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6da692.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6da692();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6da692();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6da692();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/6da692.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/6da692.wgsl.expected.wgsl
index 2497a69..ba262d3 100644
--- a/test/tint/builtins/gen/literal/textureStore/6da692.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/6da692.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_6da692();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_6da692();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_6da692();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/731349.wgsl b/test/tint/builtins/gen/literal/textureStore/731349.wgsl
index c8f81d6..b8cda8e 100644
--- a/test/tint/builtins/gen/literal/textureStore/731349.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/731349.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_731349();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_731349();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_731349();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/731349.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/731349.wgsl.expected.wgsl
index d5466d8..b659a87 100644
--- a/test/tint/builtins/gen/literal/textureStore/731349.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/731349.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_731349();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_731349();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_731349();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/752da6.wgsl b/test/tint/builtins/gen/literal/textureStore/752da6.wgsl
index e2b0734..65b7216 100644
--- a/test/tint/builtins/gen/literal/textureStore/752da6.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/752da6.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_752da6();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_752da6();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_752da6();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/752da6.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/752da6.wgsl.expected.wgsl
index e029549a8..9f4e06d 100644
--- a/test/tint/builtins/gen/literal/textureStore/752da6.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/752da6.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_752da6();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_752da6();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_752da6();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl b/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl
index a5b9a55..bf17f8d 100644
--- a/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_77c0ae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_77c0ae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_77c0ae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl.expected.wgsl
index 56cf079..ddc73be 100644
--- a/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/77c0ae.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_77c0ae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_77c0ae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_77c0ae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl b/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl
index 282ed90..368bfcd 100644
--- a/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_7cec8d();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_7cec8d();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_7cec8d();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl.expected.wgsl
index 89ed161..5218558 100644
--- a/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/7cec8d.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_7cec8d();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_7cec8d();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_7cec8d();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl b/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl
index 8b5d0dd..6a205dd 100644
--- a/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_7f7fae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_7f7fae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_7f7fae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl.expected.wgsl
index 438123c..eba8d52 100644
--- a/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/7f7fae.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_7f7fae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_7f7fae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_7f7fae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/804942.wgsl b/test/tint/builtins/gen/literal/textureStore/804942.wgsl
index 4024631..cc50e1b 100644
--- a/test/tint/builtins/gen/literal/textureStore/804942.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/804942.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_804942();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_804942();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_804942();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/804942.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/804942.wgsl.expected.wgsl
index f4a325c..06f7057 100644
--- a/test/tint/builtins/gen/literal/textureStore/804942.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/804942.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_804942();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_804942();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_804942();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/805dae.wgsl b/test/tint/builtins/gen/literal/textureStore/805dae.wgsl
index 65e4420..3ea3e2d 100644
--- a/test/tint/builtins/gen/literal/textureStore/805dae.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/805dae.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_805dae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_805dae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_805dae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/805dae.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/805dae.wgsl.expected.wgsl
index 667452d..7bc5067 100644
--- a/test/tint/builtins/gen/literal/textureStore/805dae.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/805dae.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_805dae();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_805dae();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_805dae();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl b/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl
index 55d29c6..fd88fc1 100644
--- a/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_83bcc1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_83bcc1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_83bcc1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl.expected.wgsl
index f0ee863..e7e464e 100644
--- a/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/83bcc1.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_83bcc1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_83bcc1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_83bcc1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/872747.wgsl b/test/tint/builtins/gen/literal/textureStore/872747.wgsl
index 98ccead..15b1772 100644
--- a/test/tint/builtins/gen/literal/textureStore/872747.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/872747.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_872747();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_872747();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_872747();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/872747.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/872747.wgsl.expected.wgsl
index b5c2914..50b7106 100644
--- a/test/tint/builtins/gen/literal/textureStore/872747.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/872747.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_872747();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_872747();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_872747();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl b/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl
index 271ed67..6f8b386 100644
--- a/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_8e0479();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_8e0479();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_8e0479();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl.expected.wgsl
index 07408ae..c831783 100644
--- a/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/8e0479.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_8e0479();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_8e0479();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_8e0479();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl b/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl
index 4d840ce..5f53a96 100644
--- a/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_8f71a1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_8f71a1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_8f71a1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl.expected.wgsl
index 9bacd31..5ce5918 100644
--- a/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/8f71a1.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_8f71a1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_8f71a1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_8f71a1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/969534.wgsl b/test/tint/builtins/gen/literal/textureStore/969534.wgsl
index d73a34d..809d51e 100644
--- a/test/tint/builtins/gen/literal/textureStore/969534.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/969534.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_969534();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_969534();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_969534();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/969534.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/969534.wgsl.expected.wgsl
index 467cbda..41ae906 100644
--- a/test/tint/builtins/gen/literal/textureStore/969534.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/969534.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_969534();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_969534();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_969534();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl b/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl
index 72656e1..513ef4e 100644
--- a/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9a3ecc();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9a3ecc();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9a3ecc();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl.expected.wgsl
index 6fe66cb..072ae70 100644
--- a/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9a3ecc.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9a3ecc();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9a3ecc();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9a3ecc();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl b/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl
index 484a3cd..ee264b2 100644
--- a/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9d9cd5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9d9cd5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9d9cd5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl.expected.wgsl
index 346f010..f321044 100644
--- a/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9d9cd5.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9d9cd5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9d9cd5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9d9cd5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl b/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl
index da9282e..1f6eefc 100644
--- a/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9e3ec5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9e3ec5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9e3ec5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl.expected.wgsl
index aaa8bf6..b2822b3 100644
--- a/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/9e3ec5.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_9e3ec5();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_9e3ec5();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_9e3ec5();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl b/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl
index 42d2b30..24023d5 100644
--- a/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ac67aa();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ac67aa();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ac67aa();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl.expected.wgsl
index 21d0d22..c3095f6 100644
--- a/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ac67aa.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ac67aa();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ac67aa();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ac67aa();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl b/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl
index 949c932..bcc47da 100644
--- a/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_b706b1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_b706b1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_b706b1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl.expected.wgsl
index 7b77dc8..3c379e9 100644
--- a/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/b706b1.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_b706b1();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_b706b1();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_b706b1();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl b/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl
index ee04924..f53663a 100644
--- a/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_bbcb7f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_bbcb7f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_bbcb7f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl.expected.wgsl
index 8c4bb4d..3f24147 100644
--- a/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/bbcb7f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_bbcb7f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_bbcb7f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_bbcb7f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl b/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl
index c192581..6546e46 100644
--- a/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_be6e30();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_be6e30();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_be6e30();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl.expected.wgsl
index 3a4c747..cf77394 100644
--- a/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/be6e30.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_be6e30();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_be6e30();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_be6e30();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl b/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl
index e8a7f14..bdb2d08 100644
--- a/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_bf775c();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_bf775c();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_bf775c();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl.expected.wgsl
index db5d80e..47b4f23 100644
--- a/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/bf775c.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_bf775c();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_bf775c();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_bf775c();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl b/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl
index 716f491..4c10f22 100644
--- a/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_c5af1e();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_c5af1e();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_c5af1e();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl.expected.wgsl
index 02c2ef8..003aee2 100644
--- a/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/c5af1e.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_c5af1e();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_c5af1e();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_c5af1e();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/c863be.wgsl b/test/tint/builtins/gen/literal/textureStore/c863be.wgsl
index abb287f..b927835 100644
--- a/test/tint/builtins/gen/literal/textureStore/c863be.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/c863be.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_c863be();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_c863be();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_c863be();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/c863be.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/c863be.wgsl.expected.wgsl
index 46b5a69..7c6f8ac 100644
--- a/test/tint/builtins/gen/literal/textureStore/c863be.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/c863be.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_c863be();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_c863be();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_c863be();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl b/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl
index 44468a8..5b73aef 100644
--- a/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_d73b5c();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_d73b5c();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_d73b5c();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl.expected.wgsl
index 2b09cbc..bed8187 100644
--- a/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/d73b5c.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_d73b5c();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_d73b5c();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_d73b5c();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl b/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl
index 8ea8125..b1cd4fa 100644
--- a/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_dd7d81();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_dd7d81();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_dd7d81();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl.expected.wgsl
index f7f4bc1..4c1f529 100644
--- a/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/dd7d81.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_dd7d81();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_dd7d81();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_dd7d81();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/dde364.wgsl b/test/tint/builtins/gen/literal/textureStore/dde364.wgsl
index e09a8e5..93f5a32 100644
--- a/test/tint/builtins/gen/literal/textureStore/dde364.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/dde364.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_dde364();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_dde364();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_dde364();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/dde364.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/dde364.wgsl.expected.wgsl
index 044b09d..f7dc0c1 100644
--- a/test/tint/builtins/gen/literal/textureStore/dde364.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/dde364.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_dde364();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_dde364();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_dde364();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl b/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl
index a6706d5..2689758 100644
--- a/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_e885e8();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_e885e8();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_e885e8();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl.expected.wgsl
index eac4ea6..f7b3caa 100644
--- a/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/e885e8.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_e885e8();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_e885e8();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_e885e8();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl b/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl
index 916fc32..5e80ec5 100644
--- a/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_eb702f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_eb702f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_eb702f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl.expected.wgsl
index c2b8653..0e86ad2 100644
--- a/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/eb702f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_eb702f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_eb702f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_eb702f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl b/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl
index 6d0e377..4888ed2 100644
--- a/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_eb78b9();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_eb78b9();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_eb78b9();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl.expected.wgsl
index 09472ae..c200bb2 100644
--- a/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/eb78b9.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_eb78b9();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_eb78b9();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_eb78b9();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl b/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl
index a5b15c1..5be75a6 100644
--- a/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ee6acc();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ee6acc();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ee6acc();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl.expected.wgsl
index 8b5b0ab..377310e 100644
--- a/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ee6acc.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<f32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ee6acc();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ee6acc();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ee6acc();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl b/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl
index b3effb6..d2a3935 100644
--- a/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ef9f2f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ef9f2f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ef9f2f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl.expected.wgsl
index 72cb6d0..9a0b5cb 100644
--- a/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/ef9f2f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_ef9f2f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_ef9f2f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_ef9f2f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl b/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl
index b49ebc3..f3054ef 100644
--- a/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_f8dead();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_f8dead();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_f8dead();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl.expected.wgsl
index eabfbca..450fd25 100644
--- a/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/f8dead.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec3<i32>(), vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_f8dead();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_f8dead();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_f8dead();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl b/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl
index cb34f7a..5677404 100644
--- a/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_f9be83();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_f9be83();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_f9be83();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl.expected.wgsl
index c6bb3c4..0b0ece1 100644
--- a/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/f9be83.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_f9be83();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_f9be83();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_f9be83();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl b/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl
index fa69110..3186385 100644
--- a/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_fb9a8f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_fb9a8f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_fb9a8f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl.expected.wgsl
index 27aa300..e1b6643 100644
--- a/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/fb9a8f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, 1, vec4<u32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_fb9a8f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_fb9a8f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_fb9a8f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl b/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl
index 541f45c..b65045b 100644
--- a/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl
@@ -29,18 +29,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_fbf53f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_fbf53f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_fbf53f();
 }
diff --git a/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl.expected.wgsl
index 8ee5316..dd790f2 100644
--- a/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl.expected.wgsl
+++ b/test/tint/builtins/gen/literal/textureStore/fbf53f.wgsl.expected.wgsl
@@ -4,18 +4,18 @@
   textureStore(arg_0, vec2<i32>(), 1, vec4<i32>());
 }
 
-@stage(vertex)
+@vertex
 fn vertex_main() -> @builtin(position) vec4<f32> {
   textureStore_fbf53f();
   return vec4<f32>();
 }
 
-@stage(fragment)
+@fragment
 fn fragment_main() {
   textureStore_fbf53f();
 }
 
-@stage(compute) @workgroup_size(1)
+@compute @workgroup_size(1)
 fn compute_main() {
   textureStore_fbf53f();
 }