diff --git a/src/ast/builtin.cc b/src/ast/builtin.cc index 3dac5cdf76..bdfaf887c7 100644 --- a/src/ast/builtin.cc +++ b/src/ast/builtin.cc @@ -59,6 +59,10 @@ std::ostream& operator<<(std::ostream& out, Builtin builtin) { out << "global_invocation_id"; break; } + case Builtin::kWorkgroupId: { + out << "workgroup_id"; + break; + } case Builtin::kSampleIndex: { out << "sample_index"; break; diff --git a/src/ast/builtin.h b/src/ast/builtin.h index e065d69df8..5b37aeb55b 100644 --- a/src/ast/builtin.h +++ b/src/ast/builtin.h @@ -32,6 +32,7 @@ enum class Builtin { kLocalInvocationId, kLocalInvocationIndex, kGlobalInvocationId, + kWorkgroupId, kSampleIndex, kSampleMask, kSampleMaskIn, // TODO(crbug.com/tint/715): Remove this diff --git a/src/reader/spirv/enum_converter.cc b/src/reader/spirv/enum_converter.cc index 364b596669..b07da2679f 100644 --- a/src/reader/spirv/enum_converter.cc +++ b/src/reader/spirv/enum_converter.cc @@ -86,6 +86,8 @@ ast::Builtin EnumConverter::ToBuiltin(SpvBuiltIn b) { return ast::Builtin::kLocalInvocationIndex; case SpvBuiltInGlobalInvocationId: return ast::Builtin::kGlobalInvocationId; + case SpvBuiltInWorkgroupId: + return ast::Builtin::kWorkgroupId; case SpvBuiltInSampleId: return ast::Builtin::kSampleIndex; case SpvBuiltInSampleMask: diff --git a/src/reader/spirv/enum_converter_test.cc b/src/reader/spirv/enum_converter_test.cc index 58e69241ef..19f6ad2167 100644 --- a/src/reader/spirv/enum_converter_test.cc +++ b/src/reader/spirv/enum_converter_test.cc @@ -217,6 +217,7 @@ INSTANTIATE_TEST_SUITE_P( ast::Builtin::kLocalInvocationIndex}, BuiltinCase{SpvBuiltInGlobalInvocationId, true, ast::Builtin::kGlobalInvocationId}, + BuiltinCase{SpvBuiltInWorkgroupId, true, ast::Builtin::kWorkgroupId}, BuiltinCase{SpvBuiltInSampleId, true, ast::Builtin::kSampleIndex}, BuiltinCase{SpvBuiltInSampleMask, true, ast::Builtin::kSampleMask})); diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc index dd0b884648..794cad8615 100644 --- a/src/reader/wgsl/parser_impl.cc +++ b/src/reader/wgsl/parser_impl.cc @@ -85,6 +85,9 @@ ast::Builtin ident_to_builtin(const std::string& str) { if (str == "global_invocation_id") { return ast::Builtin::kGlobalInvocationId; } + if (str == "workgroup_id") { + return ast::Builtin::kWorkgroupId; + } if (str == "sample_index") { return ast::Builtin::kSampleIndex; } diff --git a/src/reader/wgsl/parser_impl_variable_decoration_test.cc b/src/reader/wgsl/parser_impl_variable_decoration_test.cc index 9430cf1172..c324483a2b 100644 --- a/src/reader/wgsl/parser_impl_variable_decoration_test.cc +++ b/src/reader/wgsl/parser_impl_variable_decoration_test.cc @@ -121,6 +121,7 @@ INSTANTIATE_TEST_SUITE_P( BuiltinData{"local_invocation_index", ast::Builtin::kLocalInvocationIndex}, BuiltinData{"global_invocation_id", ast::Builtin::kGlobalInvocationId}, + BuiltinData{"workgroup_id", ast::Builtin::kWorkgroupId}, BuiltinData{"sample_index", ast::Builtin::kSampleIndex}, BuiltinData{"sample_mask", ast::Builtin::kSampleMask}, BuiltinData{"sample_mask_in", ast::Builtin::kSampleMaskIn}, diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 55505347b2..b7dd0d26e9 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -1949,6 +1949,8 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const { return "SV_GroupIndex"; case ast::Builtin::kGlobalInvocationId: return "SV_DispatchThreadID"; + case ast::Builtin::kWorkgroupId: + return "SV_GroupID"; case ast::Builtin::kSampleIndex: return "SV_SampleIndex"; case ast::Builtin::kSampleMask: diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc index d37ff5395a..b04d496773 100644 --- a/src/writer/hlsl/generator_impl_test.cc +++ b/src/writer/hlsl/generator_impl_test.cc @@ -82,6 +82,7 @@ INSTANTIATE_TEST_SUITE_P( HlslBuiltinData{ast::Builtin::kLocalInvocationIndex, "SV_GroupIndex"}, HlslBuiltinData{ast::Builtin::kGlobalInvocationId, "SV_DispatchThreadID"}, + HlslBuiltinData{ast::Builtin::kWorkgroupId, "SV_GroupID"}, HlslBuiltinData{ast::Builtin::kSampleIndex, "SV_SampleIndex"}, HlslBuiltinData{ast::Builtin::kSampleMask, "SV_Coverage"}, HlslBuiltinData{ast::Builtin::kSampleMaskIn, "SV_Coverage"}, diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index ffca978bb5..74e05e434d 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -1352,6 +1352,8 @@ std::string GeneratorImpl::builtin_to_attribute(ast::Builtin builtin) const { return "thread_index_in_threadgroup"; case ast::Builtin::kGlobalInvocationId: return "thread_position_in_grid"; + case ast::Builtin::kWorkgroupId: + return "threadgroup_position_in_grid"; case ast::Builtin::kSampleIndex: return "sample_id"; case ast::Builtin::kSampleMask: diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc index bc2ee5532b..9cdf4b345f 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc @@ -73,6 +73,8 @@ INSTANTIATE_TEST_SUITE_P( "thread_index_in_threadgroup"}, MslBuiltinData{ast::Builtin::kGlobalInvocationId, "thread_position_in_grid"}, + MslBuiltinData{ast::Builtin::kWorkgroupId, + "threadgroup_position_in_grid"}, MslBuiltinData{ast::Builtin::kSampleIndex, "sample_id"}, MslBuiltinData{ast::Builtin::kSampleMask, "sample_mask"}, MslBuiltinData{ast::Builtin::kSampleMaskIn, "sample_mask"}, diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index 77e3f11838..4482b7db49 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -3304,6 +3304,8 @@ SpvBuiltIn Builder::ConvertBuiltin(ast::Builtin builtin, return SpvBuiltInGlobalInvocationId; case ast::Builtin::kPointSize: return SpvBuiltInPointSize; + case ast::Builtin::kWorkgroupId: + return SpvBuiltInWorkgroupId; case ast::Builtin::kSampleIndex: push_capability(SpvCapabilitySampleRateShading); return SpvBuiltInSampleId; diff --git a/src/writer/spirv/builder_global_variable_test.cc b/src/writer/spirv/builder_global_variable_test.cc index b6bf64c827..eb0ed7886f 100644 --- a/src/writer/spirv/builder_global_variable_test.cc +++ b/src/writer/spirv/builder_global_variable_test.cc @@ -368,6 +368,8 @@ INSTANTIATE_TEST_SUITE_P( ast::StorageClass::kInput, SpvBuiltInLocalInvocationIndex}, BuiltinData{ast::Builtin::kGlobalInvocationId, ast::StorageClass::kInput, SpvBuiltInGlobalInvocationId}, + BuiltinData{ast::Builtin::kWorkgroupId, ast::StorageClass::kInput, + SpvBuiltInWorkgroupId}, BuiltinData{ast::Builtin::kSampleIndex, ast::StorageClass::kInput, SpvBuiltInSampleId}, BuiltinData{ast::Builtin::kSampleMask, ast::StorageClass::kInput, diff --git a/src/writer/wgsl/generator_impl_test.cc b/src/writer/wgsl/generator_impl_test.cc index 4a6e8c10dd..287a6a7294 100644 --- a/src/writer/wgsl/generator_impl_test.cc +++ b/src/writer/wgsl/generator_impl_test.cc @@ -74,6 +74,7 @@ INSTANTIATE_TEST_SUITE_P( "local_invocation_index"}, WgslBuiltinData{ast::Builtin::kGlobalInvocationId, "global_invocation_id"}, + WgslBuiltinData{ast::Builtin::kWorkgroupId, "workgroup_id"}, WgslBuiltinData{ast::Builtin::kSampleIndex, "sample_index"}, WgslBuiltinData{ast::Builtin::kSampleMaskIn, "sample_mask_in"}, WgslBuiltinData{ast::Builtin::kSampleMaskOut, "sample_mask_out"}));