diff --git a/src/inspector/entry_point.h b/src/inspector/entry_point.h index a97cc67cce..ca72cca0b5 100644 --- a/src/inspector/entry_point.h +++ b/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 workgroup_size() { diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc index 54da3c364d..1c1906ea9c 100644 --- a/src/inspector/inspector.cc +++ b/src/inspector/inspector.cc @@ -173,6 +173,9 @@ std::vector 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()) { diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc index 488d928685..9854f4058b 100644 --- a/src/inspector/inspector_test.cc +++ b/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(), {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(), + {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)}));