inspector: reflect num_workgroups builtin usage

Dawn needs to know if a given entry point uses this builtin, so that
it can pass this information via a root constant for HLSL.

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

src/inspector/entry_point.h

@@ -135,6 +135,8 @@ struct EntryPoint {
   bool front_facing_used = false;
   /// Does the entry point use the sample_index builtin
   bool sample_index_used = false;
+  /// Does the entry point use the num_workgroups builtin
+  bool num_workgroups_used = false;
 
   /// @returns the size of the workgroup in {x,y,z} format
   std::tuple<uint32_t, uint32_t, uint32_t> workgroup_size() {

src/inspector/inspector.cc

@@ -173,6 +173,9 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
       entry_point.input_sample_mask_used |=
           ContainsBuiltin(ast::Builtin::kSampleMask, param->Type(),
                           param->Declaration()->decorations());
+      entry_point.num_workgroups_used |=
+          ContainsBuiltin(ast::Builtin::kNumWorkgroups, param->Type(),
+                          param->Declaration()->decorations());
     }
 
     if (!sem->ReturnType()->Is<sem::Void>()) {

src/inspector/inspector_test.cc

@@ -642,6 +642,7 @@ TEST_F(InspectorGetEntryPointTest, BuiltinNotReferenced) {
   EXPECT_FALSE(result[0].input_position_used);
   EXPECT_FALSE(result[0].front_facing_used);
   EXPECT_FALSE(result[0].sample_index_used);
+  EXPECT_FALSE(result[0].num_workgroups_used);
 }
 
 TEST_F(InspectorGetEntryPointTest, InputSampleMaskSimpleReferenced) {
@@ -805,6 +806,38 @@ TEST_F(InspectorGetEntryPointTest, SampleIndexStructReferenced) {
   EXPECT_TRUE(result[0].sample_index_used);
 }
 
+TEST_F(InspectorGetEntryPointTest, NumWorkgroupsSimpleReferenced) {
+  auto* in_var =
+      Param("in_var", ty.vec3<u32>(), {Builtin(ast::Builtin::kNumWorkgroups)});
+  Func("ep_func", {in_var}, ty.void_(), {Return()},
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)}, {});
+
+  Inspector& inspector = Build();
+
+  auto result = inspector.GetEntryPoints();
+
+  ASSERT_EQ(1u, result.size());
+  EXPECT_TRUE(result[0].num_workgroups_used);
+}
+
+TEST_F(InspectorGetEntryPointTest, NumWorkgroupsStructReferenced) {
+  ast::StructMemberList members;
+  members.push_back(Member("inner_position", ty.vec3<u32>(),
+                           {Builtin(ast::Builtin::kNumWorkgroups)}));
+  Structure("in_struct", members, {});
+  auto* in_var = Param("in_var", ty.type_name("in_struct"), {});
+
+  Func("ep_func", {in_var}, ty.void_(), {Return()},
+       {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)}, {});
+
+  Inspector& inspector = Build();
+
+  auto result = inspector.GetEntryPoints();
+
+  ASSERT_EQ(1u, result.size());
+  EXPECT_TRUE(result[0].num_workgroups_used);
+}
+
 TEST_F(InspectorGetEntryPointTest, ImplicitInterpolate) {
   ast::StructMemberList members;
   members.push_back(Member("struct_inner", ty.f32(), {Location(0)}));