commit 71da3e8aa054ec1e8d36f65a07b975ad8d02de0f
author dan sinclair <dsinclair@chromium.org> Mon Nov 27 02:00:59 2023 +0000
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Mon Nov 27 02:00:59 2023 +0000
tree 86b6ec1e0e321ed28d19452d3fe185efa79ee1b7
parent 9ffc1d92cb4b6a599dc8803a9d84f28821c5b27d

[ir][msl] Add thread and threadgroup position.

Add support for the `GlobalInvocationId`, `LocalInvocationId`,
`LocalInvocationIndex`, `NumWorkgroups` and `WorkgroupId` entry point
attributes.

Bug: tint:1967
Change-Id: I74ffb8ca93d87db4c531dc4ba30375a6818de417
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/162823
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

src/tint/lang/msl/writer/printer/printer.cc

diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index d8532de..283d974 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -259,6 +259,31 @@
                 }
 
                 EmitFunctionParam(out, param);
+
+                if (param->Builtin().has_value()) {
+                    out << " [[";
+                    switch (param->Builtin().value()) {
+                        case core::ir::FunctionParam::Builtin::kGlobalInvocationId:
+                            out << "thread_position_in_grid";
+                            break;
+                        case core::ir::FunctionParam::Builtin::kLocalInvocationId:
+                            out << "thread_position_in_threadgroup";
+                            break;
+                        case core::ir::FunctionParam::Builtin::kLocalInvocationIndex:
+                            out << "thread_index_in_threadgroup";
+                            break;
+                        case core::ir::FunctionParam::Builtin::kNumWorkgroups:
+                            out << "threadgroups_per_grid";
+                            break;
+                        case core::ir::FunctionParam::Builtin::kWorkgroupId:
+                            out << "threadgroup_position_in_grid";
+                            break;
+
+                        default:
+                            break;
+                    }
+                    out << "]]";
+                }
             }
 
             out << ") {";

test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.msl

diff --git a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.msl b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.msl
index 28251dc..34ad493 100644
--- a/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.msl
+++ b/test/tint/types/functions/shader_io/compute_input_builtins.wgsl.expected.ir.msl
@@ -1,9 +1,6 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
 
-<dawn>/src/tint/lang/msl/writer/printer/printer.cc:247 internal compiler error: Switch() matched no cases. Type: tint::core::ir::Access
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+kernel void tint_symbol(uint3 local_invocation_id [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], uint3 global_invocation_id [[thread_position_in_grid]], uint3 workgroup_id [[threadgroup_position_in_grid]], uint3 num_workgroups [[threadgroups_per_grid]]) {
+  uint const foo = ((((local_invocation_id[0u] + local_invocation_index) + global_invocation_id[0u]) + workgroup_id[0u]) + num_workgroups[0u]);
+}