Add support for [[builtin(workgroup_id)]] This is a trivial mapping to/from WGSL in all cases. Bug: tint:478 Change-Id: I7f21a2392543a880906b54fddbdb8bbd149a526e Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/48140 Commit-Queue: James Price <jrprice@google.com> Commit-Queue: Ben Clayton <bclayton@google.com> Auto-Submit: James Price <jrprice@google.com> Reviewed-by: Ben Clayton <bclayton@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/ast/builtin.cc b/src/ast/builtin.cc index 3dac5cd..bdfaf88 100644 --- a/src/ast/builtin.cc +++ b/src/ast/builtin.cc
@@ -59,6 +59,10 @@ 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 e065d69..5b37aeb 100644 --- a/src/ast/builtin.h +++ b/src/ast/builtin.h
@@ -32,6 +32,7 @@ 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 364b596..b07da26 100644 --- a/src/reader/spirv/enum_converter.cc +++ b/src/reader/spirv/enum_converter.cc
@@ -86,6 +86,8 @@ 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 58e6924..19f6ad2 100644 --- a/src/reader/spirv/enum_converter_test.cc +++ b/src/reader/spirv/enum_converter_test.cc
@@ -217,6 +217,7 @@ 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 dd0b884..794cad8 100644 --- a/src/reader/wgsl/parser_impl.cc +++ b/src/reader/wgsl/parser_impl.cc
@@ -85,6 +85,9 @@ 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 9430cf1..c324483 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 @@ 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 5550534..b7dd0d2 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc
@@ -1949,6 +1949,8 @@ 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 d37ff53..b04d496 100644 --- a/src/writer/hlsl/generator_impl_test.cc +++ b/src/writer/hlsl/generator_impl_test.cc
@@ -82,6 +82,7 @@ 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 ffca978..74e05e4 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc
@@ -1352,6 +1352,8 @@ 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 bc2ee55..9cdf4b3 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc
@@ -73,6 +73,8 @@ "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 77e3f11..4482b7d 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc
@@ -3304,6 +3304,8 @@ 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 b6bf64c..eb0ed78 100644 --- a/src/writer/spirv/builder_global_variable_test.cc +++ b/src/writer/spirv/builder_global_variable_test.cc
@@ -368,6 +368,8 @@ 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 4a6e8c1..287a6a7 100644 --- a/src/writer/wgsl/generator_impl_test.cc +++ b/src/writer/wgsl/generator_impl_test.cc
@@ -74,6 +74,7 @@ "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"}));