intrinsics: Add new struct form of modf(), frexp()

Implement these for all the writers.
SPIR-V reader not implemented (the old overloads weren't implemented either).

Deprecate the old overloads.

Fixed: tint:54
Change-Id: If66d26dbac3389ff604734f31b426abe47868b91
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59302
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/test/intrinsics/frexp.wgsl b/test/intrinsics/frexp.wgsl
index 54217fe..fed55e5 100644
--- a/test/intrinsics/frexp.wgsl
+++ b/test/intrinsics/frexp.wgsl
@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-    var exponent : i32;
-    let significand : f32 = frexp(1.23, &exponent);
+    let res = frexp(1.23);
+    let exp : i32 = res.exp;
+    let sig : f32 = res.sig;
 }
diff --git a/test/intrinsics/frexp.wgsl.expected.hlsl b/test/intrinsics/frexp.wgsl.expected.hlsl
index e35149a..4429974 100644
--- a/test/intrinsics/frexp.wgsl.expected.hlsl
+++ b/test/intrinsics/frexp.wgsl.expected.hlsl
@@ -1,13 +1,18 @@
-float tint_frexp(float param_0, inout int param_1) {
-  float float_exp;
-  float significand = frexp(param_0, float_exp);
-  param_1 = int(float_exp);
-  return significand;
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  float exp;
+  float sig = frexp(param_0, exp);
+  frexp_result result = {sig, int(exp)};
+  return result;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  int exponent = 0;
-  const float significand = tint_frexp(1.230000019f, exponent);
+  const frexp_result res = tint_frexp(1.230000019f);
+  const int exp = res.exp;
+  const float sig = res.sig;
   return;
 }
diff --git a/test/intrinsics/frexp.wgsl.expected.msl b/test/intrinsics/frexp.wgsl.expected.msl
index cda2b1b..d850760 100644
--- a/test/intrinsics/frexp.wgsl.expected.msl
+++ b/test/intrinsics/frexp.wgsl.expected.msl
@@ -1,9 +1,21 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
 kernel void tint_symbol() {
-  int exponent = 0;
-  float const significand = frexp(1.230000019f, *(&(exponent)));
+  frexp_result const res = tint_frexp(1.230000019f);
+  int const exp = res.exp;
+  float const sig = res.sig;
   return;
 }
 
diff --git a/test/intrinsics/frexp.wgsl.expected.spvasm b/test/intrinsics/frexp.wgsl.expected.spvasm
index 7d6543e..f282195 100644
--- a/test/intrinsics/frexp.wgsl.expected.spvasm
+++ b/test/intrinsics/frexp.wgsl.expected.spvasm
@@ -1,25 +1,29 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 14
+; Bound: 13
 ; Schema: 0
                OpCapability Shader
-         %11 = OpExtInstImport "GLSL.std.450"
+          %9 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
                OpName %main "main"
-               OpName %exponent "exponent"
+               OpName %_frexp_result "_frexp_result"
+               OpMemberName %_frexp_result 0 "sig"
+               OpMemberName %_frexp_result 1 "exp"
+               OpMemberDecorate %_frexp_result 0 Offset 0
+               OpMemberDecorate %_frexp_result 1 Offset 4
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
-          %8 = OpConstantNull %int
       %float = OpTypeFloat 32
+        %int = OpTypeInt 32 1
+%_frexp_result = OpTypeStruct %float %int
 %float_1_23000002 = OpConstant %float 1.23000002
        %main = OpFunction %void None %1
           %4 = OpLabel
-   %exponent = OpVariable %_ptr_Function_int Function %8
-          %9 = OpExtInst %float %11 Frexp %float_1_23000002 %exponent
+          %5 = OpExtInst %_frexp_result %9 FrexpStruct %float_1_23000002
+         %11 = OpCompositeExtract %int %5 1
+         %12 = OpCompositeExtract %float %5 0
                OpReturn
                OpFunctionEnd
diff --git a/test/intrinsics/frexp.wgsl.expected.wgsl b/test/intrinsics/frexp.wgsl.expected.wgsl
index d80f8f5..519ffb6 100644
--- a/test/intrinsics/frexp.wgsl.expected.wgsl
+++ b/test/intrinsics/frexp.wgsl.expected.wgsl
@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var exponent : i32;
-  let significand : f32 = frexp(1.230000019, &(exponent));
+  let res = frexp(1.230000019);
+  let exp : i32 = res.exp;
+  let sig : f32 = res.sig;
 }
diff --git a/test/intrinsics/gen/frexp/013caa.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/013caa.wgsl.expected.hlsl
index 9e077f5..8a964ed 100644
--- a/test/intrinsics/gen/frexp/013caa.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/013caa.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/013caa.wgsl.expected.msl b/test/intrinsics/gen/frexp/013caa.wgsl.expected.msl
index 1b4ab95..bd2fdbd 100644
--- a/test/intrinsics/gen/frexp/013caa.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/013caa.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_013caa() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/013caa.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/013caa.wgsl.expected.spvasm
index 9f2dda9..d1e2268 100644
--- a/test/intrinsics/gen/frexp/013caa.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/013caa.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/013caa.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/013caa.wgsl.expected.wgsl
index 990ed6e..0fccee6 100644
--- a/test/intrinsics/gen/frexp/013caa.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/013caa.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_013caa() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl
index 06c20aa..276869f 100644
--- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.msl b/test/intrinsics/gen/frexp/0da285.wgsl.expected.msl
index 7763244..b1cbd7b 100644
--- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, threadgroup int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_0da285(threadgroup int* const tint_symbol_1) {
-  float res = frexp(1.0f, *(&(*(tint_symbol_1))));
+  float res = tint_frexp(1.0f, &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:15: error: no matching function for call to 'frexp'
-  float res = frexp(1.0f, *(&(*(tint_symbol_1))));
-              ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int'), parameter type must be 'int &'
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int'), parameter type must be 'int &'
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int2 &' (aka 'int2 &') for 2nd argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int3 &' (aka 'int3 &') for 2nd argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int4 &' (aka 'int4 &') for 2nd argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int2 &' (aka 'int2 &') for 2nd argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int3 &' (aka 'int3 &') for 2nd argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int4 &' (aka 'int4 &') for 2nd argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm
index 3ea3539..e6de8e1 100644
--- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/0da285.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/0da285.wgsl.expected.wgsl
index a3bf648..d66ac56 100644
--- a/test/intrinsics/gen/frexp/0da285.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/0da285.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 var<workgroup> arg_1 : i32;
 
 fn frexp_0da285() {
diff --git a/test/intrinsics/gen/frexp/12f1da.wgsl b/test/intrinsics/gen/frexp/12f1da.wgsl
new file mode 100644
index 0000000..4db8bd9
--- /dev/null
+++ b/test/intrinsics/gen/frexp/12f1da.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(f32) -> _frexp_result
+fn frexp_12f1da() {
+  var res = frexp(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_12f1da();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_12f1da();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_12f1da();
+}
diff --git a/test/intrinsics/gen/frexp/12f1da.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.hlsl
new file mode 100644
index 0000000..caf7961
--- /dev/null
+++ b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  float exp;
+  float sig = frexp(param_0, exp);
+  frexp_result result = {sig, int(exp)};
+  return result;
+}
+
+void frexp_12f1da() {
+  frexp_result res = tint_frexp(1.0f);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_12f1da();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_12f1da();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_12f1da();
+  return;
+}
diff --git a/test/intrinsics/gen/frexp/12f1da.wgsl.expected.msl b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.msl
new file mode 100644
index 0000000..cf7b0cb
--- /dev/null
+++ b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_12f1da() {
+  frexp_result res = tint_frexp(1.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_12f1da();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_12f1da();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_12f1da();
+  return;
+}
+
diff --git a/test/intrinsics/gen/frexp/12f1da.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.spvasm
new file mode 100644
index 0000000..1887cc1
--- /dev/null
+++ b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.spvasm
@@ -0,0 +1,75 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+         %16 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_12f1da "frexp_12f1da"
+               OpName %_frexp_result "_frexp_result"
+               OpMemberName %_frexp_result 0 "sig"
+               OpMemberName %_frexp_result 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result 0 Offset 0
+               OpMemberDecorate %_frexp_result 1 Offset 4
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_frexp_result = OpTypeStruct %float %int
+    %float_1 = OpConstant %float 1
+%_ptr_Function__frexp_result = OpTypePointer Function %_frexp_result
+         %20 = OpConstantNull %_frexp_result
+         %21 = OpTypeFunction %void %v4float
+%frexp_12f1da = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result Function %20
+         %13 = OpExtInst %_frexp_result %16 FrexpStruct %float_1
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %21
+%tint_symbol = OpFunctionParameter %v4float
+         %24 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %26 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %27 = OpFunctionCall %void %frexp_12f1da
+         %28 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %frexp_12f1da
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %frexp_12f1da
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/frexp/12f1da.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.wgsl
new file mode 100644
index 0000000..76b85c3
--- /dev/null
+++ b/test/intrinsics/gen/frexp/12f1da.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn frexp_12f1da() {
+  var res = frexp(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_12f1da();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_12f1da();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_12f1da();
+}
diff --git a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.hlsl
index cd65967..446b4b9 100644
--- a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.msl b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.msl
index 35452ee..0a3ab3e 100644
--- a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_15edf3() {
   int2 arg_1 = 0;
-  float2 res = frexp(float2(), *(&(arg_1)));
+  float2 res = tint_frexp(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.spvasm
index e624b76..702a52c 100644
--- a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.wgsl
index 742494547..66dc841 100644
--- a/test/intrinsics/gen/frexp/15edf3.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/15edf3.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_15edf3() {
   var arg_1 : vec2<i32>;
   var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.hlsl
index e35f419..40eef60 100644
--- a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.msl b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.msl
index 2759c91..b2c38fe 100644
--- a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_19ab15() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.spvasm
index e11b4a2..ed54aad 100644
--- a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.wgsl
index 4bc6335..16cfe37 100644
--- a/test/intrinsics/gen/frexp/19ab15.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/19ab15.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_19ab15() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.hlsl
index bddd0e2..32b7550 100644
--- a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.msl b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.msl
index 5568b87..c3c1293 100644
--- a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_2052e9() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.spvasm
index 9a2809d..6646081 100644
--- a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.wgsl
index 07203f9..48eb9c0 100644
--- a/test/intrinsics/gen/frexp/2052e9.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/2052e9.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_2052e9() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl
index c7452b4..90c074a 100644
--- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.msl b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.msl
index cd1d5ea..fbab4cb 100644
--- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, threadgroup int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_40fc9b(threadgroup int3* const tint_symbol_1) {
-  float3 res = frexp(float3(), *(&(*(tint_symbol_1))));
+  float3 res = tint_frexp(float3(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float3 res = frexp(float3(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int3' (vector of 3 'int' values)), parameter type must be 'metal::int3 &' (aka 'int3 &')
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm
index 2facf1b..379017c 100644
--- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.wgsl
index 7452296..05b42fd 100644
--- a/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/40fc9b.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec3<i32>;
 
 fn frexp_40fc9b() {
diff --git a/test/intrinsics/gen/frexp/41e931.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/41e931.wgsl.expected.hlsl
index 080b770..da408a3 100644
--- a/test/intrinsics/gen/frexp/41e931.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/41e931.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/41e931.wgsl.expected.msl b/test/intrinsics/gen/frexp/41e931.wgsl.expected.msl
index a562b55..9e60f17 100644
--- a/test/intrinsics/gen/frexp/41e931.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/41e931.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_41e931() {
   int arg_1 = 0;
-  float res = frexp(1.0f, *(&(arg_1)));
+  float res = tint_frexp(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/41e931.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/41e931.wgsl.expected.spvasm
index fb8b59b..b9cce46 100644
--- a/test/intrinsics/gen/frexp/41e931.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/41e931.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/41e931.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/41e931.wgsl.expected.wgsl
index 0c6e160..afe0e12 100644
--- a/test/intrinsics/gen/frexp/41e931.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/41e931.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 fn frexp_41e931() {
   var arg_1 : i32;
   var res : f32 = frexp(1.0, &(arg_1));
diff --git a/test/intrinsics/gen/frexp/481e59.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/481e59.wgsl.expected.hlsl
index 668bc6e..ef6c94e 100644
--- a/test/intrinsics/gen/frexp/481e59.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/481e59.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/481e59.wgsl.expected.msl b/test/intrinsics/gen/frexp/481e59.wgsl.expected.msl
index 23f14ce..cb21d9b 100644
--- a/test/intrinsics/gen/frexp/481e59.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/481e59.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_481e59() {
   int arg_1 = 0;
-  float res = frexp(1.0f, *(&(arg_1)));
+  float res = tint_frexp(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/481e59.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/481e59.wgsl.expected.spvasm
index 9b2cd5a..703c948 100644
--- a/test/intrinsics/gen/frexp/481e59.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/481e59.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/481e59.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/481e59.wgsl.expected.wgsl
index ec1d927..e305a6d 100644
--- a/test/intrinsics/gen/frexp/481e59.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/481e59.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 fn frexp_481e59() {
   var arg_1 : i32;
   var res : f32 = frexp(1.0, &(arg_1));
diff --git a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.hlsl
index 53500da..84a57f9 100644
--- a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.msl b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.msl
index 0f1bc32..fc06cdb 100644
--- a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_5a141e() {
   int3 arg_1 = 0;
-  float3 res = frexp(float3(), *(&(arg_1)));
+  float3 res = tint_frexp(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.spvasm
index 25145e1..46ccae0 100644
--- a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.wgsl
index 6ee87a8..697ee96 100644
--- a/test/intrinsics/gen/frexp/5a141e.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/5a141e.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_5a141e() {
   var arg_1 : vec3<i32>;
   var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.hlsl
index 10803d1..eddcbc5 100644
--- a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.msl b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.msl
index 4bb1c13..0c30348 100644
--- a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_6d0058() {
   int3 arg_1 = 0;
-  float3 res = frexp(float3(), *(&(arg_1)));
+  float3 res = tint_frexp(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.spvasm
index 51638f6..cbeae7b 100644
--- a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.wgsl
index 172066b..336437e 100644
--- a/test/intrinsics/gen/frexp/6d0058.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/6d0058.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_6d0058() {
   var arg_1 : vec3<i32>;
   var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.hlsl
index 7960528..d967d1c 100644
--- a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.msl b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.msl
index 58fc2c7..3911450 100644
--- a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_6efa09(thread int3* const tint_symbol_2) {
-  float3 res = frexp(float3(), *(&(*(tint_symbol_2))));
+  float3 res = tint_frexp(float3(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.spvasm
index 0991859..869f633 100644
--- a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.wgsl
index 8d5c133..6605ee7 100644
--- a/test/intrinsics/gen/frexp/6efa09.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/6efa09.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 var<private> arg_1 : vec3<i32>;
 
 fn frexp_6efa09() {
diff --git a/test/intrinsics/gen/frexp/a0eb3b.wgsl b/test/intrinsics/gen/frexp/a0eb3b.wgsl
new file mode 100644
index 0000000..beca50c
--- /dev/null
+++ b/test/intrinsics/gen/frexp/a0eb3b.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(vec<3, f32>) -> _frexp_result_vec<3>
+fn frexp_a0eb3b() {
+  var res = frexp(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_a0eb3b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_a0eb3b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_a0eb3b();
+}
diff --git a/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.hlsl
new file mode 100644
index 0000000..2ca0a68
--- /dev/null
+++ b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct frexp_result_vec3 {
+  float3 sig;
+  int3 exp;
+};
+frexp_result_vec3 tint_frexp(float3 param_0) {
+  float3 exp;
+  float3 sig = frexp(param_0, exp);
+  frexp_result_vec3 result = {sig, int3(exp)};
+  return result;
+}
+
+void frexp_a0eb3b() {
+  frexp_result_vec3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_a0eb3b();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_a0eb3b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_a0eb3b();
+  return;
+}
diff --git a/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.msl b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.msl
new file mode 100644
index 0000000..6ca1b90
--- /dev/null
+++ b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec3 {
+  float3 sig;
+  int3 exp;
+};
+frexp_result_vec3 tint_frexp(float3 param_0) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_a0eb3b() {
+  frexp_result_vec3 res = tint_frexp(float3());
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_a0eb3b();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_a0eb3b();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_a0eb3b();
+  return;
+}
+
diff --git a/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.spvasm
new file mode 100644
index 0000000..3049fc1
--- /dev/null
+++ b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.spvasm
@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+         %18 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_a0eb3b "frexp_a0eb3b"
+               OpName %_frexp_result_vec3 "_frexp_result_vec3"
+               OpMemberName %_frexp_result_vec3 0 "sig"
+               OpMemberName %_frexp_result_vec3 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result_vec3 0 Offset 0
+               OpMemberDecorate %_frexp_result_vec3 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_frexp_result_vec3 = OpTypeStruct %v3float %v3int
+         %19 = OpConstantNull %v3float
+%_ptr_Function__frexp_result_vec3 = OpTypePointer Function %_frexp_result_vec3
+         %22 = OpConstantNull %_frexp_result_vec3
+         %23 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%frexp_a0eb3b = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result_vec3 Function %22
+         %13 = OpExtInst %_frexp_result_vec3 %18 FrexpStruct %19
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %23
+%tint_symbol = OpFunctionParameter %v4float
+         %26 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %28 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %30 = OpFunctionCall %void %frexp_a0eb3b
+         %31 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %frexp_a0eb3b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %frexp_a0eb3b
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.wgsl
new file mode 100644
index 0000000..fba88d3
--- /dev/null
+++ b/test/intrinsics/gen/frexp/a0eb3b.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn frexp_a0eb3b() {
+  var res = frexp(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_a0eb3b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_a0eb3b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_a0eb3b();
+}
diff --git a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.hlsl
index 68180a7..df13399 100644
--- a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.msl b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.msl
index 802e9f0..5ab05d4 100644
--- a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_a2a617(thread int* const tint_symbol_2) {
-  float res = frexp(1.0f, *(&(*(tint_symbol_2))));
+  float res = tint_frexp(1.0f, &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.spvasm
index 182842e..955c165 100644
--- a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.wgsl
index a5631237..e5fdacd 100644
--- a/test/intrinsics/gen/frexp/a2a617.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/a2a617.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 var<private> arg_1 : i32;
 
 fn frexp_a2a617() {
diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl
index f540a08..0c05d53 100644
--- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.msl b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.msl
index 7f841e0..b259049 100644
--- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, threadgroup int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_a3f940(threadgroup int2* const tint_symbol_1) {
-  float2 res = frexp(float2(), *(&(*(tint_symbol_1))));
+  float2 res = tint_frexp(float2(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float2 res = frexp(float2(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int2' (vector of 2 'int' values)), parameter type must be 'metal::int2 &' (aka 'int2 &')
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm
index 11801b6..e0c3785 100644
--- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.wgsl
index 745f7d5..b40075f 100644
--- a/test/intrinsics/gen/frexp/a3f940.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/a3f940.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec2<i32>;
 
 fn frexp_a3f940() {
diff --git a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.hlsl
index f88516c..88c9565 100644
--- a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.msl b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.msl
index 7f3a315..56dd076 100644
--- a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_a951b5() {
   int2 arg_1 = 0;
-  float2 res = frexp(float2(), *(&(arg_1)));
+  float2 res = tint_frexp(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.spvasm
index 8f5353e..49ad292 100644
--- a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.wgsl
index 44449dd..431a4a7 100644
--- a/test/intrinsics/gen/frexp/a951b5.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/a951b5.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_a951b5() {
   var arg_1 : vec2<i32>;
   var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/b45525.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/b45525.wgsl.expected.hlsl
index f9e83b8..28ee0ea 100644
--- a/test/intrinsics/gen/frexp/b45525.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/b45525.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/b45525.wgsl.expected.msl b/test/intrinsics/gen/frexp/b45525.wgsl.expected.msl
index d77f3a2..537576a 100644
--- a/test/intrinsics/gen/frexp/b45525.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/b45525.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_b45525(thread int4* const tint_symbol_2) {
-  float4 res = frexp(float4(), *(&(*(tint_symbol_2))));
+  float4 res = tint_frexp(float4(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/b45525.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/b45525.wgsl.expected.spvasm
index 8280479..f3c321a 100644
--- a/test/intrinsics/gen/frexp/b45525.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/b45525.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/b45525.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/b45525.wgsl.expected.wgsl
index d9cff0a..b31806c 100644
--- a/test/intrinsics/gen/frexp/b45525.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/b45525.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 var<private> arg_1 : vec4<i32>;
 
 fn frexp_b45525() {
diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl
index ddb2f5f..a39bf1e 100644
--- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.msl b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.msl
index dbbfcb5..d3107ae 100644
--- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, threadgroup int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_b87f4e(threadgroup int4* const tint_symbol_1) {
-  float4 res = frexp(float4(), *(&(*(tint_symbol_1))));
+  float4 res = tint_frexp(float4(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float4 res = frexp(float4(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int4' (vector of 4 'int' values)), parameter type must be 'metal::int4 &' (aka 'int4 &')
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm
index f8c3322..1d9d9af 100644
--- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.wgsl
index b9ed9da..2d23758 100644
--- a/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/b87f4e.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec4<i32>;
 
 fn frexp_b87f4e() {
diff --git a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.hlsl
index 1b06a32..527404b 100644
--- a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b9e4de.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.msl b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.msl
index 55c997b..937b87f 100644
--- a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/b9e4de.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_b9e4de() {
   int3 arg_1 = 0;
-  float3 res = frexp(float3(), *(&(arg_1)));
+  float3 res = tint_frexp(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.spvasm
index 1d6285c..0687916 100644
--- a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b9e4de.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.wgsl
index 38c0091..7c137ca 100644
--- a/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/b9e4de.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b9e4de.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_b9e4de() {
   var arg_1 : vec3<i32>;
   var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.hlsl
index 7db21c8..b264ed9 100644
--- a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/c084e3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.msl b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.msl
index b2c4648..3bd387e 100644
--- a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/c084e3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_c084e3(thread int2* const tint_symbol_2) {
-  float2 res = frexp(float2(), *(&(*(tint_symbol_2))));
+  float2 res = tint_frexp(float2(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.spvasm
index 5e2d7a9..ac5c4a0 100644
--- a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/c084e3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.wgsl
index 7a2d8c6..f924d31 100644
--- a/test/intrinsics/gen/frexp/c084e3.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/c084e3.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/c084e3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 var<private> arg_1 : vec2<i32>;
 
 fn frexp_c084e3() {
diff --git a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.hlsl
index b806072..1efdb62 100644
--- a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/d06c2c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.msl b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.msl
index 4b10b99..8f580e1 100644
--- a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/d06c2c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_d06c2c() {
   int2 arg_1 = 0;
-  float2 res = frexp(float2(), *(&(arg_1)));
+  float2 res = tint_frexp(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.spvasm
index 45cf3cc..2b15722 100644
--- a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/d06c2c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.wgsl
index 1dcaaf6..1680cd9 100644
--- a/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/d06c2c.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/d06c2c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_d06c2c() {
   var arg_1 : vec2<i32>;
   var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/frexp/d80367.wgsl b/test/intrinsics/gen/frexp/d80367.wgsl
new file mode 100644
index 0000000..bd8112b
--- /dev/null
+++ b/test/intrinsics/gen/frexp/d80367.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(vec<4, f32>) -> _frexp_result_vec<4>
+fn frexp_d80367() {
+  var res = frexp(vec4<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_d80367();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_d80367();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_d80367();
+}
diff --git a/test/intrinsics/gen/frexp/d80367.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/d80367.wgsl.expected.hlsl
new file mode 100644
index 0000000..18b14d4
--- /dev/null
+++ b/test/intrinsics/gen/frexp/d80367.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct frexp_result_vec4 {
+  float4 sig;
+  int4 exp;
+};
+frexp_result_vec4 tint_frexp(float4 param_0) {
+  float4 exp;
+  float4 sig = frexp(param_0, exp);
+  frexp_result_vec4 result = {sig, int4(exp)};
+  return result;
+}
+
+void frexp_d80367() {
+  frexp_result_vec4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_d80367();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_d80367();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_d80367();
+  return;
+}
diff --git a/test/intrinsics/gen/frexp/d80367.wgsl.expected.msl b/test/intrinsics/gen/frexp/d80367.wgsl.expected.msl
new file mode 100644
index 0000000..036089b
--- /dev/null
+++ b/test/intrinsics/gen/frexp/d80367.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec4 {
+  float4 sig;
+  int4 exp;
+};
+frexp_result_vec4 tint_frexp(float4 param_0) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_d80367() {
+  frexp_result_vec4 res = tint_frexp(float4());
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_d80367();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_d80367();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_d80367();
+  return;
+}
+
diff --git a/test/intrinsics/gen/frexp/d80367.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/d80367.wgsl.expected.spvasm
new file mode 100644
index 0000000..03f674e
--- /dev/null
+++ b/test/intrinsics/gen/frexp/d80367.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+         %17 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_d80367 "frexp_d80367"
+               OpName %_frexp_result_vec4 "_frexp_result_vec4"
+               OpMemberName %_frexp_result_vec4 0 "sig"
+               OpMemberName %_frexp_result_vec4 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result_vec4 0 Offset 0
+               OpMemberDecorate %_frexp_result_vec4 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
+%_frexp_result_vec4 = OpTypeStruct %v4float %v4int
+%_ptr_Function__frexp_result_vec4 = OpTypePointer Function %_frexp_result_vec4
+         %20 = OpConstantNull %_frexp_result_vec4
+         %21 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%frexp_d80367 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result_vec4 Function %20
+         %13 = OpExtInst %_frexp_result_vec4 %17 FrexpStruct %8
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %21
+%tint_symbol = OpFunctionParameter %v4float
+         %24 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %26 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %28 = OpFunctionCall %void %frexp_d80367
+         %29 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %frexp_d80367
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %frexp_d80367
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/frexp/d80367.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/d80367.wgsl.expected.wgsl
new file mode 100644
index 0000000..0c60003
--- /dev/null
+++ b/test/intrinsics/gen/frexp/d80367.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn frexp_d80367() {
+  var res = frexp(vec4<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_d80367();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_d80367();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_d80367();
+}
diff --git a/test/intrinsics/gen/frexp/db0637.wgsl b/test/intrinsics/gen/frexp/db0637.wgsl
new file mode 100644
index 0000000..9afab3a
--- /dev/null
+++ b/test/intrinsics/gen/frexp/db0637.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(vec<2, f32>) -> _frexp_result_vec<2>
+fn frexp_db0637() {
+  var res = frexp(vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_db0637();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_db0637();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_db0637();
+}
diff --git a/test/intrinsics/gen/frexp/db0637.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/db0637.wgsl.expected.hlsl
new file mode 100644
index 0000000..8ae0b21
--- /dev/null
+++ b/test/intrinsics/gen/frexp/db0637.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct frexp_result_vec2 {
+  float2 sig;
+  int2 exp;
+};
+frexp_result_vec2 tint_frexp(float2 param_0) {
+  float2 exp;
+  float2 sig = frexp(param_0, exp);
+  frexp_result_vec2 result = {sig, int2(exp)};
+  return result;
+}
+
+void frexp_db0637() {
+  frexp_result_vec2 res = tint_frexp(float2(0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_db0637();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_db0637();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_db0637();
+  return;
+}
diff --git a/test/intrinsics/gen/frexp/db0637.wgsl.expected.msl b/test/intrinsics/gen/frexp/db0637.wgsl.expected.msl
new file mode 100644
index 0000000..26b3332
--- /dev/null
+++ b/test/intrinsics/gen/frexp/db0637.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec2 {
+  float2 sig;
+  int2 exp;
+};
+frexp_result_vec2 tint_frexp(float2 param_0) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_db0637() {
+  frexp_result_vec2 res = tint_frexp(float2());
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_db0637();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_db0637();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_db0637();
+  return;
+}
+
diff --git a/test/intrinsics/gen/frexp/db0637.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/db0637.wgsl.expected.spvasm
new file mode 100644
index 0000000..df21833
--- /dev/null
+++ b/test/intrinsics/gen/frexp/db0637.wgsl.expected.spvasm
@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+         %18 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_db0637 "frexp_db0637"
+               OpName %_frexp_result_vec2 "_frexp_result_vec2"
+               OpMemberName %_frexp_result_vec2 0 "sig"
+               OpMemberName %_frexp_result_vec2 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result_vec2 0 Offset 0
+               OpMemberDecorate %_frexp_result_vec2 1 Offset 8
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v2float = OpTypeVector %float 2
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+%_frexp_result_vec2 = OpTypeStruct %v2float %v2int
+         %19 = OpConstantNull %v2float
+%_ptr_Function__frexp_result_vec2 = OpTypePointer Function %_frexp_result_vec2
+         %22 = OpConstantNull %_frexp_result_vec2
+         %23 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%frexp_db0637 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result_vec2 Function %22
+         %13 = OpExtInst %_frexp_result_vec2 %18 FrexpStruct %19
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %23
+%tint_symbol = OpFunctionParameter %v4float
+         %26 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %28 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %30 = OpFunctionCall %void %frexp_db0637
+         %31 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %frexp_db0637
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %frexp_db0637
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/frexp/db0637.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/db0637.wgsl.expected.wgsl
new file mode 100644
index 0000000..c56e615
--- /dev/null
+++ b/test/intrinsics/gen/frexp/db0637.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn frexp_db0637() {
+  var res = frexp(vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_db0637();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_db0637();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_db0637();
+}
diff --git a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.hlsl b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.hlsl
index 8aa4d20..3fd7e82 100644
--- a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/e061dd.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);
diff --git a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.msl b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.msl
index 5b7adae..a773306 100644
--- a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.msl
+++ b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/e061dd.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_e061dd() {
   int arg_1 = 0;
-  float res = frexp(1.0f, *(&(arg_1)));
+  float res = tint_frexp(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.spvasm b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.spvasm
index b580d60..1504a09 100644
--- a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/e061dd.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.wgsl b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.wgsl
index a69a43c..feb5bad 100644
--- a/test/intrinsics/gen/frexp/e061dd.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/frexp/e061dd.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/e061dd.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 fn frexp_e061dd() {
   var arg_1 : i32;
   var res : f32 = frexp(1.0, &(arg_1));
diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl
index 5f0f807..2207d2e 100644
--- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/1d59e5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 groupshared float4 arg_1;
 
 void modf_1d59e5() {
diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.msl b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.msl
index ae1675b..a078817 100644
--- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/modf/1d59e5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_modf(float4 param_0, threadgroup float4* param_1) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 void modf_1d59e5(threadgroup float4* const tint_symbol_1) {
-  float4 res = modf(float4(), *(&(*(tint_symbol_1))));
+  float4 res = tint_modf(float4(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'modf'
-  float4 res = modf(float4(), *(&(*(tint_symbol_1))));
-               ^~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5342:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup float4' (vector of 4 'float' values)), parameter type must be 'metal::float4 &' (aka 'float4 &')
-METAL_FUNC float4 modf(float4 x, thread float4 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3142:17: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'half' for 1st argument
-METAL_FUNC half modf(half x, thread half &intval)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3386:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 modf(half2 x, thread half2 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3630:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 modf(half3 x, thread half3 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3874:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 modf(half4 x, thread half4 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4610:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'float' for 1st argument
-METAL_FUNC float modf(float x, thread float &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4854:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 modf(float2 x, thread float2 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5098:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 modf(float3 x, thread float3 &intval)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm
index b999664..717b2e2e 100644
--- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/1d59e5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.wgsl b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.wgsl
index d94d6b1..3a82932 100644
--- a/test/intrinsics/gen/modf/1d59e5.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/1d59e5.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/1d59e5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 var<workgroup> arg_1 : vec4<f32>;
 
 fn modf_1d59e5() {
diff --git a/test/intrinsics/gen/modf/2199f1.wgsl b/test/intrinsics/gen/modf/2199f1.wgsl
new file mode 100644
index 0000000..c9cecd6
--- /dev/null
+++ b/test/intrinsics/gen/modf/2199f1.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn modf(vec<3, f32>) -> _modf_result_vec<3>
+fn modf_2199f1() {
+  var res = modf(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_2199f1();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_2199f1();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_2199f1();
+}
diff --git a/test/intrinsics/gen/modf/2199f1.wgsl.expected.hlsl b/test/intrinsics/gen/modf/2199f1.wgsl.expected.hlsl
new file mode 100644
index 0000000..9b4570e
--- /dev/null
+++ b/test/intrinsics/gen/modf/2199f1.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct modf_result_vec3 {
+  float3 fract;
+  float3 whole;
+};
+modf_result_vec3 tint_modf(float3 param_0) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  modf_result_vec3 result = {fract, whole};
+  return result;
+}
+
+void modf_2199f1() {
+  modf_result_vec3 res = tint_modf(float3(0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  modf_2199f1();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  modf_2199f1();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  modf_2199f1();
+  return;
+}
diff --git a/test/intrinsics/gen/modf/2199f1.wgsl.expected.msl b/test/intrinsics/gen/modf/2199f1.wgsl.expected.msl
new file mode 100644
index 0000000..233b91c
--- /dev/null
+++ b/test/intrinsics/gen/modf/2199f1.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec3 {
+  float3 fract;
+  float3 whole;
+};
+modf_result_vec3 tint_modf(float3 param_0) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  return {fract, whole};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void modf_2199f1() {
+  modf_result_vec3 res = tint_modf(float3());
+}
+
+vertex tint_symbol vertex_main() {
+  modf_2199f1();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  modf_2199f1();
+  return;
+}
+
+kernel void compute_main() {
+  modf_2199f1();
+  return;
+}
+
diff --git a/test/intrinsics/gen/modf/2199f1.wgsl.expected.spvasm b/test/intrinsics/gen/modf/2199f1.wgsl.expected.spvasm
new file mode 100644
index 0000000..998b149
--- /dev/null
+++ b/test/intrinsics/gen/modf/2199f1.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+         %16 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %modf_2199f1 "modf_2199f1"
+               OpName %_modf_result_vec3 "_modf_result_vec3"
+               OpMemberName %_modf_result_vec3 0 "fract"
+               OpMemberName %_modf_result_vec3 1 "whole"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_modf_result_vec3 0 Offset 0
+               OpMemberDecorate %_modf_result_vec3 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+%_modf_result_vec3 = OpTypeStruct %v3float %v3float
+         %17 = OpConstantNull %v3float
+%_ptr_Function__modf_result_vec3 = OpTypePointer Function %_modf_result_vec3
+         %20 = OpConstantNull %_modf_result_vec3
+         %21 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%modf_2199f1 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__modf_result_vec3 Function %20
+         %13 = OpExtInst %_modf_result_vec3 %16 ModfStruct %17
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %21
+%tint_symbol = OpFunctionParameter %v4float
+         %24 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %26 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %28 = OpFunctionCall %void %modf_2199f1
+         %29 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %modf_2199f1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %modf_2199f1
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/modf/2199f1.wgsl.expected.wgsl b/test/intrinsics/gen/modf/2199f1.wgsl.expected.wgsl
new file mode 100644
index 0000000..8a55a22
--- /dev/null
+++ b/test/intrinsics/gen/modf/2199f1.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn modf_2199f1() {
+  var res = modf(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_2199f1();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_2199f1();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_2199f1();
+}
diff --git a/test/intrinsics/gen/modf/353f7d.wgsl.expected.hlsl b/test/intrinsics/gen/modf/353f7d.wgsl.expected.hlsl
index b677940..7661a49 100644
--- a/test/intrinsics/gen/modf/353f7d.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/353f7d.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/353f7d.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 void modf_353f7d() {
   float arg_1 = 0.0f;
   float res = modf(1.0f, arg_1);
diff --git a/test/intrinsics/gen/modf/353f7d.wgsl.expected.msl b/test/intrinsics/gen/modf/353f7d.wgsl.expected.msl
index 256dd45..67dafcf 100644
--- a/test/intrinsics/gen/modf/353f7d.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/353f7d.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/353f7d.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_modf(float param_0, thread float* param_1) {
+  float whole;
+  float fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_353f7d() {
   float arg_1 = 0.0f;
-  float res = modf(1.0f, *(&(arg_1)));
+  float res = tint_modf(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/353f7d.wgsl.expected.spvasm b/test/intrinsics/gen/modf/353f7d.wgsl.expected.spvasm
index 8252221..83763b8 100644
--- a/test/intrinsics/gen/modf/353f7d.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/353f7d.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/353f7d.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/353f7d.wgsl.expected.wgsl b/test/intrinsics/gen/modf/353f7d.wgsl.expected.wgsl
index 247d8c1..f57dfff 100644
--- a/test/intrinsics/gen/modf/353f7d.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/353f7d.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/353f7d.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 fn modf_353f7d() {
   var arg_1 : f32;
   var res : f32 = modf(1.0, &(arg_1));
diff --git a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.hlsl b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.hlsl
index 4f4e6fd..7d4c708 100644
--- a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3b79d5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 void modf_3b79d5() {
   float3 arg_1 = float3(0.0f, 0.0f, 0.0f);
   float3 res = modf(float3(0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.msl b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.msl
index 31af906..2881b41 100644
--- a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/3b79d5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_modf(float3 param_0, thread float3* param_1) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_3b79d5() {
   float3 arg_1 = 0.0f;
-  float3 res = modf(float3(), *(&(arg_1)));
+  float3 res = tint_modf(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.spvasm b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.spvasm
index 545e00f..9a2e520 100644
--- a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3b79d5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.wgsl b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.wgsl
index 78f7646..616c2b8 100644
--- a/test/intrinsics/gen/modf/3b79d5.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/3b79d5.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3b79d5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_3b79d5() {
   var arg_1 : vec3<f32>;
   var res : vec3<f32> = modf(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.hlsl b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.hlsl
index 2020461..5f97ee4 100644
--- a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3d00e2.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 static float4 arg_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
 
 void modf_3d00e2() {
diff --git a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.msl b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.msl
index d7ce378..10b1b38 100644
--- a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/modf/3d00e2.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_modf(float4 param_0, thread float4* param_1) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_3d00e2(thread float4* const tint_symbol_2) {
-  float4 res = modf(float4(), *(&(*(tint_symbol_2))));
+  float4 res = tint_modf(float4(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.spvasm b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.spvasm
index 2bc383b..35ccc6a 100644
--- a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3d00e2.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.wgsl b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.wgsl
index 6175df8..8f08f94 100644
--- a/test/intrinsics/gen/modf/3d00e2.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/3d00e2.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/3d00e2.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 var<private> arg_1 : vec4<f32>;
 
 fn modf_3d00e2() {
diff --git a/test/intrinsics/gen/modf/4bb324.wgsl.expected.hlsl b/test/intrinsics/gen/modf/4bb324.wgsl.expected.hlsl
index 3401374..caf030c 100644
--- a/test/intrinsics/gen/modf/4bb324.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/4bb324.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4bb324.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 void modf_4bb324() {
   float4 arg_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
   float4 res = modf(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/4bb324.wgsl.expected.msl b/test/intrinsics/gen/modf/4bb324.wgsl.expected.msl
index 50d7c15..e87e26c 100644
--- a/test/intrinsics/gen/modf/4bb324.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/4bb324.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/4bb324.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_modf(float4 param_0, thread float4* param_1) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_4bb324() {
   float4 arg_1 = 0.0f;
-  float4 res = modf(float4(), *(&(arg_1)));
+  float4 res = tint_modf(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/4bb324.wgsl.expected.spvasm b/test/intrinsics/gen/modf/4bb324.wgsl.expected.spvasm
index 6c2baae..e654ee6 100644
--- a/test/intrinsics/gen/modf/4bb324.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/4bb324.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4bb324.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/4bb324.wgsl.expected.wgsl b/test/intrinsics/gen/modf/4bb324.wgsl.expected.wgsl
index 4cc1f48..61a8ca8 100644
--- a/test/intrinsics/gen/modf/4bb324.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/4bb324.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4bb324.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_4bb324() {
   var arg_1 : vec4<f32>;
   var res : vec4<f32> = modf(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.hlsl b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.hlsl
index dec22ac..2731d81 100644
--- a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4fe3d9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 void modf_4fe3d9() {
   float3 arg_1 = float3(0.0f, 0.0f, 0.0f);
   float3 res = modf(float3(0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.msl b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.msl
index 49e45d3..66c8950 100644
--- a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/4fe3d9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_modf(float3 param_0, thread float3* param_1) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_4fe3d9() {
   float3 arg_1 = 0.0f;
-  float3 res = modf(float3(), *(&(arg_1)));
+  float3 res = tint_modf(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.spvasm b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.spvasm
index 9bd4558..a5dbaa0 100644
--- a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4fe3d9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.wgsl b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.wgsl
index 0f8e2a3..c145208 100644
--- a/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/4fe3d9.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/4fe3d9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_4fe3d9() {
   var arg_1 : vec3<f32>;
   var res : vec3<f32> = modf(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.hlsl b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.hlsl
index 8b28ed7..0d22edf 100644
--- a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/51e4c6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 void modf_51e4c6() {
   float2 arg_1 = float2(0.0f, 0.0f);
   float2 res = modf(float2(0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.msl b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.msl
index 7cdaac4..ad760d5 100644
--- a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/51e4c6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_modf(float2 param_0, thread float2* param_1) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_51e4c6() {
   float2 arg_1 = 0.0f;
-  float2 res = modf(float2(), *(&(arg_1)));
+  float2 res = tint_modf(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.spvasm b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.spvasm
index 8952ffa..515cd31 100644
--- a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/51e4c6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.wgsl b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.wgsl
index 209a9fb6..700ffe5 100644
--- a/test/intrinsics/gen/modf/51e4c6.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/51e4c6.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/51e4c6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_51e4c6() {
   var arg_1 : vec2<f32>;
   var res : vec2<f32> = modf(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/546e09.wgsl.expected.hlsl b/test/intrinsics/gen/modf/546e09.wgsl.expected.hlsl
index 2df7bdf..1fcd41c 100644
--- a/test/intrinsics/gen/modf/546e09.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/546e09.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/546e09.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 void modf_546e09() {
   float arg_1 = 0.0f;
   float res = modf(1.0f, arg_1);
diff --git a/test/intrinsics/gen/modf/546e09.wgsl.expected.msl b/test/intrinsics/gen/modf/546e09.wgsl.expected.msl
index 64e3475..a00cf02 100644
--- a/test/intrinsics/gen/modf/546e09.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/546e09.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/546e09.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_modf(float param_0, thread float* param_1) {
+  float whole;
+  float fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_546e09() {
   float arg_1 = 0.0f;
-  float res = modf(1.0f, *(&(arg_1)));
+  float res = tint_modf(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/546e09.wgsl.expected.spvasm b/test/intrinsics/gen/modf/546e09.wgsl.expected.spvasm
index 5a56fcf..8c70ec3 100644
--- a/test/intrinsics/gen/modf/546e09.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/546e09.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/546e09.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/546e09.wgsl.expected.wgsl b/test/intrinsics/gen/modf/546e09.wgsl.expected.wgsl
index 331c087..ab9908d 100644
--- a/test/intrinsics/gen/modf/546e09.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/546e09.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/546e09.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 fn modf_546e09() {
   var arg_1 : f32;
   var res : f32 = modf(1.0, &(arg_1));
diff --git a/test/intrinsics/gen/modf/5e8476.wgsl.expected.hlsl b/test/intrinsics/gen/modf/5e8476.wgsl.expected.hlsl
index 8fabfe2..dbe27e6 100644
--- a/test/intrinsics/gen/modf/5e8476.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/5e8476.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/5e8476.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 static float arg_1 = 0.0f;
 
 void modf_5e8476() {
diff --git a/test/intrinsics/gen/modf/5e8476.wgsl.expected.msl b/test/intrinsics/gen/modf/5e8476.wgsl.expected.msl
index 7b9a4e6..531930f 100644
--- a/test/intrinsics/gen/modf/5e8476.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/5e8476.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/modf/5e8476.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_modf(float param_0, thread float* param_1) {
+  float whole;
+  float fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_5e8476(thread float* const tint_symbol_2) {
-  float res = modf(1.0f, *(&(*(tint_symbol_2))));
+  float res = tint_modf(1.0f, &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/5e8476.wgsl.expected.spvasm b/test/intrinsics/gen/modf/5e8476.wgsl.expected.spvasm
index f3a8bdb..f9c55e2 100644
--- a/test/intrinsics/gen/modf/5e8476.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/5e8476.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/5e8476.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/5e8476.wgsl.expected.wgsl b/test/intrinsics/gen/modf/5e8476.wgsl.expected.wgsl
index 6d26235..5888c21 100644
--- a/test/intrinsics/gen/modf/5e8476.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/5e8476.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/5e8476.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 var<private> arg_1 : f32;
 
 fn modf_5e8476() {
diff --git a/test/intrinsics/gen/modf/684d46.wgsl b/test/intrinsics/gen/modf/684d46.wgsl
new file mode 100644
index 0000000..d983b72
--- /dev/null
+++ b/test/intrinsics/gen/modf/684d46.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn modf(f32) -> _modf_result
+fn modf_684d46() {
+  var res = modf(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_684d46();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_684d46();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_684d46();
+}
diff --git a/test/intrinsics/gen/modf/684d46.wgsl.expected.hlsl b/test/intrinsics/gen/modf/684d46.wgsl.expected.hlsl
new file mode 100644
index 0000000..77b05b7
--- /dev/null
+++ b/test/intrinsics/gen/modf/684d46.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct modf_result {
+  float fract;
+  float whole;
+};
+modf_result tint_modf(float param_0) {
+  float whole;
+  float fract = modf(param_0, whole);
+  modf_result result = {fract, whole};
+  return result;
+}
+
+void modf_684d46() {
+  modf_result res = tint_modf(1.0f);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  modf_684d46();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  modf_684d46();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  modf_684d46();
+  return;
+}
diff --git a/test/intrinsics/gen/modf/684d46.wgsl.expected.msl b/test/intrinsics/gen/modf/684d46.wgsl.expected.msl
new file mode 100644
index 0000000..9944617
--- /dev/null
+++ b/test/intrinsics/gen/modf/684d46.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result {
+  float fract;
+  float whole;
+};
+modf_result tint_modf(float param_0) {
+  float whole;
+  float fract = modf(param_0, whole);
+  return {fract, whole};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void modf_684d46() {
+  modf_result res = tint_modf(1.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  modf_684d46();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  modf_684d46();
+  return;
+}
+
+kernel void compute_main() {
+  modf_684d46();
+  return;
+}
+
diff --git a/test/intrinsics/gen/modf/684d46.wgsl.expected.spvasm b/test/intrinsics/gen/modf/684d46.wgsl.expected.spvasm
new file mode 100644
index 0000000..9a37b68
--- /dev/null
+++ b/test/intrinsics/gen/modf/684d46.wgsl.expected.spvasm
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+         %15 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %modf_684d46 "modf_684d46"
+               OpName %_modf_result "_modf_result"
+               OpMemberName %_modf_result 0 "fract"
+               OpMemberName %_modf_result 1 "whole"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_modf_result 0 Offset 0
+               OpMemberDecorate %_modf_result 1 Offset 4
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+%_modf_result = OpTypeStruct %float %float
+    %float_1 = OpConstant %float 1
+%_ptr_Function__modf_result = OpTypePointer Function %_modf_result
+         %19 = OpConstantNull %_modf_result
+         %20 = OpTypeFunction %void %v4float
+%modf_684d46 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__modf_result Function %19
+         %13 = OpExtInst %_modf_result %15 ModfStruct %float_1
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %20
+%tint_symbol = OpFunctionParameter %v4float
+         %23 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %25 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %26 = OpFunctionCall %void %modf_684d46
+         %27 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %modf_684d46
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %modf_684d46
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/modf/684d46.wgsl.expected.wgsl b/test/intrinsics/gen/modf/684d46.wgsl.expected.wgsl
new file mode 100644
index 0000000..3686ead
--- /dev/null
+++ b/test/intrinsics/gen/modf/684d46.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn modf_684d46() {
+  var res = modf(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_684d46();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_684d46();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_684d46();
+}
diff --git a/test/intrinsics/gen/modf/86441c.wgsl.expected.hlsl b/test/intrinsics/gen/modf/86441c.wgsl.expected.hlsl
index 21ca45c..8106372 100644
--- a/test/intrinsics/gen/modf/86441c.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/86441c.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/86441c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 void modf_86441c() {
   float2 arg_1 = float2(0.0f, 0.0f);
   float2 res = modf(float2(0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/86441c.wgsl.expected.msl b/test/intrinsics/gen/modf/86441c.wgsl.expected.msl
index 900085d..bd5a595 100644
--- a/test/intrinsics/gen/modf/86441c.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/86441c.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/86441c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_modf(float2 param_0, thread float2* param_1) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_86441c() {
   float2 arg_1 = 0.0f;
-  float2 res = modf(float2(), *(&(arg_1)));
+  float2 res = tint_modf(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/86441c.wgsl.expected.spvasm b/test/intrinsics/gen/modf/86441c.wgsl.expected.spvasm
index 7a41b14..3567057 100644
--- a/test/intrinsics/gen/modf/86441c.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/86441c.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/86441c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/86441c.wgsl.expected.wgsl b/test/intrinsics/gen/modf/86441c.wgsl.expected.wgsl
index a9c1e38..f2940fe 100644
--- a/test/intrinsics/gen/modf/86441c.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/86441c.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/86441c.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_86441c() {
   var arg_1 : vec2<f32>;
   var res : vec2<f32> = modf(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/955651.wgsl.expected.hlsl b/test/intrinsics/gen/modf/955651.wgsl.expected.hlsl
index 02cf941..de9a0a0 100644
--- a/test/intrinsics/gen/modf/955651.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/955651.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/955651.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 void modf_955651() {
   float3 arg_1 = float3(0.0f, 0.0f, 0.0f);
   float3 res = modf(float3(0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/955651.wgsl.expected.msl b/test/intrinsics/gen/modf/955651.wgsl.expected.msl
index 25b521c..937fbe1 100644
--- a/test/intrinsics/gen/modf/955651.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/955651.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/955651.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_modf(float3 param_0, thread float3* param_1) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_955651() {
   float3 arg_1 = 0.0f;
-  float3 res = modf(float3(), *(&(arg_1)));
+  float3 res = tint_modf(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/955651.wgsl.expected.spvasm b/test/intrinsics/gen/modf/955651.wgsl.expected.spvasm
index 192d324..01fe4be 100644
--- a/test/intrinsics/gen/modf/955651.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/955651.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/955651.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/955651.wgsl.expected.wgsl b/test/intrinsics/gen/modf/955651.wgsl.expected.wgsl
index bf94c38..aab1a1e 100644
--- a/test/intrinsics/gen/modf/955651.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/955651.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/955651.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_955651() {
   var arg_1 : vec3<f32>;
   var res : vec3<f32> = modf(vec3<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/9b44a9.wgsl b/test/intrinsics/gen/modf/9b44a9.wgsl
new file mode 100644
index 0000000..2556b01
--- /dev/null
+++ b/test/intrinsics/gen/modf/9b44a9.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn modf(vec<4, f32>) -> _modf_result_vec<4>
+fn modf_9b44a9() {
+  var res = modf(vec4<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_9b44a9();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_9b44a9();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_9b44a9();
+}
diff --git a/test/intrinsics/gen/modf/9b44a9.wgsl.expected.hlsl b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.hlsl
new file mode 100644
index 0000000..56faa20
--- /dev/null
+++ b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct modf_result_vec4 {
+  float4 fract;
+  float4 whole;
+};
+modf_result_vec4 tint_modf(float4 param_0) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  modf_result_vec4 result = {fract, whole};
+  return result;
+}
+
+void modf_9b44a9() {
+  modf_result_vec4 res = tint_modf(float4(0.0f, 0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  modf_9b44a9();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  modf_9b44a9();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  modf_9b44a9();
+  return;
+}
diff --git a/test/intrinsics/gen/modf/9b44a9.wgsl.expected.msl b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.msl
new file mode 100644
index 0000000..86541b3
--- /dev/null
+++ b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec4 {
+  float4 fract;
+  float4 whole;
+};
+modf_result_vec4 tint_modf(float4 param_0) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  return {fract, whole};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void modf_9b44a9() {
+  modf_result_vec4 res = tint_modf(float4());
+}
+
+vertex tint_symbol vertex_main() {
+  modf_9b44a9();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  modf_9b44a9();
+  return;
+}
+
+kernel void compute_main() {
+  modf_9b44a9();
+  return;
+}
+
diff --git a/test/intrinsics/gen/modf/9b44a9.wgsl.expected.spvasm b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.spvasm
new file mode 100644
index 0000000..c15a37c
--- /dev/null
+++ b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.spvasm
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+               OpCapability Shader
+         %15 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %modf_9b44a9 "modf_9b44a9"
+               OpName %_modf_result_vec4 "_modf_result_vec4"
+               OpMemberName %_modf_result_vec4 0 "fract"
+               OpMemberName %_modf_result_vec4 1 "whole"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_modf_result_vec4 0 Offset 0
+               OpMemberDecorate %_modf_result_vec4 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+%_modf_result_vec4 = OpTypeStruct %v4float %v4float
+%_ptr_Function__modf_result_vec4 = OpTypePointer Function %_modf_result_vec4
+         %18 = OpConstantNull %_modf_result_vec4
+         %19 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%modf_9b44a9 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__modf_result_vec4 Function %18
+         %13 = OpExtInst %_modf_result_vec4 %15 ModfStruct %8
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %19
+%tint_symbol = OpFunctionParameter %v4float
+         %22 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %24 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %26 = OpFunctionCall %void %modf_9b44a9
+         %27 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %29 = OpLabel
+         %30 = OpFunctionCall %void %modf_9b44a9
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %32 = OpLabel
+         %33 = OpFunctionCall %void %modf_9b44a9
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/modf/9b44a9.wgsl.expected.wgsl b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.wgsl
new file mode 100644
index 0000000..2009bff
--- /dev/null
+++ b/test/intrinsics/gen/modf/9b44a9.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn modf_9b44a9() {
+  var res = modf(vec4<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_9b44a9();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_9b44a9();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_9b44a9();
+}
diff --git a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.hlsl b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.hlsl
index 26b8eaf..112167d 100644
--- a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9c6a91.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 static float2 arg_1 = float2(0.0f, 0.0f);
 
 void modf_9c6a91() {
diff --git a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.msl b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.msl
index 475c8e0..b702b5c 100644
--- a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/modf/9c6a91.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_modf(float2 param_0, thread float2* param_1) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_9c6a91(thread float2* const tint_symbol_2) {
-  float2 res = modf(float2(), *(&(*(tint_symbol_2))));
+  float2 res = tint_modf(float2(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.spvasm b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.spvasm
index 73bad5d..5ebe38b 100644
--- a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9c6a91.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.wgsl b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.wgsl
index 87aae4f..dff1475 100644
--- a/test/intrinsics/gen/modf/9c6a91.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/9c6a91.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9c6a91.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 var<private> arg_1 : vec2<f32>;
 
 fn modf_9c6a91() {
diff --git a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.hlsl b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.hlsl
index 72befc3..1197885 100644
--- a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9cecfc.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 static float3 arg_1 = float3(0.0f, 0.0f, 0.0f);
 
 void modf_9cecfc() {
diff --git a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.msl b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.msl
index 617d3c7..2a8186e 100644
--- a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.msl
@@ -1,12 +1,24 @@
+intrinsics/gen/modf/9cecfc.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_modf(float3 param_0, thread float3* param_1) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_9cecfc(thread float3* const tint_symbol_2) {
-  float3 res = modf(float3(), *(&(*(tint_symbol_2))));
+  float3 res = tint_modf(float3(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.spvasm b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.spvasm
index 0700451..70dd5db 100644
--- a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9cecfc.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.wgsl b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.wgsl
index 48ea3c8..7ad6c49 100644
--- a/test/intrinsics/gen/modf/9cecfc.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/9cecfc.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/9cecfc.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 var<private> arg_1 : vec3<f32>;
 
 fn modf_9cecfc() {
diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl b/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl
index b900c7b..8bb5fd8 100644
--- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a128ab.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 groupshared float2 arg_1;
 
 void modf_a128ab() {
diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.msl b/test/intrinsics/gen/modf/a128ab.wgsl.expected.msl
index bc45ed1..175efd8 100644
--- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/modf/a128ab.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_modf(float2 param_0, threadgroup float2* param_1) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 void modf_a128ab(threadgroup float2* const tint_symbol_1) {
-  float2 res = modf(float2(), *(&(*(tint_symbol_1))));
+  float2 res = tint_modf(float2(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'modf'
-  float2 res = modf(float2(), *(&(*(tint_symbol_1))));
-               ^~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4854:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup float2' (vector of 2 'float' values)), parameter type must be 'metal::float2 &' (aka 'float2 &')
-METAL_FUNC float2 modf(float2 x, thread float2 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3142:17: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'half' for 1st argument
-METAL_FUNC half modf(half x, thread half &intval)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3386:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 modf(half2 x, thread half2 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3630:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 modf(half3 x, thread half3 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3874:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 modf(half4 x, thread half4 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4610:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'float' for 1st argument
-METAL_FUNC float modf(float x, thread float &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5098:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 modf(float3 x, thread float3 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5342:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 modf(float4 x, thread float4 &intval)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm b/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm
index ef64e6d..bb3dca7 100644
--- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a128ab.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/a128ab.wgsl.expected.wgsl b/test/intrinsics/gen/modf/a128ab.wgsl.expected.wgsl
index fffc39b..89a4a28 100644
--- a/test/intrinsics/gen/modf/a128ab.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/a128ab.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a128ab.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 var<workgroup> arg_1 : vec2<f32>;
 
 fn modf_a128ab() {
diff --git a/test/intrinsics/gen/modf/a54eca.wgsl.expected.hlsl b/test/intrinsics/gen/modf/a54eca.wgsl.expected.hlsl
index 59e0e3c..d3ed29a 100644
--- a/test/intrinsics/gen/modf/a54eca.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/a54eca.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a54eca.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 void modf_a54eca() {
   float2 arg_1 = float2(0.0f, 0.0f);
   float2 res = modf(float2(0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/a54eca.wgsl.expected.msl b/test/intrinsics/gen/modf/a54eca.wgsl.expected.msl
index 6c583ed..94ea876 100644
--- a/test/intrinsics/gen/modf/a54eca.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/a54eca.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/a54eca.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_modf(float2 param_0, thread float2* param_1) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_a54eca() {
   float2 arg_1 = 0.0f;
-  float2 res = modf(float2(), *(&(arg_1)));
+  float2 res = tint_modf(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/a54eca.wgsl.expected.spvasm b/test/intrinsics/gen/modf/a54eca.wgsl.expected.spvasm
index c8d52c5..99928e1 100644
--- a/test/intrinsics/gen/modf/a54eca.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/a54eca.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a54eca.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/a54eca.wgsl.expected.wgsl b/test/intrinsics/gen/modf/a54eca.wgsl.expected.wgsl
index 39c4a77..4e886cd 100644
--- a/test/intrinsics/gen/modf/a54eca.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/a54eca.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/a54eca.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = modf(vec2<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_a54eca() {
   var arg_1 : vec2<f32>;
   var res : vec2<f32> = modf(vec2<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl b/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl
index 22f570d..7a37367 100644
--- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/bb9088.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 groupshared float3 arg_1;
 
 void modf_bb9088() {
diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.msl b/test/intrinsics/gen/modf/bb9088.wgsl.expected.msl
index 44671b2..45dc355 100644
--- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/modf/bb9088.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_modf(float3 param_0, threadgroup float3* param_1) {
+  float3 whole;
+  float3 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 void modf_bb9088(threadgroup float3* const tint_symbol_1) {
-  float3 res = modf(float3(), *(&(*(tint_symbol_1))));
+  float3 res = tint_modf(float3(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'modf'
-  float3 res = modf(float3(), *(&(*(tint_symbol_1))));
-               ^~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5098:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup float3' (vector of 3 'float' values)), parameter type must be 'metal::float3 &' (aka 'float3 &')
-METAL_FUNC float3 modf(float3 x, thread float3 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3142:17: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'half' for 1st argument
-METAL_FUNC half modf(half x, thread half &intval)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3386:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 modf(half2 x, thread half2 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3630:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 modf(half3 x, thread half3 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3874:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 modf(half4 x, thread half4 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4610:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'float' for 1st argument
-METAL_FUNC float modf(float x, thread float &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4854:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 modf(float2 x, thread float2 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5342:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 modf(float4 x, thread float4 &intval)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm b/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm
index df3e8e6..51b1647 100644
--- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/bb9088.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/bb9088.wgsl.expected.wgsl b/test/intrinsics/gen/modf/bb9088.wgsl.expected.wgsl
index bd9e14c..50f629e 100644
--- a/test/intrinsics/gen/modf/bb9088.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/bb9088.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/bb9088.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = modf(vec3<f32>(), &arg_1);
+                       ^^^^
+
 var<workgroup> arg_1 : vec3<f32>;
 
 fn modf_bb9088() {
diff --git a/test/intrinsics/gen/modf/c87851.wgsl b/test/intrinsics/gen/modf/c87851.wgsl
new file mode 100644
index 0000000..55f7ea8
--- /dev/null
+++ b/test/intrinsics/gen/modf/c87851.wgsl
@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn modf(vec<2, f32>) -> _modf_result_vec<2>
+fn modf_c87851() {
+  var res = modf(vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_c87851();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_c87851();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_c87851();
+}
diff --git a/test/intrinsics/gen/modf/c87851.wgsl.expected.hlsl b/test/intrinsics/gen/modf/c87851.wgsl.expected.hlsl
new file mode 100644
index 0000000..bffd646
--- /dev/null
+++ b/test/intrinsics/gen/modf/c87851.wgsl.expected.hlsl
@@ -0,0 +1,35 @@
+struct modf_result_vec2 {
+  float2 fract;
+  float2 whole;
+};
+modf_result_vec2 tint_modf(float2 param_0) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  modf_result_vec2 result = {fract, whole};
+  return result;
+}
+
+void modf_c87851() {
+  modf_result_vec2 res = tint_modf(float2(0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  modf_c87851();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  modf_c87851();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  modf_c87851();
+  return;
+}
diff --git a/test/intrinsics/gen/modf/c87851.wgsl.expected.msl b/test/intrinsics/gen/modf/c87851.wgsl.expected.msl
new file mode 100644
index 0000000..4bcc460
--- /dev/null
+++ b/test/intrinsics/gen/modf/c87851.wgsl.expected.msl
@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_vec2 {
+  float2 fract;
+  float2 whole;
+};
+modf_result_vec2 tint_modf(float2 param_0) {
+  float2 whole;
+  float2 fract = modf(param_0, whole);
+  return {fract, whole};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void modf_c87851() {
+  modf_result_vec2 res = tint_modf(float2());
+}
+
+vertex tint_symbol vertex_main() {
+  modf_c87851();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  modf_c87851();
+  return;
+}
+
+kernel void compute_main() {
+  modf_c87851();
+  return;
+}
+
diff --git a/test/intrinsics/gen/modf/c87851.wgsl.expected.spvasm b/test/intrinsics/gen/modf/c87851.wgsl.expected.spvasm
new file mode 100644
index 0000000..d8cbd3d
--- /dev/null
+++ b/test/intrinsics/gen/modf/c87851.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+         %16 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %modf_c87851 "modf_c87851"
+               OpName %_modf_result_vec2 "_modf_result_vec2"
+               OpMemberName %_modf_result_vec2 0 "fract"
+               OpMemberName %_modf_result_vec2 1 "whole"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_modf_result_vec2 0 Offset 0
+               OpMemberDecorate %_modf_result_vec2 1 Offset 8
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v2float = OpTypeVector %float 2
+%_modf_result_vec2 = OpTypeStruct %v2float %v2float
+         %17 = OpConstantNull %v2float
+%_ptr_Function__modf_result_vec2 = OpTypePointer Function %_modf_result_vec2
+         %20 = OpConstantNull %_modf_result_vec2
+         %21 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%modf_c87851 = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__modf_result_vec2 Function %20
+         %13 = OpExtInst %_modf_result_vec2 %16 ModfStruct %17
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %21
+%tint_symbol = OpFunctionParameter %v4float
+         %24 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %26 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %28 = OpFunctionCall %void %modf_c87851
+         %29 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %31 = OpLabel
+         %32 = OpFunctionCall %void %modf_c87851
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %34 = OpLabel
+         %35 = OpFunctionCall %void %modf_c87851
+               OpReturn
+               OpFunctionEnd
diff --git a/test/intrinsics/gen/modf/c87851.wgsl.expected.wgsl b/test/intrinsics/gen/modf/c87851.wgsl.expected.wgsl
new file mode 100644
index 0000000..9edd1bf
--- /dev/null
+++ b/test/intrinsics/gen/modf/c87851.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn modf_c87851() {
+  var res = modf(vec2<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  modf_c87851();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  modf_c87851();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  modf_c87851();
+}
diff --git a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.hlsl b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.hlsl
index afd659d..f017b20b 100644
--- a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/d1d6f6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 void modf_d1d6f6() {
   float4 arg_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
   float4 res = modf(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.msl b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.msl
index c2dd7d2..549d90e 100644
--- a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/d1d6f6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_modf(float4 param_0, thread float4* param_1) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_d1d6f6() {
   float4 arg_1 = 0.0f;
-  float4 res = modf(float4(), *(&(arg_1)));
+  float4 res = tint_modf(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.spvasm b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.spvasm
index 1d40f64..aa8a3b8 100644
--- a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/d1d6f6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.wgsl b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.wgsl
index dd45517..113e0c1 100644
--- a/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/d1d6f6.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/d1d6f6.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_d1d6f6() {
   var arg_1 : vec4<f32>;
   var res : vec4<f32> = modf(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl
index 726d9b5..7c29406 100644
--- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e38ae6.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 groupshared float arg_1;
 
 void modf_e38ae6() {
diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.msl b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.msl
index 51877e8..5ce4bb3 100644
--- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.msl
@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/modf/e38ae6.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_modf(float param_0, threadgroup float* param_1) {
+  float whole;
+  float fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 void modf_e38ae6(threadgroup float* const tint_symbol_1) {
-  float res = modf(1.0f, *(&(*(tint_symbol_1))));
+  float res = tint_modf(1.0f, &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@
   return;
 }
 
-Compilation failed: 
-
-program_source:5:15: error: no matching function for call to 'modf'
-  float res = modf(1.0f, *(&(*(tint_symbol_1))));
-              ^~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4610:18: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup float'), parameter type must be 'float &'
-METAL_FUNC float modf(float x, thread float &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3142:17: note: candidate function not viable: no known conversion from 'threadgroup float' to 'half &' for 2nd argument
-METAL_FUNC half modf(half x, thread half &intval)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3386:18: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::half2 &' (aka 'half2 &') for 2nd argument
-METAL_FUNC half2 modf(half2 x, thread half2 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3630:18: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::half3 &' (aka 'half3 &') for 2nd argument
-METAL_FUNC half3 modf(half3 x, thread half3 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3874:18: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::half4 &' (aka 'half4 &') for 2nd argument
-METAL_FUNC half4 modf(half4 x, thread half4 &intval)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4854:19: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::float2 &' (aka 'float2 &') for 2nd argument
-METAL_FUNC float2 modf(float2 x, thread float2 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5098:19: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::float3 &' (aka 'float3 &') for 2nd argument
-METAL_FUNC float3 modf(float3 x, thread float3 &intval)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5342:19: note: candidate function not viable: no known conversion from 'threadgroup float' to 'metal::float4 &' (aka 'float4 &') for 2nd argument
-METAL_FUNC float4 modf(float4 x, thread float4 &intval)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-
diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm
index 3dfafee..b5fd9e9 100644
--- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e38ae6.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.wgsl b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.wgsl
index 6f45f8f..af186d1 100644
--- a/test/intrinsics/gen/modf/e38ae6.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/e38ae6.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e38ae6.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 var<workgroup> arg_1 : f32;
 
 fn modf_e38ae6() {
diff --git a/test/intrinsics/gen/modf/e83560.wgsl.expected.hlsl b/test/intrinsics/gen/modf/e83560.wgsl.expected.hlsl
index fbd15df..dbc6c6e 100644
--- a/test/intrinsics/gen/modf/e83560.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/e83560.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e83560.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 void modf_e83560() {
   float4 arg_1 = float4(0.0f, 0.0f, 0.0f, 0.0f);
   float4 res = modf(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
diff --git a/test/intrinsics/gen/modf/e83560.wgsl.expected.msl b/test/intrinsics/gen/modf/e83560.wgsl.expected.msl
index 922a972..28bee9f 100644
--- a/test/intrinsics/gen/modf/e83560.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/e83560.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/e83560.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_modf(float4 param_0, thread float4* param_1) {
+  float4 whole;
+  float4 fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_e83560() {
   float4 arg_1 = 0.0f;
-  float4 res = modf(float4(), *(&(arg_1)));
+  float4 res = tint_modf(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/e83560.wgsl.expected.spvasm b/test/intrinsics/gen/modf/e83560.wgsl.expected.spvasm
index 8e5622c..99f906d 100644
--- a/test/intrinsics/gen/modf/e83560.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/e83560.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e83560.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/e83560.wgsl.expected.wgsl b/test/intrinsics/gen/modf/e83560.wgsl.expected.wgsl
index 55a52e2..719781c 100644
--- a/test/intrinsics/gen/modf/e83560.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/e83560.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/e83560.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = modf(vec4<f32>(), &arg_1);
+                       ^^^^
+
 fn modf_e83560() {
   var arg_1 : vec4<f32>;
   var res : vec4<f32> = modf(vec4<f32>(), &(arg_1));
diff --git a/test/intrinsics/gen/modf/f90945.wgsl.expected.hlsl b/test/intrinsics/gen/modf/f90945.wgsl.expected.hlsl
index 5e79aa3..7255685 100644
--- a/test/intrinsics/gen/modf/f90945.wgsl.expected.hlsl
+++ b/test/intrinsics/gen/modf/f90945.wgsl.expected.hlsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/f90945.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 void modf_f90945() {
   float arg_1 = 0.0f;
   float res = modf(1.0f, arg_1);
diff --git a/test/intrinsics/gen/modf/f90945.wgsl.expected.msl b/test/intrinsics/gen/modf/f90945.wgsl.expected.msl
index eed1548..795e7ef 100644
--- a/test/intrinsics/gen/modf/f90945.wgsl.expected.msl
+++ b/test/intrinsics/gen/modf/f90945.wgsl.expected.msl
@@ -1,13 +1,25 @@
+intrinsics/gen/modf/f90945.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_modf(float param_0, thread float* param_1) {
+  float whole;
+  float fract = modf(param_0, whole);
+  *param_1 = whole;
+  return fract;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void modf_f90945() {
   float arg_1 = 0.0f;
-  float res = modf(1.0f, *(&(arg_1)));
+  float res = tint_modf(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {
diff --git a/test/intrinsics/gen/modf/f90945.wgsl.expected.spvasm b/test/intrinsics/gen/modf/f90945.wgsl.expected.spvasm
index 0872a2f..29b6948 100644
--- a/test/intrinsics/gen/modf/f90945.wgsl.expected.spvasm
+++ b/test/intrinsics/gen/modf/f90945.wgsl.expected.spvasm
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/f90945.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
diff --git a/test/intrinsics/gen/modf/f90945.wgsl.expected.wgsl b/test/intrinsics/gen/modf/f90945.wgsl.expected.wgsl
index 6665ed2..fb28341 100644
--- a/test/intrinsics/gen/modf/f90945.wgsl.expected.wgsl
+++ b/test/intrinsics/gen/modf/f90945.wgsl.expected.wgsl
@@ -1,3 +1,7 @@
+intrinsics/gen/modf/f90945.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = modf(1.0, &arg_1);
+                 ^^^^
+
 fn modf_f90945() {
   var arg_1 : f32;
   var res : f32 = modf(1.0, &(arg_1));
diff --git a/test/intrinsics/intrinsics.wgsl.tmpl b/test/intrinsics/intrinsics.wgsl.tmpl
index 1751b6d..9ed3c42 100644
--- a/test/intrinsics/intrinsics.wgsl.tmpl
+++ b/test/intrinsics/intrinsics.wgsl.tmpl
@@ -123,7 +123,7 @@
 {{- /* Make the call to the intrinsic */ -}}
 {{- /*indent*/}}  {{/*indent*/ -}}
 {{-   if .ReturnType -}}
-  var res: {{template "Type" .ReturnType}} = {{/* preserve space after = */ -}}
+  var res{{if IsDeclarable .ReturnType}}: {{template "Type" .ReturnType}}{{end}} = {{/* preserve space after = */ -}}
 {{-   end -}}
   {{$function}}(
 {{-   range $i, $p := .Parameters -}}
diff --git a/test/intrinsics/modf.wgsl b/test/intrinsics/modf.wgsl
index dae3447..f8b1cf3 100644
--- a/test/intrinsics/modf.wgsl
+++ b/test/intrinsics/modf.wgsl
@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-    var whole : f32;
-    let frac : f32 = modf(1.23, &whole);
+    let res = modf(1.23);
+    let fract : f32 = res.fract;
+    let whole : f32 = res.whole;
 }
diff --git a/test/intrinsics/modf.wgsl.expected.hlsl b/test/intrinsics/modf.wgsl.expected.hlsl
index 200d084..d2bdca4 100644
--- a/test/intrinsics/modf.wgsl.expected.hlsl
+++ b/test/intrinsics/modf.wgsl.expected.hlsl
@@ -1,6 +1,18 @@
+struct modf_result {
+  float fract;
+  float whole;
+};
+modf_result tint_modf(float param_0) {
+  float whole;
+  float fract = modf(param_0, whole);
+  modf_result result = {fract, whole};
+  return result;
+}
+
 [numthreads(1, 1, 1)]
 void main() {
-  float whole = 0.0f;
-  const float tint_symbol = modf(1.230000019f, whole);
+  const modf_result res = tint_modf(1.230000019f);
+  const float fract = res.fract;
+  const float whole = res.whole;
   return;
 }
diff --git a/test/intrinsics/modf.wgsl.expected.msl b/test/intrinsics/modf.wgsl.expected.msl
index 3f2627a..6f6cf69 100644
--- a/test/intrinsics/modf.wgsl.expected.msl
+++ b/test/intrinsics/modf.wgsl.expected.msl
@@ -1,9 +1,21 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+struct modf_result {
+  float fract;
+  float whole;
+};
+modf_result tint_modf(float param_0) {
+  float whole;
+  float fract = modf(param_0, whole);
+  return {fract, whole};
+}
+
 kernel void tint_symbol() {
-  float whole = 0.0f;
-  float const frac = modf(1.230000019f, *(&(whole)));
+  modf_result const res = tint_modf(1.230000019f);
+  float const fract = res.fract;
+  float const whole = res.whole;
   return;
 }
 
diff --git a/test/intrinsics/modf.wgsl.expected.spvasm b/test/intrinsics/modf.wgsl.expected.spvasm
index 090d8b8..3122347 100644
--- a/test/intrinsics/modf.wgsl.expected.spvasm
+++ b/test/intrinsics/modf.wgsl.expected.spvasm
@@ -1,24 +1,28 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 13
+; Bound: 12
 ; Schema: 0
                OpCapability Shader
-         %10 = OpExtInstImport "GLSL.std.450"
+          %8 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
                OpName %main "main"
-               OpName %whole "whole"
+               OpName %_modf_result "_modf_result"
+               OpMemberName %_modf_result 0 "fract"
+               OpMemberName %_modf_result 1 "whole"
+               OpMemberDecorate %_modf_result 0 Offset 0
+               OpMemberDecorate %_modf_result 1 Offset 4
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
       %float = OpTypeFloat 32
-%_ptr_Function_float = OpTypePointer Function %float
-          %8 = OpConstantNull %float
+%_modf_result = OpTypeStruct %float %float
 %float_1_23000002 = OpConstant %float 1.23000002
        %main = OpFunction %void None %1
           %4 = OpLabel
-      %whole = OpVariable %_ptr_Function_float Function %8
-          %9 = OpExtInst %float %10 Modf %float_1_23000002 %whole
+          %5 = OpExtInst %_modf_result %8 ModfStruct %float_1_23000002
+         %10 = OpCompositeExtract %float %5 0
+         %11 = OpCompositeExtract %float %5 1
                OpReturn
                OpFunctionEnd
diff --git a/test/intrinsics/modf.wgsl.expected.wgsl b/test/intrinsics/modf.wgsl.expected.wgsl
index eea8562..ba6a7c6 100644
--- a/test/intrinsics/modf.wgsl.expected.wgsl
+++ b/test/intrinsics/modf.wgsl.expected.wgsl
@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var whole : f32;
-  let frac : f32 = modf(1.230000019, &(whole));
+  let res = modf(1.230000019);
+  let fract : f32 = res.fract;
+  let whole : f32 = res.whole;
 }