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>

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;

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

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:

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}));
 

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;
   }

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},

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:

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"},

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:

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"},

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;

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,

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"}));