validation: compute shader must include 'workgroup_size' in its attributes

Bug: tint:884
Change-Id: If96c6df3247fee142a779117fa26d006afd4f7ef
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/55680
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/inspector/inspector_test.cc

@@ -787,9 +787,9 @@ TEST_F(InspectorGetEntryPointTest, MultipleEntryPoints) {
                                    Stage(ast::PipelineStage::kFragment),
                                });
 
-  MakeEmptyBodyFunction("bar", ast::DecorationList{
+  MakeEmptyBodyFunction("bar",
-                                   Stage(ast::PipelineStage::kCompute),
+                        ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                               });
+                                            WorkgroupSize(1)});
 
   // TODO(dsinclair): Update to run the namer transform when available.
 
@@ -810,10 +810,10 @@ TEST_F(InspectorGetEntryPointTest, MultipleEntryPoints) {
 TEST_F(InspectorGetEntryPointTest, MixFunctionsAndEntryPoints) {
   MakeEmptyBodyFunction("func", {});
 
-  MakeCallerBodyFunction("foo", {"func"},
+  MakeCallerBodyFunction(
-                         ast::DecorationList{
+      "foo", {"func"},
-                             Stage(ast::PipelineStage::kCompute),
+      ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                         });
+                          WorkgroupSize(1)});
 
   MakeCallerBodyFunction("bar", {"func"},
                          ast::DecorationList{
@@ -837,9 +837,9 @@ TEST_F(InspectorGetEntryPointTest, MixFunctionsAndEntryPoints) {
 }
 
 TEST_F(InspectorGetEntryPointTest, DefaultWorkgroupSize) {
-  MakeEmptyBodyFunction("foo", ast::DecorationList{
+  MakeEmptyBodyFunction("foo",
-                                   Stage(ast::PipelineStage::kCompute),
+                        ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                               });
+                                            WorkgroupSize(8, 2, 1)});
 
   Inspector& inspector = Build();
 
@@ -849,8 +849,8 @@ TEST_F(InspectorGetEntryPointTest, DefaultWorkgroupSize) {
   ASSERT_EQ(1u, result.size());
   uint32_t x, y, z;
   std::tie(x, y, z) = result[0].workgroup_size();
-  EXPECT_EQ(1u, x);
+  EXPECT_EQ(8u, x);
-  EXPECT_EQ(1u, y);
+  EXPECT_EQ(2u, y);
   EXPECT_EQ(1u, z);
 }
 
@@ -1329,10 +1329,10 @@ TEST_F(InspectorGetEntryPointTest, MultipleEntryPointsInOutVariables_Legacy) {
                                     Stage(ast::PipelineStage::kFragment),
                                 });
 
-  MakeInOutVariableBodyFunction("bar", {{"in2_var", "out_var"}},
+  MakeInOutVariableBodyFunction(
-                                ast::DecorationList{
+      "bar", {{"in2_var", "out_var"}},
-                                    Stage(ast::PipelineStage::kCompute),
+      ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                                });
+                          WorkgroupSize(1)});
 
   // TODO(dsinclair): Update to run the namer transform when
   // available.
@@ -1387,10 +1387,10 @@ TEST_F(InspectorGetEntryPointTest,
                                           Stage(ast::PipelineStage::kFragment),
                                       });
 
-  MakeCallerBodyFunction("bar", {"func"},
+  MakeCallerBodyFunction(
-                         ast::DecorationList{
+      "bar", {"func"},
-                             Stage(ast::PipelineStage::kCompute),
+      ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                         });
+                          WorkgroupSize(1)});
 
   // TODO(dsinclair): Update to run the namer transform when
   // available.
@@ -1482,7 +1482,8 @@ TEST_F(InspectorGetEntryPointTest, BuiltInsNotStageVariables_Legacy) {
 
 TEST_F(InspectorGetEntryPointTest, OverridableConstantUnreferenced) {
   AddOverridableConstantWithoutID<float>("foo", ty.f32(), nullptr);
-  MakeEmptyBodyFunction("ep_func", {Stage(ast::PipelineStage::kCompute)});
+  MakeEmptyBodyFunction(
+      "ep_func", {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   Inspector& inspector = Build();
 
@@ -1494,8 +1495,9 @@ TEST_F(InspectorGetEntryPointTest, OverridableConstantUnreferenced) {
 
 TEST_F(InspectorGetEntryPointTest, OverridableConstantReferencedByEntryPoint) {
   AddOverridableConstantWithoutID<float>("foo", ty.f32(), nullptr);
-  MakeConstReferenceBodyFunction("ep_func", "foo", ty.f32(),
+  MakeConstReferenceBodyFunction(
-                                 {Stage(ast::PipelineStage::kCompute)});
+      "ep_func", "foo", ty.f32(),
+      {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   Inspector& inspector = Build();
 
@@ -1512,8 +1514,9 @@ TEST_F(InspectorGetEntryPointTest, OverridableConstantReferencedByEntryPoint) {
 TEST_F(InspectorGetEntryPointTest, OverridableConstantReferencedByCallee) {
   AddOverridableConstantWithoutID<float>("foo", ty.f32(), nullptr);
   MakeConstReferenceBodyFunction("callee_func", "foo", ty.f32(), {});
-  MakeCallerBodyFunction("ep_func", {"callee_func"},
+  MakeCallerBodyFunction(
-                         {Stage(ast::PipelineStage::kCompute)});
+      "ep_func", {"callee_func"},
+      {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   Inspector& inspector = Build();
 
@@ -1528,8 +1531,9 @@ TEST_F(InspectorGetEntryPointTest, OverridableConstantSomeReferenced) {
   AddOverridableConstantWithID<float>("foo", 1, ty.f32(), nullptr);
   AddOverridableConstantWithID<float>("bar", 2, ty.f32(), nullptr);
   MakeConstReferenceBodyFunction("callee_func", "foo", ty.f32(), {});
-  MakeCallerBodyFunction("ep_func", {"callee_func"},
+  MakeCallerBodyFunction(
-                         {Stage(ast::PipelineStage::kCompute)});
+      "ep_func", {"callee_func"},
+      {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   Inspector& inspector = Build();
 
@@ -1606,9 +1610,9 @@ TEST_F(InspectorGetRemappedNameForEntryPointTest,
   // TODO(dsinclair): Update to run the namer transform when
   // available.
 
-  MakeEmptyBodyFunction("bar", ast::DecorationList{
+  MakeEmptyBodyFunction("bar",
-                                   Stage(ast::PipelineStage::kCompute),
+                        ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                               });
+                                            WorkgroupSize(1)});
 
   Inspector& inspector = Build();
 

src/program_builder.cc

@@ -144,7 +144,8 @@ ast::Statement* ProgramBuilder::WrapInStatement(ast::Statement* stmt) {
 
 ast::Function* ProgramBuilder::WrapInFunction(ast::StatementList stmts) {
   return Func("test_function", {}, ty.void_(), std::move(stmts),
-              {create<ast::StageDecoration>(ast::PipelineStage::kCompute)});
+              {create<ast::StageDecoration>(ast::PipelineStage::kCompute),
+               WorkgroupSize(1, 1, 1)});
 }
 
 }  // namespace tint

src/resolver/builtins_validation_test.cc

@@ -273,7 +273,7 @@ TEST_F(ResolverBuiltinsValidationTest, VertexBuiltin_Pass) {
 }
 
 TEST_F(ResolverBuiltinsValidationTest, ComputeBuiltin_Pass) {
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn main(
   //   [[builtin(local_invocationId)]] li_id: vec3<u32>,
   //   [[builtin(local_invocationIndex)]] li_index: u32,

src/resolver/decoration_validation_test.cc

@@ -271,7 +271,7 @@ TEST_P(EntryPointReturnTypeDecorationTest, IsValid) {
 
   Func("main", ast::VariableList{}, ty.vec4<f32>(),
        {Return(Construct(ty.vec4<f32>(), 1.f))},
-       {Stage(ast::PipelineStage::kCompute)},
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)},
        createDecorations({}, *this, params.kind));
 
   if (params.should_pass) {
@@ -933,6 +933,24 @@ namespace WorkgroupDecorationTests {
 namespace {
 
 using WorkgroupDecoration = ResolverTest;
+TEST_F(WorkgroupDecoration, ComputeShaderPass) {
+  Func("main", {}, ty.void_(), {},
+       {Stage(ast::PipelineStage::kCompute),
+        create<ast::WorkgroupDecoration>(Source{{12, 34}}, Expr(1))});
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(WorkgroupDecoration, Missing) {
+  Func(Source{{12, 34}}, "main", {}, ty.void_(), {},
+       {Stage(ast::PipelineStage::kCompute)});
+
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      "12:34 error: a compute shader must include 'workgroup_size' in its "
+      "attributes");
+}
 
 TEST_F(WorkgroupDecoration, NotAnEntryPoint) {
   Func("main", {}, ty.void_(), {},

src/resolver/intrinsic_validation_test.cc

@@ -49,12 +49,12 @@ using IntrinsicType = sem::IntrinsicType;
 using ResolverIntrinsicValidationTest = ResolverTest;
 
 TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageDirect) {
-  // [[stage(compute)]] fn func { return dpdx(1.0); }
+  // [[stage(compute), workgroup_size(1)]] fn func { return dpdx(1.0); }
 
   auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"),
                                            ast::ExpressionList{Expr(1.0f)});
   Func(Source{{1, 2}}, "func", ast::VariableList{}, ty.void_(), {Ignore(dpdx)},
-       {Stage(ast::PipelineStage::kCompute)});
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   EXPECT_FALSE(r()->Resolve());
   EXPECT_EQ(r()->error(),
@@ -65,7 +65,7 @@ TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageIndirect) {
   // fn f0 { return dpdx(1.0); }
   // fn f1 { f0(); }
   // fn f2 { f1(); }
-  // [[stage(compute)]] fn main { return f2(); }
+  // [[stage(compute), workgroup_size(1)]] fn main { return f2(); }
 
   auto* dpdx = create<ast::CallExpression>(Source{{3, 4}}, Expr("dpdx"),
                                            ast::ExpressionList{Expr(1.0f)});
@@ -78,7 +78,8 @@ TEST_F(ResolverIntrinsicValidationTest, InvalidPipelineStageIndirect) {
        {Ignore(Call("f1"))});
 
   Func(Source{{7, 8}}, "main", ast::VariableList{}, ty.void_(),
-       {Ignore(Call("f2"))}, {Stage(ast::PipelineStage::kCompute)});
+       {Ignore(Call("f2"))},
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   EXPECT_FALSE(r()->Resolve());
   EXPECT_EQ(r()->error(),

src/resolver/resolver.cc

@@ -1365,6 +1365,16 @@ bool Resolver::ValidateEntryPoint(const ast::Function* func,
     }
   }
 
+  if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
+    if (!ast::HasDecoration<ast::WorkgroupDecoration>(func->decorations())) {
+      AddError(
+          "a compute shader must include 'workgroup_size' in its "
+          "attributes",
+          func->source());
+      return false;
+    }
+  }
+
   // Validate there are no resource variable binding collisions
   std::unordered_map<sem::BindingPoint, const ast::Variable*> binding_points;
   for (auto* var_info : info->referenced_module_vars) {

src/resolver/resolver_test.cc

@@ -1107,7 +1107,7 @@ TEST_F(ResolverTest, Function_ReturnStatements) {
 }
 
 TEST_F(ResolverTest, Function_WorkgroupSize_NotSet) {
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn main() {}
   auto* func = Func("main", ast::VariableList{}, ty.void_(), {}, {});
 
@@ -2067,17 +2067,15 @@ TEST_F(ResolverTest, Function_EntryPoints_StageDecoration) {
                         Assign("call_a", Call("a")),
                         Assign("call_b", Call("b")),
                     },
-                    ast::DecorationList{
+                    ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                        Stage(ast::PipelineStage::kCompute),
+                                        WorkgroupSize(1)});
-                    });
 
   auto* ep_2 = Func("ep_2", params, ty.void_(),
                     {
                         Assign("call_c", Call("c")),
                     },
-                    ast::DecorationList{
+                    ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                        Stage(ast::PipelineStage::kCompute),
+                                        WorkgroupSize(1)});
-                    });
 
   ASSERT_TRUE(r()->Resolve()) << r()->error();
 
@@ -2154,9 +2152,7 @@ TEST_F(ResolverTest, Function_EntryPoints_LinearTime) {
            create<ast::CallStatement>(Call(fn_a(0))),
            create<ast::CallStatement>(Call(fn_b(0))),
        },
-       {
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
-           Stage(ast::PipelineStage::kCompute),
-       });
 
   ASSERT_TRUE(r()->Resolve()) << r()->error();
 }

src/resolver/struct_pipeline_stage_use_test.cc

@@ -126,7 +126,7 @@ TEST_F(ResolverPipelineStageUseTest, StructUsedAsComputeShaderParam) {
       {Member("a", ty.u32(), {Builtin(ast::Builtin::kLocalInvocationIndex)})});
 
   Func("main", {Param("param", ty.Of(s))}, ty.void_(), {},
-       {Stage(ast::PipelineStage::kCompute)});
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   ASSERT_TRUE(r()->Resolve()) << r()->error();
 

src/transform/array_length_from_uniform_test.cc

@@ -70,7 +70,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
@@ -92,7 +92,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
 }
@@ -123,7 +123,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
@@ -146,7 +146,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = ((tint_symbol_1.buffer_size[0u] - 8u) / 64u);
 }
@@ -184,7 +184,7 @@ struct SB2 {
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -216,7 +216,7 @@ struct SB2 {
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
   var len2 : u32 = ((tint_symbol_1.buffer_size[1u] - 16u) / 16u);
@@ -249,7 +249,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   ignore(&(sb.arr));
 }
@@ -287,7 +287,7 @@ struct SB2 {
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));

src/transform/binding_remapper_test.cc

@@ -34,7 +34,7 @@ struct S {
 
 [[group(3), binding(2)]] var<storage, read> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -59,7 +59,7 @@ struct S {
 
 [[group(3), binding(2)]] var<storage, read> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -73,7 +73,7 @@ struct S {
 
 [[group(3), binding(2)]] var<storage, read> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -103,7 +103,7 @@ struct S {
 
 [[group(4), binding(3)]] var<storage, read> c : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -119,7 +119,7 @@ struct S {
 
 [[group(4), binding(3)]] var<storage, read> c : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -157,7 +157,7 @@ type A = S;
 
 [[group(4), binding(3)]] var<storage> c : A;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -179,7 +179,7 @@ type A = S;
 
 [[group(4), binding(3)]] var<storage, write> c : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -207,7 +207,7 @@ struct S {
 
 [[group(3), binding(2)]] var<storage, read> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -221,7 +221,7 @@ struct S {
 
 [[group(6), binding(7)]] var<storage, write> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
 }
 )";
@@ -256,7 +256,7 @@ struct S {
 
 [[group(5), binding(4)]] var<storage, read> d : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
   let x : i32 = (((a.i + b.i) + c.i) + d.i);
 }
@@ -276,7 +276,7 @@ struct S {
 
 [[internal(disable_validation__binding_point_collision), group(5), binding(4)]] var<storage, read> d : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {
   let x : i32 = (((a.i + b.i) + c.i) + d.i);
 }
@@ -310,12 +310,12 @@ struct S {
 
 [[group(5), binding(4)]] var<storage, read> d : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f1() {
   let x : i32 = (a.i + c.i);
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f2() {
   let x : i32 = (b.i + d.i);
 }
@@ -335,12 +335,12 @@ struct S {
 
 [[group(5), binding(4)]] var<storage, read> d : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f1() {
   let x : i32 = (a.i + c.i);
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f2() {
   let x : i32 = (b.i + d.i);
 }
@@ -368,7 +368,7 @@ struct S {
 [[group(2), binding(1)]] var<storage, read> a : S;
 [[group(3), binding(2)]] var<storage, read> b : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn f() {}
 )";
 

src/transform/calculate_array_length_test.cc

@@ -32,7 +32,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
@@ -50,7 +50,7 @@ fn tint_symbol(buffer : SB, result : ptr<function, u32>)
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var tint_symbol_1 : u32 = 0u;
   tint_symbol(sb, &(tint_symbol_1));
@@ -74,7 +74,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var a : u32 = arrayLength(&sb.arr);
   var b : u32 = arrayLength(&sb.arr);
@@ -94,7 +94,7 @@ fn tint_symbol(buffer : SB, result : ptr<function, u32>)
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var tint_symbol_1 : u32 = 0u;
   tint_symbol(sb, &(tint_symbol_1));
@@ -121,7 +121,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len : u32 = arrayLength(&sb.arr);
 }
@@ -140,7 +140,7 @@ fn tint_symbol(buffer : SB, result : ptr<function, u32>)
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var tint_symbol_1 : u32 = 0u;
   tint_symbol(sb, &(tint_symbol_1));
@@ -164,7 +164,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   if (true) {
     var len : u32 = arrayLength(&sb.arr);
@@ -188,7 +188,7 @@ fn tint_symbol(buffer : SB, result : ptr<function, u32>)
 
 [[group(0), binding(0)]] var<storage, read> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   if (true) {
     var tint_symbol_1 : u32 = 0u;
@@ -229,7 +229,7 @@ struct SB2 {
 
 [[group(0), binding(1)]] var<storage, read> sb2 : SB2;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -260,7 +260,7 @@ fn tint_symbol_3(buffer : SB2, result : ptr<function, u32>)
 
 [[group(0), binding(1)]] var<storage, read> sb2 : SB2;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var tint_symbol_1 : u32 = 0u;
   tint_symbol(sb1, &(tint_symbol_1));

src/transform/canonicalize_entry_point_io_test.cc

@@ -1095,7 +1095,7 @@ fn vert_main1() -> [[builtin(position)]] vec4<f32> {
   return vec4<f32>();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
 }
 )";
@@ -1133,7 +1133,7 @@ fn vert_main1() -> tint_symbol_3 {
   return tint_symbol_3(vec4<f32>());
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
 }
 )";

src/transform/decompose_memory_access_test.cc

@@ -52,7 +52,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var a : i32 = sb.a;
   var b : u32 = sb.b;
@@ -184,7 +184,7 @@ fn tint_symbol_21([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var a : i32 = tint_symbol(sb, 0u);
   var b : u32 = tint_symbol_1(sb, 4u);
@@ -246,7 +246,7 @@ struct UB {
 
 [[group(0), binding(0)]] var<uniform> ub : UB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var a : i32 = ub.a;
   var b : u32 = ub.b;
@@ -378,7 +378,7 @@ fn tint_symbol_21([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<uniform> ub : UB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var a : i32 = tint_symbol(ub, 0u);
   var b : u32 = tint_symbol_1(ub, 4u);
@@ -440,7 +440,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   sb.a = i32();
   sb.b = u32();
@@ -591,7 +591,7 @@ fn tint_symbol_21([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   tint_symbol(sb, 0u, i32());
   tint_symbol_1(sb, 4u, u32());
@@ -653,7 +653,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : SB = sb;
 }
@@ -768,7 +768,7 @@ fn tint_symbol_22([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : SB = tint_symbol_22(sb, 0u);
 }
@@ -809,7 +809,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   sb = SB();
 }
@@ -964,7 +964,7 @@ fn tint_symbol_22([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   tint_symbol_22(sb, 0u, SB());
 }
@@ -998,7 +998,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : f32 = sb.b[4].b[1].b.z;
 }
@@ -1036,7 +1036,7 @@ fn tint_symbol([[internal(disable_validation__ignore_atomic_function_parameter)]
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : f32 = tint_symbol(sb, 1224u);
 }
@@ -1070,7 +1070,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : i32 = 4;
   var j : u32 = 1u;
@@ -1104,7 +1104,7 @@ fn tint_symbol([[internal(disable_validation__ignore_atomic_function_parameter)]
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : i32 = 4;
   var j : u32 = 1u;
@@ -1149,7 +1149,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : i32 = 4;
   var j : u32 = 1u;
@@ -1191,7 +1191,7 @@ fn tint_symbol([[internal(disable_validation__ignore_atomic_function_parameter)]
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : i32 = 4;
   var j : u32 = 1u;
@@ -1216,7 +1216,7 @@ struct SB {
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   atomicStore(&sb.a, 123);
   ignore(atomicLoad(&sb.a));
@@ -1312,7 +1312,7 @@ fn tint_symbol_19([[internal(disable_validation__ignore_atomic_function_paramete
 
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   tint_symbol(sb, 16u, 123);
   ignore(tint_symbol_1(sb, 16u));
@@ -1352,7 +1352,7 @@ struct S {
 
 var<workgroup> w : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   atomicStore(&(w.a), 123);
   ignore(atomicLoad(&(w.a)));

src/transform/hlsl.cc

@@ -83,7 +83,8 @@ void Hlsl::AddEmptyEntryPoint(CloneContext& ctx) const {
   }
   ctx.dst->Func(ctx.dst->Symbols().New("unused_entry_point"), {},
                 ctx.dst->ty.void_(), {},
-                {ctx.dst->Stage(ast::PipelineStage::kCompute)});
+                {ctx.dst->Stage(ast::PipelineStage::kCompute),
+                 ctx.dst->WorkgroupSize(1)});
 }
 
 }  // namespace transform

src/transform/hlsl_test.cc

@@ -26,7 +26,7 @@ TEST_F(HlslTest, AddEmptyEntryPoint) {
   auto* src = R"()";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn unused_entry_point() {
 }
 )";

src/transform/msl.cc

@@ -128,7 +128,7 @@ void Msl::HandleModuleScopeVariables(CloneContext& ctx) const {
   //   v = v + 1.0;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn main() {
   //   foo();
   // }
@@ -140,7 +140,7 @@ void Msl::HandleModuleScopeVariables(CloneContext& ctx) const {
   //   *v = *v + 1.0;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn main() {
   //   var<private> v : f32 = 2.0;
   //   foo(&v);

src/transform/msl_test.cc

@@ -27,14 +27,14 @@ TEST_F(MslTest, HandleModuleScopeVariables_Basic) {
 var<private> p : f32;
 var<workgroup> w : f32;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   w = p;
 }
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
   [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_1 : f32;
   [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_2 : f32;
@@ -70,7 +70,7 @@ fn foo(a : f32) {
   no_uses();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   foo(1.0);
 }
@@ -91,7 +91,7 @@ fn foo(a : f32, tint_symbol_3 : ptr<private, f32>, tint_symbol_4 : ptr<workgroup
   no_uses();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
   [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_5 : f32;
   [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_6 : f32;
@@ -113,14 +113,14 @@ TEST_F(MslTest, HandleModuleScopeVariables_Constructors) {
 var<private> a : f32 = 1.0;
 var<private> b : f32 = f32();
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let x : f32 = a + b;
 }
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol : f32 = 1.0;
   [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_1 : f32 = f32();
@@ -138,7 +138,7 @@ TEST_F(MslTest, HandleModuleScopeVariables_Pointers) {
 var<private> p : f32;
 var<workgroup> w : f32;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p_ptr : ptr<private, f32> = &p;
   let w_ptr : ptr<workgroup, f32> = &w;
@@ -148,7 +148,7 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
   [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_1 : f32;
   [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_2 : f32;
@@ -171,13 +171,13 @@ TEST_F(MslTest, HandleModuleScopeVariables_UnusedVariables) {
 var<private> p : f32;
 var<workgroup> w : f32;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
 }
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
 }
 )";
@@ -196,7 +196,7 @@ struct S {
 [[group(0), binding(0)]]
 var<uniform> u : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
 }
 )";
@@ -208,7 +208,7 @@ struct S {
 
 [[group(0), binding(0)]] var<uniform> u : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
 }
 )";
@@ -223,7 +223,7 @@ TEST_F(MslTest, HandleModuleScopeVariables_HandleTypes_Basic) {
 [[group(0), binding(0)]] var t : texture_2d<f32>;
 [[group(0), binding(1)]] var s : sampler;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   ignore(t);
   ignore(s);
@@ -231,7 +231,7 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter)]] tint_symbol : texture_2d<f32>, [[group(0), binding(1), internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : sampler) {
   ignore(tint_symbol);
   ignore(tint_symbol_1);
@@ -263,7 +263,7 @@ fn foo(a : f32) {
   no_uses();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   foo(1.0);
 }
@@ -285,7 +285,7 @@ fn foo(a : f32, tint_symbol_2 : texture_2d<f32>, tint_symbol_3 : sampler) {
   no_uses();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter)]] tint_symbol_4 : texture_2d<f32>, [[group(0), binding(1), internal(disable_validation__entry_point_parameter)]] tint_symbol_5 : sampler) {
   foo(1.0, tint_symbol_4, tint_symbol_5);
 }

src/transform/promote_initializers_to_const_var_test.cc

@@ -24,7 +24,7 @@ using PromoteInitializersToConstVarTest = TransformTest;
 
 TEST_F(PromoteInitializersToConstVarTest, BasicArray) {
   auto* src = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var f0 : f32 = 1.0;
   var f1 : f32 = 2.0;
@@ -35,7 +35,7 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var f0 : f32 = 1.0;
   var f1 : f32 = 2.0;
@@ -59,7 +59,7 @@ struct S {
   c : vec3<f32>;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : f32 = S(1, 2.0, vec3<f32>()).b;
 }
@@ -72,7 +72,7 @@ struct S {
   c : vec3<f32>;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let tint_symbol : S = S(1, 2.0, vec3<f32>());
   var x : f32 = tint_symbol.b;
@@ -86,14 +86,14 @@ fn main() {
 
 TEST_F(PromoteInitializersToConstVarTest, ArrayInArrayArray) {
   auto* src = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : f32 = array<array<f32, 2>, 2>(array<f32, 2>(1.0, 2.0), array<f32, 2>(3.0, 4.0))[0][1];
 }
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let tint_symbol : array<f32, 2> = array<f32, 2>(1.0, 2.0);
   let tint_symbol_1 : array<f32, 2> = array<f32, 2>(3.0, 4.0);
@@ -123,7 +123,7 @@ struct S3 {
   a : S2;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : i32 = S3(S2(1, S1(2), 3)).a.b.a;
 }
@@ -144,7 +144,7 @@ struct S3 {
   a : S2;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let tint_symbol : S1 = S1(2);
   let tint_symbol_1 : S2 = S2(1, tint_symbol, 3);
@@ -168,7 +168,7 @@ struct S2 {
   a : array<S1, 3>;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var x : i32 = S2(array<S1, 3>(S1(1), S1(2), S1(3))).a[1].a;
 }
@@ -183,7 +183,7 @@ struct S2 {
   a : array<S1, 3>;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let tint_symbol : S1 = S1(1);
   let tint_symbol_1 : S1 = S1(2);
@@ -207,7 +207,7 @@ struct S {
   c : i32;
 };
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var local_arr : array<f32, 4> = array<f32, 4>(0.0, 1.0, 2.0, 3.0);
   var local_str : S = S(1, 2.0, 3);

src/transform/single_entry_point_test.cc

@@ -86,7 +86,7 @@ fn main() {}
 
 TEST_F(SingleEntryPointTest, SingleEntryPoint) {
   auto* src = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
 }
 )";
@@ -111,17 +111,17 @@ fn vert_main() -> [[builtin(position)]] vec4<f32> {
 fn frag_main() {
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main2() {
 }
 )";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
 }
 )";
@@ -156,12 +156,12 @@ fn frag_main() {
   b = 0.0;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   c = 0.0;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main2() {
   d = 0.0;
 }
@@ -170,7 +170,7 @@ fn comp_main2() {
   auto* expect = R"(
 var<private> c : f32;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   c = 0.0;
 }
@@ -206,12 +206,12 @@ fn frag_main() {
   let local_b : f32 = b;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   let local_c : f32 = c;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main2() {
   let local_d : f32 = d;
 }
@@ -226,7 +226,7 @@ let c : f32 = 1.0;
 
 let d : f32 = 1.0;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   let local_c : f32 = c;
 }
@@ -262,12 +262,12 @@ fn outer2() {
   inner_shared();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   outer1();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main2() {
   outer2();
 }
@@ -285,7 +285,7 @@ fn outer1() {
   inner_shared();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   outer1();
 }
@@ -336,12 +336,12 @@ fn outer2() {
   outer2_var = 0.0;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   outer1();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main2() {
   outer2();
 }
@@ -368,7 +368,7 @@ fn outer1() {
   outer1_var = 0.0;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn comp_main1() {
   outer1();
 }

src/transform/spirv.cc

@@ -300,7 +300,8 @@ void Spirv::AddEmptyEntryPoint(CloneContext& ctx) const {
     }
   }
   ctx.dst->Func(ctx.dst->Sym("unused_entry_point"), {}, ctx.dst->ty.void_(), {},
-                {ctx.dst->Stage(ast::PipelineStage::kCompute)});
+                {ctx.dst->Stage(ast::PipelineStage::kCompute),
+                 ctx.dst->WorkgroupSize(1)});
 }
 
 Symbol Spirv::HoistToInputVariables(

src/transform/spirv_test.cc

@@ -30,7 +30,7 @@ fn frag_main([[builtin(position)]] coord : vec4<f32>,
   var col : f32 = (coord.x * loc1);
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(8, 1, 1)]]
 fn compute_main([[builtin(local_invocation_id)]] local_id : vec3<u32>,
                 [[builtin(local_invocation_index)]] local_index : u32) {
   var id_x : u32 = local_id.x;
@@ -51,7 +51,7 @@ fn frag_main() {
 
 [[builtin(local_invocation_index), internal(disable_validation__ignore_storage_class)]] var<in> tint_symbol_3 : u32;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(8, 1, 1)]]
 fn compute_main() {
   var id_x : u32 = tint_symbol_2.x;
 }
@@ -891,7 +891,7 @@ fn main3() {
 
 TEST_F(SpirvTest, EmitVertexPointSize_NoVertexShaders) {
   auto* src = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(8, 1, 1)]]
 fn main() {
 }
 )";
@@ -907,7 +907,7 @@ TEST_F(SpirvTest, AddEmptyEntryPoint) {
   auto* src = R"()";
 
   auto* expect = R"(
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn unused_entry_point() {
 }
 )";

src/writer/hlsl/generator_impl_function_test.cc

@@ -660,9 +660,7 @@ TEST_F(HlslGeneratorImplTest_Function, Emit_Decoration_EntryPoint_Compute) {
        {
            Return(),
        },
-       {
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
-           Stage(ast::PipelineStage::kCompute),
-       });
 
   GeneratorImpl& gen = Build();
 
@@ -800,13 +798,13 @@ TEST_F(HlslGeneratorImplTest_Function,
   // };
   // [[binding(0), group(0)]] var<storage> data : Data;
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn a() {
   //   var v = data.d;
   //   return;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn b() {
   //   var v = data.d;
   //   return;
@@ -830,9 +828,7 @@ TEST_F(HlslGeneratorImplTest_Function,
              Decl(var),
              Return(),
          },
-         {
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
-             Stage(ast::PipelineStage::kCompute),
-         });
   }
 
   {
@@ -844,9 +840,7 @@ TEST_F(HlslGeneratorImplTest_Function,
              Decl(var),
              Return(),
          },
-         {
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
-             Stage(ast::PipelineStage::kCompute),
-         });
   }
 
   GeneratorImpl& gen = SanitizeAndBuild();

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -535,7 +535,10 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Unpack2x16Float) {
 TEST_F(HlslGeneratorImplTest_Intrinsic, StorageBarrier) {
   Func("main", {}, ty.void_(),
        {create<ast::CallStatement>(Call("storageBarrier"))},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
 
   GeneratorImpl& gen = Build();
 
@@ -551,7 +554,10 @@ void main() {
 TEST_F(HlslGeneratorImplTest_Intrinsic, WorkgroupBarrier) {
   Func("main", {}, ty.void_(),
        {create<ast::CallStatement>(Call("workgroupBarrier"))},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
 
   GeneratorImpl& gen = Build();
 
@@ -570,7 +576,10 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Ignore) {
 
   Func("main", {}, ty.void_(),
        {create<ast::CallStatement>(Call("ignore", Call("f", 1, 2, 3)))},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
 
   GeneratorImpl& gen = Build();
 

src/writer/hlsl/generator_impl_workgroup_var_test.cc

@@ -29,7 +29,10 @@ TEST_F(HlslGeneratorImplTest_WorkgroupVar, Basic) {
   Global("wg", ty.f32(), ast::StorageClass::kWorkgroup);
 
   Func("main", {}, ty.void_(), {Assign("wg", 1.2f)},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
   GeneratorImpl& gen = Build();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();
@@ -42,7 +45,10 @@ TEST_F(HlslGeneratorImplTest_WorkgroupVar, Aliased) {
   Global("wg", ty.Of(alias), ast::StorageClass::kWorkgroup);
 
   Func("main", {}, ty.void_(), {Assign("wg", 1.2f)},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
   GeneratorImpl& gen = Build();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();

src/writer/msl/generator_impl_function_test.cc

@@ -612,12 +612,12 @@ TEST_F(MslGeneratorImplTest,
   // };
   // [[binding(0), group(0)]] var<storage> data : Data;
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn a() {
   //   return;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn b() {
   //   return;
   // }
@@ -642,6 +642,7 @@ TEST_F(MslGeneratorImplTest,
          },
          {
              Stage(ast::PipelineStage::kCompute),
+             WorkgroupSize(1),
          });
   }
 
@@ -651,7 +652,10 @@ TEST_F(MslGeneratorImplTest,
 
     Func("b", ast::VariableList{}, ty.void_(),
          ast::StatementList{Decl(var), Return()},
-         {Stage(ast::PipelineStage::kCompute)});
+         {
+             Stage(ast::PipelineStage::kCompute),
+             WorkgroupSize(1),
+         });
   }
 
   GeneratorImpl& gen = Build();

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -347,7 +347,10 @@ TEST_F(MslGeneratorImplTest, Ignore) {
        ty.i32(), {Return(Mul(Add("a", "b"), "c"))});
 
   Func("func", {}, ty.void_(), {Ignore(Call("f", 1, 2, 3))},
-       {Stage(ast::PipelineStage::kCompute)});
+       {
+           Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
+       });
 
   GeneratorImpl& gen = Build();
 

src/writer/msl/generator_impl_test.cc

@@ -36,6 +36,7 @@ TEST_F(MslGeneratorImplTest, Generate) {
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::DecorationList{
            Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
        });
 
   GeneratorImpl& gen = Build();

src/writer/spirv/builder_entry_point_test.cc

@@ -296,7 +296,7 @@ OpFunctionEnd
 TEST_F(BuilderTest, SampleIndex_SampleRateShadingCapability) {
   Func("main",
        {Param("sample_index", ty.u32(), {Builtin(ast::Builtin::kSampleIndex)})},
-       ty.void_(), {}, {Stage(ast::PipelineStage::kCompute)});
+       ty.void_(), {}, {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
 
   spirv::Builder& b = SanitizeAndBuild();
 

src/writer/spirv/builder_function_decoration_test.cc

@@ -62,11 +62,12 @@ TEST_P(Decoration_StageTest, Emit) {
     ret_type = ty.void_();
   }
 
-  auto* func = Func("main", {}, ret_type, body,
+  auto deco_list = ast::DecorationList{Stage(params.stage)};
-                    ast::DecorationList{
+  if (params.stage == ast::PipelineStage::kCompute) {
-                        Stage(params.stage),
+    deco_list.push_back(WorkgroupSize(1));
-                    },
+  }
-                    ret_type_decos);
+
+  auto* func = Func("main", {}, ret_type, body, deco_list, ret_type_decos);
 
   spirv::Builder& b = Build();
 
@@ -109,9 +110,8 @@ TEST_F(BuilderTest, Decoration_ExecutionMode_Fragment_OriginUpperLeft) {
 
 TEST_F(BuilderTest, Decoration_ExecutionMode_WorkgroupSize_Default) {
   auto* func = Func("main", {}, ty.void_(), ast::StatementList{},
-                    ast::DecorationList{
+                    ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-                        Stage(ast::PipelineStage::kCompute),
+                                        WorkgroupSize(1)});
-                    });
 
   spirv::Builder& b = Build();
 

src/writer/spirv/builder_function_test.cc

@@ -191,12 +191,12 @@ TEST_F(BuilderTest, Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
   // };
   // [[binding(0), group(0)]] var<storage> data : Data;
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn a() {
   //   return;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn b() {
   //   return;
   // }
@@ -219,9 +219,8 @@ TEST_F(BuilderTest, Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
              Decl(var),
              Return(),
          },
-         ast::DecorationList{
+         ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-             Stage(ast::PipelineStage::kCompute),
+                             WorkgroupSize(1)});
-         });
   }
 
   {
@@ -233,9 +232,8 @@ TEST_F(BuilderTest, Emit_Multiple_EntryPoint_With_Same_ModuleVar) {
              Decl(var),
              Return(),
          },
-         ast::DecorationList{
+         ast::DecorationList{Stage(ast::PipelineStage::kCompute),
-             Stage(ast::PipelineStage::kCompute),
+                             WorkgroupSize(1)});
-         });
   }
 
   spirv::Builder& b = Build();

src/writer/spirv/builder_intrinsic_test.cc

@@ -2366,6 +2366,7 @@ TEST_F(IntrinsicBuilderTest, Call_WorkgroupBarrier) {
        },
        ast::DecorationList{
            Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
        });
 
   spirv::Builder& b = Build();
@@ -2398,6 +2399,7 @@ TEST_F(IntrinsicBuilderTest, Call_StorageBarrier) {
        },
        ast::DecorationList{
            Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
        });
 
   spirv::Builder& b = Build();
@@ -2434,6 +2436,7 @@ TEST_F(IntrinsicBuilderTest, Call_Ignore) {
        },
        {
            Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
        });
 
   spirv::Builder& b = Build();

src/writer/wgsl/generator_impl_function_test.cc

@@ -162,12 +162,12 @@ TEST_F(WgslGeneratorImplTest,
   // };
   // [[binding(0), group(0)]] var<storage> data : Data;
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn a() {
   //   return;
   // }
   //
-  // [[stage(compute)]]
+  // [[stage(compute), workgroup_size(1)]]
   // fn b() {
   //   return;
   // }
@@ -192,6 +192,7 @@ TEST_F(WgslGeneratorImplTest,
          },
          ast::DecorationList{
              Stage(ast::PipelineStage::kCompute),
+             WorkgroupSize(1),
          });
   }
 
@@ -206,6 +207,7 @@ TEST_F(WgslGeneratorImplTest,
          },
          ast::DecorationList{
              Stage(ast::PipelineStage::kCompute),
+             WorkgroupSize(1),
          });
   }
 
@@ -219,13 +221,13 @@ struct Data {
 
 [[binding(0), group(0)]] var<storage, read_write> data : Data;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn a() {
   var v : f32 = data.d;
   return;
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn b() {
   var v : f32 = data.d;
   return;

src/writer/wgsl/generator_impl_global_decl_test.cc

@@ -35,7 +35,7 @@ TEST_F(WgslGeneratorImplTest, Emit_GlobalDeclAfterFunction) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  [[stage(compute)]]
+  EXPECT_EQ(gen.result(), R"(  [[stage(compute), workgroup_size(1, 1, 1)]]
   fn test_function() {
     var a : f32;
   }
@@ -67,6 +67,7 @@ TEST_F(WgslGeneratorImplTest, Emit_GlobalsInterleaved) {
        },
        ast::DecorationList{
            Stage(ast::PipelineStage::kCompute),
+           WorkgroupSize(1),
        });
 
   GeneratorImpl& gen = Build();
@@ -90,7 +91,7 @@ TEST_F(WgslGeneratorImplTest, Emit_GlobalsInterleaved) {
     a : i32;
   };
 
-  [[stage(compute)]]
+  [[stage(compute), workgroup_size(1)]]
   fn main() {
     var s0 : S0;
     var s1 : S1;

test/access/let/matrix.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute)]]
+[[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)]]
+[[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.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute)]]
+[[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)]]
+[[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.wgsl

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

test/access/var/vector.wgsl

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

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)]]
+[[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)]]
+[[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)]]
+[[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)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let a1 : array<f32, 4> = f1();
   let a2 : array<array<f32, 4>, 3> = f2();

test/array/type_constructor.wgsl

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

test/array/type_constructor.wgsl.expected.wgsl

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

test/bug/tint/221.wgsl

@@ -8,7 +8,7 @@ struct Buf{
 
 [[group(0), binding (0)]] var<storage, read_write> b : Buf;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : u32 = 0u;
   loop {

test/bug/tint/221.wgsl.expected.wgsl

@@ -8,7 +8,7 @@ struct Buf {
 
 [[group(0), binding(0)]] var<storage, read_write> b : Buf;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var i : u32 = 0u;
   loop {

test/bug/tint/453.wgsl

@@ -1,7 +1,7 @@
 [[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
 [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var srcValue : vec4<u32>;
   let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));

test/bug/tint/453.wgsl.expected.wgsl

@@ -2,7 +2,7 @@
 
 [[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var srcValue : vec4<u32>;
   let x_22 : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));

test/bug/tint/492.wgsl

@@ -1,7 +1,7 @@
 [[block]] struct S { a : i32; };
 [[group(0), binding(0)]] var<storage, read_write> buf : S;
 
-[[stage(compute)]] fn main() {
+[[stage(compute), workgroup_size(1)]] fn main() {
   let p : ptr<storage, i32, read_write> = &buf.a;
   *p = 12;
 }

test/bug/tint/492.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read_write> buf : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p : ptr<storage, i32> = &(buf.a);
   *(p) = 12;

test/bug/tint/757.wgsl

@@ -11,7 +11,7 @@
 };
 [[group(0), binding(3)]] var<storage, read_write> result : Result;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
   var flatIndex : u32 =
     2u * 2u * GlobalInvocationID.z +

test/bug/tint/757.wgsl.expected.wgsl

@@ -14,7 +14,7 @@ struct Result {
 
 [[group(0), binding(3)]] var<storage, read_write> result : Result;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
   var flatIndex : u32 = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
   flatIndex = (flatIndex * 1u);

test/bug/tint/827.wgsl

@@ -7,7 +7,7 @@ let width : u32 = 128u;
 [[group(0), binding(0)]] var tex : texture_depth_2d;
 [[group(0), binding(1)]] var<storage, read_write> result : Result;
 
-[[stage(compute)]] fn main(
+[[stage(compute), workgroup_size(1)]] fn main(
     [[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>
 ) {
     result.values[GlobalInvocationId.y * width + GlobalInvocationId.x] = textureLoad(

test/bug/tint/827.wgsl.expected.wgsl

@@ -9,7 +9,7 @@ let width : u32 = 128u;
 
 [[group(0), binding(1)]] var<storage, read_write> result : Result;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main([[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>) {
   result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = textureLoad(tex, vec2<i32>(i32(GlobalInvocationId.x), i32(GlobalInvocationId.y)), 0);
 }

test/bug/tint/926.wgsl

@@ -5,7 +5,7 @@
 
 var<private> cubeVerts : u32 = 0u;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn computeMain([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
   // Increment cubeVerts based on some criteria...
 

test/bug/tint/926.wgsl.expected.msl

@@ -8,7 +8,7 @@ struct DrawIndirectArgs {
 
 [[group(0), binding(5)]] var<storage, read_write> drawOut : DrawIndirectArgs;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn computeMain([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
   [[internal(disable_validation__function_var_storage_class)]] var<private> tint_symbol_1 : u32 = 0u;
   let firstVertex : u32 = atomicAdd(&(drawOut.vertexCount), tint_symbol_1);

test/bug/tint/926.wgsl.expected.wgsl

@@ -7,7 +7,7 @@ struct DrawIndirectArgs {
 
 var<private> cubeVerts : u32 = 0u;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn computeMain([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
   let firstVertex : u32 = atomicAdd(&(drawOut.vertexCount), cubeVerts);
 }

test/intrinsics/arrayLength/complex_via_let.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
     let p = &G;
     let p2 = &((*p).a);

test/intrinsics/arrayLength/complex_via_let.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p = &(G);
   let p2 = &((*(p)).a);

test/intrinsics/arrayLength/deprecated.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
     let l1 : u32 = arrayLength(&G.a);
 

test/intrinsics/arrayLength/deprecated.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let l1 : u32 = arrayLength(&(G.a));
   let p = &(G.a);

test/intrinsics/arrayLength/simple.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
     let l1 : u32 = arrayLength(&G.a);
 }

test/intrinsics/arrayLength/simple.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let l1 : u32 = arrayLength(&(G.a));
 }

test/intrinsics/arrayLength/via_let.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
     let p = &G.a;
     let p2 = p;

test/intrinsics/arrayLength/via_let.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p = &(G.a);
   let p2 = p;

test/intrinsics/arrayLength/via_let_complex.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p = &*&G;
   let p2 = &*p;

test/intrinsics/arrayLength/via_let_complex.wgsl.expected.wgsl

@@ -5,7 +5,7 @@ struct S {
 
 [[group(0), binding(0)]] var<storage, read> G : S;
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   let p = &(*(&(G)));
   let p2 = &(*(p));

test/intrinsics/frexp.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
     var exponent : i32;
     let significand : f32 = frexp(1.23, &exponent);

test/intrinsics/frexp.wgsl.expected.wgsl

@@ -1,4 +1,4 @@
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn main() {
   var exponent : i32;
   let significand : f32 = frexp(1.230000019, &(exponent));

test/intrinsics/gen/abs/002533.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_002533();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_002533();
 }

test/intrinsics/gen/abs/002533.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_002533();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_002533();
 }

test/intrinsics/gen/abs/005174.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_005174();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_005174();
 }

test/intrinsics/gen/abs/005174.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_005174();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_005174();
 }

test/intrinsics/gen/abs/1ce782.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_1ce782();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_1ce782();
 }

test/intrinsics/gen/abs/1ce782.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_1ce782();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_1ce782();
 }

test/intrinsics/gen/abs/1e9d53.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_1e9d53();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_1e9d53();
 }

test/intrinsics/gen/abs/1e9d53.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_1e9d53();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_1e9d53();
 }

test/intrinsics/gen/abs/467cd1.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_467cd1();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_467cd1();
 }

test/intrinsics/gen/abs/467cd1.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_467cd1();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_467cd1();
 }

test/intrinsics/gen/abs/4ad288.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_4ad288();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_4ad288();
 }

test/intrinsics/gen/abs/4ad288.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_4ad288();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_4ad288();
 }

test/intrinsics/gen/abs/5ad50a.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_5ad50a();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_5ad50a();
 }

test/intrinsics/gen/abs/5ad50a.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_5ad50a();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_5ad50a();
 }

test/intrinsics/gen/abs/7326de.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_7326de();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7326de();
 }

test/intrinsics/gen/abs/7326de.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_7326de();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7326de();
 }

test/intrinsics/gen/abs/7f28e6.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_7f28e6();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7f28e6();
 }

test/intrinsics/gen/abs/7f28e6.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_7f28e6();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7f28e6();
 }

test/intrinsics/gen/abs/7faa9e.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_7faa9e();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7faa9e();
 }

test/intrinsics/gen/abs/7faa9e.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_7faa9e();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_7faa9e();
 }

test/intrinsics/gen/abs/9c80a6.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_9c80a6();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_9c80a6();
 }

test/intrinsics/gen/abs/9c80a6.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_9c80a6();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_9c80a6();
 }

test/intrinsics/gen/abs/b96037.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   abs_b96037();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_b96037();
 }

test/intrinsics/gen/abs/b96037.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   abs_b96037();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   abs_b96037();
 }

test/intrinsics/gen/acos/489247.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   acos_489247();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   acos_489247();
 }

test/intrinsics/gen/acos/489247.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   acos_489247();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   acos_489247();
 }

test/intrinsics/gen/acos/8e2acf.wgsl

@@ -39,7 +39,7 @@ fn fragment_main() {
   acos_8e2acf();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   acos_8e2acf();
 }

test/intrinsics/gen/acos/8e2acf.wgsl.expected.wgsl

@@ -13,7 +13,7 @@ fn fragment_main() {
   acos_8e2acf();
 }
 
-[[stage(compute)]]
+[[stage(compute), workgroup_size(1)]]
 fn compute_main() {
   acos_8e2acf();
 }