intrinsics: Remove deprecated modf & frexp overloads

These have been deprecated for multiple chrome releases.

Change-Id: I4cc05a74ff8f085e6d13f93aefb93077480e52f5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/66261
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

docs/origin-trial-changes.md

@@ -1,5 +1,11 @@
 # Tint changes during Origin Trial
 
+## Changes for M96
+
+### Breaking Changes
+
+* Deprecated `modf()` and `frexp()` builtin overloads that take a pointer second parameter have been removed.
+
 ## Changes for M95
 
 ### New Features
@@ -10,5 +16,5 @@
 ### Fixes
 
 * Hex floats: now correctly errors when the magnitude is non-zero, and the exponent would cause overflow. [tint:1150](https://crbug.com/tint/1150), [tint:1166](https://crbug.com/tint/1166)
-* Identifers beginning with an underscore are now correctly rejected.  [tint:1179](https://crbug.com/tint/1179)
+* Identifiers beginning with an underscore are now correctly rejected.  [tint:1179](https://crbug.com/tint/1179)
 * `abs()` fixed for unsigned integers on SPIR-V backend   [tint:1179](https://crbug.com/tint/1194)

src/intrinsic_table_test.cc

@@ -16,6 +16,7 @@
 
 #include "gmock/gmock.h"
 #include "src/program_builder.h"
+#include "src/sem/atomic_type.h"
 #include "src/sem/depth_multisampled_texture_type.h"
 #include "src/sem/depth_texture_type.h"
 #include "src/sem/external_texture_type.h"
@@ -211,22 +212,24 @@ TEST_F(IntrinsicTableTest, MismatchBool) {
 }
 
 TEST_F(IntrinsicTableTest, MatchPointer) {
-  auto* f32 = create<sem::F32>();
+  auto* i32 = create<sem::I32>();
-  auto* ptr = create<sem::Pointer>(f32, ast::StorageClass::kFunction,
+  auto* atomicI32 = create<sem::Atomic>(i32);
+  auto* ptr = create<sem::Pointer>(atomicI32, ast::StorageClass::kWorkgroup,
                                    ast::Access::kReadWrite);
-  auto* result = table->Lookup(IntrinsicType::kModf, {f32, ptr}, Source{});
+  auto* result = table->Lookup(IntrinsicType::kAtomicLoad, {ptr}, Source{});
   ASSERT_NE(result, nullptr) << Diagnostics().str();
   ASSERT_EQ(Diagnostics().str(), "");
-  EXPECT_THAT(result->Type(), IntrinsicType::kModf);
+  EXPECT_THAT(result->Type(), IntrinsicType::kAtomicLoad);
-  EXPECT_THAT(result->ReturnType(), f32);
+  EXPECT_THAT(result->ReturnType(), i32);
-  ASSERT_EQ(result->Parameters().size(), 2u);
+  ASSERT_EQ(result->Parameters().size(), 1u);
-  EXPECT_EQ(result->Parameters()[0]->Type(), f32);
+  EXPECT_EQ(result->Parameters()[0]->Type(), ptr);
-  EXPECT_EQ(result->Parameters()[1]->Type(), ptr);
 }
 
 TEST_F(IntrinsicTableTest, MismatchPointer) {
-  auto* f32 = create<sem::F32>();
+  auto* i32 = create<sem::I32>();
-  auto* result = table->Lookup(IntrinsicType::kModf, {f32, f32}, Source{});
+  auto* atomicI32 = create<sem::Atomic>(i32);
+  auto* result =
+      table->Lookup(IntrinsicType::kAtomicLoad, {atomicI32}, Source{});
   ASSERT_EQ(result, nullptr);
   ASSERT_THAT(Diagnostics().str(), HasSubstr("no matching call"));
 }

src/intrinsics.def

@@ -314,8 +314,6 @@ fn fma(f32, f32, f32) -> f32
 fn fma<N: num>(vec<N, f32>, vec<N, f32>, vec<N, f32>) -> vec<N, f32>
 fn fract(f32) -> f32
 fn fract<N: num>(vec<N, f32>) -> vec<N, f32>
-[[deprecated]] fn frexp<S: function_private_workgroup, A: access>(f32, ptr<S, i32, A>) -> f32
-[[deprecated]] fn frexp<N: num, S: function_private_workgroup, A: access>(vec<N, f32>, ptr<S, vec<N, i32>, A>) -> vec<N, f32>
 fn frexp(f32) -> _frexp_result
 fn frexp<N: num>(vec<N, f32>) -> _frexp_result_vec<N>
 [[stage("fragment")]] fn fwidth(f32) -> f32
@@ -350,8 +348,6 @@ fn min<N: num, T: fiu32>(vec<N, T>, vec<N, T>) -> vec<N, T>
 fn mix(f32, f32, f32) -> f32
 fn mix<N: num>(vec<N, f32>, vec<N, f32>, vec<N, f32>) -> vec<N, f32>
 fn mix<N: num>(vec<N, f32>, vec<N, f32>, f32) -> vec<N, f32>
-[[deprecated]] fn modf<S: function_private_workgroup, A: access>(f32, ptr<S, f32, A>) -> f32
-[[deprecated]] fn modf<N: num, S: function_private_workgroup, A: access>(vec<N, f32>, ptr<S, vec<N, f32>, A>) -> vec<N, f32>
 fn modf(f32) -> _modf_result
 fn modf<N: num>(vec<N, f32>) -> _modf_result_vec<N>
 fn normalize<N: num>(vec<N, f32>) -> vec<N, f32>

src/resolver/builtins_validation_test.cc

@@ -766,87 +766,131 @@ TEST_F(ResolverBuiltinsValidationTest, Determinant_Mat4x4) {
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Frexp_Scalar) {
-  auto* a = Var("a", ty.i32());
+  auto* builtin = Call("frexp", 1.0f);
-  auto* builtin = Call("frexp", 1.0f, AddressOf(Expr("a")));
+  WrapInFunction(builtin);
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->Is<sem::F32>());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  EXPECT_TRUE(members[0]->Type()->Is<sem::F32>());
+  EXPECT_TRUE(members[1]->Type()->Is<sem::I32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Frexp_Vec2) {
-  auto* a = Var("a", ty.vec2<int>());
+  auto* builtin = Call("frexp", vec2<f32>(1.0f, 1.0f));
-  auto* builtin = Call("frexp", vec2<f32>(1.0f, 1.0f), AddressOf(Expr("a")));
+  WrapInFunction(builtin);
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 2u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 2u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::I32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Frexp_Vec3) {
-  auto* a = Var("a", ty.vec3<int>());
+  auto* builtin = Call("frexp", vec3<f32>(1.0f, 1.0f, 1.0f));
-  auto* builtin =
+  WrapInFunction(builtin);
-      Call("frexp", vec3<f32>(1.0f, 1.0f, 1.0f), AddressOf(Expr("a")));
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::I32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Frexp_Vec4) {
-  auto* a = Var("a", ty.vec4<int>());
+  auto* builtin = Call("frexp", vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f));
-  auto* builtin =
+  WrapInFunction(builtin);
-      Call("frexp", vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f), AddressOf(Expr("a")));
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 4u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 4u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::I32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Modf_Scalar) {
-  auto* a = Var("a", ty.f32());
+  auto* builtin = Call("modf", 1.0f);
-  auto* builtin = Call("modf", 1.0f, AddressOf(Expr("a")));
+  WrapInFunction(builtin);
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->Is<sem::F32>());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  EXPECT_TRUE(members[0]->Type()->Is<sem::F32>());
+  EXPECT_TRUE(members[1]->Type()->Is<sem::F32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Modf_Vec2) {
-  auto* a = Var("a", ty.vec2<f32>());
+  auto* builtin = Call("modf", vec2<f32>(1.0f, 1.0f));
-  auto* builtin = Call("modf", vec2<f32>(1.0f, 1.0f), AddressOf(Expr("a")));
+  WrapInFunction(builtin);
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 2u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 2u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Modf_Vec3) {
-  auto* a = Var("a", ty.vec3<f32>());
+  auto* builtin = Call("modf", vec3<f32>(1.0f, 1.0f, 1.0f));
-  auto* builtin =
+  WrapInFunction(builtin);
-      Call("modf", vec3<f32>(1.0f, 1.0f, 1.0f), AddressOf(Expr("a")));
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Modf_Vec4) {
-  auto* a = Var("a", ty.vec4<f32>());
+  auto* builtin = Call("modf", vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f));
-  auto* builtin =
+  WrapInFunction(builtin);
-      Call("modf", vec4<f32>(1.0f, 1.0f, 1.0f, 1.0f), AddressOf(Expr("a")));
-  WrapInFunction(Decl(a), builtin);
 
   EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_TRUE(TypeOf(builtin)->is_float_vector());
+  auto* res_ty = TypeOf(builtin)->As<sem::Struct>();
-  EXPECT_TRUE(TypeOf(builtin->params()[1])->Is<sem::Pointer>());
+  ASSERT_TRUE(res_ty != nullptr);
+  auto& members = res_ty->Members();
+  ASSERT_EQ(members.size(), 2u);
+  ASSERT_TRUE(members[0]->Type()->Is<sem::Vector>());
+  ASSERT_TRUE(members[1]->Type()->Is<sem::Vector>());
+  EXPECT_EQ(members[0]->Type()->As<sem::Vector>()->Width(), 4u);
+  EXPECT_TRUE(members[0]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(members[1]->Type()->As<sem::Vector>()->Width(), 4u);
+  EXPECT_TRUE(members[1]->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
 }
 
 TEST_F(ResolverBuiltinsValidationTest, Cross_Float_Vec3) {

src/resolver/intrinsic_test.cc

@@ -828,29 +828,6 @@ TEST_F(ResolverIntrinsicDataTest, Normalize_Error_NoParams) {
 )");
 }
 
-TEST_F(ResolverIntrinsicDataTest, DEPRECATED_FrexpScalar) {
-  Global("exp", ty.i32(), ast::StorageClass::kWorkgroup);
-  auto* call = Call("frexp", 1.0f, AddressOf("exp"));
-  WrapInFunction(call);
-
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
-
-  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(call)->Is<sem::F32>());
-}
-
-TEST_F(ResolverIntrinsicDataTest, DEPRECATED_FrexpVector) {
-  Global("exp", ty.vec3<i32>(), ast::StorageClass::kWorkgroup);
-  auto* call = Call("frexp", vec3<f32>(1.0f, 2.0f, 3.0f), AddressOf("exp"));
-  WrapInFunction(call);
-
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
-
-  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(call)->Is<sem::Vector>());
-  EXPECT_TRUE(TypeOf(call)->As<sem::Vector>()->type()->Is<sem::F32>());
-}
-
 TEST_F(ResolverIntrinsicDataTest, FrexpScalar) {
   auto* call = Call("frexp", 1.0f);
   WrapInFunction(call);
@@ -924,9 +901,7 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_FirstParamInt) {
       r()->error(),
       R"(error: no matching call to frexp(i32, ptr<workgroup, i32, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  frexp(f32, ptr<S, i32, A>) -> f32  where: S is function, private or workgroup
-  frexp(vecN<f32>, ptr<S, vecN<i32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   frexp(f32) -> _frexp_result
   frexp(vecN<f32>) -> _frexp_result_vecN
 )");
@@ -943,10 +918,8 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_SecondParamFloatPtr) {
       r()->error(),
       R"(error: no matching call to frexp(f32, ptr<workgroup, f32, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  frexp(f32, ptr<S, i32, A>) -> f32  where: S is function, private or workgroup
   frexp(f32) -> _frexp_result
-  frexp(vecN<f32>, ptr<S, vecN<i32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   frexp(vecN<f32>) -> _frexp_result_vecN
 )");
 }
@@ -959,10 +932,8 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_SecondParamNotAPointer) {
 
   EXPECT_EQ(r()->error(), R"(error: no matching call to frexp(f32, i32)
 
-4 candidate functions:
+2 candidate functions:
-  frexp(f32, ptr<S, i32, A>) -> f32  where: S is function, private or workgroup
   frexp(f32) -> _frexp_result
-  frexp(vecN<f32>, ptr<S, vecN<i32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   frexp(vecN<f32>) -> _frexp_result_vecN
 )");
 }
@@ -978,37 +949,12 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_VectorSizesDontMatch) {
       r()->error(),
       R"(error: no matching call to frexp(vec2<f32>, ptr<workgroup, vec4<i32>, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  frexp(vecN<f32>, ptr<S, vecN<i32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   frexp(vecN<f32>) -> _frexp_result_vecN
-  frexp(f32, ptr<S, i32, A>) -> f32  where: S is function, private or workgroup
   frexp(f32) -> _frexp_result
 )");
 }
 
-TEST_F(ResolverIntrinsicDataTest, DEPRECATED_ModfScalar) {
-  Global("whole", ty.f32(), ast::StorageClass::kWorkgroup);
-  auto* call = Call("modf", 1.0f, AddressOf("whole"));
-  WrapInFunction(call);
-
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
-
-  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(call)->Is<sem::F32>());
-}
-
-TEST_F(ResolverIntrinsicDataTest, DEPRECATED_ModfVector) {
-  Global("whole", ty.vec3<f32>(), ast::StorageClass::kWorkgroup);
-  auto* call = Call("modf", vec3<f32>(1.0f, 2.0f, 3.0f), AddressOf("whole"));
-  WrapInFunction(call);
-
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
-
-  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(call)->Is<sem::Vector>());
-  EXPECT_TRUE(TypeOf(call)->As<sem::Vector>()->type()->Is<sem::F32>());
-}
-
 TEST_F(ResolverIntrinsicDataTest, ModfScalar) {
   auto* call = Call("modf", 1.0f);
   WrapInFunction(call);
@@ -1082,9 +1028,7 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_FirstParamInt) {
       r()->error(),
       R"(error: no matching call to modf(i32, ptr<workgroup, f32, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  modf(f32, ptr<S, f32, A>) -> f32  where: S is function, private or workgroup
-  modf(vecN<f32>, ptr<S, vecN<f32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   modf(f32) -> _modf_result
   modf(vecN<f32>) -> _modf_result_vecN
 )");
@@ -1101,10 +1045,8 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_SecondParamIntPtr) {
       r()->error(),
       R"(error: no matching call to modf(f32, ptr<workgroup, i32, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  modf(f32, ptr<S, f32, A>) -> f32  where: S is function, private or workgroup
   modf(f32) -> _modf_result
-  modf(vecN<f32>, ptr<S, vecN<f32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   modf(vecN<f32>) -> _modf_result_vecN
 )");
 }
@@ -1117,10 +1059,8 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_SecondParamNotAPointer) {
 
   EXPECT_EQ(r()->error(), R"(error: no matching call to modf(f32, f32)
 
-4 candidate functions:
+2 candidate functions:
-  modf(f32, ptr<S, f32, A>) -> f32  where: S is function, private or workgroup
   modf(f32) -> _modf_result
-  modf(vecN<f32>, ptr<S, vecN<f32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   modf(vecN<f32>) -> _modf_result_vecN
 )");
 }
@@ -1136,10 +1076,8 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_VectorSizesDontMatch) {
       r()->error(),
       R"(error: no matching call to modf(vec2<f32>, ptr<workgroup, vec4<f32>, read_write>)
 
-4 candidate functions:
+2 candidate functions:
-  modf(vecN<f32>, ptr<S, vecN<f32>, A>) -> vecN<f32>  where: S is function, private or workgroup
   modf(vecN<f32>) -> _modf_result_vecN
-  modf(f32, ptr<S, f32, A>) -> f32  where: S is function, private or workgroup
   modf(f32) -> _modf_result
 )");
 }

src/writer/glsl/generator_impl_intrinsic_test.cc

@@ -297,6 +297,7 @@ TEST_F(GlslGeneratorImplTest_Intrinsic, Select_Vector) {
   EXPECT_EQ(out.str(), "(bvec2(true, false) ? ivec2(3, 4) : ivec2(1, 2))");
 }
 
+#if 0
 TEST_F(GlslGeneratorImplTest_Intrinsic, Modf_Scalar) {
   auto* res = Var("res", ty.f32());
   auto* call = Call("modf", 1.0f, AddressOf(res));
@@ -319,7 +320,6 @@ TEST_F(GlslGeneratorImplTest_Intrinsic, Modf_Vector) {
   EXPECT_THAT(gen.result(), HasSubstr("modf(vec3(0.0f, 0.0f, 0.0f), res)"));
 }
 
-#if 0
 TEST_F(GlslGeneratorImplTest_Intrinsic, Frexp_Scalar_i32) {
   auto* exp = Var("exp", ty.i32());
   auto* call = Call("frexp", 1.0f, AddressOf(exp));

src/writer/hlsl/generator_impl.cc

@@ -1366,7 +1366,6 @@ bool GeneratorImpl::EmitSelectCall(std::ostream& out,
 bool GeneratorImpl::EmitModfCall(std::ostream& out,
                                  ast::CallExpression* expr,
                                  const sem::Intrinsic* intrinsic) {
-  if (expr->params().size() == 1) {
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
@@ -1400,23 +1399,9 @@ bool GeneratorImpl::EmitModfCall(std::ostream& out,
       });
 }
 
-  // DEPRECATED
-  out << "modf";
-  ScopedParen sp(out);
-  if (!EmitExpression(out, expr->params()[0])) {
-    return false;
-  }
-  out << ", ";
-  if (!EmitExpression(out, expr->params()[1])) {
-    return false;
-  }
-  return true;
-}
-
 bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
                                   ast::CallExpression* expr,
                                   const sem::Intrinsic* intrinsic) {
-  if (expr->params().size() == 1) {
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
@@ -1449,40 +1434,6 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
         return true;
       });
 }
-  // DEPRECATED
-  // Exponent is an integer in WGSL, but HLSL wants a float.
-  // We need to make the call with a temporary float, and then cast.
-  return CallIntrinsicHelper(
-      out, expr, intrinsic,
-      [&](TextBuffer* b, const std::vector<std::string>& params) {
-        auto* significand_ty = intrinsic->Parameters()[0]->Type();
-        auto significand = params[0];
-        auto* exponent_ty = intrinsic->Parameters()[1]->Type();
-        auto exponent = params[1];
-
-        std::string width;
-        if (auto* vec = significand_ty->As<sem::Vector>()) {
-          width = std::to_string(vec->Width());
-        }
-
-        // Exponent is an integer, which HLSL does not have an overload for.
-        // We need to cast from a float.
-        line(b) << "float" << width << " float_exp;";
-        line(b) << "float" << width << " significand = frexp(" << significand
-                << ", float_exp);";
-        {
-          auto l = line(b);
-          l << exponent << " = ";
-          if (!EmitType(l, exponent_ty->UnwrapPtr(), ast::StorageClass::kNone,
-                        ast::Access::kUndefined, "")) {
-            return false;
-          }
-          l << "(float_exp);";
-        }
-        line(b) << "return significand;";
-        return true;
-      });
-}
 
 bool GeneratorImpl::EmitIsNormalCall(std::ostream& out,
                                      ast::CallExpression* expr,

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -298,72 +298,104 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Select_Vector) {
 }
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Modf_Scalar) {
-  auto* res = Var("res", ty.f32());
+  auto* call = Call("modf", 1.0f);
-  auto* call = Call("modf", 1.0f, AddressOf(res));
+  WrapInFunction(call);
-  WrapInFunction(res, call);
 
   GeneratorImpl& gen = SanitizeAndBuild();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_THAT(gen.result(), HasSubstr("modf(1.0f, res)"));
+  EXPECT_EQ(gen.result(), R"(struct modf_result {
-}
+  float fract;
-
+  float whole;
-TEST_F(HlslGeneratorImplTest_Intrinsic, Modf_Vector) {
+};
-  auto* res = Var("res", ty.vec3<f32>());
+modf_result tint_modf(float param_0) {
-  auto* call = Call("modf", vec3<f32>(), AddressOf(res));
+  float whole;
-  WrapInFunction(res, call);
+  float fract = modf(param_0, whole);
-
+  modf_result result = {fract, whole};
-  GeneratorImpl& gen = SanitizeAndBuild();
+  return result;
-
-  ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_THAT(gen.result(), HasSubstr("modf(float3(0.0f, 0.0f, 0.0f), res)"));
-}
-
-TEST_F(HlslGeneratorImplTest_Intrinsic, Frexp_Scalar_i32) {
-  auto* exp = Var("exp", ty.i32());
-  auto* call = Call("frexp", 1.0f, AddressOf(exp));
-  WrapInFunction(exp, call);
-
-  GeneratorImpl& gen = SanitizeAndBuild();
-
-  ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_EQ(gen.result(),
-            R"(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;
 }
 
 [numthreads(1, 1, 1)]
 void test_function() {
-  int exp = 0;
+  tint_modf(1.0f);
-  tint_frexp(1.0f, exp);
+  return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Intrinsic, Modf_Vector) {
+  auto* call = Call("modf", vec3<f32>());
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(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;
+}
+
+[numthreads(1, 1, 1)]
+void test_function() {
+  tint_modf(float3(0.0f, 0.0f, 0.0f));
+  return;
+}
+)");
+}
+
+TEST_F(HlslGeneratorImplTest_Intrinsic, Frexp_Scalar_i32) {
+  auto* call = Call("frexp", 1.0f);
+  WrapInFunction(call);
+
+  GeneratorImpl& gen = SanitizeAndBuild();
+
+  ASSERT_TRUE(gen.Generate()) << gen.error();
+  EXPECT_EQ(gen.result(), R"(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 test_function() {
+  tint_frexp(1.0f);
   return;
 }
 )");
 }
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, Frexp_Vector_i32) {
-  auto* res = Var("res", ty.vec3<i32>());
+  auto* call = Call("frexp", vec3<f32>());
-  auto* call = Call("frexp", vec3<f32>(), AddressOf(res));
+  WrapInFunction(call);
-  WrapInFunction(res, call);
 
   GeneratorImpl& gen = SanitizeAndBuild();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_EQ(gen.result(),
+  EXPECT_EQ(gen.result(), R"(struct frexp_result_vec3 {
-            R"(float3 tint_frexp(float3 param_0, inout int3 param_1) {
+  float3 sig;
-  float3 float_exp;
+  int3 exp;
-  float3 significand = frexp(param_0, float_exp);
+};
-  param_1 = int3(float_exp);
+frexp_result_vec3 tint_frexp(float3 param_0) {
-  return significand;
+  float3 exp;
+  float3 sig = frexp(param_0, exp);
+  frexp_result_vec3 result = {sig, int3(exp)};
+  return result;
 }
 
 [numthreads(1, 1, 1)]
 void test_function() {
-  int3 res = int3(0, 0, 0);
+  tint_frexp(float3(0.0f, 0.0f, 0.0f));
-  tint_frexp(float3(0.0f, 0.0f, 0.0f), res);
   return;
 }
 )");

src/writer/msl/generator_impl.cc

@@ -1034,7 +1034,6 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
 bool GeneratorImpl::EmitModfCall(std::ostream& out,
                                  ast::CallExpression* expr,
                                  const sem::Intrinsic* intrinsic) {
-  if (expr->params().size() == 1) {
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
@@ -1059,31 +1058,10 @@ bool GeneratorImpl::EmitModfCall(std::ostream& out,
         return true;
       });
 }
-  // DEPRECATED
-  return CallIntrinsicHelper(
-      out, expr, intrinsic,
-      [&](TextBuffer* b, const std::vector<std::string>& params) {
-        auto* ty = intrinsic->Parameters()[0]->Type();
-        auto in = params[0];
-        auto out_whole = params[1];
-
-        std::string width;
-        if (auto* vec = ty->As<sem::Vector>()) {
-          width = std::to_string(vec->Width());
-        }
-
-        line(b) << "float" << width << " whole;";
-        line(b) << "float" << width << " fract = modf(" << in << ", whole);";
-        line(b) << "*" << out_whole << " = whole;";
-        line(b) << "return fract;";
-        return true;
-      });
-}
 
 bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
                                   ast::CallExpression* expr,
                                   const sem::Intrinsic* intrinsic) {
-  if (expr->params().size() == 1) {
   return CallIntrinsicHelper(
       out, expr, intrinsic,
       [&](TextBuffer* b, const std::vector<std::string>& params) {
@@ -1109,27 +1087,6 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
       });
 }
 
-  // DEPRECATED
-  return CallIntrinsicHelper(
-      out, expr, intrinsic,
-      [&](TextBuffer* b, const std::vector<std::string>& params) {
-        auto* ty = intrinsic->Parameters()[0]->Type();
-        auto in = params[0];
-        auto out_exp = params[1];
-
-        std::string width;
-        if (auto* vec = ty->As<sem::Vector>()) {
-          width = std::to_string(vec->Width());
-        }
-
-        line(b) << "int" << width << " exp;";
-        line(b) << "float" << width << " sig = frexp(" << in << ", exp);";
-        line(b) << "*" << out_exp << " = exp;";
-        line(b) << "return sig;";
-        return true;
-      });
-}
-
 std::string GeneratorImpl::generate_builtin_name(
     const sem::Intrinsic* intrinsic) {
   std::string out = "";

src/writer/spirv/builder.cc

@@ -165,8 +165,7 @@ uint32_t intrinsic_to_glsl_method(const sem::Intrinsic* intrinsic) {
     case IntrinsicType::kFract:
       return GLSLstd450Fract;
     case IntrinsicType::kFrexp:
-      return (intrinsic->Parameters().size() == 1) ? GLSLstd450FrexpStruct
+      return GLSLstd450FrexpStruct;
-                                                   : GLSLstd450Frexp;
     case IntrinsicType::kInverseSqrt:
       return GLSLstd450InverseSqrt;
     case IntrinsicType::kLdexp:
@@ -196,8 +195,7 @@ uint32_t intrinsic_to_glsl_method(const sem::Intrinsic* intrinsic) {
     case IntrinsicType::kMix:
       return GLSLstd450FMix;
     case IntrinsicType::kModf:
-      return (intrinsic->Parameters().size() == 1) ? GLSLstd450ModfStruct
+      return GLSLstd450ModfStruct;
-                                                   : GLSLstd450Modf;
     case IntrinsicType::kNormalize:
       return GLSLstd450Normalize;
     case IntrinsicType::kPack4x8snorm:

src/writer/spirv/builder_intrinsic_test.cc

@@ -1452,16 +1452,48 @@ INSTANTIATE_TEST_SUITE_P(IntrinsicBuilderTest,
                          testing::Values(IntrinsicData{"clamp", "UClamp"}));
 
 TEST_F(IntrinsicBuilderTest, Call_Modf) {
-  auto* out = Var("out", ty.vec2<f32>());
+  auto* expr = Call("modf", vec2<f32>(1.0f, 2.0f));
-  auto* expr = Call("modf", vec2<f32>(1.0f, 2.0f), AddressOf("out"));
+  Func("a_func", {}, ty.void_(), {Ignore(expr)},
-  Func("a_func", ast::VariableList{}, ty.void_(),
+       {Stage(ast::PipelineStage::kFragment)});
-       ast::StatementList{
+
-           Decl(out),
+  spirv::Builder& b = Build();
-           Ignore(expr),
+
-       },
+  ASSERT_TRUE(b.Build()) << b.error();
-       ast::DecorationList{
+  auto got = DumpBuilder(b);
-           Stage(ast::PipelineStage::kFragment),
+  auto* expect = R"(OpCapability Shader
-       });
+%10 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint Fragment %3 "a_func"
+OpExecutionMode %3 OriginUpperLeft
+OpName %3 "a_func"
+OpName %7 "_modf_result_vec2"
+OpMemberName %7 0 "fract"
+OpMemberName %7 1 "whole"
+OpMemberDecorate %7 0 Offset 0
+OpMemberDecorate %7 1 Offset 8
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%9 = OpTypeFloat 32
+%8 = OpTypeVector %9 2
+%7 = OpTypeStruct %8 %8
+%11 = OpConstant %9 1
+%12 = OpConstant %9 2
+%13 = OpConstantComposite %8 %11 %12
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+%6 = OpExtInst %7 %10 ModfStruct %13
+OpReturn
+OpFunctionEnd
+)";
+  EXPECT_EQ(expect, got);
+
+  Validate(b);
+}
+
+TEST_F(IntrinsicBuilderTest, Call_Frexp) {
+  auto* expr = Call("frexp", vec2<f32>(1.0f, 2.0f));
+  Func("a_func", {}, ty.void_(), {Ignore(expr)},
+       {Stage(ast::PipelineStage::kFragment)});
 
   spirv::Builder& b = Build();
 
@@ -1473,66 +1505,24 @@ OpMemoryModel Logical GLSL450
 OpEntryPoint Fragment %3 "a_func"
 OpExecutionMode %3 OriginUpperLeft
 OpName %3 "a_func"
-OpName %5 "out"
+OpName %7 "_frexp_result_vec2"
+OpMemberName %7 0 "sig"
+OpMemberName %7 1 "exp"
+OpMemberDecorate %7 0 Offset 0
+OpMemberDecorate %7 1 Offset 8
 %2 = OpTypeVoid
 %1 = OpTypeFunction %2
-%8 = OpTypeFloat 32
+%9 = OpTypeFloat 32
-%7 = OpTypeVector %8 2
+%8 = OpTypeVector %9 2
-%6 = OpTypePointer Function %7
+%11 = OpTypeInt 32 1
-%9 = OpConstantNull %7
+%10 = OpTypeVector %11 2
-%13 = OpConstant %8 1
+%7 = OpTypeStruct %8 %10
-%14 = OpConstant %8 2
+%13 = OpConstant %9 1
-%15 = OpConstantComposite %7 %13 %14
+%14 = OpConstant %9 2
+%15 = OpConstantComposite %8 %13 %14
 %3 = OpFunction %2 None %1
 %4 = OpLabel
-%5 = OpVariable %6 Function %9
+%6 = OpExtInst %7 %12 FrexpStruct %15
-%11 = OpExtInst %7 %12 Modf %15 %5
-OpReturn
-OpFunctionEnd
-)";
-  EXPECT_EQ(expect, got);
-
-  Validate(b);
-}
-
-TEST_F(IntrinsicBuilderTest, Call_Frexp) {
-  auto* out = Var("out", ty.vec2<i32>());
-  auto* expr = Call("frexp", vec2<f32>(1.0f, 2.0f), AddressOf("out"));
-  Func("a_func", ast::VariableList{}, ty.void_(),
-       ast::StatementList{
-           Decl(out),
-           Ignore(expr),
-       },
-       ast::DecorationList{
-           Stage(ast::PipelineStage::kFragment),
-       });
-
-  spirv::Builder& b = Build();
-
-  ASSERT_TRUE(b.Build()) << b.error();
-  auto got = DumpBuilder(b);
-  auto* expect = R"(OpCapability Shader
-%14 = OpExtInstImport "GLSL.std.450"
-OpMemoryModel Logical GLSL450
-OpEntryPoint Fragment %3 "a_func"
-OpExecutionMode %3 OriginUpperLeft
-OpName %3 "a_func"
-OpName %5 "out"
-%2 = OpTypeVoid
-%1 = OpTypeFunction %2
-%8 = OpTypeInt 32 1
-%7 = OpTypeVector %8 2
-%6 = OpTypePointer Function %7
-%9 = OpConstantNull %7
-%13 = OpTypeFloat 32
-%12 = OpTypeVector %13 2
-%15 = OpConstant %13 1
-%16 = OpConstant %13 2
-%17 = OpConstantComposite %12 %15 %16
-%3 = OpFunction %2 None %1
-%4 = OpLabel
-%5 = OpVariable %6 Function %9
-%11 = OpExtInst %12 %14 Frexp %17 %5
 OpReturn
 OpFunctionEnd
 )";

test/intrinsics/gen/frexp/013caa.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<4, i32>, write>) -> vec<4, f32>
-fn frexp_013caa() {
-  var arg_1: vec4<i32>;
-  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_013caa();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_013caa();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_013caa();
-}

test/intrinsics/gen/frexp/013caa.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int4(float_exp);
-  return significand;
-}
-
-void frexp_013caa() {
-  int4 arg_1 = int4(0, 0, 0, 0);
-  float4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_013caa();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_013caa();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_013caa();
-  return;
-}

test/intrinsics/gen/frexp/013caa.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float4(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_013caa();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_013caa();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_013caa();
-  return;
-}
-

test/intrinsics/gen/frexp/013caa.wgsl.expected.spvasm

@@ -1,75 +0,0 @@
-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
-; Bound: 37
-; Schema: 0
-               OpCapability Shader
-         %19 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_013caa "frexp_013caa"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v4int = OpTypeVector %int 4
-%_ptr_Function_v4int = OpTypePointer Function %v4int
-         %17 = OpConstantNull %v4int
-%_ptr_Function_v4float = OpTypePointer Function %v4float
-         %23 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_013caa = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v4int Function %17
-        %res = OpVariable %_ptr_Function_v4float Function %5
-         %18 = OpExtInst %v4float %19 Frexp %5 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_013caa
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %frexp_013caa
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %35 = OpLabel
-         %36 = OpFunctionCall %void %frexp_013caa
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/013caa.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_013caa();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_013caa();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_013caa();
-}

test/intrinsics/gen/frexp/0da285.wgsl

@@ -1,35 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<workgroup> arg_1: i32;
-
-// fn frexp(f32, ptr<workgroup, i32, read_write>) -> f32
-fn frexp_0da285() {
-  var res: f32 = frexp(1.0, &arg_1);
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_0da285();
-}

test/intrinsics/gen/frexp/0da285.wgsl.expected.hlsl

@@ -1,34 +0,0 @@
-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);
-  param_1 = int(float_exp);
-  return significand;
-}
-
-groupshared int arg_1;
-
-void frexp_0da285() {
-  float res = tint_frexp(1.0f, arg_1);
-}
-
-struct tint_symbol_1 {
-  uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
-  {
-    arg_1 = 0;
-  }
-  GroupMemoryBarrierWithGroupSync();
-  frexp_0da285();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner(tint_symbol.local_invocation_index);
-  return;
-}

test/intrinsics/gen/frexp/0da285.wgsl.expected.msl

@@ -1,33 +0,0 @@
-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) {
-  float res = tint_frexp(1.0f, tint_symbol);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup int* const tint_symbol_1) {
-  {
-    *(tint_symbol_1) = int();
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  frexp_0da285(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/intrinsics/gen/frexp/0da285.wgsl.expected.spvasm

@@ -1,59 +0,0 @@
-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
-; Bound: 32
-; Schema: 0
-               OpCapability Shader
-         %13 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %local_invocation_index_1 "local_invocation_index_1"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_0da285 "frexp_0da285"
-               OpName %res "res"
-               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
-               OpName %compute_main "compute_main"
-               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
-       %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
-        %int = OpTypeInt 32 1
-%_ptr_Workgroup_int = OpTypePointer Workgroup %int
-      %arg_1 = OpVariable %_ptr_Workgroup_int Workgroup
-       %void = OpTypeVoid
-          %7 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %float_1 = OpConstant %float 1
-%_ptr_Function_float = OpTypePointer Function %float
-         %18 = OpConstantNull %float
-         %19 = OpTypeFunction %void %uint
-         %23 = OpConstantNull %int
-     %uint_2 = OpConstant %uint 2
-   %uint_264 = OpConstant %uint 264
-%frexp_0da285 = OpFunction %void None %7
-         %10 = OpLabel
-        %res = OpVariable %_ptr_Function_float Function %18
-         %11 = OpExtInst %float %13 Frexp %float_1 %arg_1
-               OpStore %res %11
-               OpReturn
-               OpFunctionEnd
-%compute_main_inner = OpFunction %void None %19
-%local_invocation_index = OpFunctionParameter %uint
-         %22 = OpLabel
-               OpStore %arg_1 %23
-               OpControlBarrier %uint_2 %uint_2 %uint_264
-         %27 = OpFunctionCall %void %frexp_0da285
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %7
-         %29 = OpLabel
-         %31 = OpLoad %uint %local_invocation_index_1
-         %30 = OpFunctionCall %void %compute_main_inner %31
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/0da285.wgsl.expected.wgsl

@@ -1,14 +0,0 @@
-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() {
-  var res : f32 = frexp(1.0, &(arg_1));
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_0da285();
-}

test/intrinsics/gen/frexp/15edf3.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<2, i32>, read_write>) -> vec<2, f32>
-fn frexp_15edf3() {
-  var arg_1: vec2<i32>;
-  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_15edf3();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_15edf3();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_15edf3();
-}

test/intrinsics/gen/frexp/15edf3.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int2(float_exp);
-  return significand;
-}
-
-void frexp_15edf3() {
-  int2 arg_1 = int2(0, 0);
-  float2 res = tint_frexp(float2(0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_15edf3();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_15edf3();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_15edf3();
-  return;
-}

test/intrinsics/gen/frexp/15edf3.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float2(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_15edf3();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_15edf3();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_15edf3();
-  return;
-}
-

test/intrinsics/gen/frexp/15edf3.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_15edf3 "frexp_15edf3"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v2int = OpTypeVector %int 2
-%_ptr_Function_v2int = OpTypePointer Function %v2int
-         %17 = OpConstantNull %v2int
-    %v2float = OpTypeVector %float 2
-         %21 = OpConstantNull %v2float
-%_ptr_Function_v2float = OpTypePointer Function %v2float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_15edf3 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v2int Function %17
-        %res = OpVariable %_ptr_Function_v2float Function %21
-         %18 = OpExtInst %v2float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_15edf3
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_15edf3
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_15edf3
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/15edf3.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_15edf3();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_15edf3();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_15edf3();
-}

test/intrinsics/gen/frexp/19ab15.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<4, i32>, read_write>) -> vec<4, f32>
-fn frexp_19ab15() {
-  var arg_1: vec4<i32>;
-  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_19ab15();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_19ab15();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_19ab15();
-}

test/intrinsics/gen/frexp/19ab15.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int4(float_exp);
-  return significand;
-}
-
-void frexp_19ab15() {
-  int4 arg_1 = int4(0, 0, 0, 0);
-  float4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_19ab15();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_19ab15();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_19ab15();
-  return;
-}

test/intrinsics/gen/frexp/19ab15.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float4(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_19ab15();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_19ab15();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_19ab15();
-  return;
-}
-

test/intrinsics/gen/frexp/19ab15.wgsl.expected.spvasm

@@ -1,75 +0,0 @@
-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
-; Bound: 37
-; Schema: 0
-               OpCapability Shader
-         %19 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_19ab15 "frexp_19ab15"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v4int = OpTypeVector %int 4
-%_ptr_Function_v4int = OpTypePointer Function %v4int
-         %17 = OpConstantNull %v4int
-%_ptr_Function_v4float = OpTypePointer Function %v4float
-         %23 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_19ab15 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v4int Function %17
-        %res = OpVariable %_ptr_Function_v4float Function %5
-         %18 = OpExtInst %v4float %19 Frexp %5 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_19ab15
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %frexp_19ab15
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %35 = OpLabel
-         %36 = OpFunctionCall %void %frexp_19ab15
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/19ab15.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_19ab15();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_19ab15();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_19ab15();
-}

test/intrinsics/gen/frexp/2052e9.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<4, i32>, read>) -> vec<4, f32>
-fn frexp_2052e9() {
-  var arg_1: vec4<i32>;
-  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_2052e9();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_2052e9();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_2052e9();
-}

test/intrinsics/gen/frexp/2052e9.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int4(float_exp);
-  return significand;
-}
-
-void frexp_2052e9() {
-  int4 arg_1 = int4(0, 0, 0, 0);
-  float4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_2052e9();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_2052e9();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_2052e9();
-  return;
-}

test/intrinsics/gen/frexp/2052e9.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float4(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_2052e9();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_2052e9();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_2052e9();
-  return;
-}
-

test/intrinsics/gen/frexp/2052e9.wgsl.expected.spvasm

@@ -1,75 +0,0 @@
-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
-; Bound: 37
-; Schema: 0
-               OpCapability Shader
-         %19 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_2052e9 "frexp_2052e9"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v4int = OpTypeVector %int 4
-%_ptr_Function_v4int = OpTypePointer Function %v4int
-         %17 = OpConstantNull %v4int
-%_ptr_Function_v4float = OpTypePointer Function %v4float
-         %23 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_2052e9 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v4int Function %17
-        %res = OpVariable %_ptr_Function_v4float Function %5
-         %18 = OpExtInst %v4float %19 Frexp %5 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_2052e9
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %frexp_2052e9
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %35 = OpLabel
-         %36 = OpFunctionCall %void %frexp_2052e9
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/2052e9.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_2052e9();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_2052e9();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_2052e9();
-}

test/intrinsics/gen/frexp/40fc9b.wgsl

@@ -1,35 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<workgroup> arg_1: vec3<i32>;
-
-// fn frexp(vec<3, f32>, ptr<workgroup, vec<3, i32>, read_write>) -> vec<3, f32>
-fn frexp_40fc9b() {
-  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_40fc9b();
-}

test/intrinsics/gen/frexp/40fc9b.wgsl.expected.hlsl

@@ -1,34 +0,0 @@
-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);
-  param_1 = int3(float_exp);
-  return significand;
-}
-
-groupshared int3 arg_1;
-
-void frexp_40fc9b() {
-  float3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol_1 {
-  uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
-  {
-    arg_1 = int3(0, 0, 0);
-  }
-  GroupMemoryBarrierWithGroupSync();
-  frexp_40fc9b();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner(tint_symbol.local_invocation_index);
-  return;
-}

test/intrinsics/gen/frexp/40fc9b.wgsl.expected.msl

@@ -1,33 +0,0 @@
-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) {
-  float3 res = tint_frexp(float3(), tint_symbol);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup int3* const tint_symbol_1) {
-  {
-    *(tint_symbol_1) = int3();
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  frexp_40fc9b(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int3 tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/intrinsics/gen/frexp/40fc9b.wgsl.expected.spvasm

@@ -1,60 +0,0 @@
-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
-; Bound: 33
-; Schema: 0
-               OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %local_invocation_index_1 "local_invocation_index_1"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_40fc9b "frexp_40fc9b"
-               OpName %res "res"
-               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
-               OpName %compute_main "compute_main"
-               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
-       %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
-        %int = OpTypeInt 32 1
-      %v3int = OpTypeVector %int 3
-%_ptr_Workgroup_v3int = OpTypePointer Workgroup %v3int
-      %arg_1 = OpVariable %_ptr_Workgroup_v3int Workgroup
-       %void = OpTypeVoid
-          %8 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v3float = OpTypeVector %float 3
-         %16 = OpConstantNull %v3float
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %20 = OpTypeFunction %void %uint
-         %24 = OpConstantNull %v3int
-     %uint_2 = OpConstant %uint 2
-   %uint_264 = OpConstant %uint 264
-%frexp_40fc9b = OpFunction %void None %8
-         %11 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %16
-         %12 = OpExtInst %v3float %15 Frexp %16 %arg_1
-               OpStore %res %12
-               OpReturn
-               OpFunctionEnd
-%compute_main_inner = OpFunction %void None %20
-%local_invocation_index = OpFunctionParameter %uint
-         %23 = OpLabel
-               OpStore %arg_1 %24
-               OpControlBarrier %uint_2 %uint_2 %uint_264
-         %28 = OpFunctionCall %void %frexp_40fc9b
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %8
-         %30 = OpLabel
-         %32 = OpLoad %uint %local_invocation_index_1
-         %31 = OpFunctionCall %void %compute_main_inner %32
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/40fc9b.wgsl.expected.wgsl

@@ -1,14 +0,0 @@
-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() {
-  var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_40fc9b();
-}

test/intrinsics/gen/frexp/41e931.wgsl

@@ -1,46 +0,0 @@
-// 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, ptr<function, i32, read_write>) -> f32
-fn frexp_41e931() {
-  var arg_1: i32;
-  var res: f32 = frexp(1.0, &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_41e931();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_41e931();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_41e931();
-}

test/intrinsics/gen/frexp/41e931.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int(float_exp);
-  return significand;
-}
-
-void frexp_41e931() {
-  int arg_1 = 0;
-  float res = tint_frexp(1.0f, arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_41e931();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_41e931();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_41e931();
-  return;
-}

test/intrinsics/gen/frexp/41e931.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(1.0f, &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_41e931();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_41e931();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_41e931();
-  return;
-}
-

test/intrinsics/gen/frexp/41e931.wgsl.expected.spvasm

@@ -1,74 +0,0 @@
-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
-; Bound: 36
-; Schema: 0
-               OpCapability Shader
-         %18 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_41e931 "frexp_41e931"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
-         %16 = OpConstantNull %int
-    %float_1 = OpConstant %float 1
-%_ptr_Function_float = OpTypePointer Function %float
-         %23 = OpTypeFunction %v4float
-%frexp_41e931 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_int Function %16
-        %res = OpVariable %_ptr_Function_float Function %8
-         %17 = OpExtInst %float %18 Frexp %float_1 %arg_1
-               OpStore %res %17
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_41e931
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %frexp_41e931
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_41e931
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/41e931.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_41e931();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_41e931();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_41e931();
-}

test/intrinsics/gen/frexp/481e59.wgsl

@@ -1,46 +0,0 @@
-// 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, ptr<function, i32, read>) -> f32
-fn frexp_481e59() {
-  var arg_1: i32;
-  var res: f32 = frexp(1.0, &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_481e59();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_481e59();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_481e59();
-}

test/intrinsics/gen/frexp/481e59.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int(float_exp);
-  return significand;
-}
-
-void frexp_481e59() {
-  int arg_1 = 0;
-  float res = tint_frexp(1.0f, arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_481e59();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_481e59();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_481e59();
-  return;
-}

test/intrinsics/gen/frexp/481e59.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(1.0f, &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_481e59();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_481e59();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_481e59();
-  return;
-}
-

test/intrinsics/gen/frexp/481e59.wgsl.expected.spvasm

@@ -1,74 +0,0 @@
-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
-; Bound: 36
-; Schema: 0
-               OpCapability Shader
-         %18 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_481e59 "frexp_481e59"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
-         %16 = OpConstantNull %int
-    %float_1 = OpConstant %float 1
-%_ptr_Function_float = OpTypePointer Function %float
-         %23 = OpTypeFunction %v4float
-%frexp_481e59 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_int Function %16
-        %res = OpVariable %_ptr_Function_float Function %8
-         %17 = OpExtInst %float %18 Frexp %float_1 %arg_1
-               OpStore %res %17
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_481e59
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %frexp_481e59
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_481e59
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/481e59.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_481e59();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_481e59();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_481e59();
-}

test/intrinsics/gen/frexp/5a141e.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<3, i32>, read>) -> vec<3, f32>
-fn frexp_5a141e() {
-  var arg_1: vec3<i32>;
-  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_5a141e();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_5a141e();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_5a141e();
-}

test/intrinsics/gen/frexp/5a141e.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int3(float_exp);
-  return significand;
-}
-
-void frexp_5a141e() {
-  int3 arg_1 = int3(0, 0, 0);
-  float3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_5a141e();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_5a141e();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_5a141e();
-  return;
-}

test/intrinsics/gen/frexp/5a141e.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float3(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_5a141e();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_5a141e();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_5a141e();
-  return;
-}
-

test/intrinsics/gen/frexp/5a141e.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_5a141e "frexp_5a141e"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v3int = OpTypeVector %int 3
-%_ptr_Function_v3int = OpTypePointer Function %v3int
-         %17 = OpConstantNull %v3int
-    %v3float = OpTypeVector %float 3
-         %21 = OpConstantNull %v3float
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_5a141e = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v3int Function %17
-        %res = OpVariable %_ptr_Function_v3float Function %21
-         %18 = OpExtInst %v3float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_5a141e
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_5a141e
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_5a141e
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/5a141e.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_5a141e();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_5a141e();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_5a141e();
-}

test/intrinsics/gen/frexp/6d0058.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<3, i32>, read_write>) -> vec<3, f32>
-fn frexp_6d0058() {
-  var arg_1: vec3<i32>;
-  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_6d0058();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_6d0058();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_6d0058();
-}

test/intrinsics/gen/frexp/6d0058.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int3(float_exp);
-  return significand;
-}
-
-void frexp_6d0058() {
-  int3 arg_1 = int3(0, 0, 0);
-  float3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_6d0058();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_6d0058();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_6d0058();
-  return;
-}

test/intrinsics/gen/frexp/6d0058.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float3(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_6d0058();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_6d0058();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_6d0058();
-  return;
-}
-

test/intrinsics/gen/frexp/6d0058.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_6d0058 "frexp_6d0058"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v3int = OpTypeVector %int 3
-%_ptr_Function_v3int = OpTypePointer Function %v3int
-         %17 = OpConstantNull %v3int
-    %v3float = OpTypeVector %float 3
-         %21 = OpConstantNull %v3float
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_6d0058 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v3int Function %17
-        %res = OpVariable %_ptr_Function_v3float Function %21
-         %18 = OpExtInst %v3float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_6d0058
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_6d0058
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_6d0058
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/6d0058.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_6d0058();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_6d0058();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_6d0058();
-}

test/intrinsics/gen/frexp/6efa09.wgsl

@@ -1,46 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<private> arg_1: vec3<i32>;
-
-// fn frexp(vec<3, f32>, ptr<private, vec<3, i32>, read_write>) -> vec<3, f32>
-fn frexp_6efa09() {
-  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_6efa09();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_6efa09();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_6efa09();
-}

test/intrinsics/gen/frexp/6efa09.wgsl.expected.hlsl

@@ -1,43 +0,0 @@
-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);
-  param_1 = int3(float_exp);
-  return significand;
-}
-
-static int3 arg_1 = int3(0, 0, 0);
-
-void frexp_6efa09() {
-  float3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_6efa09();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_6efa09();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_6efa09();
-  return;
-}

test/intrinsics/gen/frexp/6efa09.wgsl.expected.msl

@@ -1,48 +0,0 @@
-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_1) {
-  float3 res = tint_frexp(float3(), tint_symbol_1);
-}
-
-float4 vertex_main_inner(thread int3* const tint_symbol_2) {
-  frexp_6efa09(tint_symbol_2);
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  thread int3 tint_symbol_3 = 0;
-  float4 const inner_result = vertex_main_inner(&(tint_symbol_3));
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  thread int3 tint_symbol_4 = 0;
-  frexp_6efa09(&(tint_symbol_4));
-  return;
-}
-
-kernel void compute_main() {
-  thread int3 tint_symbol_5 = 0;
-  frexp_6efa09(&(tint_symbol_5));
-  return;
-}
-

test/intrinsics/gen/frexp/6efa09.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_6efa09 "frexp_6efa09"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-        %int = OpTypeInt 32 1
-      %v3int = OpTypeVector %int 3
-%_ptr_Private_v3int = OpTypePointer Private %v3int
-         %13 = OpConstantNull %v3int
-      %arg_1 = OpVariable %_ptr_Private_v3int Private %13
-       %void = OpTypeVoid
-         %14 = OpTypeFunction %void
-    %v3float = OpTypeVector %float 3
-         %21 = OpConstantNull %v3float
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_6efa09 = OpFunction %void None %14
-         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v3float Function %21
-         %18 = OpExtInst %v3float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_6efa09
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %14
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %14
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_6efa09
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %14
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_6efa09
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/6efa09.wgsl.expected.wgsl

@@ -1,25 +0,0 @@
-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() {
-  var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_6efa09();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_6efa09();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_6efa09();
-}

test/intrinsics/gen/frexp/a2a617.wgsl

@@ -1,46 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<private> arg_1: i32;
-
-// fn frexp(f32, ptr<private, i32, read_write>) -> f32
-fn frexp_a2a617() {
-  var res: f32 = frexp(1.0, &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_a2a617();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_a2a617();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a2a617();
-}

test/intrinsics/gen/frexp/a2a617.wgsl.expected.hlsl

@@ -1,43 +0,0 @@
-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);
-  param_1 = int(float_exp);
-  return significand;
-}
-
-static int arg_1 = 0;
-
-void frexp_a2a617() {
-  float res = tint_frexp(1.0f, arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_a2a617();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_a2a617();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_a2a617();
-  return;
-}

test/intrinsics/gen/frexp/a2a617.wgsl.expected.msl

@@ -1,48 +0,0 @@
-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_1) {
-  float res = tint_frexp(1.0f, tint_symbol_1);
-}
-
-float4 vertex_main_inner(thread int* const tint_symbol_2) {
-  frexp_a2a617(tint_symbol_2);
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  thread int tint_symbol_3 = 0;
-  float4 const inner_result = vertex_main_inner(&(tint_symbol_3));
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  thread int tint_symbol_4 = 0;
-  frexp_a2a617(&(tint_symbol_4));
-  return;
-}
-
-kernel void compute_main() {
-  thread int tint_symbol_5 = 0;
-  frexp_a2a617(&(tint_symbol_5));
-  return;
-}
-

test/intrinsics/gen/frexp/a2a617.wgsl.expected.spvasm

@@ -1,74 +0,0 @@
-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
-; Bound: 36
-; Schema: 0
-               OpCapability Shader
-         %18 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_a2a617 "frexp_a2a617"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-        %int = OpTypeInt 32 1
-%_ptr_Private_int = OpTypePointer Private %int
-         %12 = OpConstantNull %int
-      %arg_1 = OpVariable %_ptr_Private_int Private %12
-       %void = OpTypeVoid
-         %13 = OpTypeFunction %void
-    %float_1 = OpConstant %float 1
-%_ptr_Function_float = OpTypePointer Function %float
-         %23 = OpTypeFunction %v4float
-%frexp_a2a617 = OpFunction %void None %13
-         %16 = OpLabel
-        %res = OpVariable %_ptr_Function_float Function %8
-         %17 = OpExtInst %float %18 Frexp %float_1 %arg_1
-               OpStore %res %17
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_a2a617
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %13
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %13
-         %31 = OpLabel
-         %32 = OpFunctionCall %void %frexp_a2a617
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %13
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_a2a617
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/a2a617.wgsl.expected.wgsl

@@ -1,25 +0,0 @@
-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() {
-  var res : f32 = frexp(1.0, &(arg_1));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_a2a617();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_a2a617();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a2a617();
-}

test/intrinsics/gen/frexp/a3f940.wgsl

@@ -1,35 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<workgroup> arg_1: vec2<i32>;
-
-// fn frexp(vec<2, f32>, ptr<workgroup, vec<2, i32>, read_write>) -> vec<2, f32>
-fn frexp_a3f940() {
-  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a3f940();
-}

test/intrinsics/gen/frexp/a3f940.wgsl.expected.hlsl

@@ -1,34 +0,0 @@
-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);
-  param_1 = int2(float_exp);
-  return significand;
-}
-
-groupshared int2 arg_1;
-
-void frexp_a3f940() {
-  float2 res = tint_frexp(float2(0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol_1 {
-  uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
-  {
-    arg_1 = int2(0, 0);
-  }
-  GroupMemoryBarrierWithGroupSync();
-  frexp_a3f940();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner(tint_symbol.local_invocation_index);
-  return;
-}

test/intrinsics/gen/frexp/a3f940.wgsl.expected.msl

@@ -1,33 +0,0 @@
-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) {
-  float2 res = tint_frexp(float2(), tint_symbol);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup int2* const tint_symbol_1) {
-  {
-    *(tint_symbol_1) = int2();
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  frexp_a3f940(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int2 tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/intrinsics/gen/frexp/a3f940.wgsl.expected.spvasm

@@ -1,60 +0,0 @@
-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
-; Bound: 33
-; Schema: 0
-               OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %local_invocation_index_1 "local_invocation_index_1"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_a3f940 "frexp_a3f940"
-               OpName %res "res"
-               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
-               OpName %compute_main "compute_main"
-               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
-       %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
-        %int = OpTypeInt 32 1
-      %v2int = OpTypeVector %int 2
-%_ptr_Workgroup_v2int = OpTypePointer Workgroup %v2int
-      %arg_1 = OpVariable %_ptr_Workgroup_v2int Workgroup
-       %void = OpTypeVoid
-          %8 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v2float = OpTypeVector %float 2
-         %16 = OpConstantNull %v2float
-%_ptr_Function_v2float = OpTypePointer Function %v2float
-         %20 = OpTypeFunction %void %uint
-         %24 = OpConstantNull %v2int
-     %uint_2 = OpConstant %uint 2
-   %uint_264 = OpConstant %uint 264
-%frexp_a3f940 = OpFunction %void None %8
-         %11 = OpLabel
-        %res = OpVariable %_ptr_Function_v2float Function %16
-         %12 = OpExtInst %v2float %15 Frexp %16 %arg_1
-               OpStore %res %12
-               OpReturn
-               OpFunctionEnd
-%compute_main_inner = OpFunction %void None %20
-%local_invocation_index = OpFunctionParameter %uint
-         %23 = OpLabel
-               OpStore %arg_1 %24
-               OpControlBarrier %uint_2 %uint_2 %uint_264
-         %28 = OpFunctionCall %void %frexp_a3f940
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %8
-         %30 = OpLabel
-         %32 = OpLoad %uint %local_invocation_index_1
-         %31 = OpFunctionCall %void %compute_main_inner %32
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/a3f940.wgsl.expected.wgsl

@@ -1,14 +0,0 @@
-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() {
-  var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a3f940();
-}

test/intrinsics/gen/frexp/a951b5.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<2, i32>, write>) -> vec<2, f32>
-fn frexp_a951b5() {
-  var arg_1: vec2<i32>;
-  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_a951b5();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_a951b5();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a951b5();
-}

test/intrinsics/gen/frexp/a951b5.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int2(float_exp);
-  return significand;
-}
-
-void frexp_a951b5() {
-  int2 arg_1 = int2(0, 0);
-  float2 res = tint_frexp(float2(0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_a951b5();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_a951b5();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_a951b5();
-  return;
-}

test/intrinsics/gen/frexp/a951b5.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float2(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_a951b5();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_a951b5();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_a951b5();
-  return;
-}
-

test/intrinsics/gen/frexp/a951b5.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_a951b5 "frexp_a951b5"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v2int = OpTypeVector %int 2
-%_ptr_Function_v2int = OpTypePointer Function %v2int
-         %17 = OpConstantNull %v2int
-    %v2float = OpTypeVector %float 2
-         %21 = OpConstantNull %v2float
-%_ptr_Function_v2float = OpTypePointer Function %v2float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_a951b5 = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v2int Function %17
-        %res = OpVariable %_ptr_Function_v2float Function %21
-         %18 = OpExtInst %v2float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_a951b5
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_a951b5
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_a951b5
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/a951b5.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_a951b5();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_a951b5();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_a951b5();
-}

test/intrinsics/gen/frexp/b45525.wgsl

@@ -1,46 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<private> arg_1: vec4<i32>;
-
-// fn frexp(vec<4, f32>, ptr<private, vec<4, i32>, read_write>) -> vec<4, f32>
-fn frexp_b45525() {
-  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_b45525();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_b45525();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b45525();
-}

test/intrinsics/gen/frexp/b45525.wgsl.expected.hlsl

@@ -1,43 +0,0 @@
-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);
-  param_1 = int4(float_exp);
-  return significand;
-}
-
-static int4 arg_1 = int4(0, 0, 0, 0);
-
-void frexp_b45525() {
-  float4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_b45525();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_b45525();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_b45525();
-  return;
-}

test/intrinsics/gen/frexp/b45525.wgsl.expected.msl

@@ -1,48 +0,0 @@
-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_1) {
-  float4 res = tint_frexp(float4(), tint_symbol_1);
-}
-
-float4 vertex_main_inner(thread int4* const tint_symbol_2) {
-  frexp_b45525(tint_symbol_2);
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  thread int4 tint_symbol_3 = 0;
-  float4 const inner_result = vertex_main_inner(&(tint_symbol_3));
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  thread int4 tint_symbol_4 = 0;
-  frexp_b45525(&(tint_symbol_4));
-  return;
-}
-
-kernel void compute_main() {
-  thread int4 tint_symbol_5 = 0;
-  frexp_b45525(&(tint_symbol_5));
-  return;
-}
-

test/intrinsics/gen/frexp/b45525.wgsl.expected.spvasm

@@ -1,75 +0,0 @@
-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
-; Bound: 37
-; Schema: 0
-               OpCapability Shader
-         %19 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_b45525 "frexp_b45525"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-        %int = OpTypeInt 32 1
-      %v4int = OpTypeVector %int 4
-%_ptr_Private_v4int = OpTypePointer Private %v4int
-         %13 = OpConstantNull %v4int
-      %arg_1 = OpVariable %_ptr_Private_v4int Private %13
-       %void = OpTypeVoid
-         %14 = OpTypeFunction %void
-%_ptr_Function_v4float = OpTypePointer Function %v4float
-         %23 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_b45525 = OpFunction %void None %14
-         %17 = OpLabel
-        %res = OpVariable %_ptr_Function_v4float Function %5
-         %18 = OpExtInst %v4float %19 Frexp %5 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %23
-         %25 = OpLabel
-         %26 = OpFunctionCall %void %frexp_b45525
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %14
-         %28 = OpLabel
-         %29 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %29
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %14
-         %32 = OpLabel
-         %33 = OpFunctionCall %void %frexp_b45525
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %14
-         %35 = OpLabel
-         %36 = OpFunctionCall %void %frexp_b45525
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/b45525.wgsl.expected.wgsl

@@ -1,25 +0,0 @@
-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() {
-  var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_b45525();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_b45525();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b45525();
-}

test/intrinsics/gen/frexp/b87f4e.wgsl

@@ -1,35 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<workgroup> arg_1: vec4<i32>;
-
-// fn frexp(vec<4, f32>, ptr<workgroup, vec<4, i32>, read_write>) -> vec<4, f32>
-fn frexp_b87f4e() {
-  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b87f4e();
-}

test/intrinsics/gen/frexp/b87f4e.wgsl.expected.hlsl

@@ -1,34 +0,0 @@
-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);
-  param_1 = int4(float_exp);
-  return significand;
-}
-
-groupshared int4 arg_1;
-
-void frexp_b87f4e() {
-  float4 res = tint_frexp(float4(0.0f, 0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol_1 {
-  uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
-  {
-    arg_1 = int4(0, 0, 0, 0);
-  }
-  GroupMemoryBarrierWithGroupSync();
-  frexp_b87f4e();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
-  compute_main_inner(tint_symbol.local_invocation_index);
-  return;
-}

test/intrinsics/gen/frexp/b87f4e.wgsl.expected.msl

@@ -1,33 +0,0 @@
-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) {
-  float4 res = tint_frexp(float4(), tint_symbol);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup int4* const tint_symbol_1) {
-  {
-    *(tint_symbol_1) = int4();
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  frexp_b87f4e(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup int4 tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/intrinsics/gen/frexp/b87f4e.wgsl.expected.spvasm

@@ -1,60 +0,0 @@
-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
-; Bound: 33
-; Schema: 0
-               OpCapability Shader
-         %15 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %local_invocation_index_1 "local_invocation_index_1"
-               OpName %arg_1 "arg_1"
-               OpName %frexp_b87f4e "frexp_b87f4e"
-               OpName %res "res"
-               OpName %compute_main_inner "compute_main_inner"
-               OpName %local_invocation_index "local_invocation_index"
-               OpName %compute_main "compute_main"
-               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
-       %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
-        %int = OpTypeInt 32 1
-      %v4int = OpTypeVector %int 4
-%_ptr_Workgroup_v4int = OpTypePointer Workgroup %v4int
-      %arg_1 = OpVariable %_ptr_Workgroup_v4int Workgroup
-       %void = OpTypeVoid
-          %8 = OpTypeFunction %void
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-         %16 = OpConstantNull %v4float
-%_ptr_Function_v4float = OpTypePointer Function %v4float
-         %20 = OpTypeFunction %void %uint
-         %24 = OpConstantNull %v4int
-     %uint_2 = OpConstant %uint 2
-   %uint_264 = OpConstant %uint 264
-%frexp_b87f4e = OpFunction %void None %8
-         %11 = OpLabel
-        %res = OpVariable %_ptr_Function_v4float Function %16
-         %12 = OpExtInst %v4float %15 Frexp %16 %arg_1
-               OpStore %res %12
-               OpReturn
-               OpFunctionEnd
-%compute_main_inner = OpFunction %void None %20
-%local_invocation_index = OpFunctionParameter %uint
-         %23 = OpLabel
-               OpStore %arg_1 %24
-               OpControlBarrier %uint_2 %uint_2 %uint_264
-         %28 = OpFunctionCall %void %frexp_b87f4e
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %8
-         %30 = OpLabel
-         %32 = OpLoad %uint %local_invocation_index_1
-         %31 = OpFunctionCall %void %compute_main_inner %32
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/b87f4e.wgsl.expected.wgsl

@@ -1,14 +0,0 @@
-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() {
-  var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b87f4e();
-}

test/intrinsics/gen/frexp/b9e4de.wgsl

@@ -1,46 +0,0 @@
-// 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>, ptr<function, vec<3, i32>, write>) -> vec<3, f32>
-fn frexp_b9e4de() {
-  var arg_1: vec3<i32>;
-  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_b9e4de();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_b9e4de();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b9e4de();
-}

test/intrinsics/gen/frexp/b9e4de.wgsl.expected.hlsl

@@ -1,42 +0,0 @@
-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);
-  param_1 = int3(float_exp);
-  return significand;
-}
-
-void frexp_b9e4de() {
-  int3 arg_1 = int3(0, 0, 0);
-  float3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_b9e4de();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_b9e4de();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_b9e4de();
-  return;
-}

test/intrinsics/gen/frexp/b9e4de.wgsl.expected.msl

@@ -1,46 +0,0 @@
-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 = tint_frexp(float3(), &(arg_1));
-}
-
-float4 vertex_main_inner() {
-  frexp_b9e4de();
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  float4 const inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  frexp_b9e4de();
-  return;
-}
-
-kernel void compute_main() {
-  frexp_b9e4de();
-  return;
-}
-

test/intrinsics/gen/frexp/b9e4de.wgsl.expected.spvasm

@@ -1,77 +0,0 @@
-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
-; Bound: 39
-; Schema: 0
-               OpCapability Shader
-         %20 = OpExtInstImport "GLSL.std.450"
-               OpMemoryModel Logical GLSL450
-               OpEntryPoint Vertex %vertex_main "vertex_main" %value %vertex_point_size
-               OpEntryPoint Fragment %fragment_main "fragment_main"
-               OpEntryPoint GLCompute %compute_main "compute_main"
-               OpExecutionMode %fragment_main OriginUpperLeft
-               OpExecutionMode %compute_main LocalSize 1 1 1
-               OpName %value "value"
-               OpName %vertex_point_size "vertex_point_size"
-               OpName %frexp_b9e4de "frexp_b9e4de"
-               OpName %arg_1 "arg_1"
-               OpName %res "res"
-               OpName %vertex_main_inner "vertex_main_inner"
-               OpName %vertex_main "vertex_main"
-               OpName %fragment_main "fragment_main"
-               OpName %compute_main "compute_main"
-               OpDecorate %value BuiltIn Position
-               OpDecorate %vertex_point_size BuiltIn PointSize
-      %float = OpTypeFloat 32
-    %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
-          %5 = OpConstantNull %v4float
-      %value = OpVariable %_ptr_Output_v4float Output %5
-%_ptr_Output_float = OpTypePointer Output %float
-          %8 = OpConstantNull %float
-%vertex_point_size = OpVariable %_ptr_Output_float Output %8
-       %void = OpTypeVoid
-          %9 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-      %v3int = OpTypeVector %int 3
-%_ptr_Function_v3int = OpTypePointer Function %v3int
-         %17 = OpConstantNull %v3int
-    %v3float = OpTypeVector %float 3
-         %21 = OpConstantNull %v3float
-%_ptr_Function_v3float = OpTypePointer Function %v3float
-         %25 = OpTypeFunction %v4float
-    %float_1 = OpConstant %float 1
-%frexp_b9e4de = OpFunction %void None %9
-         %12 = OpLabel
-      %arg_1 = OpVariable %_ptr_Function_v3int Function %17
-        %res = OpVariable %_ptr_Function_v3float Function %21
-         %18 = OpExtInst %v3float %20 Frexp %21 %arg_1
-               OpStore %res %18
-               OpReturn
-               OpFunctionEnd
-%vertex_main_inner = OpFunction %v4float None %25
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %frexp_b9e4de
-               OpReturnValue %5
-               OpFunctionEnd
-%vertex_main = OpFunction %void None %9
-         %30 = OpLabel
-         %31 = OpFunctionCall %v4float %vertex_main_inner
-               OpStore %value %31
-               OpStore %vertex_point_size %float_1
-               OpReturn
-               OpFunctionEnd
-%fragment_main = OpFunction %void None %9
-         %34 = OpLabel
-         %35 = OpFunctionCall %void %frexp_b9e4de
-               OpReturn
-               OpFunctionEnd
-%compute_main = OpFunction %void None %9
-         %37 = OpLabel
-         %38 = OpFunctionCall %void %frexp_b9e4de
-               OpReturn
-               OpFunctionEnd

test/intrinsics/gen/frexp/b9e4de.wgsl.expected.wgsl

@@ -1,24 +0,0 @@
-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));
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_b9e4de();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_b9e4de();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_b9e4de();
-}

test/intrinsics/gen/frexp/c084e3.wgsl

@@ -1,46 +0,0 @@
-// 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
-////////////////////////////////////////////////////////////////////////////////
-
-var<private> arg_1: vec2<i32>;
-
-// fn frexp(vec<2, f32>, ptr<private, vec<2, i32>, read_write>) -> vec<2, f32>
-fn frexp_c084e3() {
-  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
-}
-
-[[stage(vertex)]]
-fn vertex_main() -> [[builtin(position)]] vec4<f32> {
-  frexp_c084e3();
-  return vec4<f32>();
-}
-
-[[stage(fragment)]]
-fn fragment_main() {
-  frexp_c084e3();
-}
-
-[[stage(compute), workgroup_size(1)]]
-fn compute_main() {
-  frexp_c084e3();
-}

test/intrinsics/gen/frexp/c084e3.wgsl.expected.hlsl

@@ -1,43 +0,0 @@
-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);
-  param_1 = int2(float_exp);
-  return significand;
-}
-
-static int2 arg_1 = int2(0, 0);
-
-void frexp_c084e3() {
-  float2 res = tint_frexp(float2(0.0f, 0.0f), arg_1);
-}
-
-struct tint_symbol {
-  float4 value : SV_Position;
-};
-
-float4 vertex_main_inner() {
-  frexp_c084e3();
-  return float4(0.0f, 0.0f, 0.0f, 0.0f);
-}
-
-tint_symbol vertex_main() {
-  const float4 inner_result = vertex_main_inner();
-  tint_symbol wrapper_result = (tint_symbol)0;
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-void fragment_main() {
-  frexp_c084e3();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  frexp_c084e3();
-  return;
-}

test/intrinsics/gen/frexp/c084e3.wgsl.expected.msl

@@ -1,48 +0,0 @@
-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_1) {
-  float2 res = tint_frexp(float2(), tint_symbol_1);
-}
-
-float4 vertex_main_inner(thread int2* const tint_symbol_2) {
-  frexp_c084e3(tint_symbol_2);
-  return float4();
-}
-
-vertex tint_symbol vertex_main() {
-  thread int2 tint_symbol_3 = 0;
-  float4 const inner_result = vertex_main_inner(&(tint_symbol_3));
-  tint_symbol wrapper_result = {};
-  wrapper_result.value = inner_result;
-  return wrapper_result;
-}
-
-fragment void fragment_main() {
-  thread int2 tint_symbol_4 = 0;
-  frexp_c084e3(&(tint_symbol_4));
-  return;
-}
-
-kernel void compute_main() {
-  thread int2 tint_symbol_5 = 0;
-  frexp_c084e3(&(tint_symbol_5));
-  return;
-}
-