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