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

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

Deprecate the old overloads.

Fixed: tint:54
Change-Id: If66d26dbac3389ff604734f31b426abe47868b91
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59302
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>

src/intrinsic_table.cc

@@ -29,6 +29,7 @@
 #include "src/sem/storage_texture_type.h"
 #include "src/utils/get_or_create.h"
 #include "src/utils/hash.h"
+#include "src/utils/math.h"
 #include "src/utils/scoped_assignment.h"
 
 namespace tint {
@@ -637,6 +638,91 @@ const sem::ExternalTexture* build_texture_external(MatchState& state) {
   return state.builder.create<sem::ExternalTexture>();
 }
 
+// Builtin types starting with a _ prefix cannot be declared in WGSL, so they
+// can only be used as return types. Because of this, they must only match Any,
+// which is used as the return type matcher.
+bool match_modf_result(const sem::Type* ty) {
+  return ty->Is<Any>();
+}
+bool match_modf_result_vec(const sem::Type* ty, Number& N) {
+  if (!ty->Is<Any>()) {
+    return false;
+  }
+  N = Number::any;
+  return true;
+}
+bool match_frexp_result(const sem::Type* ty) {
+  return ty->Is<Any>();
+}
+bool match_frexp_result_vec(const sem::Type* ty, Number& N) {
+  if (!ty->Is<Any>()) {
+    return false;
+  }
+  N = Number::any;
+  return true;
+}
+
+struct NameAndType {
+  std::string name;
+  sem::Type* type;
+};
+const sem::Struct* build_struct(
+    MatchState& state,
+    std::string name,
+    std::initializer_list<NameAndType> member_names_and_types) {
+  uint32_t offset = 0;
+  uint32_t max_align = 0;
+  sem::StructMemberList members;
+  for (auto& m : member_names_and_types) {
+    uint32_t align = m.type->Align();
+    uint32_t size = m.type->Size();
+    offset = utils::RoundUp(align, offset);
+    max_align = std::max(max_align, align);
+    members.emplace_back(state.builder.create<sem::StructMember>(
+        /* declaration */ nullptr,
+        /* name */ state.builder.Sym(m.name),
+        /* type */ m.type,
+        /* index */ static_cast<uint32_t>(members.size()),
+        /* offset */ offset,
+        /* align */ align,
+        /* size */ size));
+    offset += size;
+  }
+  uint32_t size_without_padding = offset;
+  uint32_t size_with_padding = utils::RoundUp(max_align, offset);
+  return state.builder.create<sem::Struct>(
+      /* declaration */ nullptr,
+      /* name */ state.builder.Sym(name),
+      /* members */ members,
+      /* align */ max_align,
+      /* size */ size_with_padding,
+      /* size_no_padding */ size_without_padding);
+}
+
+const sem::Struct* build_modf_result(MatchState& state) {
+  auto* f32 = state.builder.create<sem::F32>();
+  return build_struct(state, "_modf_result", {{"fract", f32}, {"whole", f32}});
+}
+const sem::Struct* build_modf_result_vec(MatchState& state, Number& n) {
+  auto* vec_f32 = state.builder.create<sem::Vector>(
+      state.builder.create<sem::F32>(), n.Value());
+  return build_struct(state, "_modf_result_vec" + std::to_string(n.Value()),
+                      {{"fract", vec_f32}, {"whole", vec_f32}});
+}
+const sem::Struct* build_frexp_result(MatchState& state) {
+  auto* f32 = state.builder.create<sem::F32>();
+  auto* i32 = state.builder.create<sem::I32>();
+  return build_struct(state, "_frexp_result", {{"sig", f32}, {"exp", i32}});
+}
+const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n) {
+  auto* vec_f32 = state.builder.create<sem::Vector>(
+      state.builder.create<sem::F32>(), n.Value());
+  auto* vec_i32 = state.builder.create<sem::Vector>(
+      state.builder.create<sem::I32>(), n.Value());
+  return build_struct(state, "_frexp_result_vec" + std::to_string(n.Value()),
+                      {{"sig", vec_f32}, {"exp", vec_i32}});
+}
+
 /// ParameterInfo describes a parameter
 struct ParameterInfo {
   /// The parameter usage (parameter name in definition file)

src/intrinsic_table.inl.tmpl

@@ -148,7 +148,7 @@ const sem::Type* {{$class}}::Match(MatchState& state, const sem::Type* ty) const
 {{- range .TemplateParams }}
 {{-   template "DeclareLocalTemplateParam" . }}
 {{- end  }}
-  if (!match_{{.Name}}(ty{{range .TemplateParams}}, {{.GetName}}{{end}})) {
+  if (!match_{{TrimPrefix .Name "_"}}(ty{{range .TemplateParams}}, {{.GetName}}{{end}})) {
     return nullptr;
   }
 {{- range .TemplateParams }}
@@ -157,7 +157,7 @@ const sem::Type* {{$class}}::Match(MatchState& state, const sem::Type* ty) const
     return nullptr;
   }
 {{- end  }}
-  return build_{{.Name}}(state{{range .TemplateParams}}, {{.GetName}}{{end}});
+  return build_{{TrimPrefix .Name "_"}}(state{{range .TemplateParams}}, {{.GetName}}{{end}});
 }
 
 std::string {{$class}}::String(MatchState&{{if .TemplateParams}} state{{end}}) const {

src/intrinsics.def

@@ -96,6 +96,11 @@ type texture_storage_2d_array<F: texel_format, A: access>
 type texture_storage_3d<F: texel_format, A: access>
 type texture_external
 
+type _modf_result
+[[display("_modf_result_vec{N}")]] type _modf_result_vec<N: num>
+type _frexp_result
+[[display("_frexp_result_vec{N}")]] type _frexp_result_vec<N: num>
+
 ////////////////////////////////////////////////////////////////////////////////
 // Type matchers                                                              //
 //                                                                            //
@@ -308,8 +313,10 @@ 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>
-fn frexp<S: function_private_workgroup, A: access>(f32, ptr<S, i32, A>) -> f32
+[[deprecated]] fn frexp<S: function_private_workgroup, A: access>(f32, ptr<S, i32, A>) -> f32
-fn frexp<N: num, S: function_private_workgroup, A: access>(vec<N, f32>, ptr<S, vec<N, i32>, A>) -> vec<N, 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
 [[stage("fragment")]] fn fwidth<N: num>(vec<N, f32>) -> vec<N, f32>
 [[stage("fragment")]] fn fwidthCoarse(f32) -> f32
@@ -341,8 +348,10 @@ fn min<T: fiu32>(T, T) -> T
 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 modf<S: function_private_workgroup, A: access>(f32, ptr<S, f32, A>) -> f32
+[[deprecated]] fn modf<S: function_private_workgroup, A: access>(f32, ptr<S, f32, A>) -> f32
-fn modf<N: num, S: function_private_workgroup, A: access>(vec<N, f32>, ptr<S, vec<N, f32>, A>) -> vec<N, 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>
 fn pack2x16float(vec2<f32>) -> u32
 fn pack2x16snorm(vec2<f32>) -> u32

src/resolver/inferred_type_test.cc

@@ -156,11 +156,11 @@ TEST_F(ResolverInferredTypeTest, InferStruct_Pass) {
   auto* member = Member("x", ty.i32());
   auto* str = Structure("S", {member}, {create<ast::StructBlockDecoration>()});
 
-  auto* expected_type =
+  auto* expected_type = create<sem::Struct>(
-      create<sem::Struct>(str,
+      str, str->name(),
-                          sem::StructMemberList{create<sem::StructMember>(
+      sem::StructMemberList{create<sem::StructMember>(
-                              member, create<sem::I32>(), 0, 0, 0, 4)},
+          member, member->symbol(), create<sem::I32>(), 0, 0, 0, 4)},
-                          0, 4, 4);
+      0, 4, 4);
 
   auto* ctor_expr = Construct(ty.Of(str));
 

src/resolver/intrinsic_test.cc

@@ -828,7 +828,7 @@ TEST_F(ResolverIntrinsicDataTest, Normalize_Error_NoParams) {
 )");
 }
 
-TEST_F(ResolverIntrinsicDataTest, FrexpScalar) {
+TEST_F(ResolverIntrinsicDataTest, DEPRECATED_FrexpScalar) {
   Global("exp", ty.i32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", 1.0f, AddressOf("exp"));
   WrapInFunction(call);
@@ -839,7 +839,7 @@ TEST_F(ResolverIntrinsicDataTest, FrexpScalar) {
   EXPECT_TRUE(TypeOf(call)->Is<sem::F32>());
 }
 
-TEST_F(ResolverIntrinsicDataTest, FrexpVector) {
+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);
@@ -851,6 +851,68 @@ TEST_F(ResolverIntrinsicDataTest, FrexpVector) {
   EXPECT_TRUE(TypeOf(call)->As<sem::Vector>()->type()->Is<sem::F32>());
 }
 
+TEST_F(ResolverIntrinsicDataTest, FrexpScalar) {
+  auto* call = Call("frexp", 1.0f);
+  WrapInFunction(call);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(call), nullptr);
+  auto* ty = TypeOf(call)->As<sem::Struct>();
+  ASSERT_NE(ty, nullptr);
+  ASSERT_EQ(ty->Members().size(), 2u);
+
+  auto* sig = ty->Members()[0];
+  EXPECT_TRUE(sig->Type()->Is<sem::F32>());
+  EXPECT_EQ(sig->Offset(), 0u);
+  EXPECT_EQ(sig->Size(), 4u);
+  EXPECT_EQ(sig->Align(), 4u);
+  EXPECT_EQ(sig->Name(), Sym("sig"));
+
+  auto* exp = ty->Members()[1];
+  EXPECT_TRUE(exp->Type()->Is<sem::I32>());
+  EXPECT_EQ(exp->Offset(), 4u);
+  EXPECT_EQ(exp->Size(), 4u);
+  EXPECT_EQ(exp->Align(), 4u);
+  EXPECT_EQ(exp->Name(), Sym("exp"));
+
+  EXPECT_EQ(ty->Size(), 8u);
+  EXPECT_EQ(ty->SizeNoPadding(), 8u);
+}
+
+TEST_F(ResolverIntrinsicDataTest, FrexpVector) {
+  auto* call = Call("frexp", vec3<f32>());
+  WrapInFunction(call);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(call), nullptr);
+  auto* ty = TypeOf(call)->As<sem::Struct>();
+  ASSERT_NE(ty, nullptr);
+  ASSERT_EQ(ty->Members().size(), 2u);
+
+  auto* sig = ty->Members()[0];
+  ASSERT_TRUE(sig->Type()->Is<sem::Vector>());
+  EXPECT_EQ(sig->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(sig->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(sig->Offset(), 0u);
+  EXPECT_EQ(sig->Size(), 12u);
+  EXPECT_EQ(sig->Align(), 16u);
+  EXPECT_EQ(sig->Name(), Sym("sig"));
+
+  auto* exp = ty->Members()[1];
+  ASSERT_TRUE(exp->Type()->Is<sem::Vector>());
+  EXPECT_EQ(exp->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(exp->Type()->As<sem::Vector>()->type()->Is<sem::I32>());
+  EXPECT_EQ(exp->Offset(), 16u);
+  EXPECT_EQ(exp->Size(), 12u);
+  EXPECT_EQ(exp->Align(), 16u);
+  EXPECT_EQ(exp->Name(), Sym("exp"));
+
+  EXPECT_EQ(ty->Size(), 32u);
+  EXPECT_EQ(ty->SizeNoPadding(), 28u);
+}
+
 TEST_F(ResolverIntrinsicDataTest, Frexp_Error_FirstParamInt) {
   Global("exp", ty.i32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("frexp", 1, AddressOf("exp"));
@@ -862,9 +924,11 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_FirstParamInt) {
       r()->error(),
       R"(error: no matching call to frexp(i32, ptr<workgroup, i32, read_write>)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -879,9 +943,11 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_SecondParamFloatPtr) {
       r()->error(),
       R"(error: no matching call to frexp(f32, ptr<workgroup, f32, read_write>)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -893,9 +959,11 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_SecondParamNotAPointer) {
 
   EXPECT_EQ(r()->error(), R"(error: no matching call to frexp(f32, i32)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -910,13 +978,15 @@ TEST_F(ResolverIntrinsicDataTest, Frexp_Error_VectorSizesDontMatch) {
       r()->error(),
       R"(error: no matching call to frexp(vec2<f32>, ptr<workgroup, vec4<i32>, read_write>)
 
-2 candidate functions:
+4 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, ModfScalar) {
+TEST_F(ResolverIntrinsicDataTest, DEPRECATED_ModfScalar) {
   Global("whole", ty.f32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", 1.0f, AddressOf("whole"));
   WrapInFunction(call);
@@ -927,7 +997,7 @@ TEST_F(ResolverIntrinsicDataTest, ModfScalar) {
   EXPECT_TRUE(TypeOf(call)->Is<sem::F32>());
 }
 
-TEST_F(ResolverIntrinsicDataTest, ModfVector) {
+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);
@@ -939,6 +1009,68 @@ TEST_F(ResolverIntrinsicDataTest, ModfVector) {
   EXPECT_TRUE(TypeOf(call)->As<sem::Vector>()->type()->Is<sem::F32>());
 }
 
+TEST_F(ResolverIntrinsicDataTest, ModfScalar) {
+  auto* call = Call("modf", 1.0f);
+  WrapInFunction(call);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(call), nullptr);
+  auto* ty = TypeOf(call)->As<sem::Struct>();
+  ASSERT_NE(ty, nullptr);
+  ASSERT_EQ(ty->Members().size(), 2u);
+
+  auto* fract = ty->Members()[0];
+  EXPECT_TRUE(fract->Type()->Is<sem::F32>());
+  EXPECT_EQ(fract->Offset(), 0u);
+  EXPECT_EQ(fract->Size(), 4u);
+  EXPECT_EQ(fract->Align(), 4u);
+  EXPECT_EQ(fract->Name(), Sym("fract"));
+
+  auto* whole = ty->Members()[1];
+  EXPECT_TRUE(whole->Type()->Is<sem::F32>());
+  EXPECT_EQ(whole->Offset(), 4u);
+  EXPECT_EQ(whole->Size(), 4u);
+  EXPECT_EQ(whole->Align(), 4u);
+  EXPECT_EQ(whole->Name(), Sym("whole"));
+
+  EXPECT_EQ(ty->Size(), 8u);
+  EXPECT_EQ(ty->SizeNoPadding(), 8u);
+}
+
+TEST_F(ResolverIntrinsicDataTest, ModfVector) {
+  auto* call = Call("modf", vec3<f32>());
+  WrapInFunction(call);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(call), nullptr);
+  auto* ty = TypeOf(call)->As<sem::Struct>();
+  ASSERT_NE(ty, nullptr);
+  ASSERT_EQ(ty->Members().size(), 2u);
+
+  auto* fract = ty->Members()[0];
+  ASSERT_TRUE(fract->Type()->Is<sem::Vector>());
+  EXPECT_EQ(fract->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(fract->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(fract->Offset(), 0u);
+  EXPECT_EQ(fract->Size(), 12u);
+  EXPECT_EQ(fract->Align(), 16u);
+  EXPECT_EQ(fract->Name(), Sym("fract"));
+
+  auto* whole = ty->Members()[1];
+  ASSERT_TRUE(whole->Type()->Is<sem::Vector>());
+  EXPECT_EQ(whole->Type()->As<sem::Vector>()->Width(), 3u);
+  EXPECT_TRUE(whole->Type()->As<sem::Vector>()->type()->Is<sem::F32>());
+  EXPECT_EQ(whole->Offset(), 16u);
+  EXPECT_EQ(whole->Size(), 12u);
+  EXPECT_EQ(whole->Align(), 16u);
+  EXPECT_EQ(whole->Name(), Sym("whole"));
+
+  EXPECT_EQ(ty->Size(), 32u);
+  EXPECT_EQ(ty->SizeNoPadding(), 28u);
+}
+
 TEST_F(ResolverIntrinsicDataTest, Modf_Error_FirstParamInt) {
   Global("whole", ty.f32(), ast::StorageClass::kWorkgroup);
   auto* call = Call("modf", 1, AddressOf("whole"));
@@ -950,9 +1082,11 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_FirstParamInt) {
       r()->error(),
       R"(error: no matching call to modf(i32, ptr<workgroup, f32, read_write>)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -967,9 +1101,11 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_SecondParamIntPtr) {
       r()->error(),
       R"(error: no matching call to modf(f32, ptr<workgroup, i32, read_write>)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -981,9 +1117,11 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_SecondParamNotAPointer) {
 
   EXPECT_EQ(r()->error(), R"(error: no matching call to modf(f32, f32)
 
-2 candidate functions:
+4 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
 )");
 }
 
@@ -998,9 +1136,11 @@ TEST_F(ResolverIntrinsicDataTest, Modf_Error_VectorSizesDontMatch) {
       r()->error(),
       R"(error: no matching call to modf(vec2<f32>, ptr<workgroup, vec4<f32>, read_write>)
 
-2 candidate functions:
+4 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/resolver/resolver.cc

@@ -2998,7 +2998,7 @@ bool Resolver::MemberAccessor(ast::MemberAccessorExpression* expr) {
 
     const sem::StructMember* member = nullptr;
     for (auto* m : str->Members()) {
-      if (m->Declaration()->symbol() == symbol) {
+      if (m->Name() == symbol) {
         ret = m->Type();
         member = m;
         break;
@@ -4088,7 +4088,7 @@ sem::Struct* Resolver::Structure(const ast::Struct* str) {
     offset = utils::RoundUp(align, offset);
 
     auto* sem_member = builder_->create<sem::StructMember>(
-        member, const_cast<sem::Type*>(type),
+        member, member->symbol(), const_cast<sem::Type*>(type),
         static_cast<uint32_t>(sem_members.size()), offset, align, size);
     builder_->Sem().Add(member, sem_member);
     sem_members.emplace_back(sem_member);
@@ -4100,8 +4100,9 @@ sem::Struct* Resolver::Structure(const ast::Struct* str) {
   auto size_no_padding = struct_size;
   struct_size = utils::RoundUp(struct_align, struct_size);
 
-  auto* out = builder_->create<sem::Struct>(str, sem_members, struct_align,
+  auto* out =
-                                            struct_size, size_no_padding);
+      builder_->create<sem::Struct>(str, str->name(), sem_members, struct_align,
+                                    struct_size, size_no_padding);
 
   // Keep track of atomic members for validation after all usages have been
   // determined.

src/sem/sem_struct_test.cc

@@ -27,8 +27,9 @@ TEST_F(StructTest, Creation) {
   auto* impl =
       create<ast::Struct>(name, ast::StructMemberList{}, ast::DecorationList{});
   auto* ptr = impl;
-  auto* s = create<sem::Struct>(impl, StructMemberList{}, 4 /* align */,
+  auto* s =
-                                8 /* size */, 16 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name(), StructMemberList{}, 4 /* align */,
+                          8 /* size */, 16 /* size_no_padding */);
   EXPECT_EQ(s->Declaration(), ptr);
   EXPECT_EQ(s->Align(), 4u);
   EXPECT_EQ(s->Size(), 8u);
@@ -39,8 +40,9 @@ TEST_F(StructTest, TypeName) {
   auto name = Sym("my_struct");
   auto* impl =
       create<ast::Struct>(name, ast::StructMemberList{}, ast::DecorationList{});
-  auto* s = create<sem::Struct>(impl, StructMemberList{}, 4 /* align */,
+  auto* s =
-                                4 /* size */, 4 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name(), StructMemberList{}, 4 /* align */,
+                          4 /* size */, 4 /* size_no_padding */);
   EXPECT_EQ(s->type_name(), "__struct_$1");
 }
 
@@ -48,8 +50,9 @@ TEST_F(StructTest, FriendlyName) {
   auto name = Sym("my_struct");
   auto* impl =
       create<ast::Struct>(name, ast::StructMemberList{}, ast::DecorationList{});
-  auto* s = create<sem::Struct>(impl, StructMemberList{}, 4 /* align */,
+  auto* s =
-                                4 /* size */, 4 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name(), StructMemberList{}, 4 /* align */,
+                          4 /* size */, 4 /* size_no_padding */);
   EXPECT_EQ(s->FriendlyName(Symbols()), "my_struct");
 }
 

src/sem/struct.cc

@@ -27,11 +27,13 @@ namespace tint {
 namespace sem {
 
 Struct::Struct(const ast::Struct* declaration,
+               Symbol name,
                StructMemberList members,
                uint32_t align,
                uint32_t size,
                uint32_t size_no_padding)
     : declaration_(declaration),
+      name_(name),
       members_(std::move(members)),
       align_(align),
       size_(size),
@@ -57,7 +59,7 @@ const StructMember* Struct::FindMember(Symbol name) const {
 }
 
 std::string Struct::type_name() const {
-  return declaration_->type_name();
+  return "__struct_" + name_.to_str();
 }
 
 uint32_t Struct::Align() const {
@@ -69,7 +71,7 @@ uint32_t Struct::Size() const {
 }
 
 std::string Struct::FriendlyName(const SymbolTable& symbols) const {
-  return symbols.NameFor(declaration_->name());
+  return symbols.NameFor(name_);
 }
 
 bool Struct::IsConstructible() const {
@@ -77,12 +79,14 @@ bool Struct::IsConstructible() const {
 }
 
 StructMember::StructMember(ast::StructMember* declaration,
+                           Symbol name,
                            sem::Type* type,
                            uint32_t index,
                            uint32_t offset,
                            uint32_t align,
                            uint32_t size)
     : declaration_(declaration),
+      name_(name),
       type_(type),
       index_(index),
       offset_(offset),

src/sem/struct.h

@@ -58,12 +58,14 @@ class Struct : public Castable<Struct, Type> {
  public:
   /// Constructor
   /// @param declaration the AST structure declaration
+  /// @param name the name of the structure
   /// @param members the structure members
   /// @param align the byte alignment of the structure
   /// @param size the byte size of the structure
   /// @param size_no_padding size of the members without the end of structure
   /// alignment padding
   Struct(const ast::Struct* declaration,
+         Symbol name,
          StructMemberList members,
          uint32_t align,
          uint32_t size,
@@ -75,6 +77,9 @@ class Struct : public Castable<Struct, Type> {
   /// @returns the struct
   const ast::Struct* Declaration() const { return declaration_; }
 
+  /// @returns the name of the structure
+  Symbol Name() const { return name_; }
+
   /// @returns the members of the structure
   const StructMemberList& Members() const { return members_; }
 
@@ -156,6 +161,7 @@ class Struct : public Castable<Struct, Type> {
   uint64_t LargestMemberBaseAlignment(MemoryLayout mem_layout) const;
 
   ast::Struct const* const declaration_;
+  Symbol const name_;
   StructMemberList const members_;
   uint32_t const align_;
   uint32_t const size_;
@@ -170,12 +176,14 @@ class StructMember : public Castable<StructMember, Node> {
  public:
   /// Constructor
   /// @param declaration the AST declaration node
+  /// @param name the name of the structure
   /// @param type the type of the member
   /// @param index the index of the member in the structure
   /// @param offset the byte offset from the base of the structure
   /// @param align the byte alignment of the member
   /// @param size the byte size of the member
   StructMember(ast::StructMember* declaration,
+               Symbol name,
                sem::Type* type,
                uint32_t index,
                uint32_t offset,
@@ -188,6 +196,9 @@ class StructMember : public Castable<StructMember, Node> {
   /// @returns the AST declaration node
   ast::StructMember* Declaration() const { return declaration_; }
 
+  /// @returns the name of the structure
+  Symbol Name() const { return name_; }
+
   /// @returns the type of the member
   sem::Type* Type() const { return type_; }
 
@@ -205,6 +216,7 @@ class StructMember : public Castable<StructMember, Node> {
 
  private:
   ast::StructMember* const declaration_;
+  Symbol const name_;
   sem::Type* const type_;
   uint32_t const index_;
   uint32_t const offset_;

src/transform/transform_test.cc

@@ -106,8 +106,9 @@ TEST_F(CreateASTTypeForTest, ArrayNonImplicitStride) {
 TEST_F(CreateASTTypeForTest, Struct) {
   auto* str = create([](ProgramBuilder& b) {
     auto* decl = b.Structure("S", {}, {});
-    return b.create<sem::Struct>(decl, sem::StructMemberList{}, 4 /* align */,
+    return b.create<sem::Struct>(decl, decl->name(), sem::StructMemberList{},
-                                 4 /* size */, 4 /* size_no_padding */);
+                                 4 /* align */, 4 /* size */,
+                                 4 /* size_no_padding */);
   });
   ASSERT_TRUE(str->Is<ast::TypeName>());
   EXPECT_EQ(

src/writer/hlsl/generator_impl.cc

@@ -147,7 +147,7 @@ bool GeneratorImpl::Generate() {
         return false;
       }
     } else if (auto* str = decl->As<ast::Struct>()) {
-      if (!EmitStructType(builder_.Sem().Get(str))) {
+      if (!EmitStructType(current_buffer_, builder_.Sem().Get(str))) {
         return false;
       }
     } else if (auto* func = decl->As<ast::Function>()) {
@@ -521,6 +521,8 @@ bool GeneratorImpl::EmitCall(std::ostream& out, ast::CallExpression* expr) {
       return EmitTextureCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kSelect) {
       return EmitSelectCall(out, expr);
+    } else if (intrinsic->Type() == sem::IntrinsicType::kModf) {
+      return EmitModfCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kFrexp) {
       return EmitFrexpCall(out, expr, intrinsic);
     } else if (intrinsic->Type() == sem::IntrinsicType::kIsNormal) {
@@ -1270,9 +1272,93 @@ bool GeneratorImpl::EmitSelectCall(std::ostream& out,
   return true;
 }
 
+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) {
+          auto* ty = intrinsic->Parameters()[0]->Type();
+          auto in = params[0];
+
+          std::string width;
+          if (auto* vec = ty->As<sem::Vector>()) {
+            width = std::to_string(vec->Width());
+          }
+
+          // Emit the builtin return type unique to this overload. This does not
+          // exist in the AST, so it will not be generated in Generate().
+          if (!EmitStructType(&helpers_,
+                              intrinsic->ReturnType()->As<sem::Struct>())) {
+            return false;
+          }
+
+          line(b) << "float" << width << " whole;";
+          line(b) << "float" << width << " fract = modf(" << in << ", whole);";
+          {
+            auto l = line(b);
+            if (!EmitType(l, intrinsic->ReturnType(), ast::StorageClass::kNone,
+                          ast::Access::kUndefined, "")) {
+              return false;
+            }
+            l << " result = {fract, whole};";
+          }
+          line(b) << "return result;";
+          return true;
+        });
+  }
+
+  // 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) {
+          auto* ty = intrinsic->Parameters()[0]->Type();
+          auto in = params[0];
+
+          std::string width;
+          if (auto* vec = ty->As<sem::Vector>()) {
+            width = std::to_string(vec->Width());
+          }
+
+          // Emit the builtin return type unique to this overload. This does not
+          // exist in the AST, so it will not be generated in Generate().
+          if (!EmitStructType(&helpers_,
+                              intrinsic->ReturnType()->As<sem::Struct>())) {
+            return false;
+          }
+
+          line(b) << "float" << width << " exp;";
+          line(b) << "float" << width << " sig = frexp(" << in << ", exp);";
+          {
+            auto l = line(b);
+            if (!EmitType(l, intrinsic->ReturnType(), ast::StorageClass::kNone,
+                          ast::Access::kUndefined, "")) {
+              return false;
+            }
+            l << " result = {sig, int" << width << "(exp)};";
+          }
+          line(b) << "return result;";
+          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(
@@ -2906,7 +2992,7 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     }
     out << "State";
   } else if (auto* str = type->As<sem::Struct>()) {
-    out << builder_.Symbols().NameFor(str->Declaration()->name());
+    out << StructName(str);
   } else if (auto* tex = type->As<sem::Texture>()) {
     auto* storage = tex->As<sem::StorageTexture>();
     auto* multism = tex->As<sem::MultisampledTexture>();
@@ -3015,7 +3101,7 @@ bool GeneratorImpl::EmitTypeAndName(std::ostream& out,
   return true;
 }
 
-bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
+bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
   auto storage_class_uses = str->StorageClassUsage();
   if (storage_class_uses.size() ==
       (storage_class_uses.count(ast::StorageClass::kStorage) +
@@ -3029,71 +3115,74 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
     return true;
   }
 
-  auto struct_name = builder_.Symbols().NameFor(str->Declaration()->name());
+  line(b) << "struct " << StructName(str) << " {";
-  line() << "struct " << struct_name << " {";
   {
-    ScopedIndent si(this);
+    ScopedIndent si(b);
     for (auto* mem : str->Members()) {
-      auto name = builder_.Symbols().NameFor(mem->Declaration()->symbol());
+      auto name = builder_.Symbols().NameFor(mem->Name());
 
       auto* ty = mem->Type();
 
-      auto out = line();
+      auto out = line(b);
 
       std::string pre, post;
 
-      for (auto* deco : mem->Declaration()->decorations()) {
+      if (auto* decl = mem->Declaration()) {
-        if (auto* location = deco->As<ast::LocationDecoration>()) {
+        for (auto* deco : decl->decorations()) {
-          auto& pipeline_stage_uses = str->PipelineStageUses();
+          if (auto* location = deco->As<ast::LocationDecoration>()) {
-          if (pipeline_stage_uses.size() != 1) {
+            auto& pipeline_stage_uses = str->PipelineStageUses();
-            TINT_ICE(Writer, diagnostics_)
+            if (pipeline_stage_uses.size() != 1) {
-                << "invalid entry point IO struct uses";
+              TINT_ICE(Writer, diagnostics_)
-          }
+                  << "invalid entry point IO struct uses";
+            }
 
-          if (pipeline_stage_uses.count(
+            if (pipeline_stage_uses.count(
-                  sem::PipelineStageUsage::kVertexInput)) {
+                    sem::PipelineStageUsage::kVertexInput)) {
-            post += " : TEXCOORD" + std::to_string(location->value());
+              post += " : TEXCOORD" + std::to_string(location->value());
-          } else if (pipeline_stage_uses.count(
+            } else if (pipeline_stage_uses.count(
-                         sem::PipelineStageUsage::kVertexOutput)) {
+                           sem::PipelineStageUsage::kVertexOutput)) {
-            post += " : TEXCOORD" + std::to_string(location->value());
+              post += " : TEXCOORD" + std::to_string(location->value());
-          } else if (pipeline_stage_uses.count(
+            } else if (pipeline_stage_uses.count(
-                         sem::PipelineStageUsage::kFragmentInput)) {
+                           sem::PipelineStageUsage::kFragmentInput)) {
-            post += " : TEXCOORD" + std::to_string(location->value());
+              post += " : TEXCOORD" + std::to_string(location->value());
-          } else if (pipeline_stage_uses.count(
+            } else if (pipeline_stage_uses.count(
-                         sem::PipelineStageUsage::kFragmentOutput)) {
+                           sem::PipelineStageUsage::kFragmentOutput)) {
-            post += " : SV_Target" + std::to_string(location->value());
+              post += " : SV_Target" + std::to_string(location->value());
-          } else {
+            } else {
+              TINT_ICE(Writer, diagnostics_)
+                  << "invalid use of location decoration";
+            }
+          } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
+            auto attr = builtin_to_attribute(builtin->value());
+            if (attr.empty()) {
+              diagnostics_.add_error(diag::System::Writer,
+                                     "unsupported builtin");
+              return false;
+            }
+            post += " : " + attr;
+          } else if (auto* interpolate =
+                         deco->As<ast::InterpolateDecoration>()) {
+            auto mod = interpolation_to_modifiers(interpolate->type(),
+                                                  interpolate->sampling());
+            if (mod.empty()) {
+              diagnostics_.add_error(diag::System::Writer,
+                                     "unsupported interpolation");
+              return false;
+            }
+            pre += mod;
+
+          } else if (deco->Is<ast::InvariantDecoration>()) {
+            // Note: `precise` is not exactly the same as `invariant`, but is
+            // stricter and therefore provides the necessary guarantees.
+            // See discussion here: https://github.com/gpuweb/gpuweb/issues/893
+            pre += "precise ";
+          } else if (!deco->IsAnyOf<ast::StructMemberAlignDecoration,
+                                    ast::StructMemberOffsetDecoration,
+                                    ast::StructMemberSizeDecoration>()) {
             TINT_ICE(Writer, diagnostics_)
-                << "invalid use of location decoration";
+                << "unhandled struct member attribute: " << deco->name();
-          }
-        } else if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
-          auto attr = builtin_to_attribute(builtin->value());
-          if (attr.empty()) {
-            diagnostics_.add_error(diag::System::Writer, "unsupported builtin");
             return false;
           }
-          post += " : " + attr;
-        } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) {
-          auto mod = interpolation_to_modifiers(interpolate->type(),
-                                                interpolate->sampling());
-          if (mod.empty()) {
-            diagnostics_.add_error(diag::System::Writer,
-                                   "unsupported interpolation");
-            return false;
-          }
-          pre += mod;
-
-        } else if (deco->Is<ast::InvariantDecoration>()) {
-          // Note: `precise` is not exactly the same as `invariant`, but is
-          // stricter and therefore provides the necessary guarantees.
-          // See discussion here: https://github.com/gpuweb/gpuweb/issues/893
-          pre += "precise ";
-        } else if (!deco->IsAnyOf<ast::StructMemberAlignDecoration,
-                                  ast::StructMemberOffsetDecoration,
-                                  ast::StructMemberSizeDecoration>()) {
-          TINT_ICE(Writer, diagnostics_)
-              << "unhandled struct member attribute: " << deco->name();
-          return false;
         }
       }
 
@@ -3106,7 +3195,7 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
     }
   }
 
-  line() << "};";
+  line(b) << "};";
 
   return true;
 }
@@ -3239,11 +3328,14 @@ bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out,
                                         F&& build) {
   // Generate the helper function if it hasn't been created already
   auto fn = utils::GetOrCreate(intrinsics_, intrinsic, [&]() -> std::string {
+    TextBuffer b;
+    TINT_DEFER(helpers_.Append(b));
+
     auto fn_name =
         UniqueIdentifier(std::string("tint_") + sem::str(intrinsic->Type()));
     std::vector<std::string> parameter_names;
     {
-      auto decl = line(&helpers_);
+      auto decl = line(&b);
       if (!EmitTypeAndName(decl, intrinsic->ReturnType(),
                            ast::StorageClass::kNone, ast::Access::kUndefined,
                            fn_name)) {
@@ -3271,13 +3363,13 @@ bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out,
       decl << " {";
     }
     {
-      ScopedIndent si(&helpers_);
+      ScopedIndent si(&b);
-      if (!build(&helpers_, parameter_names)) {
+      if (!build(&b, parameter_names)) {
         return "";
       }
     }
-    line(&helpers_) << "}";
+    line(&b) << "}";
-    line(&helpers_);
+    line(&b);
     return fn_name;
   });
 

src/writer/hlsl/generator_impl.h

@@ -155,6 +155,14 @@ class GeneratorImpl : public TextGenerator {
   /// @param expr the call expression
   /// @returns true if the call expression is emitted
   bool EmitSelectCall(std::ostream& out, ast::CallExpression* expr);
+  /// Handles generating a call to the `modf()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitModfCall(std::ostream& out,
+                    ast::CallExpression* expr,
+                    const sem::Intrinsic* intrinsic);
   /// Handles generating a call to the `frexp()` intrinsic
   /// @param out the output of the expression stream
   /// @param expr the call expression
@@ -320,7 +328,7 @@ class GeneratorImpl : public TextGenerator {
   /// @param type the type to generate
   /// @param storage_class the storage class of the variable
   /// @param access the access control type of the variable
-  /// @param name the name of the variable, used for array emission.
+  /// @param name the name to emit
   /// @returns true if the type is emitted
   bool EmitTypeAndName(std::ostream& out,
                        const sem::Type* type,
@@ -328,9 +336,10 @@ class GeneratorImpl : public TextGenerator {
                        ast::Access access,
                        const std::string& name);
   /// Handles generating a structure declaration
+  /// @param buffer the text buffer that the type declaration will be written to
   /// @param ty the struct to generate
   /// @returns true if the struct is emitted
-  bool EmitStructType(const sem::Struct* ty);
+  bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
   /// Handles a unary op expression
   /// @param out the output of the expression stream
   /// @param expr the expression to emit
@@ -413,7 +422,6 @@ class GeneratorImpl : public TextGenerator {
   ///          `buffer` is the body of the generated function
   ///          `params` is the name of all the generated function parameters
   /// @returns true if the call expression is emitted
-
   template <typename F>
   bool CallIntrinsicHelper(std::ostream& out,
                            ast::CallExpression* call,

src/writer/hlsl/generator_impl_type_test.cc

@@ -173,9 +173,10 @@ TEST_F(HlslGeneratorImplTest_Type, EmitType_StructDecl) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(struct S {
+  EXPECT_EQ(buf.String(), R"(struct S {
   int a;
   float b;
 };
@@ -197,9 +198,10 @@ TEST_F(HlslGeneratorImplTest_Type, EmitType_StructDecl_OmittedIfStorageBuffer) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
-  EXPECT_EQ(gen.result(), "");
+  EXPECT_EQ(buf.String(), "");
 }
 
 TEST_F(HlslGeneratorImplTest_Type, EmitType_Struct) {
@@ -247,9 +249,10 @@ TEST_F(HlslGeneratorImplTest_Type, EmitType_Struct_WithOffsetAttributes) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(struct S {
+  EXPECT_EQ(buf.String(), R"(struct S {
   int a;
   float b;
 };

src/writer/msl/generator_impl.cc

@@ -150,7 +150,7 @@ bool GeneratorImpl::Generate() {
 
 bool GeneratorImpl::EmitTypeDecl(const sem::Type* ty) {
   if (auto* str = ty->As<sem::Struct>()) {
-    if (!EmitStructType(str)) {
+    if (!EmitStructType(current_buffer_, str)) {
       return false;
     }
   } else {
@@ -392,6 +392,11 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
   auto name = generate_builtin_name(intrinsic);
 
   switch (intrinsic->Type()) {
+    case sem::IntrinsicType::kModf:
+      return EmitModfCall(out, expr, intrinsic);
+    case sem::IntrinsicType::kFrexp:
+      return EmitFrexpCall(out, expr, intrinsic);
+
     case sem::IntrinsicType::kPack2x16float:
     case sem::IntrinsicType::kUnpack2x16float: {
       if (intrinsic->Type() == sem::IntrinsicType::kPack2x16float) {
@@ -423,28 +428,6 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
       return true;
     }
 
-    case sem::IntrinsicType::kFrexp:
-    case sem::IntrinsicType::kModf: {
-      // TODO(bclayton): These intrinsics are likely to change signature.
-      // See: crbug.com/tint/54 and https://github.com/gpuweb/gpuweb/issues/1480
-      out << name;
-      ScopedParen sp(out);
-      if (!EmitExpression(out, expr->params()[0])) {
-        return false;
-      }
-      out << ", ";
-      {
-        // MSL has a reference for the second parameter, but WGSL has a pointer.
-        // Dereference the argument.
-        out << "*";
-        ScopedParen sp2(out);
-        if (!EmitExpression(out, expr->params()[1])) {
-          return false;
-        }
-      }
-      return true;
-    }
-
     case sem::IntrinsicType::kLength: {
       auto* sem = builder_.Sem().Get(expr->params()[0]);
       if (sem->Type()->UnwrapRef()->is_scalar()) {
@@ -869,6 +852,105 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
   return true;
 }
 
+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) {
+          auto* ty = intrinsic->Parameters()[0]->Type();
+          auto in = params[0];
+
+          std::string width;
+          if (auto* vec = ty->As<sem::Vector>()) {
+            width = std::to_string(vec->Width());
+          }
+
+          // Emit the builtin return type unique to this overload. This does not
+          // exist in the AST, so it will not be generated in Generate().
+          if (!EmitStructType(&helpers_,
+                              intrinsic->ReturnType()->As<sem::Struct>())) {
+            return false;
+          }
+
+          line(b) << "float" << width << " whole;";
+          line(b) << "float" << width << " fract = modf(" << in << ", whole);";
+          line(b) << "return {fract, whole};";
+          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) {
+          auto* ty = intrinsic->Parameters()[0]->Type();
+          auto in = params[0];
+
+          std::string width;
+          if (auto* vec = ty->As<sem::Vector>()) {
+            width = std::to_string(vec->Width());
+          }
+
+          // Emit the builtin return type unique to this overload. This does not
+          // exist in the AST, so it will not be generated in Generate().
+          if (!EmitStructType(&helpers_,
+                              intrinsic->ReturnType()->As<sem::Struct>())) {
+            return false;
+          }
+
+          line(b) << "int" << width << " exp;";
+          line(b) << "float" << width << " sig = frexp(" << in << ", exp);";
+          line(b) << "return {sig, exp};";
+          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_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 = "";
@@ -1846,7 +1928,8 @@ bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
 
 bool GeneratorImpl::EmitType(std::ostream& out,
                              const sem::Type* type,
-                             const std::string& name) {
+                             const std::string& name,
+                             bool* name_printed /* = nullptr */) {
   if (auto* atomic = type->As<sem::Atomic>()) {
     if (atomic->Type()->Is<sem::I32>()) {
       out << "atomic_int";
@@ -1877,6 +1960,9 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     }
     if (!name.empty()) {
       out << " " << name;
+      if (name_printed) {
+        *name_printed = true;
+      }
     }
     for (uint32_t size : sizes) {
       out << "[" << size << "]";
@@ -1914,6 +2000,9 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     out << " ";
     if (ptr->StoreType()->Is<sem::Array>()) {
       std::string inner = "(*" + name + ")";
+      if (name_printed) {
+        *name_printed = true;
+      }
       if (!EmitType(out, ptr->StoreType(), inner)) {
         return false;
       }
@@ -1922,6 +2011,9 @@ bool GeneratorImpl::EmitType(std::ostream& out,
         return false;
       }
       out << "* " << name;
+      if (name_printed) {
+        *name_printed = true;
+      }
     }
     return true;
   }
@@ -1934,7 +2026,7 @@ bool GeneratorImpl::EmitType(std::ostream& out,
   if (auto* str = type->As<sem::Struct>()) {
     // The struct type emits as just the name. The declaration would be emitted
     // as part of emitting the declared types.
-    out << program_->Symbols().NameFor(str->Declaration()->name());
+    out << StructName(str);
     return true;
   }
 
@@ -2031,6 +2123,19 @@ bool GeneratorImpl::EmitType(std::ostream& out,
   return false;
 }
 
+bool GeneratorImpl::EmitTypeAndName(std::ostream& out,
+                                    const sem::Type* type,
+                                    const std::string& name) {
+  bool name_printed = false;
+  if (!EmitType(out, type, name, &name_printed)) {
+    return false;
+  }
+  if (!name_printed) {
+    out << " " << name;
+  }
+  return true;
+}
+
 bool GeneratorImpl::EmitStorageClass(std::ostream& out, ast::StorageClass sc) {
   switch (sc) {
     case ast::StorageClass::kFunction:
@@ -2069,9 +2174,8 @@ bool GeneratorImpl::EmitPackedType(std::ostream& out,
   return EmitType(out, type, name);
 }
 
-bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
+bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
-  line() << "struct " << program_->Symbols().NameFor(str->Declaration()->name())
+  line(b) << "struct " << StructName(str) << " {";
-         << " {";
 
   bool is_host_shareable = str->IsHostShareable();
 
@@ -2089,16 +2193,17 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
       name = UniqueIdentifier("tint_pad");
     } while (str->FindMember(program_->Symbols().Get(name)));
 
-    auto out = line();
+    auto out = line(b);
     add_byte_offset_comment(out, msl_offset);
     out << "int8_t " << name << "[" << size << "];";
   };
 
-  increment_indent();
+  b->IncrementIndent();
+
   uint32_t msl_offset = 0;
   for (auto* mem : str->Members()) {
-    auto out = line();
+    auto out = line(b);
-    auto name = program_->Symbols().NameFor(mem->Declaration()->symbol());
+    auto name = program_->Symbols().NameFor(mem->Name());
     auto wgsl_offset = mem->Offset();
 
     if (is_host_shareable) {
@@ -2135,53 +2240,56 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
     }
 
     // Emit decorations
-    for (auto* deco : mem->Declaration()->decorations()) {
+    if (auto* decl = mem->Declaration()) {
-      if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
+      for (auto* deco : decl->decorations()) {
-        auto attr = builtin_to_attribute(builtin->value());
+        if (auto* builtin = deco->As<ast::BuiltinDecoration>()) {
-        if (attr.empty()) {
+          auto attr = builtin_to_attribute(builtin->value());
-          diagnostics_.add_error(diag::System::Writer, "unknown builtin");
+          if (attr.empty()) {
-          return false;
+            diagnostics_.add_error(diag::System::Writer, "unknown builtin");
-        }
+            return false;
-        out << " [[" << attr << "]]";
+          }
-      } else if (auto* loc = deco->As<ast::LocationDecoration>()) {
+          out << " [[" << attr << "]]";
-        auto& pipeline_stage_uses = str->PipelineStageUses();
+        } else if (auto* loc = deco->As<ast::LocationDecoration>()) {
-        if (pipeline_stage_uses.size() != 1) {
+          auto& pipeline_stage_uses = str->PipelineStageUses();
-          TINT_ICE(Writer, diagnostics_)
+          if (pipeline_stage_uses.size() != 1) {
-              << "invalid entry point IO struct uses";
+            TINT_ICE(Writer, diagnostics_)
-        }
+                << "invalid entry point IO struct uses";
+          }
 
-        if (pipeline_stage_uses.count(sem::PipelineStageUsage::kVertexInput)) {
+          if (pipeline_stage_uses.count(
-          out << " [[attribute(" + std::to_string(loc->value()) + ")]]";
+                  sem::PipelineStageUsage::kVertexInput)) {
-        } else if (pipeline_stage_uses.count(
+            out << " [[attribute(" + std::to_string(loc->value()) + ")]]";
-                       sem::PipelineStageUsage::kVertexOutput)) {
+          } else if (pipeline_stage_uses.count(
-          out << " [[user(locn" + std::to_string(loc->value()) + ")]]";
+                         sem::PipelineStageUsage::kVertexOutput)) {
-        } else if (pipeline_stage_uses.count(
+            out << " [[user(locn" + std::to_string(loc->value()) + ")]]";
-                       sem::PipelineStageUsage::kFragmentInput)) {
+          } else if (pipeline_stage_uses.count(
-          out << " [[user(locn" + std::to_string(loc->value()) + ")]]";
+                         sem::PipelineStageUsage::kFragmentInput)) {
-        } else if (pipeline_stage_uses.count(
+            out << " [[user(locn" + std::to_string(loc->value()) + ")]]";
-                       sem::PipelineStageUsage::kFragmentOutput)) {
+          } else if (pipeline_stage_uses.count(
-          out << " [[color(" + std::to_string(loc->value()) + ")]]";
+                         sem::PipelineStageUsage::kFragmentOutput)) {
-        } else {
+            out << " [[color(" + std::to_string(loc->value()) + ")]]";
+          } else {
+            TINT_ICE(Writer, diagnostics_)
+                << "invalid use of location decoration";
+          }
+        } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) {
+          auto attr = interpolation_to_attribute(interpolate->type(),
+                                                 interpolate->sampling());
+          if (attr.empty()) {
+            diagnostics_.add_error(diag::System::Writer,
+                                   "unknown interpolation attribute");
+            return false;
+          }
+          out << " [[" << attr << "]]";
+        } else if (deco->Is<ast::InvariantDecoration>()) {
+          out << " [[invariant]]";
+          has_invariant_ = true;
+        } else if (!deco->IsAnyOf<ast::StructMemberOffsetDecoration,
+                                  ast::StructMemberAlignDecoration,
+                                  ast::StructMemberSizeDecoration>()) {
           TINT_ICE(Writer, diagnostics_)
-              << "invalid use of location decoration";
+              << "unhandled struct member attribute: " << deco->name();
         }
-      } else if (auto* interpolate = deco->As<ast::InterpolateDecoration>()) {
-        auto attr = interpolation_to_attribute(interpolate->type(),
-                                               interpolate->sampling());
-        if (attr.empty()) {
-          diagnostics_.add_error(diag::System::Writer,
-                                 "unknown interpolation attribute");
-          return false;
-        }
-        out << " [[" << attr << "]]";
-      } else if (deco->Is<ast::InvariantDecoration>()) {
-        out << " [[invariant]]";
-        has_invariant_ = true;
-      } else if (!deco->IsAnyOf<ast::StructMemberOffsetDecoration,
-                                ast::StructMemberAlignDecoration,
-                                ast::StructMemberSizeDecoration>()) {
-        TINT_ICE(Writer, diagnostics_)
-            << "unhandled struct member attribute: " << deco->name();
       }
     }
 
@@ -2204,9 +2312,9 @@ bool GeneratorImpl::EmitStructType(const sem::Struct* str) {
     add_padding(str->Size() - msl_offset, msl_offset);
   }
 
-  decrement_indent();
+  b->DecrementIndent();
 
-  line() << "};";
+  line(b) << "};";
   return true;
 }
 
@@ -2406,6 +2514,72 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
   return {};
 }
 
+template <typename F>
+bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out,
+                                        ast::CallExpression* call,
+                                        const sem::Intrinsic* intrinsic,
+                                        F&& build) {
+  // Generate the helper function if it hasn't been created already
+  auto fn = utils::GetOrCreate(intrinsics_, intrinsic, [&]() -> std::string {
+    TextBuffer b;
+    TINT_DEFER(helpers_.Append(b));
+
+    auto fn_name =
+        UniqueIdentifier(std::string("tint_") + sem::str(intrinsic->Type()));
+    std::vector<std::string> parameter_names;
+    {
+      auto decl = line(&b);
+      if (!EmitTypeAndName(decl, intrinsic->ReturnType(), fn_name)) {
+        return "";
+      }
+      {
+        ScopedParen sp(decl);
+        for (auto* param : intrinsic->Parameters()) {
+          if (!parameter_names.empty()) {
+            decl << ", ";
+          }
+          auto param_name = "param_" + std::to_string(parameter_names.size());
+          if (!EmitTypeAndName(decl, param->Type(), param_name)) {
+            return "";
+          }
+          parameter_names.emplace_back(std::move(param_name));
+        }
+      }
+      decl << " {";
+    }
+    {
+      ScopedIndent si(&b);
+      if (!build(&b, parameter_names)) {
+        return "";
+      }
+    }
+    line(&b) << "}";
+    line(&b);
+    return fn_name;
+  });
+
+  if (fn.empty()) {
+    return false;
+  }
+
+  // Call the helper
+  out << fn;
+  {
+    ScopedParen sp(out);
+    bool first = true;
+    for (auto* arg : call->params()) {
+      if (!first) {
+        out << ", ";
+      }
+      first = false;
+      if (!EmitExpression(out, arg)) {
+        return false;
+      }
+    }
+  }
+  return true;
+}
+
 }  // namespace msl
 }  // namespace writer
 }  // namespace tint

src/writer/msl/generator_impl.h

@@ -126,6 +126,22 @@ class GeneratorImpl : public TextGenerator {
   bool EmitTextureCall(std::ostream& out,
                        ast::CallExpression* expr,
                        const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `modf()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitModfCall(std::ostream& out,
+                    ast::CallExpression* expr,
+                    const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to the `frexp()` intrinsic
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitFrexpCall(std::ostream& out,
+                     ast::CallExpression* expr,
+                     const sem::Intrinsic* intrinsic);
   /// Handles a case statement
   /// @param stmt the statement
   /// @returns true if the statement was emitted successfully
@@ -218,10 +234,20 @@ class GeneratorImpl : public TextGenerator {
   /// @param out the output of the type stream
   /// @param type the type to generate
   /// @param name the name of the variable, only used for array emission
+  /// @param name_printed (optional) if not nullptr and an array was printed
   /// @returns true if the type is emitted
   bool EmitType(std::ostream& out,
                 const sem::Type* type,
-                const std::string& name);
+                const std::string& name,
+                bool* name_printed = nullptr);
+  /// Handles generating type and name
+  /// @param out the output stream
+  /// @param type the type to generate
+  /// @param name the name to emit
+  /// @returns true if the type is emitted
+  bool EmitTypeAndName(std::ostream& out,
+                       const sem::Type* type,
+                       const std::string& name);
   /// Handles generating a storage class
   /// @param out the output of the type stream
   /// @param sc the storage class to generate
@@ -238,9 +264,10 @@ class GeneratorImpl : public TextGenerator {
                       const sem::Type* type,
                       const std::string& name);
   /// Handles generating a struct declaration
+  /// @param buffer the text buffer that the type declaration will be written to
   /// @param str the struct to generate
   /// @returns true if the struct is emitted
-  bool EmitStructType(const sem::Struct* str);
+  bool EmitStructType(TextBuffer* buffer, const sem::Struct* str);
   /// Handles emitting a type constructor
   /// @param out the output of the expression stream
   /// @param expr the type constructor expression
@@ -291,6 +318,25 @@ class GeneratorImpl : public TextGenerator {
     uint32_t align;
   };
 
+  /// CallIntrinsicHelper will call the intrinsic helper function, creating it
+  /// if it hasn't been built already. If the intrinsic needs to be built then
+  /// CallIntrinsicHelper will generate the function signature and will call
+  /// `build` to emit the body of the function.
+  /// @param out the output of the expression stream
+  /// @param call the call expression
+  /// @param intrinsic the semantic information for the intrinsic
+  /// @param build a function with the signature:
+  ///        `bool(TextBuffer* buffer, const std::vector<std::string>& params)`
+  ///        Where:
+  ///          `buffer` is the body of the generated function
+  ///          `params` is the name of all the generated function parameters
+  /// @returns true if the call expression is emitted
+  template <typename F>
+  bool CallIntrinsicHelper(std::ostream& out,
+                           ast::CallExpression* call,
+                           const sem::Intrinsic* intrinsic,
+                           F&& build);
+
   TextBuffer helpers_;  // Helper functions emitted at the top of the output
 
   /// @returns the MSL packed type size and alignment in bytes for the given
@@ -308,6 +354,8 @@ class GeneratorImpl : public TextGenerator {
 
   /// True if an invariant attribute has been generated.
   bool has_invariant_ = false;
+
+  std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
 };
 
 }  // namespace msl

src/writer/msl/generator_impl_type_test.cc

@@ -219,9 +219,10 @@ TEST_F(MslGeneratorImplTest, EmitType_StructDecl) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(struct S {
+  EXPECT_EQ(buf.String(), R"(struct S {
   int a;
   float b;
 };
@@ -269,8 +270,9 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_NonComposites) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
 
   // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, NAME, SUFFIX)
   // for each field of the structure s.
@@ -320,7 +322,7 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_NonComposites) {
   "  /* " #ADDR " */ " #TYPE " " #NAME #SUFFIX ";\n"
   auto* expect = "struct S {\n" ALL_FIELDS() "};\n";
 #undef FIELD
-  EXPECT_EQ(gen.result(), expect);
+  EXPECT_EQ(buf.String(), expect);
 
   // 1.4 Metal and C++14
   // The Metal programming language is a C++14-based Specification with
@@ -378,8 +380,9 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_Structures) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
 
   // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, NAME, SUFFIX)
   // for each field of the structure s.
@@ -397,7 +400,7 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_Structures) {
   "  /* " #ADDR " */ " #TYPE " " #NAME #SUFFIX ";\n"
   auto* expect = "struct S {\n" ALL_FIELDS() "};\n";
 #undef FIELD
-  EXPECT_EQ(gen.result(), expect);
+  EXPECT_EQ(buf.String(), expect);
 
   // 1.4 Metal and C++14
   // The Metal programming language is a C++14-based Specification with
@@ -472,8 +475,9 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayDefaultStride) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
 
   // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, NAME, SUFFIX)
   // for each field of the structure s.
@@ -492,7 +496,7 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayDefaultStride) {
   "  /* " #ADDR " */ " #TYPE " " #NAME #SUFFIX ";\n"
   auto* expect = "struct S {\n" ALL_FIELDS() "};\n";
 #undef FIELD
-  EXPECT_EQ(gen.result(), expect);
+  EXPECT_EQ(buf.String(), expect);
 
   // 1.4 Metal and C++14
   // The Metal programming language is a C++14-based Specification with
@@ -557,8 +561,9 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
 
   // ALL_FIELDS() calls the macro FIELD(ADDR, TYPE, NAME, SUFFIX)
   // for each field of the structure s.
@@ -574,7 +579,7 @@ TEST_F(MslGeneratorImplTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
   "  /* " #ADDR " */ " #TYPE " " #NAME #SUFFIX ";\n"
   auto* expect = "struct S {\n" ALL_FIELDS() "};\n";
 #undef FIELD
-  EXPECT_EQ(gen.result(), expect);
+  EXPECT_EQ(buf.String(), expect);
 }
 
 TEST_F(MslGeneratorImplTest, AttemptTintPadSymbolCollision) {
@@ -619,9 +624,10 @@ TEST_F(MslGeneratorImplTest, AttemptTintPadSymbolCollision) {
 
   GeneratorImpl& gen = Build();
 
+  TextGenerator::TextBuffer buf;
   auto* sem_s = program->TypeOf(s)->As<sem::Struct>();
-  ASSERT_TRUE(gen.EmitStructType(sem_s)) << gen.error();
+  ASSERT_TRUE(gen.EmitStructType(&buf, sem_s)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(struct S {
+  EXPECT_EQ(buf.String(), R"(struct S {
   /* 0x0000 */ int tint_pad_2;
   /* 0x0004 */ int8_t tint_pad_10[124];
   /* 0x0080 */ float tint_pad_20;

src/writer/spirv/builder.cc

@@ -162,7 +162,8 @@ uint32_t intrinsic_to_glsl_method(const sem::Intrinsic* intrinsic) {
     case IntrinsicType::kFract:
       return GLSLstd450Fract;
     case IntrinsicType::kFrexp:
-      return GLSLstd450Frexp;
+      return (intrinsic->Parameters().size() == 1) ? GLSLstd450FrexpStruct
+                                                   : GLSLstd450Frexp;
     case IntrinsicType::kInverseSqrt:
       return GLSLstd450InverseSqrt;
     case IntrinsicType::kLdexp:
@@ -192,7 +193,8 @@ uint32_t intrinsic_to_glsl_method(const sem::Intrinsic* intrinsic) {
     case IntrinsicType::kMix:
       return GLSLstd450FMix;
     case IntrinsicType::kModf:
-      return GLSLstd450Modf;
+      return (intrinsic->Parameters().size() == 1) ? GLSLstd450ModfStruct
+                                                   : GLSLstd450Modf;
     case IntrinsicType::kNormalize:
       return GLSLstd450Normalize;
     case IntrinsicType::kPack4x8snorm:
@@ -3932,25 +3934,25 @@ bool Builder::GenerateReferenceType(const sem::Reference* ref,
 bool Builder::GenerateStructType(const sem::Struct* struct_type,
                                  const Operand& result) {
   auto struct_id = result.to_i();
-  auto* impl = struct_type->Declaration();
 
-  if (impl->name().IsValid()) {
+  if (struct_type->Name().IsValid()) {
-    push_debug(spv::Op::OpName,
+    push_debug(
-               {Operand::Int(struct_id),
+        spv::Op::OpName,
-                Operand::String(builder_.Symbols().NameFor(impl->name()))});
+        {Operand::Int(struct_id),
+         Operand::String(builder_.Symbols().NameFor(struct_type->Name()))});
   }
 
   OperandList ops;
   ops.push_back(result);
 
-  if (impl->IsBlockDecorated()) {
+  auto* decl = struct_type->Declaration();
+  if (decl && decl->IsBlockDecorated()) {
     push_annot(spv::Op::OpDecorate,
                {Operand::Int(struct_id), Operand::Int(SpvDecorationBlock)});
   }
 
-  auto& members = impl->members();
+  for (uint32_t i = 0; i < struct_type->Members().size(); ++i) {
-  for (uint32_t i = 0; i < members.size(); ++i) {
+    auto mem_id = GenerateStructMember(struct_id, i, struct_type->Members()[i]);
-    auto mem_id = GenerateStructMember(struct_id, i, members[i]);
     if (mem_id == 0) {
       return false;
     }
@@ -3964,10 +3966,10 @@ bool Builder::GenerateStructType(const sem::Struct* struct_type,
 
 uint32_t Builder::GenerateStructMember(uint32_t struct_id,
                                        uint32_t idx,
-                                       ast::StructMember* member) {
+                                       const sem::StructMember* member) {
   push_debug(spv::Op::OpMemberName,
              {Operand::Int(struct_id), Operand::Int(idx),
-              Operand::String(builder_.Symbols().NameFor(member->symbol()))});
+              Operand::String(builder_.Symbols().NameFor(member->Name()))});
 
   // Note: This will generate layout annotations for *all* structs, whether or
   // not they are used in host-shareable variables. This is officially ok in
@@ -3975,19 +3977,13 @@ uint32_t Builder::GenerateStructMember(uint32_t struct_id,
   // to only generate the layout info for structs used for certain storage
   // classes.
 
-  auto* sem_member = builder_.Sem().Get(member);
-  if (!sem_member) {
-    error_ = "Struct member has no semantic information";
-    return 0;
-  }
-
   push_annot(
       spv::Op::OpMemberDecorate,
       {Operand::Int(struct_id), Operand::Int(idx),
-       Operand::Int(SpvDecorationOffset), Operand::Int(sem_member->Offset())});
+       Operand::Int(SpvDecorationOffset), Operand::Int(member->Offset())});
 
   // Infer and emit matrix layout.
-  auto* matrix_type = GetNestedMatrixType(sem_member->Type());
+  auto* matrix_type = GetNestedMatrixType(member->Type());
   if (matrix_type) {
     push_annot(spv::Op::OpMemberDecorate,
                {Operand::Int(struct_id), Operand::Int(idx),
@@ -4004,7 +4000,7 @@ uint32_t Builder::GenerateStructMember(uint32_t struct_id,
                 Operand::Int(effective_row_count * scalar_elem_size)});
   }
 
-  return GenerateTypeIfNeeded(sem_member->Type());
+  return GenerateTypeIfNeeded(member->Type());
 }
 
 bool Builder::GenerateVectorType(const sem::Vector* vec,

src/writer/spirv/builder.h

@@ -486,7 +486,7 @@ class Builder {
   /// @returns the id of the struct member or 0 on error.
   uint32_t GenerateStructMember(uint32_t struct_id,
                                 uint32_t idx,
-                                ast::StructMember* member);
+                                const sem::StructMember* member);
   /// Generates a variable declaration statement
   /// @param stmt the statement to generate
   /// @returns true on successfull generation

src/writer/text_generator.cc

@@ -17,6 +17,8 @@
 #include <algorithm>
 #include <limits>
 
+#include "src/utils/get_or_create.h"
+
 namespace tint {
 namespace writer {
 
@@ -29,6 +31,15 @@ std::string TextGenerator::UniqueIdentifier(const std::string& prefix) {
   return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
 }
 
+std::string TextGenerator::StructName(const sem::Struct* s) {
+  auto name = builder_.Symbols().NameFor(s->Name());
+  if (name.size() > 0 && name[0] == '_') {
+    name = utils::GetOrCreate(builtin_struct_names_, s,
+                              [&] { return UniqueIdentifier(name.substr(1)); });
+  }
+  return name;
+}
+
 std::string TextGenerator::TrimSuffix(std::string str,
                                       const std::string& suffix) {
   if (str.size() >= suffix.size()) {

src/writer/text_generator.h

@@ -17,6 +17,7 @@
 
 #include <sstream>
 #include <string>
+#include <unordered_map>
 #include <utility>
 #include <vector>
 
@@ -29,37 +30,6 @@ namespace writer {
 /// Helper methods for generators which are creating text output
 class TextGenerator {
  public:
-  /// Constructor
-  /// @param program the program used by the generator
-  explicit TextGenerator(const Program* program);
-  ~TextGenerator();
-
-  /// Increment the emitter indent level
-  void increment_indent() { current_buffer_->IncrementIndent(); }
-  /// Decrement the emitter indent level
-  void decrement_indent() { current_buffer_->DecrementIndent(); }
-
-  /// @returns the result data
-  std::string result() const { return main_buffer_.String(); }
-
-  /// @returns the list of diagnostics raised by the generator.
-  const diag::List& Diagnostics() const { return diagnostics_; }
-
-  /// @returns the error
-  std::string error() const { return diagnostics_.str(); }
-
-  /// @return a new, unique identifier with the given prefix.
-  /// @param prefix optional prefix to apply to the generated identifier. If
-  /// empty "tint_symbol" will be used.
-  std::string UniqueIdentifier(const std::string& prefix = "");
-
-  /// @param str the string
-  /// @param suffix the suffix to remove
-  /// @return returns str without the provided trailing suffix string. If str
-  /// doesn't end with suffix, str is returned unchanged.
-  std::string TrimSuffix(std::string str, const std::string& suffix);
-
- protected:
   /// Line holds a single line of text
   struct Line {
     /// The indentation of the line in whitespaces
@@ -117,6 +87,44 @@ class TextGenerator {
     std::vector<Line> lines;
   };
 
+  /// Constructor
+  /// @param program the program used by the generator
+  explicit TextGenerator(const Program* program);
+  ~TextGenerator();
+
+  /// Increment the emitter indent level
+  void increment_indent() { current_buffer_->IncrementIndent(); }
+  /// Decrement the emitter indent level
+  void decrement_indent() { current_buffer_->DecrementIndent(); }
+
+  /// @returns the result data
+  std::string result() const { return main_buffer_.String(); }
+
+  /// @returns the list of diagnostics raised by the generator.
+  const diag::List& Diagnostics() const { return diagnostics_; }
+
+  /// @returns the error
+  std::string error() const { return diagnostics_.str(); }
+
+  /// @return a new, unique identifier with the given prefix.
+  /// @param prefix optional prefix to apply to the generated identifier. If
+  /// empty "tint_symbol" will be used.
+  std::string UniqueIdentifier(const std::string& prefix = "");
+
+  /// @param s the semantic structure
+  /// @returns the name of the structure, taking special care of builtin
+  /// structures that start with a leading underscore. If the structure is a
+  /// builtin, then the returned name will be a unique name without the leading
+  /// underscore.
+  std::string StructName(const sem::Struct* s);
+
+  /// @param str the string
+  /// @param suffix the suffix to remove
+  /// @return returns str without the provided trailing suffix string. If str
+  /// doesn't end with suffix, str is returned unchanged.
+  std::string TrimSuffix(std::string str, const std::string& suffix);
+
+ protected:
   /// LineWriter is a helper that acts as a string buffer, who's content is
   /// emitted to the TextBuffer as a single line on destruction.
   struct LineWriter {
@@ -224,6 +232,8 @@ class TextGenerator {
  private:
   /// The primary text buffer that the generator will emit
   TextBuffer main_buffer_;
+  /// Map of builtin structure to unique generated name
+  std::unordered_map<const sem::Struct*, std::string> builtin_struct_names_;
 };
 
 }  // namespace writer

test/intrinsics/frexp.wgsl

@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-    var exponent : i32;
+    let res = frexp(1.23);
-    let significand : f32 = frexp(1.23, &exponent);
+    let exp : i32 = res.exp;
+    let sig : f32 = res.sig;
 }

test/intrinsics/frexp.wgsl.expected.hlsl

@@ -1,13 +1,18 @@
-float tint_frexp(float param_0, inout int param_1) {
+struct frexp_result {
-  float float_exp;
+  float sig;
-  float significand = frexp(param_0, float_exp);
+  int exp;
-  param_1 = int(float_exp);
+};
-  return significand;
+frexp_result tint_frexp(float param_0) {
+  float exp;
+  float sig = frexp(param_0, exp);
+  frexp_result result = {sig, int(exp)};
+  return result;
 }
 
 [numthreads(1, 1, 1)]
 void main() {
-  int exponent = 0;
+  const frexp_result res = tint_frexp(1.230000019f);
-  const float significand = tint_frexp(1.230000019f, exponent);
+  const int exp = res.exp;
+  const float sig = res.sig;
   return;
 }

test/intrinsics/frexp.wgsl.expected.msl

@@ -1,9 +1,21 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
 kernel void tint_symbol() {
-  int exponent = 0;
+  frexp_result const res = tint_frexp(1.230000019f);
-  float const significand = frexp(1.230000019f, *(&(exponent)));
+  int const exp = res.exp;
+  float const sig = res.sig;
   return;
 }
 

test/intrinsics/frexp.wgsl.expected.spvasm

@@ -1,25 +1,29 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 14
+; Bound: 13
 ; Schema: 0
                OpCapability Shader
-         %11 = OpExtInstImport "GLSL.std.450"
+          %9 = OpExtInstImport "GLSL.std.450"
                OpMemoryModel Logical GLSL450
                OpEntryPoint GLCompute %main "main"
                OpExecutionMode %main LocalSize 1 1 1
                OpName %main "main"
-               OpName %exponent "exponent"
+               OpName %_frexp_result "_frexp_result"
+               OpMemberName %_frexp_result 0 "sig"
+               OpMemberName %_frexp_result 1 "exp"
+               OpMemberDecorate %_frexp_result 0 Offset 0
+               OpMemberDecorate %_frexp_result 1 Offset 4
        %void = OpTypeVoid
           %1 = OpTypeFunction %void
-        %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
-          %8 = OpConstantNull %int
       %float = OpTypeFloat 32
+        %int = OpTypeInt 32 1
+%_frexp_result = OpTypeStruct %float %int
 %float_1_23000002 = OpConstant %float 1.23000002
        %main = OpFunction %void None %1
           %4 = OpLabel
-   %exponent = OpVariable %_ptr_Function_int Function %8
+          %5 = OpExtInst %_frexp_result %9 FrexpStruct %float_1_23000002
-          %9 = OpExtInst %float %11 Frexp %float_1_23000002 %exponent
+         %11 = OpCompositeExtract %int %5 1
+         %12 = OpCompositeExtract %float %5 0
                OpReturn
                OpFunctionEnd

test/intrinsics/frexp.wgsl.expected.wgsl

@@ -1,5 +1,6 @@
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var exponent : i32;
+  let res = frexp(1.230000019);
-  let significand : f32 = frexp(1.230000019, &(exponent));
+  let exp : i32 = res.exp;
+  let sig : f32 = res.sig;
 }

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_013caa() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/013caa.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_013caa() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);

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

@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, threadgroup int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_0da285(threadgroup int* const tint_symbol_1) {
-  float res = frexp(1.0f, *(&(*(tint_symbol_1))));
+  float res = tint_frexp(1.0f, &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@ kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgro
   return;
 }
 
-Compilation failed: 
-
-program_source:5:15: error: no matching function for call to 'frexp'
-  float res = frexp(1.0f, *(&(*(tint_symbol_1))));
-              ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int'), parameter type must be 'int &'
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int'), parameter type must be 'int &'
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int2 &' (aka 'int2 &') for 2nd argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int3 &' (aka 'int3 &') for 2nd argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int4 &' (aka 'int4 &') for 2nd argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int2 &' (aka 'int2 &') for 2nd argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int3 &' (aka 'int3 &') for 2nd argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'threadgroup int' to 'metal::int4 &' (aka 'int4 &') for 2nd argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/0da285.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 var<workgroup> arg_1 : i32;
 
 fn frexp_0da285() {

test/intrinsics/gen/frexp/12f1da.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(f32) -> _frexp_result
+fn frexp_12f1da() {
+  var res = frexp(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_12f1da();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_12f1da();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_12f1da();
+}

test/intrinsics/gen/frexp/12f1da.wgsl.expected.hlsl

@@ -0,0 +1,35 @@
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  float exp;
+  float sig = frexp(param_0, exp);
+  frexp_result result = {sig, int(exp)};
+  return result;
+}
+
+void frexp_12f1da() {
+  frexp_result res = tint_frexp(1.0f);
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_12f1da();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_12f1da();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_12f1da();
+  return;
+}

test/intrinsics/gen/frexp/12f1da.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result {
+  float sig;
+  int exp;
+};
+frexp_result tint_frexp(float param_0) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_12f1da() {
+  frexp_result res = tint_frexp(1.0f);
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_12f1da();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_12f1da();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_12f1da();
+  return;
+}
+

test/intrinsics/gen/frexp/12f1da.wgsl.expected.spvasm

@@ -0,0 +1,75 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 35
+; Schema: 0
+               OpCapability Shader
+         %16 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_12f1da "frexp_12f1da"
+               OpName %_frexp_result "_frexp_result"
+               OpMemberName %_frexp_result 0 "sig"
+               OpMemberName %_frexp_result 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result 0 Offset 0
+               OpMemberDecorate %_frexp_result 1 Offset 4
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_frexp_result = OpTypeStruct %float %int
+    %float_1 = OpConstant %float 1
+%_ptr_Function__frexp_result = OpTypePointer Function %_frexp_result
+         %20 = OpConstantNull %_frexp_result
+         %21 = OpTypeFunction %void %v4float
+%frexp_12f1da = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result Function %20
+         %13 = OpExtInst %_frexp_result %16 FrexpStruct %float_1
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %21
+%tint_symbol = OpFunctionParameter %v4float
+         %24 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %26 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %27 = OpFunctionCall %void %frexp_12f1da
+         %28 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %30 = OpLabel
+         %31 = OpFunctionCall %void %frexp_12f1da
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %frexp_12f1da
+               OpReturn
+               OpFunctionEnd

test/intrinsics/gen/frexp/12f1da.wgsl.expected.wgsl

@@ -0,0 +1,19 @@
+fn frexp_12f1da() {
+  var res = frexp(1.0);
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_12f1da();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_12f1da();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_12f1da();
+}

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_15edf3() {
   int2 arg_1 = 0;
-  float2 res = frexp(float2(), *(&(arg_1)));
+  float2 res = tint_frexp(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/15edf3.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_15edf3() {
   var arg_1 : vec2<i32>;
   var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_19ab15() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/19ab15.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_19ab15() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_2052e9() {
   int4 arg_1 = 0;
-  float4 res = frexp(float4(), *(&(arg_1)));
+  float4 res = tint_frexp(float4(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/2052e9.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_2052e9() {
   var arg_1 : vec4<i32>;
   var res : vec4<f32> = frexp(vec4<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);

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

@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, threadgroup int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_40fc9b(threadgroup int3* const tint_symbol_1) {
-  float3 res = frexp(float3(), *(&(*(tint_symbol_1))));
+  float3 res = tint_frexp(float3(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@ kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgro
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float3 res = frexp(float3(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int3' (vector of 3 'int' values)), parameter type must be 'metal::int3 &' (aka 'int3 &')
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'float3' (vector of 3 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/40fc9b.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec3<i32>;
 
 fn frexp_40fc9b() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_41e931() {
   int arg_1 = 0;
-  float res = frexp(1.0f, *(&(arg_1)));
+  float res = tint_frexp(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/41e931.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 fn frexp_41e931() {
   var arg_1 : i32;
   var res : f32 = frexp(1.0, &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_481e59() {
   int arg_1 = 0;
-  float res = frexp(1.0f, *(&(arg_1)));
+  float res = tint_frexp(1.0f, &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/481e59.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 fn frexp_481e59() {
   var arg_1 : i32;
   var res : f32 = frexp(1.0, &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_5a141e() {
   int3 arg_1 = 0;
-  float3 res = frexp(float3(), *(&(arg_1)));
+  float3 res = tint_frexp(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/5a141e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_5a141e() {
   var arg_1 : vec3<i32>;
   var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_6d0058() {
   int3 arg_1 = 0;
-  float3 res = frexp(float3(), *(&(arg_1)));
+  float3 res = tint_frexp(float3(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6d0058.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_6d0058() {
   var arg_1 : vec3<i32>;
   var res : vec3<f32> = frexp(vec3<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 float3 tint_frexp(float3 param_0, inout int3 param_1) {
   float3 float_exp;
   float3 significand = frexp(param_0, float_exp);

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

@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float3 tint_frexp(float3 param_0, thread int3* param_1) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_6efa09(thread int3* const tint_symbol_2) {
-  float3 res = frexp(float3(), *(&(*(tint_symbol_2))));
+  float3 res = tint_frexp(float3(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/6efa09.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec3<f32> = frexp(vec3<f32>(), &arg_1);
+                       ^^^^^
+
 var<private> arg_1 : vec3<i32>;
 
 fn frexp_6efa09() {

test/intrinsics/gen/frexp/a0eb3b.wgsl

@@ -0,0 +1,45 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/intrinsic-gen
+// using the template:
+//   test/intrinsics/intrinsics.wgsl.tmpl
+// and the intrinsic defintion file:
+//   src/intrinsics.def
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+// fn frexp(vec<3, f32>) -> _frexp_result_vec<3>
+fn frexp_a0eb3b() {
+  var res = frexp(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_a0eb3b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_a0eb3b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_a0eb3b();
+}

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

@@ -0,0 +1,35 @@
+struct frexp_result_vec3 {
+  float3 sig;
+  int3 exp;
+};
+frexp_result_vec3 tint_frexp(float3 param_0) {
+  float3 exp;
+  float3 sig = frexp(param_0, exp);
+  frexp_result_vec3 result = {sig, int3(exp)};
+  return result;
+}
+
+void frexp_a0eb3b() {
+  frexp_result_vec3 res = tint_frexp(float3(0.0f, 0.0f, 0.0f));
+}
+
+struct tint_symbol {
+  float4 value : SV_Position;
+};
+
+tint_symbol vertex_main() {
+  frexp_a0eb3b();
+  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
+  return tint_symbol_1;
+}
+
+void fragment_main() {
+  frexp_a0eb3b();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  frexp_a0eb3b();
+  return;
+}

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

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct frexp_result_vec3 {
+  float3 sig;
+  int3 exp;
+};
+frexp_result_vec3 tint_frexp(float3 param_0) {
+  int3 exp;
+  float3 sig = frexp(param_0, exp);
+  return {sig, exp};
+}
+
+struct tint_symbol {
+  float4 value [[position]];
+};
+
+void frexp_a0eb3b() {
+  frexp_result_vec3 res = tint_frexp(float3());
+}
+
+vertex tint_symbol vertex_main() {
+  frexp_a0eb3b();
+  tint_symbol const tint_symbol_1 = {.value=float4()};
+  return tint_symbol_1;
+}
+
+fragment void fragment_main() {
+  frexp_a0eb3b();
+  return;
+}
+
+kernel void compute_main() {
+  frexp_a0eb3b();
+  return;
+}
+

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

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+         %18 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vertex_main "vertex_main" %tint_pointsize %tint_symbol_1
+               OpEntryPoint Fragment %fragment_main "fragment_main"
+               OpEntryPoint GLCompute %compute_main "compute_main"
+               OpExecutionMode %fragment_main OriginUpperLeft
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %tint_pointsize "tint_pointsize"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %frexp_a0eb3b "frexp_a0eb3b"
+               OpName %_frexp_result_vec3 "_frexp_result_vec3"
+               OpMemberName %_frexp_result_vec3 0 "sig"
+               OpMemberName %_frexp_result_vec3 1 "exp"
+               OpName %res "res"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol "tint_symbol"
+               OpName %vertex_main "vertex_main"
+               OpName %fragment_main "fragment_main"
+               OpName %compute_main "compute_main"
+               OpDecorate %tint_pointsize BuiltIn PointSize
+               OpDecorate %tint_symbol_1 BuiltIn Position
+               OpMemberDecorate %_frexp_result_vec3 0 Offset 0
+               OpMemberDecorate %_frexp_result_vec3 1 Offset 16
+      %float = OpTypeFloat 32
+%_ptr_Output_float = OpTypePointer Output %float
+          %4 = OpConstantNull %float
+%tint_pointsize = OpVariable %_ptr_Output_float Output %4
+    %v4float = OpTypeVector %float 4
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %8 = OpConstantNull %v4float
+%tint_symbol_1 = OpVariable %_ptr_Output_v4float Output %8
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+    %v3float = OpTypeVector %float 3
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_frexp_result_vec3 = OpTypeStruct %v3float %v3int
+         %19 = OpConstantNull %v3float
+%_ptr_Function__frexp_result_vec3 = OpTypePointer Function %_frexp_result_vec3
+         %22 = OpConstantNull %_frexp_result_vec3
+         %23 = OpTypeFunction %void %v4float
+    %float_1 = OpConstant %float 1
+%frexp_a0eb3b = OpFunction %void None %9
+         %12 = OpLabel
+        %res = OpVariable %_ptr_Function__frexp_result_vec3 Function %22
+         %13 = OpExtInst %_frexp_result_vec3 %18 FrexpStruct %19
+               OpStore %res %13
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_2 = OpFunction %void None %23
+%tint_symbol = OpFunctionParameter %v4float
+         %26 = OpLabel
+               OpStore %tint_symbol_1 %tint_symbol
+               OpReturn
+               OpFunctionEnd
+%vertex_main = OpFunction %void None %9
+         %28 = OpLabel
+               OpStore %tint_pointsize %float_1
+         %30 = OpFunctionCall %void %frexp_a0eb3b
+         %31 = OpFunctionCall %void %tint_symbol_2 %8
+               OpReturn
+               OpFunctionEnd
+%fragment_main = OpFunction %void None %9
+         %33 = OpLabel
+         %34 = OpFunctionCall %void %frexp_a0eb3b
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %9
+         %36 = OpLabel
+         %37 = OpFunctionCall %void %frexp_a0eb3b
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,19 @@
+fn frexp_a0eb3b() {
+  var res = frexp(vec3<f32>());
+}
+
+[[stage(vertex)]]
+fn vertex_main() -> [[builtin(position)]] vec4<f32> {
+  frexp_a0eb3b();
+  return vec4<f32>();
+}
+
+[[stage(fragment)]]
+fn fragment_main() {
+  frexp_a0eb3b();
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn compute_main() {
+  frexp_a0eb3b();
+}

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 float tint_frexp(float param_0, inout int param_1) {
   float float_exp;
   float significand = frexp(param_0, float_exp);

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

@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float tint_frexp(float param_0, thread int* param_1) {
+  int exp;
+  float sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_a2a617(thread int* const tint_symbol_2) {
-  float res = frexp(1.0f, *(&(*(tint_symbol_2))));
+  float res = tint_frexp(1.0f, &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a2a617.wgsl:29:18 warning: use of deprecated intrinsic
+  var res: f32 = frexp(1.0, &arg_1);
+                 ^^^^^
+
 var<private> arg_1 : i32;
 
 fn frexp_a2a617() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);

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

@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, threadgroup int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_a3f940(threadgroup int2* const tint_symbol_1) {
-  float2 res = frexp(float2(), *(&(*(tint_symbol_1))));
+  float2 res = tint_frexp(float2(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@ kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgro
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float2 res = frexp(float2(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int2' (vector of 2 'int' values)), parameter type must be 'metal::int2 &' (aka 'int2 &')
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: no known conversion from 'float2' (vector of 2 'float' values) to 'metal::float4' (aka 'float4') for 1st argument
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a3f940.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec2<i32>;
 
 fn frexp_a3f940() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 float2 tint_frexp(float2 param_0, inout int2 param_1) {
   float2 float_exp;
   float2 significand = frexp(param_0, float_exp);

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

@@ -1,13 +1,25 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float2 tint_frexp(float2 param_0, thread int2* param_1) {
+  int2 exp;
+  float2 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_a951b5() {
   int2 arg_1 = 0;
-  float2 res = frexp(float2(), *(&(arg_1)));
+  float2 res = tint_frexp(float2(), &(arg_1));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/a951b5.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec2<f32> = frexp(vec2<f32>(), &arg_1);
+                       ^^^^^
+
 fn frexp_a951b5() {
   var arg_1 : vec2<i32>;
   var res : vec2<f32> = frexp(vec2<f32>(), &(arg_1));

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);

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

@@ -1,12 +1,24 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, thread int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 struct tint_symbol {
   float4 value [[position]];
 };
 
 void frexp_b45525(thread int4* const tint_symbol_2) {
-  float4 res = frexp(float4(), *(&(*(tint_symbol_2))));
+  float4 res = tint_frexp(float4(), &(*(tint_symbol_2)));
 }
 
 vertex tint_symbol vertex_main() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b45525.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 var<private> arg_1 : vec4<i32>;
 
 fn frexp_b45525() {

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 float4 tint_frexp(float4 param_0, inout int4 param_1) {
   float4 float_exp;
   float4 significand = frexp(param_0, float_exp);

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

@@ -1,10 +1,20 @@
-SKIP: FAILED
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
 
 #include <metal_stdlib>
 
 using namespace metal;
+
+float4 tint_frexp(float4 param_0, threadgroup int4* param_1) {
+  int4 exp;
+  float4 sig = frexp(param_0, exp);
+  *param_1 = exp;
+  return sig;
+}
+
 void frexp_b87f4e(threadgroup int4* const tint_symbol_1) {
-  float4 res = frexp(float4(), *(&(*(tint_symbol_1))));
+  float4 res = tint_frexp(float4(), &(*(tint_symbol_1)));
 }
 
 kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
@@ -17,44 +27,3 @@ kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgro
   return;
 }
 
-Compilation failed: 
-
-program_source:5:16: error: no matching function for call to 'frexp'
-  float4 res = frexp(float4(), *(&(*(tint_symbol_1))));
-               ^~~~~
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5292:19: note: candidate function not viable: address space mismatch in 2nd argument ('threadgroup int4' (vector of 4 'int' values)), parameter type must be 'metal::int4 &' (aka 'int4 &')
-METAL_FUNC float4 frexp(float4 x, thread int4 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3092:17: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'half' for 1st argument
-METAL_FUNC half frexp(half x, thread int &exp)
-                ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3336:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half2' (aka 'half2') for 1st argument
-METAL_FUNC half2 frexp(half2 x, thread int2 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3580:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half3' (aka 'half3') for 1st argument
-METAL_FUNC half3 frexp(half3 x, thread int3 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:3824:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::half4' (aka 'half4') for 1st argument
-METAL_FUNC half4 frexp(half4 x, thread int4 &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4560:18: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'float' for 1st argument
-METAL_FUNC float frexp(float x, thread int &exp)
-                 ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:4804:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float2' (aka 'float2') for 1st argument
-METAL_FUNC float2 frexp(float2 x, thread int2 &exp)
-                  ^
-/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_math:5048:19: note: candidate function not viable: no known conversion from 'float4' (vector of 4 'float' values) to 'metal::float3' (aka 'float3') for 1st argument
-METAL_FUNC float3 frexp(float3 x, thread int3 &exp)
-                  ^
-program_source:10:31: warning: equality comparison with extraneous parentheses
-  if ((local_invocation_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:31: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((local_invocation_index == 0u)) {
-      ~                       ^    ~
-program_source:10:31: note: use '=' to turn this equality comparison into an assignment
-  if ((local_invocation_index == 0u)) {
-                              ^~
-                              =
-
-

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0

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

@@ -1,3 +1,7 @@
+intrinsics/gen/frexp/b87f4e.wgsl:29:24 warning: use of deprecated intrinsic
+  var res: vec4<f32> = frexp(vec4<f32>(), &arg_1);
+                       ^^^^^
+
 var<workgroup> arg_1 : vec4<i32>;
 
 fn frexp_b87f4e() {