Add semantic::Call, IntrinsicCall, TextureIntrinsicCall, use them.

semantic::Call derives from semantic::Expression, and Type() is the return type of the function
Pull the mutable semantic field from ast::Identifier and into a new semantic nodes.
Have the TypeDeterminer create these new semantic nodes.

Note: This change also fixes the node that holds the semantic information for a call.
Previously this was on the identifier, and this is now correctly on the CallExpression.
The identifier of the CallExpression should resolve to the target function, not the return type.
Functions can currently be represented as a type, and the identifier of a CallExpression now has no semantic information.

Bug: tint:390
Change-Id: I03521da5634815d35022f45ba521372cbbdb6bc7
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/40065
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>

BUILD.gn

@@ -286,8 +286,6 @@ source_set("libtint_core_src") {
     "src/ast/if_statement.h",
     "src/ast/int_literal.cc",
     "src/ast/int_literal.h",
-    "src/ast/intrinsic.cc",
-    "src/ast/intrinsic.h",
     "src/ast/literal.cc",
     "src/ast/literal.h",
     "src/ast/location_decoration.cc",
@@ -378,9 +376,13 @@ source_set("libtint_core_src") {
     "src/reader/reader.cc",
     "src/reader/reader.h",
     "src/scope_stack.h",
+    "src/semantic/call.h",
     "src/semantic/expression.h",
     "src/semantic/info.h",
+    "src/semantic/intrinsic.cc",
+    "src/semantic/intrinsic.h",
     "src/semantic/node.h",
+    "src/semantic/sem_call.cc",
     "src/semantic/sem_expression.cc",
     "src/semantic/sem_function.cc",
     "src/semantic/sem_info.cc",

src/CMakeLists.txt

@@ -100,8 +100,6 @@ set(TINT_LIB_SRCS
   ast/if_statement.h
   ast/int_literal.cc
   ast/int_literal.h
-  ast/intrinsic.cc
-  ast/intrinsic.h
   ast/literal.cc
   ast/literal.h
   ast/location_decoration.cc
@@ -192,9 +190,13 @@ set(TINT_LIB_SRCS
   reader/reader.cc
   reader/reader.h
   scope_stack.h
+  semantic/call.h
   semantic/expression.h
   semantic/info.h
+  semantic/intrinsic.cc
+  semantic/intrinsic.h
   semantic/node.h
+  semantic/sem_call.cc
   semantic/sem_expression.cc
   semantic/sem_function.cc
   semantic/sem_info.cc

src/ast/identifier_expression.h

@@ -19,7 +19,7 @@
 #include <utility>
 
 #include "src/ast/expression.h"
-#include "src/ast/intrinsic.h"
+#include "src/semantic/intrinsic.h"
 #include "src/symbol.h"
 
 namespace tint {
@@ -39,25 +39,6 @@ class IdentifierExpression : public Castable<IdentifierExpression, Expression> {
   /// @returns the symbol for the identifier
   Symbol symbol() const { return sym_; }
 
-  /// Sets the intrinsic for this identifier
-  /// @param i the intrinsic to set
-  void set_intrinsic(Intrinsic i) { intrinsic_ = i; }
-  /// @returns the intrinsic this identifier represents
-  Intrinsic intrinsic() const { return intrinsic_; }
-
-  /// Sets the intrinsic signature
-  /// @param s the intrinsic signature to set
-  void set_intrinsic_signature(std::unique_ptr<intrinsic::Signature> s) {
-    intrinsic_sig_ = std::move(s);
-  }
-  /// @returns the intrinsic signature for this identifier.
-  const intrinsic::Signature* intrinsic_signature() const {
-    return intrinsic_sig_.get();
-  }
-
-  /// @returns true if this identifier is for an intrinsic
-  bool IsIntrinsic() const { return intrinsic_ != Intrinsic::kNone; }
-
   /// Sets the identifier as a swizzle
   void SetIsSwizzle() { is_swizzle_ = true; }
 
@@ -88,8 +69,6 @@ class IdentifierExpression : public Castable<IdentifierExpression, Expression> {
 
   Symbol const sym_;
 
-  Intrinsic intrinsic_ = Intrinsic::kNone;               // Semantic info
-  std::unique_ptr<intrinsic::Signature> intrinsic_sig_;  // Semantic info
   bool is_swizzle_ = false;  // Semantic info
 };
 

src/program.cc

@@ -104,7 +104,7 @@ bool Program::IsValid() const {
   return is_valid_;
 }
 
-type::Type* Program::TypeOf(ast::Expression* expr) const {
+type::Type* Program::TypeOf(const ast::Expression* expr) const {
   auto* sem = Sem().Get(expr);
   return sem ? sem->Type() : nullptr;
 }

src/program.h

@@ -119,7 +119,7 @@ class Program {
   /// @param expr the AST expression
   /// @return the resolved semantic type for the expression, or nullptr if the
   /// expression has no resolved type.
-  type::Type* TypeOf(ast::Expression* expr) const;
+  type::Type* TypeOf(const ast::Expression* expr) const;
 
   /// @param demangle whether to automatically demangle the symbols in the
   /// returned string

src/reader/spirv/function.cc

@@ -44,7 +44,6 @@
 #include "src/ast/float_literal.h"
 #include "src/ast/identifier_expression.h"
 #include "src/ast/if_statement.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/loop_statement.h"
 #include "src/ast/member_accessor_expression.h"
 #include "src/ast/return_statement.h"
@@ -62,6 +61,7 @@
 #include "src/reader/spirv/construct.h"
 #include "src/reader/spirv/fail_stream.h"
 #include "src/reader/spirv/parser_impl.h"
+#include "src/semantic/intrinsic.h"
 #include "src/type/bool_type.h"
 #include "src/type/depth_texture_type.h"
 #include "src/type/f32_type.h"
@@ -464,20 +464,20 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
   return "";
 }
 
-// Returns the WGSL standard library function instrinsic for the
+// Returns the WGSL standard library function intrinsic for the
-// given instruction, or ast::Intrinsic::kNone
+// given instruction, or semantic::Intrinsic::kNone
-ast::Intrinsic GetIntrinsic(SpvOp opcode) {
+semantic::Intrinsic GetIntrinsic(SpvOp opcode) {
   switch (opcode) {
     case SpvOpBitCount:
-      return ast::Intrinsic::kCountOneBits;
+      return semantic::Intrinsic::kCountOneBits;
     case SpvOpBitReverse:
-      return ast::Intrinsic::kReverseBits;
+      return semantic::Intrinsic::kReverseBits;
     case SpvOpDot:
-      return ast::Intrinsic::kDot;
+      return semantic::Intrinsic::kDot;
     default:
       break;
   }
-  return ast::Intrinsic::kNone;
+  return semantic::Intrinsic::kNone;
 }
 
 // @param opcode a SPIR-V opcode
@@ -3102,7 +3102,7 @@ TypedExpression FunctionEmitter::MaybeEmitCombinatorialValue(
   }
 
   const auto intrinsic = GetIntrinsic(opcode);
-  if (intrinsic != ast::Intrinsic::kNone) {
+  if (intrinsic != semantic::Intrinsic::kNone) {
     return MakeIntrinsicCall(inst);
   }
 
@@ -3990,7 +3990,6 @@ TypedExpression FunctionEmitter::MakeIntrinsicCall(
   auto name = ss.str();
   auto* ident = create<ast::IdentifierExpression>(
       Source{}, builder_.Symbols().Register(name));
-  ident->set_intrinsic(intrinsic);
 
   ast::ExpressionList params;
   type::Type* first_operand_type = nullptr;
@@ -4679,7 +4678,6 @@ TypedExpression FunctionEmitter::MakeArrayLength(
   std::string call_ident_str = "arrayLength";
   auto* call_ident = create<ast::IdentifierExpression>(
       Source{}, builder_.Symbols().Register(call_ident_str));
-  call_ident->set_intrinsic(ast::Intrinsic::kArrayLength);
 
   ast::ExpressionList params{member_access};
   auto* call_expr =

src/semantic/call.h

@@ -0,0 +1,124 @@
+// 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.
+
+#ifndef SRC_SEMANTIC_CALL_H_
+#define SRC_SEMANTIC_CALL_H_
+
+#include "src/semantic/expression.h"
+#include "src/semantic/intrinsic.h"
+
+namespace tint {
+namespace semantic {
+
+/// Call is the base class for semantic nodes that hold semantic information for
+/// ast::CallExpression nodes.
+class Call : public Castable<Call, Expression> {
+ public:
+  /// Constructor
+  /// @param return_type the return type of the call
+  explicit Call(type::Type* return_type);
+
+  /// Destructor
+  ~Call() override;
+};
+
+/// IntrinsicCall holds semantic information for ast::CallExpression nodes that
+/// call intrinsic functions.
+class IntrinsicCall : public Castable<IntrinsicCall, Call> {
+ public:
+  /// Constructor
+  /// @param return_type the return type of the call
+  /// @param intrinsic the call target intrinsic
+  IntrinsicCall(type::Type* return_type, Intrinsic intrinsic);
+
+  /// Destructor
+  ~IntrinsicCall() override;
+
+  /// @returns the target intrinsic for the call
+  Intrinsic intrinsic() const { return intrinsic_; }
+
+ private:
+  Intrinsic const intrinsic_;
+};
+
+/// TextureIntrinsicCall holds semantic information for ast::CallExpression
+/// nodes that call intrinsic texture functions.
+class TextureIntrinsicCall
+    : public Castable<TextureIntrinsicCall, IntrinsicCall> {
+ public:
+  /// Parameters describes the parameters for the texture function.
+  struct Parameters {
+    /// kNotUsed is the constant that indicates the given parameter is not part
+    /// of the texture function signature.
+    static constexpr const size_t kNotUsed = ~static_cast<size_t>(0u);
+    /// Index holds each of the possible parameter indices. If a parameter index
+    /// is equal to `kNotUsed` then this parameter is not used by the function.
+    struct Index {
+      /// Constructor
+      Index();
+      /// Copy constructor
+      Index(const Index&);
+      /// `array_index` parameter index.
+      size_t array_index = kNotUsed;
+      /// `bias` parameter index.
+      size_t bias = kNotUsed;
+      /// `coords` parameter index.
+      size_t coords = kNotUsed;
+      /// `depth_ref` parameter index.
+      size_t depth_ref = kNotUsed;
+      /// `ddx` parameter index.
+      size_t ddx = kNotUsed;
+      /// `ddy` parameter index.
+      size_t ddy = kNotUsed;
+      /// `level` parameter index.
+      size_t level = kNotUsed;
+      /// `offset` parameter index.
+      size_t offset = kNotUsed;
+      /// `sampler` parameter index.
+      size_t sampler = kNotUsed;
+      /// `sample_index` parameter index.
+      size_t sample_index = kNotUsed;
+      /// `texture` parameter index.
+      size_t texture = kNotUsed;
+      /// `value` parameter index.
+      size_t value = kNotUsed;
+    };
+    /// The indices of all possible parameters.
+    Index idx;
+    /// Total number of parameters.
+    size_t count = 0;
+  };
+
+  /// Constructor
+  /// @param return_type the return type of the call
+  /// @param intrinsic the call target intrinsic
+  /// @param params the overload parameter info
+  TextureIntrinsicCall(type::Type* return_type,
+                       Intrinsic intrinsic,
+                       const Parameters& params);
+
+  /// Destructor
+  ~TextureIntrinsicCall() override;
+
+  /// @return the texture call's parameters
+  const Parameters& Params() const { return params_; }
+
+ private:
+  const Parameters params_;
+};
+
+}  // namespace semantic
+}  // namespace tint
+
+#endif  // SRC_SEMANTIC_CALL_H_

src/semantic/intrinsic.cc

@@ -12,245 +12,171 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/ast/intrinsic.h"
+#include "src/semantic/intrinsic.h"
 
 namespace tint {
-namespace ast {
+namespace semantic {
 
 std::ostream& operator<<(std::ostream& out, Intrinsic i) {
-  /// The emitted name matches the spelling in the WGSL spec.
+  out << intrinsic::str(i);
-  /// including case.
-  switch (i) {
-    case Intrinsic::kNone:
-      return out;
-    case Intrinsic::kAbs:
-      out << "abs";
-      return out;
-    case Intrinsic::kAcos:
-      out << "acos";
-      return out;
-    case Intrinsic::kAll:
-      out << "all";
-      return out;
-    case Intrinsic::kAny:
-      out << "any";
-      return out;
-    case Intrinsic::kArrayLength:
-      out << "arrayLength";
-      return out;
-    case Intrinsic::kAsin:
-      out << "asin";
-      return out;
-    case Intrinsic::kAtan:
-      out << "atan";
-      return out;
-    case Intrinsic::kAtan2:
-      out << "atan2";
-      return out;
-    case Intrinsic::kCeil:
-      out << "ceil";
-      return out;
-    case Intrinsic::kClamp:
-      out << "clamp";
-      return out;
-    case Intrinsic::kCos:
-      out << "cos";
-      return out;
-    case Intrinsic::kCosh:
-      out << "cosh";
-      return out;
-    case Intrinsic::kCountOneBits:
-      out << "countOneBits";
-      return out;
-    case Intrinsic::kCross:
-      out << "cross";
-      return out;
-    case Intrinsic::kDeterminant:
-      out << "determinant";
-      return out;
-    case Intrinsic::kDistance:
-      out << "distance";
-      return out;
-    case Intrinsic::kDot:
-      out << "dot";
-      return out;
-    case Intrinsic::kDpdx:
-      out << "dpdx";
-      return out;
-    case Intrinsic::kDpdxCoarse:
-      out << "dpdxCoarse";
-      return out;
-    case Intrinsic::kDpdxFine:
-      out << "dpdxFine";
-      return out;
-    case Intrinsic::kDpdy:
-      out << "dpdy";
-      return out;
-    case Intrinsic::kDpdyCoarse:
-      out << "dpdyCoarse";
-      return out;
-    case Intrinsic::kDpdyFine:
-      out << "dpdyFine";
-      return out;
-    case Intrinsic::kExp:
-      out << "exp";
-      return out;
-    case Intrinsic::kExp2:
-      out << "exp2";
-      return out;
-    case Intrinsic::kFaceForward:
-      out << "faceForward";
-      return out;
-    case Intrinsic::kFloor:
-      out << "floor";
-      return out;
-    case Intrinsic::kFma:
-      out << "fma";
-      return out;
-    case Intrinsic::kFract:
-      out << "fract";
-      return out;
-    case Intrinsic::kFrexp:
-      out << "frexp";
-      return out;
-    case Intrinsic::kFwidth:
-      out << "fwidth";
-      return out;
-    case Intrinsic::kFwidthCoarse:
-      out << "fwidthCoarse";
-      return out;
-    case Intrinsic::kFwidthFine:
-      out << "fwidthFine";
-      return out;
-    case Intrinsic::kInverseSqrt:
-      out << "inverseSqrt";
-      return out;
-    case Intrinsic::kIsFinite:
-      out << "isFinite";
-      return out;
-    case Intrinsic::kIsInf:
-      out << "isInf";
-      return out;
-    case Intrinsic::kIsNan:
-      out << "isNan";
-      return out;
-    case Intrinsic::kIsNormal:
-      out << "isNormal";
-      return out;
-    case Intrinsic::kLdexp:
-      out << "ldexp";
-      return out;
-    case Intrinsic::kLength:
-      out << "length";
-      return out;
-    case Intrinsic::kLog:
-      out << "log";
-      return out;
-    case Intrinsic::kLog2:
-      out << "log2";
-      return out;
-    case Intrinsic::kMax:
-      out << "max";
-      return out;
-    case Intrinsic::kMin:
-      out << "min";
-      return out;
-    case Intrinsic::kMix:
-      out << "mix";
-      return out;
-    case Intrinsic::kModf:
-      out << "modf";
-      return out;
-    case Intrinsic::kNormalize:
-      out << "normalize";
-      return out;
-    case Intrinsic::kPow:
-      out << "pow";
-      return out;
-    case Intrinsic::kReflect:
-      out << "reflect";
-      return out;
-    case Intrinsic::kReverseBits:
-      out << "reverseBits";
-      return out;
-    case Intrinsic::kRound:
-      out << "round";
-      return out;
-    case Intrinsic::kSelect:
-      out << "select";
-      return out;
-    case Intrinsic::kSign:
-      out << "sign";
-      return out;
-    case Intrinsic::kSin:
-      out << "sin";
-      return out;
-    case Intrinsic::kSinh:
-      out << "sinh";
-      return out;
-    case Intrinsic::kSmoothStep:
-      out << "smoothStep";
-      return out;
-    case Intrinsic::kSqrt:
-      out << "sqrt";
-      return out;
-    case Intrinsic::kStep:
-      out << "step";
-      return out;
-    case Intrinsic::kTan:
-      out << "tan";
-      return out;
-    case Intrinsic::kTanh:
-      out << "tanh";
-      return out;
-    case Intrinsic::kTextureDimensions:
-      out << "textureDimensions";
-      return out;
-    case Intrinsic::kTextureLoad:
-      out << "textureLoad";
-      return out;
-    case Intrinsic::kTextureNumLayers:
-      out << "textureNumLayers";
-      return out;
-    case Intrinsic::kTextureNumLevels:
-      out << "textureNumLevels";
-      return out;
-    case Intrinsic::kTextureNumSamples:
-      out << "textureNumSamples";
-      return out;
-    case Intrinsic::kTextureSample:
-      out << "textureSample";
-      return out;
-    case Intrinsic::kTextureSampleBias:
-      out << "textureSampleBias";
-      return out;
-    case Intrinsic::kTextureSampleCompare:
-      out << "textureSampleCompare";
-      return out;
-    case Intrinsic::kTextureSampleGrad:
-      out << "textureSampleGrad";
-      return out;
-    case Intrinsic::kTextureSampleLevel:
-      out << "textureSampleLevel";
-      return out;
-    case Intrinsic::kTextureStore:
-      out << "textureStore";
-      return out;
-    case Intrinsic::kTrunc:
-      out << "trunc";
-      return out;
-  }
-  out << "Unknown";
   return out;
 }
 
 namespace intrinsic {
 
-Signature::~Signature() = default;
+const char* str(Intrinsic i) {
-TextureSignature::~TextureSignature() = default;
+  /// The emitted name matches the spelling in the WGSL spec.
-
+  /// including case.
-TextureSignature::Parameters::Index::Index() = default;
+  switch (i) {
-TextureSignature::Parameters::Index::Index(const Index&) = default;
+    case Intrinsic::kNone:
+      return "<not-an-intrinsic>";
+    case Intrinsic::kAbs:
+      return "abs";
+    case Intrinsic::kAcos:
+      return "acos";
+    case Intrinsic::kAll:
+      return "all";
+    case Intrinsic::kAny:
+      return "any";
+    case Intrinsic::kArrayLength:
+      return "arrayLength";
+    case Intrinsic::kAsin:
+      return "asin";
+    case Intrinsic::kAtan:
+      return "atan";
+    case Intrinsic::kAtan2:
+      return "atan2";
+    case Intrinsic::kCeil:
+      return "ceil";
+    case Intrinsic::kClamp:
+      return "clamp";
+    case Intrinsic::kCos:
+      return "cos";
+    case Intrinsic::kCosh:
+      return "cosh";
+    case Intrinsic::kCountOneBits:
+      return "countOneBits";
+    case Intrinsic::kCross:
+      return "cross";
+    case Intrinsic::kDeterminant:
+      return "determinant";
+    case Intrinsic::kDistance:
+      return "distance";
+    case Intrinsic::kDot:
+      return "dot";
+    case Intrinsic::kDpdx:
+      return "dpdx";
+    case Intrinsic::kDpdxCoarse:
+      return "dpdxCoarse";
+    case Intrinsic::kDpdxFine:
+      return "dpdxFine";
+    case Intrinsic::kDpdy:
+      return "dpdy";
+    case Intrinsic::kDpdyCoarse:
+      return "dpdyCoarse";
+    case Intrinsic::kDpdyFine:
+      return "dpdyFine";
+    case Intrinsic::kExp:
+      return "exp";
+    case Intrinsic::kExp2:
+      return "exp2";
+    case Intrinsic::kFaceForward:
+      return "faceForward";
+    case Intrinsic::kFloor:
+      return "floor";
+    case Intrinsic::kFma:
+      return "fma";
+    case Intrinsic::kFract:
+      return "fract";
+    case Intrinsic::kFrexp:
+      return "frexp";
+    case Intrinsic::kFwidth:
+      return "fwidth";
+    case Intrinsic::kFwidthCoarse:
+      return "fwidthCoarse";
+    case Intrinsic::kFwidthFine:
+      return "fwidthFine";
+    case Intrinsic::kInverseSqrt:
+      return "inverseSqrt";
+    case Intrinsic::kIsFinite:
+      return "isFinite";
+    case Intrinsic::kIsInf:
+      return "isInf";
+    case Intrinsic::kIsNan:
+      return "isNan";
+    case Intrinsic::kIsNormal:
+      return "isNormal";
+    case Intrinsic::kLdexp:
+      return "ldexp";
+    case Intrinsic::kLength:
+      return "length";
+    case Intrinsic::kLog:
+      return "log";
+    case Intrinsic::kLog2:
+      return "log2";
+    case Intrinsic::kMax:
+      return "max";
+    case Intrinsic::kMin:
+      return "min";
+    case Intrinsic::kMix:
+      return "mix";
+    case Intrinsic::kModf:
+      return "modf";
+    case Intrinsic::kNormalize:
+      return "normalize";
+    case Intrinsic::kPow:
+      return "pow";
+    case Intrinsic::kReflect:
+      return "reflect";
+    case Intrinsic::kReverseBits:
+      return "reverseBits";
+    case Intrinsic::kRound:
+      return "round";
+    case Intrinsic::kSelect:
+      return "select";
+    case Intrinsic::kSign:
+      return "sign";
+    case Intrinsic::kSin:
+      return "sin";
+    case Intrinsic::kSinh:
+      return "sinh";
+    case Intrinsic::kSmoothStep:
+      return "smoothStep";
+    case Intrinsic::kSqrt:
+      return "sqrt";
+    case Intrinsic::kStep:
+      return "step";
+    case Intrinsic::kTan:
+      return "tan";
+    case Intrinsic::kTanh:
+      return "tanh";
+    case Intrinsic::kTextureDimensions:
+      return "textureDimensions";
+    case Intrinsic::kTextureLoad:
+      return "textureLoad";
+    case Intrinsic::kTextureNumLayers:
+      return "textureNumLayers";
+    case Intrinsic::kTextureNumLevels:
+      return "textureNumLevels";
+    case Intrinsic::kTextureNumSamples:
+      return "textureNumSamples";
+    case Intrinsic::kTextureSample:
+      return "textureSample";
+    case Intrinsic::kTextureSampleBias:
+      return "textureSampleBias";
+    case Intrinsic::kTextureSampleCompare:
+      return "textureSampleCompare";
+    case Intrinsic::kTextureSampleGrad:
+      return "textureSampleGrad";
+    case Intrinsic::kTextureSampleLevel:
+      return "textureSampleLevel";
+    case Intrinsic::kTextureStore:
+      return "textureStore";
+    case Intrinsic::kTrunc:
+      return "trunc";
+  }
+  return "<unknown>";
+}
 
 bool IsCoarseDerivative(Intrinsic i) {
   return i == Intrinsic::kDpdxCoarse || i == Intrinsic::kDpdyCoarse ||
@@ -283,12 +209,12 @@ bool IsTextureIntrinsic(Intrinsic i) {
 }
 
 bool IsImageQueryIntrinsic(Intrinsic i) {
-  return i == ast::Intrinsic::kTextureDimensions ||
+  return i == semantic::Intrinsic::kTextureDimensions ||
          i == Intrinsic::kTextureNumLayers ||
          i == Intrinsic::kTextureNumLevels ||
          i == Intrinsic::kTextureNumSamples;
 }
 
 }  // namespace intrinsic
-}  // namespace ast
+}  // namespace semantic
 }  // namespace tint

src/semantic/intrinsic.h

@@ -12,13 +12,13 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef SRC_AST_INTRINSIC_H_
+#ifndef SRC_SEMANTIC_INTRINSIC_H_
-#define SRC_AST_INTRINSIC_H_
+#define SRC_SEMANTIC_INTRINSIC_H_
 
 #include <ostream>
 
 namespace tint {
-namespace ast {
+namespace semantic {
 
 enum class Intrinsic {
   kNone = -1,
@@ -103,68 +103,6 @@ std::ostream& operator<<(std::ostream& out, Intrinsic i);
 
 namespace intrinsic {
 
-/// Signature is the base struct for all intrinsic signature types.
-/// Signatures are used to identify the particular overload for intrinsics that
-/// have different signatures with the same function name.
-struct Signature {
-  virtual ~Signature();
-};
-
-/// TextureSignature describes the signature of a texture intrinsic function.
-struct TextureSignature : public Signature {
-  /// Parameters describes the parameters for the texture function.
-  struct Parameters {
-    /// kNotUsed is the constant that indicates the given parameter is not part
-    /// of the texture function signature.
-    static constexpr const size_t kNotUsed = ~static_cast<size_t>(0u);
-    /// Index holds each of the possible parameter indices. If a parameter index
-    /// is equal to `kNotUsed` then this parameter is not used by the function.
-    struct Index {
-      /// Constructor
-      Index();
-      /// Copy constructor
-      Index(const Index&);
-      /// `array_index` parameter index.
-      size_t array_index = kNotUsed;
-      /// `bias` parameter index.
-      size_t bias = kNotUsed;
-      /// `coords` parameter index.
-      size_t coords = kNotUsed;
-      /// `depth_ref` parameter index.
-      size_t depth_ref = kNotUsed;
-      /// `ddx` parameter index.
-      size_t ddx = kNotUsed;
-      /// `ddy` parameter index.
-      size_t ddy = kNotUsed;
-      /// `level` parameter index.
-      size_t level = kNotUsed;
-      /// `offset` parameter index.
-      size_t offset = kNotUsed;
-      /// `sampler` parameter index.
-      size_t sampler = kNotUsed;
-      /// `sample_index` parameter index.
-      size_t sample_index = kNotUsed;
-      /// `texture` parameter index.
-      size_t texture = kNotUsed;
-      /// `value` parameter index.
-      size_t value = kNotUsed;
-    };
-    /// The indices of all possible parameters.
-    Index idx;
-    /// Total number of parameters.
-    size_t count = 0;
-  };
-
-  /// Construct an immutable `TextureSignature`.
-  /// @param p the texture intrinsic parameter signature.
-  explicit TextureSignature(const Parameters& p) : params(p) {}
-
-  ~TextureSignature() override;
-
-  /// The texture intrinsic parameter signature.
-  const Parameters params;
-};
-
 /// Determines if the given `i` is a coarse derivative
 /// @param i the intrinsic
 /// @returns true if the given derivative is coarse.
@@ -195,8 +133,12 @@ bool IsTextureIntrinsic(Intrinsic i);
 /// @returns true if the given `i` is a image query intrinsic
 bool IsImageQueryIntrinsic(Intrinsic i);
 
+/// @returns the name of the intrinsic function. The spelling, including case,
+/// matches the name in the WGSL spec.
+const char* str(Intrinsic i);
+
 }  // namespace intrinsic
-}  // namespace ast
+}  // namespace semantic
 }  // namespace tint
 
-#endif  // SRC_AST_INTRINSIC_H_
+#endif  // SRC_SEMANTIC_INTRINSIC_H_

src/semantic/sem_call.cc

@@ -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.
+
+#include "src/semantic/call.h"
+
+TINT_INSTANTIATE_CLASS_ID(tint::semantic::Call);
+TINT_INSTANTIATE_CLASS_ID(tint::semantic::IntrinsicCall);
+TINT_INSTANTIATE_CLASS_ID(tint::semantic::TextureIntrinsicCall);
+
+namespace tint {
+namespace semantic {
+
+Call::Call(type::Type* return_type) : Base(return_type) {}
+
+Call::~Call() = default;
+
+IntrinsicCall::IntrinsicCall(type::Type* return_type, Intrinsic intrinsic)
+    : Base(return_type), intrinsic_(intrinsic) {}
+
+IntrinsicCall::~IntrinsicCall() = default;
+
+TextureIntrinsicCall::TextureIntrinsicCall(type::Type* return_type,
+                                           Intrinsic intrinsic,
+                                           const Parameters& params)
+    : Base(return_type, intrinsic), params_(params) {}
+
+TextureIntrinsicCall::~TextureIntrinsicCall() = default;
+
+TextureIntrinsicCall::Parameters::Index::Index() = default;
+TextureIntrinsicCall::Parameters::Index::Index(const Index&) = default;
+
+}  // namespace semantic
+}  // namespace tint

src/semantic/type_mappings.h

@@ -22,6 +22,7 @@ namespace tint {
 // Forward declarations
 namespace ast {
 
+class CallExpression;
 class Expression;
 class Function;
 class Variable;
@@ -30,6 +31,7 @@ class Variable;
 
 namespace semantic {
 
+class Call;
 class Expression;
 class Function;
 class Variable;
@@ -44,6 +46,7 @@ struct TypeMappings {
   semantic::Expression* operator()(ast::Expression*);
   semantic::Function* operator()(ast::Function*);
   semantic::Variable* operator()(ast::Variable*);
+  semantic::Call* operator()(ast::CallExpression*);
   //! @endcond
 };
 

src/type_determiner.cc

@@ -33,7 +33,6 @@
 #include "src/ast/fallthrough_statement.h"
 #include "src/ast/identifier_expression.h"
 #include "src/ast/if_statement.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/loop_statement.h"
 #include "src/ast/member_accessor_expression.h"
 #include "src/ast/return_statement.h"
@@ -43,8 +42,10 @@
 #include "src/ast/unary_op_expression.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/program_builder.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/intrinsic.h"
 #include "src/semantic/variable.h"
 #include "src/type/array_type.h"
 #include "src/type/bool_type.h"
@@ -393,20 +394,28 @@ bool TypeDeterminer::DetermineBitcast(ast::BitcastExpression* expr) {
   return true;
 }
 
-bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
+bool TypeDeterminer::DetermineCall(ast::CallExpression* call) {
-  if (!DetermineResultType(expr->func())) {
+  if (!DetermineResultType(call->func())) {
     return false;
   }
-  if (!DetermineResultType(expr->params())) {
+  if (!DetermineResultType(call->params())) {
     return false;
   }
 
   // The expression has to be an identifier as you can't store function pointers
   // but, if it isn't we'll just use the normal result determination to be on
   // the safe side.
-  if (auto* ident = expr->func()->As<ast::IdentifierExpression>()) {
+  auto* ident = call->func()->As<ast::IdentifierExpression>();
-    if (ident->IsIntrinsic()) {
+  if (!ident) {
-      if (!DetermineIntrinsic(ident, expr)) {
+    set_error(call->source(), "call target is not an identifier");
+    return false;
+  }
+
+  auto name = builder_->Symbols().NameFor(ident->symbol());
+
+  auto intrinsic = MatchIntrinsic(name);
+  if (intrinsic != semantic::Intrinsic::kNone) {
+    if (!DetermineIntrinsicCall(call, intrinsic)) {
       return false;
     }
   } else {
@@ -416,9 +425,7 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
 
       auto callee_func_it = symbol_to_function_.find(ident->symbol());
       if (callee_func_it == symbol_to_function_.end()) {
-          set_error(expr->source(),
+        set_error(call->source(), "unable to find called function: " + name);
-                    "unable to find called function: " +
-                        builder_->Symbols().NameFor(ident->symbol()));
         return false;
       }
       auto* callee_func = callee_func_it->second;
@@ -428,21 +435,18 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
         set_referenced_from_function_if_needed(var, false);
       }
     }
-    }
+
-  } else {
+    auto iter = symbol_to_function_.find(ident->symbol());
-    if (!DetermineResultType(expr->func())) {
+    if (iter == symbol_to_function_.end()) {
+      set_error(call->source(),
+                "v-0005: function must be declared before use: '" + name + "'");
       return false;
     }
-  }
 
-  if (auto* type = TypeOf(expr->func())) {
+    auto* function = iter->second;
-    SetType(expr, type);
+    auto* return_ty = function->declaration->return_type();
-  } else {
+    auto* sem = builder_->create<semantic::Call>(return_ty);
-    auto func_sym = expr->func()->As<ast::IdentifierExpression>()->symbol();
+    builder_->Sem().Add(call, sem);
-    set_error(expr->source(),
-              "v-0005: function must be declared before use: '" +
-                  builder_->Symbols().NameFor(func_sym) + "'");
-    return false;
   }
 
   return true;
@@ -459,7 +463,7 @@ enum class IntrinsicDataType {
 };
 
 struct IntrinsicData {
-  ast::Intrinsic intrinsic;
+  semantic::Intrinsic intrinsic;
   IntrinsicDataType result_type;
   uint8_t result_vector_width;
   uint8_t param_for_result_type;
@@ -468,63 +472,64 @@ struct IntrinsicData {
 // Note, this isn't all the intrinsics. Some are handled specially before
 // we get to the generic code. See the DetermineIntrinsic code below.
 constexpr const IntrinsicData kIntrinsicData[] = {
-    {ast::Intrinsic::kAbs, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kAbs, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kAcos, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kAcos, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kAll, IntrinsicDataType::kBool, 1, 0},
+    {semantic::Intrinsic::kAll, IntrinsicDataType::kBool, 1, 0},
-    {ast::Intrinsic::kAny, IntrinsicDataType::kBool, 1, 0},
+    {semantic::Intrinsic::kAny, IntrinsicDataType::kBool, 1, 0},
-    {ast::Intrinsic::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1, 0},
+    {semantic::Intrinsic::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1,
-    {ast::Intrinsic::kAsin, IntrinsicDataType::kDependent, 0, 0},
+     0},
-    {ast::Intrinsic::kAtan, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kAsin, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kAtan2, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kAtan, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kCeil, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kAtan2, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kClamp, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kCeil, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kCos, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kClamp, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kCosh, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kCos, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kCosh, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kCross, IntrinsicDataType::kFloat, 3, 0},
+    {semantic::Intrinsic::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
+    {semantic::Intrinsic::kCross, IntrinsicDataType::kFloat, 3, 0},
-    {ast::Intrinsic::kDistance, IntrinsicDataType::kFloat, 1, 0},
+    {semantic::Intrinsic::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
-    {ast::Intrinsic::kDpdx, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDistance, IntrinsicDataType::kFloat, 1, 0},
-    {ast::Intrinsic::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDpdx, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDpdy, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDpdy, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kDot, IntrinsicDataType::kFloat, 1, 0},
+    {semantic::Intrinsic::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kExp, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kDot, IntrinsicDataType::kFloat, 1, 0},
-    {ast::Intrinsic::kExp2, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kExp, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kExp2, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFloor, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFwidth, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFloor, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFwidth, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFma, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFract, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFma, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kFrexp, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFract, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kFrexp, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kLdexp, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kLength, IntrinsicDataType::kFloat, 1, 0},
+    {semantic::Intrinsic::kLdexp, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kLog, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kLength, IntrinsicDataType::kFloat, 1, 0},
-    {ast::Intrinsic::kLog2, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kLog, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kMax, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kLog2, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kMin, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kMax, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kMin, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kRound, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSelect, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kRound, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSign, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSelect, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSin, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSign, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSinh, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSin, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSinh, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kSqrt, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kStep, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kSqrt, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kTan, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kStep, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kTanh, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kTan, IntrinsicDataType::kDependent, 0, 0},
-    {ast::Intrinsic::kTrunc, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kTanh, IntrinsicDataType::kDependent, 0, 0},
+    {semantic::Intrinsic::kTrunc, IntrinsicDataType::kDependent, 0, 0},
 };
 
 constexpr const uint32_t kIntrinsicDataCount =
@@ -532,35 +537,36 @@ constexpr const uint32_t kIntrinsicDataCount =
 
 }  // namespace
 
-bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
+bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
-                                        ast::CallExpression* expr) {
+                                            semantic::Intrinsic intrinsic) {
-  if (ast::intrinsic::IsFloatClassificationIntrinsic(ident->intrinsic())) {
+  auto create_sem = [&](type::Type* result) {
-    if (expr->params().size() != 1) {
+    auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic);
-      set_error(expr->source(),
+    builder_->Sem().Add(call, sem);
-                "incorrect number of parameters for " +
+  };
-                    builder_->Symbols().NameFor(ident->symbol()));
+
+  std::string name = semantic::intrinsic::str(intrinsic);
+  if (semantic::intrinsic::IsFloatClassificationIntrinsic(intrinsic)) {
+    if (call->params().size() != 1) {
+      set_error(call->source(), "incorrect number of parameters for " + name);
       return false;
     }
 
     auto* bool_type = builder_->create<type::Bool>();
 
-    auto* param_type = TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
+    auto* param_type = TypeOf(call->params()[0])->UnwrapPtrIfNeeded();
     if (auto* vec = param_type->As<type::Vector>()) {
-      SetType(expr->func(),
+      create_sem(builder_->create<type::Vector>(bool_type, vec->size()));
-              builder_->create<type::Vector>(bool_type, vec->size()));
     } else {
-      SetType(expr->func(), bool_type);
+      create_sem(bool_type);
     }
     return true;
   }
-  if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
+  if (semantic::intrinsic::IsTextureIntrinsic(intrinsic)) {
-    ast::intrinsic::TextureSignature::Parameters param;
+    semantic::TextureIntrinsicCall::Parameters param;
 
-    auto* texture_param = expr->params()[0];
+    auto* texture_param = call->params()[0];
     if (!TypeOf(texture_param)->UnwrapAll()->Is<type::Texture>()) {
-      set_error(expr->source(),
+      set_error(call->source(), "invalid first argument for " + name);
-                "invalid first argument for " +
-                    builder_->Symbols().NameFor(ident->symbol()));
       return false;
     }
     type::Texture* texture =
@@ -568,25 +574,25 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
 
     bool is_array = type::IsTextureArray(texture->dim());
     bool is_multisampled = texture->Is<type::MultisampledTexture>();
-    switch (ident->intrinsic()) {
+    switch (intrinsic) {
-      case ast::Intrinsic::kTextureDimensions:
+      case semantic::Intrinsic::kTextureDimensions:
         param.idx.texture = param.count++;
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.level = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureNumLayers:
+      case semantic::Intrinsic::kTextureNumLayers:
-      case ast::Intrinsic::kTextureNumLevels:
+      case semantic::Intrinsic::kTextureNumLevels:
-      case ast::Intrinsic::kTextureNumSamples:
+      case semantic::Intrinsic::kTextureNumSamples:
         param.idx.texture = param.count++;
         break;
-      case ast::Intrinsic::kTextureLoad:
+      case semantic::Intrinsic::kTextureLoad:
         param.idx.texture = param.count++;
         param.idx.coords = param.count++;
         if (is_array) {
           param.idx.array_index = param.count++;
         }
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           if (is_multisampled) {
             param.idx.sample_index = param.count++;
           } else {
@@ -594,18 +600,18 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
           }
         }
         break;
-      case ast::Intrinsic::kTextureSample:
+      case semantic::Intrinsic::kTextureSample:
         param.idx.texture = param.count++;
         param.idx.sampler = param.count++;
         param.idx.coords = param.count++;
         if (is_array) {
           param.idx.array_index = param.count++;
         }
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.offset = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureSampleBias:
+      case semantic::Intrinsic::kTextureSampleBias:
         param.idx.texture = param.count++;
         param.idx.sampler = param.count++;
         param.idx.coords = param.count++;
@@ -613,11 +619,11 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
           param.idx.array_index = param.count++;
         }
         param.idx.bias = param.count++;
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.offset = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureSampleLevel:
+      case semantic::Intrinsic::kTextureSampleLevel:
         param.idx.texture = param.count++;
         param.idx.sampler = param.count++;
         param.idx.coords = param.count++;
@@ -625,11 +631,11 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
           param.idx.array_index = param.count++;
         }
         param.idx.level = param.count++;
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.offset = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureSampleCompare:
+      case semantic::Intrinsic::kTextureSampleCompare:
         param.idx.texture = param.count++;
         param.idx.sampler = param.count++;
         param.idx.coords = param.count++;
@@ -637,11 +643,11 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
           param.idx.array_index = param.count++;
         }
         param.idx.depth_ref = param.count++;
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.offset = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureSampleGrad:
+      case semantic::Intrinsic::kTextureSampleGrad:
         param.idx.texture = param.count++;
         param.idx.sampler = param.count++;
         param.idx.coords = param.count++;
@@ -650,11 +656,11 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
         }
         param.idx.ddx = param.count++;
         param.idx.ddy = param.count++;
-        if (expr->params().size() > param.count) {
+        if (call->params().size() > param.count) {
           param.idx.offset = param.count++;
         }
         break;
-      case ast::Intrinsic::kTextureStore:
+      case semantic::Intrinsic::kTextureStore:
         param.idx.texture = param.count++;
         param.idx.coords = param.count++;
         if (is_array) {
@@ -663,32 +669,28 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
         param.idx.value = param.count++;
         break;
       default:
-        set_error(expr->source(),
+        set_error(call->source(),
                   "Internal compiler error: Unreachable intrinsic " +
-                      std::to_string(static_cast<int>(ident->intrinsic())));
+                      std::to_string(static_cast<int>(intrinsic)));
         return false;
     }
 
-    if (expr->params().size() != param.count) {
+    if (call->params().size() != param.count) {
-      set_error(expr->source(),
+      set_error(call->source(),
-                "incorrect number of parameters for " +
+                "incorrect number of parameters for " + name + ", got " +
-                    builder_->Symbols().NameFor(ident->symbol()) + ", got " +
+                    std::to_string(call->params().size()) + " and expected " +
-                    std::to_string(expr->params().size()) + " and expected " +
                     std::to_string(param.count));
       return false;
     }
 
-    ident->set_intrinsic_signature(
-        std::make_unique<ast::intrinsic::TextureSignature>(param));
-
     // Set the function return type
     type::Type* return_type = nullptr;
-    switch (ident->intrinsic()) {
+    switch (intrinsic) {
-      case ast::Intrinsic::kTextureDimensions: {
+      case semantic::Intrinsic::kTextureDimensions: {
         auto* i32 = builder_->create<type::I32>();
         switch (texture->dim()) {
           default:
-            set_error(expr->source(), "invalid texture dimensions");
+            set_error(call->source(), "invalid texture dimensions");
             break;
           case type::TextureDimension::k1d:
           case type::TextureDimension::k1dArray:
@@ -706,12 +708,12 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
         }
         break;
       }
-      case ast::Intrinsic::kTextureNumLayers:
+      case semantic::Intrinsic::kTextureNumLayers:
-      case ast::Intrinsic::kTextureNumLevels:
+      case semantic::Intrinsic::kTextureNumLevels:
-      case ast::Intrinsic::kTextureNumSamples:
+      case semantic::Intrinsic::kTextureNumSamples:
         return_type = builder_->create<type::I32>();
         break;
-      case ast::Intrinsic::kTextureStore:
+      case semantic::Intrinsic::kTextureStore:
         return_type = builder_->create<type::Void>();
         break;
       default: {
@@ -727,7 +729,7 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
                          texture->As<type::MultisampledTexture>()) {
             type = msampled->type();
           } else {
-            set_error(expr->source(),
+            set_error(call->source(),
                       "unknown texture type for texture sampling");
             return false;
           }
@@ -735,35 +737,35 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
         }
       }
     }
-    SetType(expr->func(), return_type);
+
+    auto* sem = builder_->create<semantic::TextureIntrinsicCall>(
+        return_type, intrinsic, param);
+    builder_->Sem().Add(call, sem);
 
     return true;
   }
 
   const IntrinsicData* data = nullptr;
   for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) {
-    if (ident->intrinsic() == kIntrinsicData[i].intrinsic) {
+    if (intrinsic == kIntrinsicData[i].intrinsic) {
       data = &kIntrinsicData[i];
       break;
     }
   }
   if (data == nullptr) {
-    error_ = "unable to find intrinsic " +
+    error_ = "unable to find intrinsic " + name;
-             builder_->Symbols().NameFor(ident->symbol());
     return false;
   }
 
   if (data->result_type == IntrinsicDataType::kDependent) {
     const auto param_idx = data->param_for_result_type;
-    if (expr->params().size() <= param_idx) {
+    if (call->params().size() <= param_idx) {
-      set_error(expr->source(),
+      set_error(call->source(),
                 "missing parameter " + std::to_string(param_idx) +
-                    " required for type determination in builtin " +
+                    " required for type determination in builtin " + name);
-                    builder_->Symbols().NameFor(ident->symbol()));
       return false;
     }
-    SetType(expr->func(),
+    create_sem(TypeOf(call->params()[param_idx])->UnwrapPtrIfNeeded());
-            TypeOf(expr->params()[param_idx])->UnwrapPtrIfNeeded());
   } else {
     // The result type is not dependent on the parameter types.
     type::Type* type = nullptr;
@@ -781,15 +783,14 @@ bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
         type = builder_->create<type::Bool>();
         break;
       default:
-        error_ = "unhandled intrinsic data type for " +
+        error_ = "unhandled intrinsic data type for " + name;
-                 builder_->Symbols().NameFor(ident->symbol());
         return false;
     }
 
     if (data->result_vector_width > 1) {
       type = builder_->create<type::Vector>(type, data->result_vector_width);
     }
-    SetType(expr->func(), type);
+    create_sem(type);
   }
 
   return true;
@@ -831,169 +832,168 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
 
   auto iter = symbol_to_function_.find(symbol);
   if (iter != symbol_to_function_.end()) {
-    SetType(expr, iter->second->declaration->return_type());
+    // Identifier is to a function, which has no type (currently).
     return true;
   }
 
-  if (!SetIntrinsicIfNeeded(expr)) {
+  std::string name = builder_->Symbols().NameFor(symbol);
-    set_error(expr->source(),
+  if (MatchIntrinsic(name) != semantic::Intrinsic::kNone) {
-              "v-0006: identifier must be declared before use: " +
+    // Identifier is to an intrinsic function, which has no type (currently).
-                  builder_->Symbols().NameFor(symbol));
-    return false;
-  }
     return true;
+  }
+
+  set_error(expr->source(),
+            "v-0006: identifier must be declared before use: " + name);
+  return false;
 }
 
-bool TypeDeterminer::SetIntrinsicIfNeeded(ast::IdentifierExpression* ident) {
+semantic::Intrinsic TypeDeterminer::MatchIntrinsic(const std::string& name) {
-  auto name = builder_->Symbols().NameFor(ident->symbol());
   if (name == "abs") {
-    ident->set_intrinsic(ast::Intrinsic::kAbs);
+    return semantic::Intrinsic::kAbs;
   } else if (name == "acos") {
-    ident->set_intrinsic(ast::Intrinsic::kAcos);
+    return semantic::Intrinsic::kAcos;
   } else if (name == "all") {
-    ident->set_intrinsic(ast::Intrinsic::kAll);
+    return semantic::Intrinsic::kAll;
   } else if (name == "any") {
-    ident->set_intrinsic(ast::Intrinsic::kAny);
+    return semantic::Intrinsic::kAny;
   } else if (name == "arrayLength") {
-    ident->set_intrinsic(ast::Intrinsic::kArrayLength);
+    return semantic::Intrinsic::kArrayLength;
   } else if (name == "asin") {
-    ident->set_intrinsic(ast::Intrinsic::kAsin);
+    return semantic::Intrinsic::kAsin;
   } else if (name == "atan") {
-    ident->set_intrinsic(ast::Intrinsic::kAtan);
+    return semantic::Intrinsic::kAtan;
   } else if (name == "atan2") {
-    ident->set_intrinsic(ast::Intrinsic::kAtan2);
+    return semantic::Intrinsic::kAtan2;
   } else if (name == "ceil") {
-    ident->set_intrinsic(ast::Intrinsic::kCeil);
+    return semantic::Intrinsic::kCeil;
   } else if (name == "clamp") {
-    ident->set_intrinsic(ast::Intrinsic::kClamp);
+    return semantic::Intrinsic::kClamp;
   } else if (name == "cos") {
-    ident->set_intrinsic(ast::Intrinsic::kCos);
+    return semantic::Intrinsic::kCos;
   } else if (name == "cosh") {
-    ident->set_intrinsic(ast::Intrinsic::kCosh);
+    return semantic::Intrinsic::kCosh;
   } else if (name == "countOneBits") {
-    ident->set_intrinsic(ast::Intrinsic::kCountOneBits);
+    return semantic::Intrinsic::kCountOneBits;
   } else if (name == "cross") {
-    ident->set_intrinsic(ast::Intrinsic::kCross);
+    return semantic::Intrinsic::kCross;
   } else if (name == "determinant") {
-    ident->set_intrinsic(ast::Intrinsic::kDeterminant);
+    return semantic::Intrinsic::kDeterminant;
   } else if (name == "distance") {
-    ident->set_intrinsic(ast::Intrinsic::kDistance);
+    return semantic::Intrinsic::kDistance;
   } else if (name == "dot") {
-    ident->set_intrinsic(ast::Intrinsic::kDot);
+    return semantic::Intrinsic::kDot;
   } else if (name == "dpdx") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdx);
+    return semantic::Intrinsic::kDpdx;
   } else if (name == "dpdxCoarse") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdxCoarse);
+    return semantic::Intrinsic::kDpdxCoarse;
   } else if (name == "dpdxFine") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdxFine);
+    return semantic::Intrinsic::kDpdxFine;
   } else if (name == "dpdy") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdy);
+    return semantic::Intrinsic::kDpdy;
   } else if (name == "dpdyCoarse") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdyCoarse);
+    return semantic::Intrinsic::kDpdyCoarse;
   } else if (name == "dpdyFine") {
-    ident->set_intrinsic(ast::Intrinsic::kDpdyFine);
+    return semantic::Intrinsic::kDpdyFine;
   } else if (name == "exp") {
-    ident->set_intrinsic(ast::Intrinsic::kExp);
+    return semantic::Intrinsic::kExp;
   } else if (name == "exp2") {
-    ident->set_intrinsic(ast::Intrinsic::kExp2);
+    return semantic::Intrinsic::kExp2;
   } else if (name == "faceForward") {
-    ident->set_intrinsic(ast::Intrinsic::kFaceForward);
+    return semantic::Intrinsic::kFaceForward;
   } else if (name == "floor") {
-    ident->set_intrinsic(ast::Intrinsic::kFloor);
+    return semantic::Intrinsic::kFloor;
   } else if (name == "fma") {
-    ident->set_intrinsic(ast::Intrinsic::kFma);
+    return semantic::Intrinsic::kFma;
   } else if (name == "fract") {
-    ident->set_intrinsic(ast::Intrinsic::kFract);
+    return semantic::Intrinsic::kFract;
   } else if (name == "frexp") {
-    ident->set_intrinsic(ast::Intrinsic::kFrexp);
+    return semantic::Intrinsic::kFrexp;
   } else if (name == "fwidth") {
-    ident->set_intrinsic(ast::Intrinsic::kFwidth);
+    return semantic::Intrinsic::kFwidth;
   } else if (name == "fwidthCoarse") {
-    ident->set_intrinsic(ast::Intrinsic::kFwidthCoarse);
+    return semantic::Intrinsic::kFwidthCoarse;
   } else if (name == "fwidthFine") {
-    ident->set_intrinsic(ast::Intrinsic::kFwidthFine);
+    return semantic::Intrinsic::kFwidthFine;
   } else if (name == "inverseSqrt") {
-    ident->set_intrinsic(ast::Intrinsic::kInverseSqrt);
+    return semantic::Intrinsic::kInverseSqrt;
   } else if (name == "isFinite") {
-    ident->set_intrinsic(ast::Intrinsic::kIsFinite);
+    return semantic::Intrinsic::kIsFinite;
   } else if (name == "isInf") {
-    ident->set_intrinsic(ast::Intrinsic::kIsInf);
+    return semantic::Intrinsic::kIsInf;
   } else if (name == "isNan") {
-    ident->set_intrinsic(ast::Intrinsic::kIsNan);
+    return semantic::Intrinsic::kIsNan;
   } else if (name == "isNormal") {
-    ident->set_intrinsic(ast::Intrinsic::kIsNormal);
+    return semantic::Intrinsic::kIsNormal;
   } else if (name == "ldexp") {
-    ident->set_intrinsic(ast::Intrinsic::kLdexp);
+    return semantic::Intrinsic::kLdexp;
   } else if (name == "length") {
-    ident->set_intrinsic(ast::Intrinsic::kLength);
+    return semantic::Intrinsic::kLength;
   } else if (name == "log") {
-    ident->set_intrinsic(ast::Intrinsic::kLog);
+    return semantic::Intrinsic::kLog;
   } else if (name == "log2") {
-    ident->set_intrinsic(ast::Intrinsic::kLog2);
+    return semantic::Intrinsic::kLog2;
   } else if (name == "max") {
-    ident->set_intrinsic(ast::Intrinsic::kMax);
+    return semantic::Intrinsic::kMax;
   } else if (name == "min") {
-    ident->set_intrinsic(ast::Intrinsic::kMin);
+    return semantic::Intrinsic::kMin;
   } else if (name == "mix") {
-    ident->set_intrinsic(ast::Intrinsic::kMix);
+    return semantic::Intrinsic::kMix;
   } else if (name == "modf") {
-    ident->set_intrinsic(ast::Intrinsic::kModf);
+    return semantic::Intrinsic::kModf;
   } else if (name == "normalize") {
-    ident->set_intrinsic(ast::Intrinsic::kNormalize);
+    return semantic::Intrinsic::kNormalize;
   } else if (name == "pow") {
-    ident->set_intrinsic(ast::Intrinsic::kPow);
+    return semantic::Intrinsic::kPow;
   } else if (name == "reflect") {
-    ident->set_intrinsic(ast::Intrinsic::kReflect);
+    return semantic::Intrinsic::kReflect;
   } else if (name == "reverseBits") {
-    ident->set_intrinsic(ast::Intrinsic::kReverseBits);
+    return semantic::Intrinsic::kReverseBits;
   } else if (name == "round") {
-    ident->set_intrinsic(ast::Intrinsic::kRound);
+    return semantic::Intrinsic::kRound;
   } else if (name == "select") {
-    ident->set_intrinsic(ast::Intrinsic::kSelect);
+    return semantic::Intrinsic::kSelect;
   } else if (name == "sign") {
-    ident->set_intrinsic(ast::Intrinsic::kSign);
+    return semantic::Intrinsic::kSign;
   } else if (name == "sin") {
-    ident->set_intrinsic(ast::Intrinsic::kSin);
+    return semantic::Intrinsic::kSin;
   } else if (name == "sinh") {
-    ident->set_intrinsic(ast::Intrinsic::kSinh);
+    return semantic::Intrinsic::kSinh;
   } else if (name == "smoothStep") {
-    ident->set_intrinsic(ast::Intrinsic::kSmoothStep);
+    return semantic::Intrinsic::kSmoothStep;
   } else if (name == "sqrt") {
-    ident->set_intrinsic(ast::Intrinsic::kSqrt);
+    return semantic::Intrinsic::kSqrt;
   } else if (name == "step") {
-    ident->set_intrinsic(ast::Intrinsic::kStep);
+    return semantic::Intrinsic::kStep;
   } else if (name == "tan") {
-    ident->set_intrinsic(ast::Intrinsic::kTan);
+    return semantic::Intrinsic::kTan;
   } else if (name == "tanh") {
-    ident->set_intrinsic(ast::Intrinsic::kTanh);
+    return semantic::Intrinsic::kTanh;
   } else if (name == "textureDimensions") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureDimensions);
+    return semantic::Intrinsic::kTextureDimensions;
   } else if (name == "textureNumLayers") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureNumLayers);
+    return semantic::Intrinsic::kTextureNumLayers;
   } else if (name == "textureNumLevels") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureNumLevels);
+    return semantic::Intrinsic::kTextureNumLevels;
   } else if (name == "textureNumSamples") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureNumSamples);
+    return semantic::Intrinsic::kTextureNumSamples;
   } else if (name == "textureLoad") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureLoad);
+    return semantic::Intrinsic::kTextureLoad;
   } else if (name == "textureStore") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureStore);
+    return semantic::Intrinsic::kTextureStore;
   } else if (name == "textureSample") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureSample);
+    return semantic::Intrinsic::kTextureSample;
   } else if (name == "textureSampleBias") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureSampleBias);
+    return semantic::Intrinsic::kTextureSampleBias;
   } else if (name == "textureSampleCompare") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureSampleCompare);
+    return semantic::Intrinsic::kTextureSampleCompare;
   } else if (name == "textureSampleGrad") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureSampleGrad);
+    return semantic::Intrinsic::kTextureSampleGrad;
   } else if (name == "textureSampleLevel") {
-    ident->set_intrinsic(ast::Intrinsic::kTextureSampleLevel);
+    return semantic::Intrinsic::kTextureSampleLevel;
   } else if (name == "trunc") {
-    ident->set_intrinsic(ast::Intrinsic::kTrunc);
+    return semantic::Intrinsic::kTrunc;
-  } else {
-    return false;
   }
-  return true;
+  return semantic::Intrinsic::kNone;
 }
 
 bool TypeDeterminer::DetermineMemberAccessor(

src/type_determiner.h

@@ -66,10 +66,10 @@ class TypeDeterminer {
   /// @returns true if the type determiner was successful
   bool Determine();
 
-  /// Sets the intrinsic data information for the identifier if needed
+  /// @param name the function name to try and match as an intrinsic.
-  /// @param ident the identifier expression
+  /// @return the semantic::Intrinsic for the given name. If `name` does not
-  /// @returns true if an intrinsic was set
+  /// match an intrinsic, returns semantic::Intrinsic::kNone
-  bool SetIntrinsicIfNeeded(ast::IdentifierExpression* ident);
+  static semantic::Intrinsic MatchIntrinsic(const std::string& name);
 
  private:
   template <typename T>
@@ -176,8 +176,8 @@ class TypeDeterminer {
   bool DetermineCall(ast::CallExpression* expr);
   bool DetermineConstructor(ast::ConstructorExpression* expr);
   bool DetermineIdentifier(ast::IdentifierExpression* expr);
-  bool DetermineIntrinsic(ast::IdentifierExpression* name,
+  bool DetermineIntrinsicCall(ast::CallExpression* call,
-                          ast::CallExpression* expr);
+                              semantic::Intrinsic intrinsic);
   bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr);
   bool DetermineUnaryOp(ast::UnaryOpExpression* expr);
 

src/type_determiner_test.cc

@@ -51,6 +51,7 @@
 #include "src/ast/unary_op_expression.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/program_builder.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
 #include "src/semantic/variable.h"
@@ -646,17 +647,17 @@ TEST_F(TypeDeterminerTest, Expr_Identifier_Function_Ptr) {
   EXPECT_TRUE(TypeOf(my_var)->As<type::Pointer>()->type()->Is<type::F32>());
 }
 
-TEST_F(TypeDeterminerTest, Expr_Identifier_Function) {
+TEST_F(TypeDeterminerTest, Expr_Call_Function) {
   Func("my_func", ast::VariableList{}, ty.f32(), ast::StatementList{},
        ast::FunctionDecorationList{});
 
-  auto* ident = Expr("my_func");
+  auto* call = Call("my_func");
-  WrapInFunction(ident);
+  WrapInFunction(call);
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
 }
 
 TEST_F(TypeDeterminerTest, Expr_Identifier_Unknown) {
@@ -1521,7 +1522,7 @@ TEST_F(TypeDeterminerTest, StorageClass_NonFunctionClassError) {
 
 struct IntrinsicData {
   const char* name;
-  ast::Intrinsic intrinsic;
+  semantic::Intrinsic intrinsic;
 };
 inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
   out << data.name;
@@ -1531,94 +1532,95 @@ using IntrinsicDataTest = TypeDeterminerTestWithParam<IntrinsicData>;
 TEST_P(IntrinsicDataTest, Lookup) {
   auto param = GetParam();
 
-  auto* ident = Expr(param.name);
+  EXPECT_EQ(TypeDeterminer::MatchIntrinsic(param.name), param.intrinsic);
-  EXPECT_TRUE(td()->SetIntrinsicIfNeeded(ident));
-  EXPECT_EQ(ident->intrinsic(), param.intrinsic);
-  EXPECT_TRUE(ident->IsIntrinsic());
 }
 INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     IntrinsicDataTest,
     testing::Values(
-        IntrinsicData{"abs", ast::Intrinsic::kAbs},
+        IntrinsicData{"abs", semantic::Intrinsic::kAbs},
-        IntrinsicData{"acos", ast::Intrinsic::kAcos},
+        IntrinsicData{"acos", semantic::Intrinsic::kAcos},
-        IntrinsicData{"all", ast::Intrinsic::kAll},
+        IntrinsicData{"all", semantic::Intrinsic::kAll},
-        IntrinsicData{"any", ast::Intrinsic::kAny},
+        IntrinsicData{"any", semantic::Intrinsic::kAny},
-        IntrinsicData{"arrayLength", ast::Intrinsic::kArrayLength},
+        IntrinsicData{"arrayLength", semantic::Intrinsic::kArrayLength},
-        IntrinsicData{"asin", ast::Intrinsic::kAsin},
+        IntrinsicData{"asin", semantic::Intrinsic::kAsin},
-        IntrinsicData{"atan", ast::Intrinsic::kAtan},
+        IntrinsicData{"atan", semantic::Intrinsic::kAtan},
-        IntrinsicData{"atan2", ast::Intrinsic::kAtan2},
+        IntrinsicData{"atan2", semantic::Intrinsic::kAtan2},
-        IntrinsicData{"ceil", ast::Intrinsic::kCeil},
+        IntrinsicData{"ceil", semantic::Intrinsic::kCeil},
-        IntrinsicData{"clamp", ast::Intrinsic::kClamp},
+        IntrinsicData{"clamp", semantic::Intrinsic::kClamp},
-        IntrinsicData{"cos", ast::Intrinsic::kCos},
+        IntrinsicData{"cos", semantic::Intrinsic::kCos},
-        IntrinsicData{"cosh", ast::Intrinsic::kCosh},
+        IntrinsicData{"cosh", semantic::Intrinsic::kCosh},
-        IntrinsicData{"countOneBits", ast::Intrinsic::kCountOneBits},
+        IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits},
-        IntrinsicData{"cross", ast::Intrinsic::kCross},
+        IntrinsicData{"cross", semantic::Intrinsic::kCross},
-        IntrinsicData{"determinant", ast::Intrinsic::kDeterminant},
+        IntrinsicData{"determinant", semantic::Intrinsic::kDeterminant},
-        IntrinsicData{"distance", ast::Intrinsic::kDistance},
+        IntrinsicData{"distance", semantic::Intrinsic::kDistance},
-        IntrinsicData{"dot", ast::Intrinsic::kDot},
+        IntrinsicData{"dot", semantic::Intrinsic::kDot},
-        IntrinsicData{"dpdx", ast::Intrinsic::kDpdx},
+        IntrinsicData{"dpdx", semantic::Intrinsic::kDpdx},
-        IntrinsicData{"dpdxCoarse", ast::Intrinsic::kDpdxCoarse},
+        IntrinsicData{"dpdxCoarse", semantic::Intrinsic::kDpdxCoarse},
-        IntrinsicData{"dpdxFine", ast::Intrinsic::kDpdxFine},
+        IntrinsicData{"dpdxFine", semantic::Intrinsic::kDpdxFine},
-        IntrinsicData{"dpdy", ast::Intrinsic::kDpdy},
+        IntrinsicData{"dpdy", semantic::Intrinsic::kDpdy},
-        IntrinsicData{"dpdyCoarse", ast::Intrinsic::kDpdyCoarse},
+        IntrinsicData{"dpdyCoarse", semantic::Intrinsic::kDpdyCoarse},
-        IntrinsicData{"dpdyFine", ast::Intrinsic::kDpdyFine},
+        IntrinsicData{"dpdyFine", semantic::Intrinsic::kDpdyFine},
-        IntrinsicData{"exp", ast::Intrinsic::kExp},
+        IntrinsicData{"exp", semantic::Intrinsic::kExp},
-        IntrinsicData{"exp2", ast::Intrinsic::kExp2},
+        IntrinsicData{"exp2", semantic::Intrinsic::kExp2},
-        IntrinsicData{"faceForward", ast::Intrinsic::kFaceForward},
+        IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward},
-        IntrinsicData{"floor", ast::Intrinsic::kFloor},
+        IntrinsicData{"floor", semantic::Intrinsic::kFloor},
-        IntrinsicData{"fma", ast::Intrinsic::kFma},
+        IntrinsicData{"fma", semantic::Intrinsic::kFma},
-        IntrinsicData{"fract", ast::Intrinsic::kFract},
+        IntrinsicData{"fract", semantic::Intrinsic::kFract},
-        IntrinsicData{"frexp", ast::Intrinsic::kFrexp},
+        IntrinsicData{"frexp", semantic::Intrinsic::kFrexp},
-        IntrinsicData{"fwidth", ast::Intrinsic::kFwidth},
+        IntrinsicData{"fwidth", semantic::Intrinsic::kFwidth},
-        IntrinsicData{"fwidthCoarse", ast::Intrinsic::kFwidthCoarse},
+        IntrinsicData{"fwidthCoarse", semantic::Intrinsic::kFwidthCoarse},
-        IntrinsicData{"fwidthFine", ast::Intrinsic::kFwidthFine},
+        IntrinsicData{"fwidthFine", semantic::Intrinsic::kFwidthFine},
-        IntrinsicData{"inverseSqrt", ast::Intrinsic::kInverseSqrt},
+        IntrinsicData{"inverseSqrt", semantic::Intrinsic::kInverseSqrt},
-        IntrinsicData{"isFinite", ast::Intrinsic::kIsFinite},
+        IntrinsicData{"isFinite", semantic::Intrinsic::kIsFinite},
-        IntrinsicData{"isInf", ast::Intrinsic::kIsInf},
+        IntrinsicData{"isInf", semantic::Intrinsic::kIsInf},
-        IntrinsicData{"isNan", ast::Intrinsic::kIsNan},
+        IntrinsicData{"isNan", semantic::Intrinsic::kIsNan},
-        IntrinsicData{"isNormal", ast::Intrinsic::kIsNormal},
+        IntrinsicData{"isNormal", semantic::Intrinsic::kIsNormal},
-        IntrinsicData{"ldexp", ast::Intrinsic::kLdexp},
+        IntrinsicData{"ldexp", semantic::Intrinsic::kLdexp},
-        IntrinsicData{"length", ast::Intrinsic::kLength},
+        IntrinsicData{"length", semantic::Intrinsic::kLength},
-        IntrinsicData{"log", ast::Intrinsic::kLog},
+        IntrinsicData{"log", semantic::Intrinsic::kLog},
-        IntrinsicData{"log2", ast::Intrinsic::kLog2},
+        IntrinsicData{"log2", semantic::Intrinsic::kLog2},
-        IntrinsicData{"max", ast::Intrinsic::kMax},
+        IntrinsicData{"max", semantic::Intrinsic::kMax},
-        IntrinsicData{"min", ast::Intrinsic::kMin},
+        IntrinsicData{"min", semantic::Intrinsic::kMin},
-        IntrinsicData{"mix", ast::Intrinsic::kMix},
+        IntrinsicData{"mix", semantic::Intrinsic::kMix},
-        IntrinsicData{"modf", ast::Intrinsic::kModf},
+        IntrinsicData{"modf", semantic::Intrinsic::kModf},
-        IntrinsicData{"normalize", ast::Intrinsic::kNormalize},
+        IntrinsicData{"normalize", semantic::Intrinsic::kNormalize},
-        IntrinsicData{"pow", ast::Intrinsic::kPow},
+        IntrinsicData{"pow", semantic::Intrinsic::kPow},
-        IntrinsicData{"reflect", ast::Intrinsic::kReflect},
+        IntrinsicData{"reflect", semantic::Intrinsic::kReflect},
-        IntrinsicData{"reverseBits", ast::Intrinsic::kReverseBits},
+        IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits},
-        IntrinsicData{"round", ast::Intrinsic::kRound},
+        IntrinsicData{"round", semantic::Intrinsic::kRound},
-        IntrinsicData{"select", ast::Intrinsic::kSelect},
+        IntrinsicData{"select", semantic::Intrinsic::kSelect},
-        IntrinsicData{"sign", ast::Intrinsic::kSign},
+        IntrinsicData{"sign", semantic::Intrinsic::kSign},
-        IntrinsicData{"sin", ast::Intrinsic::kSin},
+        IntrinsicData{"sin", semantic::Intrinsic::kSin},
-        IntrinsicData{"sinh", ast::Intrinsic::kSinh},
+        IntrinsicData{"sinh", semantic::Intrinsic::kSinh},
-        IntrinsicData{"smoothStep", ast::Intrinsic::kSmoothStep},
+        IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep},
-        IntrinsicData{"sqrt", ast::Intrinsic::kSqrt},
+        IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt},
-        IntrinsicData{"step", ast::Intrinsic::kStep},
+        IntrinsicData{"step", semantic::Intrinsic::kStep},
-        IntrinsicData{"tan", ast::Intrinsic::kTan},
+        IntrinsicData{"tan", semantic::Intrinsic::kTan},
-        IntrinsicData{"tanh", ast::Intrinsic::kTanh},
+        IntrinsicData{"tanh", semantic::Intrinsic::kTanh},
-        IntrinsicData{"textureDimensions", ast::Intrinsic::kTextureDimensions},
+        IntrinsicData{"textureDimensions",
-        IntrinsicData{"textureLoad", ast::Intrinsic::kTextureLoad},
+                      semantic::Intrinsic::kTextureDimensions},
-        IntrinsicData{"textureNumLayers", ast::Intrinsic::kTextureNumLayers},
+        IntrinsicData{"textureLoad", semantic::Intrinsic::kTextureLoad},
-        IntrinsicData{"textureNumLevels", ast::Intrinsic::kTextureNumLevels},
+        IntrinsicData{"textureNumLayers",
-        IntrinsicData{"textureNumSamples", ast::Intrinsic::kTextureNumSamples},
+                      semantic::Intrinsic::kTextureNumLayers},
-        IntrinsicData{"textureSample", ast::Intrinsic::kTextureSample},
+        IntrinsicData{"textureNumLevels",
-        IntrinsicData{"textureSampleBias", ast::Intrinsic::kTextureSampleBias},
+                      semantic::Intrinsic::kTextureNumLevels},
+        IntrinsicData{"textureNumSamples",
+                      semantic::Intrinsic::kTextureNumSamples},
+        IntrinsicData{"textureSample", semantic::Intrinsic::kTextureSample},
+        IntrinsicData{"textureSampleBias",
+                      semantic::Intrinsic::kTextureSampleBias},
         IntrinsicData{"textureSampleCompare",
-                      ast::Intrinsic::kTextureSampleCompare},
+                      semantic::Intrinsic::kTextureSampleCompare},
-        IntrinsicData{"textureSampleGrad", ast::Intrinsic::kTextureSampleGrad},
+        IntrinsicData{"textureSampleGrad",
+                      semantic::Intrinsic::kTextureSampleGrad},
         IntrinsicData{"textureSampleLevel",
-                      ast::Intrinsic::kTextureSampleLevel},
+                      semantic::Intrinsic::kTextureSampleLevel},
-        IntrinsicData{"trunc", ast::Intrinsic::kTrunc}));
+        IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
 
-TEST_F(TypeDeterminerTest, IntrinsicNotSetIfNotMatched) {
+TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) {
-  auto* ident = Expr("not_intrinsic");
+  EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"),
-  EXPECT_FALSE(td()->SetIntrinsicIfNeeded(ident));
+            semantic::Intrinsic::kNone);
-  EXPECT_EQ(ident->intrinsic(), ast::Intrinsic::kNone);
-  EXPECT_FALSE(ident->IsIntrinsic());
 }
 
 using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1631,8 +1633,8 @@ TEST_P(ImportData_SingleParamTest, Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_P(ImportData_SingleParamTest, Vector) {
@@ -1644,9 +1646,9 @@ TEST_P(ImportData_SingleParamTest, Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_SingleParamTest, Error_NoParams) {
@@ -1665,28 +1667,29 @@ TEST_P(ImportData_SingleParamTest, Error_NoParams) {
 INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     ImportData_SingleParamTest,
-    testing::Values(IntrinsicData{"acos", ast::Intrinsic::kAcos},
+    testing::Values(IntrinsicData{"acos", semantic::Intrinsic::kAcos},
-                    IntrinsicData{"asin", ast::Intrinsic::kAsin},
+                    IntrinsicData{"asin", semantic::Intrinsic::kAsin},
-                    IntrinsicData{"atan", ast::Intrinsic::kAtan},
+                    IntrinsicData{"atan", semantic::Intrinsic::kAtan},
-                    IntrinsicData{"ceil", ast::Intrinsic::kCeil},
+                    IntrinsicData{"ceil", semantic::Intrinsic::kCeil},
-                    IntrinsicData{"cos", ast::Intrinsic::kCos},
+                    IntrinsicData{"cos", semantic::Intrinsic::kCos},
-                    IntrinsicData{"cosh", ast::Intrinsic::kCosh},
+                    IntrinsicData{"cosh", semantic::Intrinsic::kCosh},
-                    IntrinsicData{"exp", ast::Intrinsic::kExp},
+                    IntrinsicData{"exp", semantic::Intrinsic::kExp},
-                    IntrinsicData{"exp2", ast::Intrinsic::kExp2},
+                    IntrinsicData{"exp2", semantic::Intrinsic::kExp2},
-                    IntrinsicData{"floor", ast::Intrinsic::kFloor},
+                    IntrinsicData{"floor", semantic::Intrinsic::kFloor},
-                    IntrinsicData{"fract", ast::Intrinsic::kFract},
+                    IntrinsicData{"fract", semantic::Intrinsic::kFract},
-                    IntrinsicData{"inverseSqrt", ast::Intrinsic::kInverseSqrt},
+                    IntrinsicData{"inverseSqrt",
-                    IntrinsicData{"log", ast::Intrinsic::kLog},
+                                  semantic::Intrinsic::kInverseSqrt},
-                    IntrinsicData{"log2", ast::Intrinsic::kLog2},
+                    IntrinsicData{"log", semantic::Intrinsic::kLog},
-                    IntrinsicData{"normalize", ast::Intrinsic::kNormalize},
+                    IntrinsicData{"log2", semantic::Intrinsic::kLog2},
-                    IntrinsicData{"round", ast::Intrinsic::kRound},
+                    IntrinsicData{"normalize", semantic::Intrinsic::kNormalize},
-                    IntrinsicData{"sign", ast::Intrinsic::kSign},
+                    IntrinsicData{"round", semantic::Intrinsic::kRound},
-                    IntrinsicData{"sin", ast::Intrinsic::kSin},
+                    IntrinsicData{"sign", semantic::Intrinsic::kSign},
-                    IntrinsicData{"sinh", ast::Intrinsic::kSinh},
+                    IntrinsicData{"sin", semantic::Intrinsic::kSin},
-                    IntrinsicData{"sqrt", ast::Intrinsic::kSqrt},
+                    IntrinsicData{"sinh", semantic::Intrinsic::kSinh},
-                    IntrinsicData{"tan", ast::Intrinsic::kTan},
+                    IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt},
-                    IntrinsicData{"tanh", ast::Intrinsic::kTanh},
+                    IntrinsicData{"tan", semantic::Intrinsic::kTan},
-                    IntrinsicData{"trunc", ast::Intrinsic::kTrunc}));
+                    IntrinsicData{"tanh", semantic::Intrinsic::kTanh},
+                    IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
 
 using ImportData_SingleParam_FloatOrInt_Test =
     TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1699,8 +1702,8 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Float_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Float_Vector) {
@@ -1712,9 +1715,9 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Float_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Scalar) {
@@ -1726,8 +1729,8 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Vector) {
@@ -1747,9 +1750,9 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Scalar) {
@@ -1764,8 +1767,8 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Vector) {
@@ -1777,9 +1780,9 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_SingleParam_FloatOrInt_Test, Error_NoParams) {
@@ -1797,8 +1800,8 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Error_NoParams) {
 
 INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
                          ImportData_SingleParam_FloatOrInt_Test,
-                         testing::Values(IntrinsicData{"abs",
+                         testing::Values(IntrinsicData{
-                                                       ast::Intrinsic::kAbs}));
+                             "abs", semantic::Intrinsic::kAbs}));
 
 TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) {
   auto* ident = Expr("length");
@@ -1808,8 +1811,8 @@ TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_F(TypeDeterminerTest, ImportData_Length_FloatVector) {
@@ -1823,8 +1826,8 @@ TEST_F(TypeDeterminerTest, ImportData_Length_FloatVector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 using ImportData_TwoParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1837,8 +1840,8 @@ TEST_P(ImportData_TwoParamTest, Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_P(ImportData_TwoParamTest, Vector) {
@@ -1851,9 +1854,9 @@ TEST_P(ImportData_TwoParamTest, Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 TEST_P(ImportData_TwoParamTest, Error_NoParams) {
   auto param = GetParam();
@@ -1870,10 +1873,10 @@ TEST_P(ImportData_TwoParamTest, Error_NoParams) {
 INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     ImportData_TwoParamTest,
-    testing::Values(IntrinsicData{"atan2", ast::Intrinsic::kAtan2},
+    testing::Values(IntrinsicData{"atan2", semantic::Intrinsic::kAtan2},
-                    IntrinsicData{"pow", ast::Intrinsic::kPow},
+                    IntrinsicData{"pow", semantic::Intrinsic::kPow},
-                    IntrinsicData{"step", ast::Intrinsic::kStep},
+                    IntrinsicData{"step", semantic::Intrinsic::kStep},
-                    IntrinsicData{"reflect", ast::Intrinsic::kReflect}));
+                    IntrinsicData{"reflect", semantic::Intrinsic::kReflect}));
 
 TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) {
   auto* ident = Expr("distance");
@@ -1883,8 +1886,8 @@ TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_F(TypeDeterminerTest, ImportData_Distance_Vector) {
@@ -1896,8 +1899,8 @@ TEST_F(TypeDeterminerTest, ImportData_Distance_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
 }
 
 TEST_F(TypeDeterminerTest, ImportData_Cross) {
@@ -1909,9 +1912,9 @@ TEST_F(TypeDeterminerTest, ImportData_Cross) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_F(TypeDeterminerTest, ImportData_Cross_AutoType) {
@@ -1922,9 +1925,9 @@ TEST_F(TypeDeterminerTest, ImportData_Cross_AutoType) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 using ImportData_ThreeParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1937,8 +1940,8 @@ TEST_P(ImportData_ThreeParamTest, Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_P(ImportData_ThreeParamTest, Vector) {
@@ -1951,9 +1954,9 @@ TEST_P(ImportData_ThreeParamTest, Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 TEST_P(ImportData_ThreeParamTest, Error_NoParams) {
   auto param = GetParam();
@@ -1971,11 +1974,11 @@ TEST_P(ImportData_ThreeParamTest, Error_NoParams) {
 INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     ImportData_ThreeParamTest,
-    testing::Values(IntrinsicData{"mix", ast::Intrinsic::kMix},
+    testing::Values(
-                    IntrinsicData{"smoothStep", ast::Intrinsic::kSmoothStep},
+        IntrinsicData{"mix", semantic::Intrinsic::kMix},
-                    IntrinsicData{"fma", ast::Intrinsic::kFma},
+        IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep},
-                    IntrinsicData{"faceForward",
+        IntrinsicData{"fma", semantic::Intrinsic::kFma},
-                                  ast::Intrinsic::kFaceForward}));
+        IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward}));
 
 using ImportData_ThreeParam_FloatOrInt_Test =
     TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1988,8 +1991,8 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Float_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_float_scalar());
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Float_Vector) {
@@ -2002,9 +2005,9 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Float_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Scalar) {
@@ -2016,8 +2019,8 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Vector) {
@@ -2030,9 +2033,9 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Scalar) {
@@ -2044,8 +2047,8 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Vector) {
@@ -2058,9 +2061,9 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Error_NoParams) {
@@ -2079,7 +2082,7 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Error_NoParams) {
 INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
                          ImportData_ThreeParam_FloatOrInt_Test,
                          testing::Values(IntrinsicData{
-                             "clamp", ast::Intrinsic::kClamp}));
+                             "clamp", semantic::Intrinsic::kClamp}));
 
 using ImportData_Int_SingleParamTest =
     TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2092,8 +2095,8 @@ TEST_P(ImportData_Int_SingleParamTest, Scalar) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_integer_scalar());
+  EXPECT_TRUE(TypeOf(call)->is_integer_scalar());
 }
 
 TEST_P(ImportData_Int_SingleParamTest, Vector) {
@@ -2105,9 +2108,9 @@ TEST_P(ImportData_Int_SingleParamTest, Vector) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_Int_SingleParamTest, Error_NoParams) {
@@ -2127,8 +2130,8 @@ INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     ImportData_Int_SingleParamTest,
     testing::Values(
-        IntrinsicData{"countOneBits", ast::Intrinsic::kCountOneBits},
+        IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits},
-        IntrinsicData{"reverseBits", ast::Intrinsic::kReverseBits}));
+        IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits}));
 
 using ImportData_FloatOrInt_TwoParamTest =
     TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2141,8 +2144,8 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Signed) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Unsigned) {
@@ -2154,8 +2157,8 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Unsigned) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Float) {
@@ -2167,8 +2170,8 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Float) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Signed) {
@@ -2180,9 +2183,9 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Signed) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Unsigned) {
@@ -2194,9 +2197,9 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Unsigned) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
+  EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Float) {
@@ -2208,9 +2211,9 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Float) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->is_float_vector());
+  EXPECT_TRUE(TypeOf(call)->is_float_vector());
-  EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+  EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
 }
 
 TEST_P(ImportData_FloatOrInt_TwoParamTest, Error_NoParams) {
@@ -2229,8 +2232,8 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Error_NoParams) {
 INSTANTIATE_TEST_SUITE_P(
     TypeDeterminerTest,
     ImportData_FloatOrInt_TwoParamTest,
-    testing::Values(IntrinsicData{"min", ast::Intrinsic::kMin},
+    testing::Values(IntrinsicData{"min", semantic::Intrinsic::kMin},
-                    IntrinsicData{"max", ast::Intrinsic::kMax}));
+                    IntrinsicData{"max", semantic::Intrinsic::kMax}));
 
 TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) {
   Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
@@ -2242,8 +2245,8 @@ TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) {
 
   EXPECT_TRUE(td()->Determine()) << td()->error();
 
-  ASSERT_NE(TypeOf(ident), nullptr);
+  ASSERT_NE(TypeOf(call), nullptr);
-  EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+  EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
 }
 
 using ImportData_Matrix_OneParam_Test =
@@ -2263,7 +2266,8 @@ TEST_P(ImportData_Matrix_OneParam_Test, NoParams) {
 INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
                          ImportData_Matrix_OneParam_Test,
                          testing::Values(IntrinsicData{
-                             "determinant", ast::Intrinsic::kDeterminant}));
+                             "determinant",
+                             semantic::Intrinsic::kDeterminant}));
 
 TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) {
   // fn b() {}
@@ -2360,37 +2364,38 @@ INSTANTIATE_TEST_SUITE_P(
     testing::ValuesIn(ast::intrinsic::test::TextureOverloadCase::ValidCases()));
 
 std::string to_str(const std::string& function,
-                   const ast::intrinsic::TextureSignature* sig) {
+                   const semantic::TextureIntrinsicCall::Parameters& params) {
   struct Parameter {
     size_t idx;
     std::string name;
   };
-  std::vector<Parameter> params;
+  std::vector<Parameter> list;
-  auto maybe_add_param = [&params](size_t idx, const char* name) {
+  auto maybe_add_param = [&list](size_t idx, const char* name) {
-    if (idx != ast::intrinsic::TextureSignature::Parameters::kNotUsed) {
+    if (idx !=
-      params.emplace_back(Parameter{idx, name});
+        semantic::TextureIntrinsicCall::Parameters::Parameters::kNotUsed) {
+      list.emplace_back(Parameter{idx, name});
     }
   };
-  maybe_add_param(sig->params.idx.array_index, "array_index");
+  maybe_add_param(params.idx.array_index, "array_index");
-  maybe_add_param(sig->params.idx.bias, "bias");
+  maybe_add_param(params.idx.bias, "bias");
-  maybe_add_param(sig->params.idx.coords, "coords");
+  maybe_add_param(params.idx.coords, "coords");
-  maybe_add_param(sig->params.idx.depth_ref, "depth_ref");
+  maybe_add_param(params.idx.depth_ref, "depth_ref");
-  maybe_add_param(sig->params.idx.ddx, "ddx");
+  maybe_add_param(params.idx.ddx, "ddx");
-  maybe_add_param(sig->params.idx.ddy, "ddy");
+  maybe_add_param(params.idx.ddy, "ddy");
-  maybe_add_param(sig->params.idx.level, "level");
+  maybe_add_param(params.idx.level, "level");
-  maybe_add_param(sig->params.idx.offset, "offset");
+  maybe_add_param(params.idx.offset, "offset");
-  maybe_add_param(sig->params.idx.sampler, "sampler");
+  maybe_add_param(params.idx.sampler, "sampler");
-  maybe_add_param(sig->params.idx.sample_index, "sample_index");
+  maybe_add_param(params.idx.sample_index, "sample_index");
-  maybe_add_param(sig->params.idx.texture, "texture");
+  maybe_add_param(params.idx.texture, "texture");
-  maybe_add_param(sig->params.idx.value, "value");
+  maybe_add_param(params.idx.value, "value");
   std::sort(
-      params.begin(), params.end(),
+      list.begin(), list.end(),
       [](const Parameter& a, const Parameter& b) { return a.idx < b.idx; });
 
   std::stringstream out;
   out << function << "(";
   bool first = true;
-  for (auto& param : params) {
+  for (auto& param : list) {
     if (!first) {
       out << ", ";
     }
@@ -2727,11 +2732,12 @@ TEST_P(TypeDeterminerTextureIntrinsicTest, Call) {
     }
   }
 
-  auto* sig = static_cast<const ast::intrinsic::TextureSignature*>(
+  auto* sem = Sem().Get(call);
-      ident->intrinsic_signature());
+  ASSERT_NE(sem, nullptr);
-  ASSERT_NE(sig, nullptr);
+  auto* intrinsic = sem->As<semantic::TextureIntrinsicCall>();
+  ASSERT_NE(intrinsic, nullptr);
 
-  auto got = ::tint::to_str(param.function, sig);
+  auto got = ::tint::to_str(param.function, intrinsic->Params());
   auto* expected = expected_texture_overload(param.overload);
   EXPECT_EQ(got, expected);
 }

src/validator/validator_impl.cc

@@ -22,7 +22,6 @@
 #include "src/ast/fallthrough_statement.h"
 #include "src/ast/function.h"
 #include "src/ast/int_literal.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/module.h"
 #include "src/ast/sint_literal.h"
 #include "src/ast/stage_decoration.h"
@@ -30,7 +29,9 @@
 #include "src/ast/switch_statement.h"
 #include "src/ast/uint_literal.h"
 #include "src/ast/variable_decl_statement.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
+#include "src/semantic/intrinsic.h"
 #include "src/semantic/variable.h"
 #include "src/type/alias_type.h"
 #include "src/type/array_type.h"
@@ -62,7 +63,7 @@ enum class IntrinsicDataType {
 };
 
 struct IntrinsicData {
-  ast::Intrinsic intrinsic;
+  semantic::Intrinsic intrinsic;
   uint32_t param_count;
   IntrinsicDataType data_type;
   uint32_t vector_size;
@@ -72,102 +73,112 @@ struct IntrinsicData {
 // Note, this isn't all the intrinsics. Some are handled specially before
 // we get to the generic code. See the ValidateCallExpr code below.
 constexpr const IntrinsicData kIntrinsicData[] = {
-    {ast::Intrinsic::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+    {semantic::Intrinsic::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector,
+     0, true},
+    {semantic::Intrinsic::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kAll, 1, IntrinsicDataType::kBoolVector, 0, false},
+    {semantic::Intrinsic::kAny, 1, IntrinsicDataType::kBoolVector, 0, false},
+    {semantic::Intrinsic::kArrayLength, 1, IntrinsicDataType::kMixed, 0, false},
+    {semantic::Intrinsic::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kAll, 1, IntrinsicDataType::kBoolVector, 0, false},
+    {semantic::Intrinsic::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-    {ast::Intrinsic::kAny, 1, IntrinsicDataType::kBoolVector, 0, false},
-    {ast::Intrinsic::kArrayLength, 1, IntrinsicDataType::kMixed, 0, false},
-    {ast::Intrinsic::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kAtan2, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kAtan2, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kClamp, 3,
+     IntrinsicDataType::kFloatOrIntScalarOrVector, 0, true},
+    {semantic::Intrinsic::kCos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kClamp, 3, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+    {semantic::Intrinsic::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kCos, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kCountOneBits, 1,
-    {ast::Intrinsic::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     IntrinsicDataType::kIntScalarOrVector, 0, true},
-     true},
+    {semantic::Intrinsic::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
-    {ast::Intrinsic::kCountOneBits, 1, IntrinsicDataType::kIntScalarOrVector, 0,
+    {semantic::Intrinsic::kDeterminant, 1, IntrinsicDataType::kMatrix, 0,
-     true},
-    {ast::Intrinsic::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
-    {ast::Intrinsic::kDeterminant, 1, IntrinsicDataType::kMatrix, 0, false},
-    {ast::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      false},
-    {ast::Intrinsic::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
+    {semantic::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector,
-    {ast::Intrinsic::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     0, false},
+    {semantic::Intrinsic::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
+    {semantic::Intrinsic::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kDpdxCoarse, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kDpdxCoarse, 1,
-     true},
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
-    {ast::Intrinsic::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector,
-     true},
-    {ast::Intrinsic::kDpdy, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-     true},
-    {ast::Intrinsic::kDpdyCoarse, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-     true},
-    {ast::Intrinsic::kDpdyFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-     true},
-    {ast::Intrinsic::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
-    {ast::Intrinsic::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-     true},
-    {ast::Intrinsic::kFaceForward, 3, IntrinsicDataType::kFloatScalarOrVector,
      0, true},
-    {ast::Intrinsic::kFloor, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kDpdy, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kFma, 3, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kDpdyCoarse, 1,
-    {ast::Intrinsic::kFract, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
-     true},
+    {semantic::Intrinsic::kDpdyFine, 1, IntrinsicDataType::kFloatScalarOrVector,
-    {ast::Intrinsic::kFrexp, 2, IntrinsicDataType::kMixed, 0, false},
-    {ast::Intrinsic::kFwidth, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-     true},
-    {ast::Intrinsic::kFwidthCoarse, 1, IntrinsicDataType::kFloatScalarOrVector,
      0, true},
-    {ast::Intrinsic::kFwidthFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kInverseSqrt, 1, IntrinsicDataType::kFloatScalarOrVector,
+    {semantic::Intrinsic::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     true},
+    {semantic::Intrinsic::kFaceForward, 3,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kFloor, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     true},
+    {semantic::Intrinsic::kFma, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
+     true},
+    {semantic::Intrinsic::kFract, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     true},
+    {semantic::Intrinsic::kFrexp, 2, IntrinsicDataType::kMixed, 0, false},
+    {semantic::Intrinsic::kFwidth, 1, IntrinsicDataType::kFloatScalarOrVector,
      0, true},
-    {ast::Intrinsic::kLdexp, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kFwidthCoarse, 1,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kFwidthFine, 1,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kInverseSqrt, 1,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kLdexp, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kLength, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kLength, 1, IntrinsicDataType::kFloatScalarOrVector,
-     false},
+     0, false},
-    {ast::Intrinsic::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-    {ast::Intrinsic::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kMax, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+    {semantic::Intrinsic::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kMin, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+    {semantic::Intrinsic::kMax, 2, IntrinsicDataType::kFloatOrIntScalarOrVector,
+     0, true},
+    {semantic::Intrinsic::kMin, 2, IntrinsicDataType::kFloatOrIntScalarOrVector,
+     0, true},
+    {semantic::Intrinsic::kMix, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kMix, 3, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
-    {ast::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 0, true},
+    {semantic::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 0,
-    {ast::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
-    {ast::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kReverseBits, 1, IntrinsicDataType::kIntScalarOrVector, 0,
+    {semantic::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector,
+     0, true},
+    {semantic::Intrinsic::kReverseBits, 1,
+     IntrinsicDataType::kIntScalarOrVector, 0, true},
+    {semantic::Intrinsic::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kSelect, 3, IntrinsicDataType::kMixed, 0, false},
+    {semantic::Intrinsic::kSelect, 3, IntrinsicDataType::kMixed, 0, false},
-    {ast::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-    {ast::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kSmoothStep, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kSmoothStep, 3,
+     IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+    {semantic::Intrinsic::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
-    {ast::Intrinsic::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
-    {ast::Intrinsic::kTrunc, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+    {semantic::Intrinsic::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+     true},
+    {semantic::Intrinsic::kTrunc, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
      true},
 };
 
@@ -646,19 +657,24 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
     return false;
   }
 
-  if (auto* ident = expr->func()->As<ast::IdentifierExpression>()) {
+  auto* call_sem = program_->Sem().Get(expr);
-    auto symbol = ident->symbol();
+  if (call_sem == nullptr) {
-    if (ident->IsIntrinsic()) {
+    add_error(expr->source(), "CallExpression is missing semantic information");
+    return false;
+  }
+
+  if (auto* intrinsic_sem = call_sem->As<semantic::IntrinsicCall>()) {
     const IntrinsicData* data = nullptr;
     for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) {
-        if (ident->intrinsic() == kIntrinsicData[i].intrinsic) {
+      if (intrinsic_sem->intrinsic() == kIntrinsicData[i].intrinsic) {
         data = &kIntrinsicData[i];
         break;
       }
     }
 
     if (data != nullptr) {
-        const auto builtin = program_->Symbols().NameFor(symbol);
+      std::string builtin =
+          semantic::intrinsic::str(intrinsic_sem->intrinsic());
       if (expr->params().size() != data->param_count) {
         add_error(expr->source(),
                   "incorrect number of parameters for " + builtin +
@@ -669,14 +685,14 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
 
       if (data->all_types_match) {
         // Check that the type is an acceptable one.
-          if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
+        if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
-                           builtin, data->data_type, data->vector_size, this)) {
+                         data->data_type, data->vector_size, this)) {
           return false;
         }
 
         // Check that all params match the result type.
         for (uint32_t i = 0; i < data->param_count; ++i) {
-            if (program_->TypeOf(expr->func())->UnwrapPtrIfNeeded() !=
+          if (program_->TypeOf(expr)->UnwrapPtrIfNeeded() !=
               program_->TypeOf(expr->params()[i])->UnwrapPtrIfNeeded()) {
             add_error(expr->params()[i]->source(),
                       "expected parameter " + std::to_string(i) +
@@ -697,8 +713,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
           for (uint32_t i = 1; i < expr->params().size(); ++i) {
             if (program_->TypeOf(p0)->UnwrapPtrIfNeeded() !=
                 program_->TypeOf(expr->params()[i])->UnwrapPtrIfNeeded()) {
-                add_error(
+              add_error(expr->source(),
-                    expr->source(),
                         "parameter " + std::to_string(i) +
                             "'s unwrapped type must match parameter 0's type");
               return false;
@@ -706,7 +721,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
           }
         } else {
           // Special cases.
-            if (data->intrinsic == ast::Intrinsic::kFrexp) {
+          if (data->intrinsic == semantic::Intrinsic::kFrexp) {
             auto* p0 = expr->params()[0];
             auto* p1 = expr->params()[1];
             auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded();
@@ -717,8 +732,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
               return false;
             }
             if (!IsValidType(t1, p1->source(), builtin,
-                               IntrinsicDataType::kIntScalarOrVector, 0,
+                             IntrinsicDataType::kIntScalarOrVector, 0, this)) {
-                               this)) {
               return false;
             }
 
@@ -749,14 +763,11 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
             }
           }
 
-            if (data->intrinsic == ast::Intrinsic::kSelect) {
+          if (data->intrinsic == semantic::Intrinsic::kSelect) {
-              auto* type = program_->TypeOf(expr->func());
+            auto* type = program_->TypeOf(expr);
-              auto* t0 =
+            auto* t0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
-                  program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
+            auto* t1 = program_->TypeOf(expr->params()[1])->UnwrapPtrIfNeeded();
-              auto* t1 =
+            auto* t2 = program_->TypeOf(expr->params()[2])->UnwrapPtrIfNeeded();
-                  program_->TypeOf(expr->params()[1])->UnwrapPtrIfNeeded();
-              auto* t2 =
-                  program_->TypeOf(expr->params()[2])->UnwrapPtrIfNeeded();
             if (!type->is_scalar() && !type->Is<type::Vector>()) {
               add_error(expr->source(),
                         "incorrect type for " + builtin +
@@ -772,8 +783,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
             }
 
             if (!t2->is_bool_scalar_or_vector()) {
-                add_error(
+              add_error(expr->params()[2]->source(),
-                    expr->params()[2]->source(),
                         "incorrect type for " + builtin +
                             ". Selector must be a bool scalar or vector value");
               return false;
@@ -799,10 +809,8 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
             }
           }
 
-            if (data->intrinsic == ast::Intrinsic::kArrayLength) {
+          if (data->intrinsic == semantic::Intrinsic::kArrayLength) {
-              if (!program_->TypeOf(expr->func())
+            if (!program_->TypeOf(expr)->UnwrapPtrIfNeeded()->Is<type::U32>()) {
-                       ->UnwrapPtrIfNeeded()
-                       ->Is<type::U32>()) {
               add_error(
                   expr->source(),
                   "incorrect type for " + builtin +
@@ -810,8 +818,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
               return false;
             }
 
-              auto* p0 =
+            auto* p0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
-                  program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
             if (!p0->Is<type::Array>() ||
                 !p0->As<type::Array>()->IsRuntimeArray()) {
               add_error(expr->params()[0]->source(),
@@ -823,49 +830,46 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
         }
 
         // Result types don't match parameter types.
-          if (data->intrinsic == ast::Intrinsic::kAll ||
+        if (data->intrinsic == semantic::Intrinsic::kAll ||
-              data->intrinsic == ast::Intrinsic::kAny) {
+            data->intrinsic == semantic::Intrinsic::kAny) {
-            if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
+          if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
-                             builtin, IntrinsicDataType::kBoolScalar, 0,
+                           IntrinsicDataType::kBoolScalar, 0, this)) {
-                             this)) {
             return false;
           }
         }
 
-          if (data->intrinsic == ast::Intrinsic::kDot) {
+        if (data->intrinsic == semantic::Intrinsic::kDot) {
-            if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
+          if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
-                             builtin, IntrinsicDataType::kFloatScalar, 0,
+                           IntrinsicDataType::kFloatScalar, 0, this)) {
-                             this)) {
             return false;
           }
         }
 
-          if (data->intrinsic == ast::Intrinsic::kLength ||
+        if (data->intrinsic == semantic::Intrinsic::kLength ||
-              data->intrinsic == ast::Intrinsic::kDistance ||
+            data->intrinsic == semantic::Intrinsic::kDistance ||
-              data->intrinsic == ast::Intrinsic::kDeterminant) {
+            data->intrinsic == semantic::Intrinsic::kDeterminant) {
-            if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
+          if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
-                             builtin, IntrinsicDataType::kFloatScalar, 0,
+                           IntrinsicDataType::kFloatScalar, 0, this)) {
-                             this)) {
             return false;
           }
         }
 
         // Must be a square matrix.
-          if (data->intrinsic == ast::Intrinsic::kDeterminant) {
+        if (data->intrinsic == semantic::Intrinsic::kDeterminant) {
           const auto* matrix =
               program_->TypeOf(expr->params()[0])->As<type::Matrix>();
           if (matrix->rows() != matrix->columns()) {
-              add_error(expr->params()[0]->source(),
+            add_error(
-                        "incorrect type for " + builtin +
+                expr->params()[0]->source(),
-                            ". Requires a square matrix");
+                "incorrect type for " + builtin + ". Requires a square matrix");
             return false;
           }
         }
       }
 
       // Last parameter must be a pointer.
-        if (data->intrinsic == ast::Intrinsic::kFrexp ||
+      if (data->intrinsic == semantic::Intrinsic::kFrexp ||
-            data->intrinsic == ast::Intrinsic::kModf) {
+          data->intrinsic == semantic::Intrinsic::kModf) {
         auto* last_param = expr->params()[data->param_count - 1];
         if (!program_->TypeOf(last_param)->Is<type::Pointer>()) {
           add_error(last_param->source(), "incorrect type for " + builtin +
@@ -874,7 +878,11 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
         }
       }
     }
-    } else {
+    return true;
+  }
+
+  if (auto* ident = expr->func()->As<ast::IdentifierExpression>()) {
+    auto symbol = ident->symbol();
     if (!function_stack_.has(symbol)) {
       add_error(expr->source(), "v-0005",
                 "function must be declared before use: '" +
@@ -887,7 +895,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
                     program_->Symbols().NameFor(symbol) + "'");
       return false;
     }
-    }
+
   } else {
     add_error(expr->source(), "Invalid function call expression");
     return false;

src/writer/hlsl/generator_impl.cc

@@ -45,6 +45,7 @@
 #include "src/ast/variable.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/program_builder.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
 #include "src/semantic/variable.h"
@@ -540,19 +541,20 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
     return 0;
   }
 
-  if (ident->IsIntrinsic()) {
+  auto* call_sem = builder_.Sem().Get(expr);
+  if (auto* sem = call_sem->As<semantic::TextureIntrinsicCall>()) {
+    return EmitTextureCall(pre, out, expr, sem);
+  }
+  if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
     const auto& params = expr->params();
-    if (ident->intrinsic() == ast::Intrinsic::kSelect) {
+    if (sem->intrinsic() == semantic::Intrinsic::kSelect) {
       error_ = "select not supported in HLSL backend yet";
       return false;
-    } else if (ident->intrinsic() == ast::Intrinsic::kIsNormal) {
+    } else if (sem->intrinsic() == semantic::Intrinsic::kIsNormal) {
       error_ = "is_normal not supported in HLSL backend yet";
       return false;
     } else {
-      if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
+      auto name = generate_builtin_name(sem);
-        return EmitTextureCall(pre, out, expr);
-      }
-      auto name = generate_builtin_name(ident);
       if (name.empty()) {
         return false;
       }
@@ -634,30 +636,29 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
 
 bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
                                     std::ostream& out,
-                                    ast::CallExpression* expr) {
+                                    ast::CallExpression* expr,
+                                    const semantic::TextureIntrinsicCall* sem) {
   auto* ident = expr->func()->As<ast::IdentifierExpression>();
 
   auto params = expr->params();
-  auto* signature = static_cast<const ast::intrinsic::TextureSignature*>(
+  auto& pidx = sem->Params().idx;
-      ident->intrinsic_signature());
+  auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
-  auto& pidx = signature->params.idx;
-  auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
 
   auto* texture = params[pidx.texture];
   auto* texture_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>();
 
-  switch (ident->intrinsic()) {
+  switch (sem->intrinsic()) {
-    case ast::Intrinsic::kTextureDimensions:
+    case semantic::Intrinsic::kTextureDimensions:
-    case ast::Intrinsic::kTextureNumLayers:
+    case semantic::Intrinsic::kTextureNumLayers:
-    case ast::Intrinsic::kTextureNumLevels:
+    case semantic::Intrinsic::kTextureNumLevels:
-    case ast::Intrinsic::kTextureNumSamples: {
+    case semantic::Intrinsic::kTextureNumSamples: {
       // All of these intrinsics use the GetDimensions() method on the texture
       int num_dimensions = 0;
       const char* swizzle = "";
       bool add_mip_level_in = false;
 
-      switch (ident->intrinsic()) {
+      switch (sem->intrinsic()) {
-        case ast::Intrinsic::kTextureDimensions:
+        case semantic::Intrinsic::kTextureDimensions:
           switch (texture_type->dim()) {
             case type::TextureDimension::kNone:
               error_ = "texture dimension is kNone";
@@ -693,7 +694,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
               break;
           }
           break;
-        case ast::Intrinsic::kTextureNumLayers:
+        case semantic::Intrinsic::kTextureNumLayers:
           switch (texture_type->dim()) {
             default:
               error_ = "texture dimension is not arrayed";
@@ -709,7 +710,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
               break;
           }
           break;
-        case ast::Intrinsic::kTextureNumLevels:
+        case semantic::Intrinsic::kTextureNumLevels:
           add_mip_level_in = true;
           switch (texture_type->dim()) {
             default:
@@ -728,7 +729,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
               break;
           }
           break;
-        case ast::Intrinsic::kTextureNumSamples:
+        case semantic::Intrinsic::kTextureNumSamples:
           switch (texture_type->dim()) {
             default:
               error_ = "texture dimension does not support multisampling";
@@ -797,29 +798,29 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
 
   bool pack_mip_in_coords = false;
 
-  switch (ident->intrinsic()) {
+  switch (sem->intrinsic()) {
-    case ast::Intrinsic::kTextureSample:
+    case semantic::Intrinsic::kTextureSample:
       out << ".Sample(";
       break;
-    case ast::Intrinsic::kTextureSampleBias:
+    case semantic::Intrinsic::kTextureSampleBias:
       out << ".SampleBias(";
       break;
-    case ast::Intrinsic::kTextureSampleLevel:
+    case semantic::Intrinsic::kTextureSampleLevel:
       out << ".SampleLevel(";
       break;
-    case ast::Intrinsic::kTextureSampleGrad:
+    case semantic::Intrinsic::kTextureSampleGrad:
       out << ".SampleGrad(";
       break;
-    case ast::Intrinsic::kTextureSampleCompare:
+    case semantic::Intrinsic::kTextureSampleCompare:
       out << ".SampleCmp(";
       break;
-    case ast::Intrinsic::kTextureLoad:
+    case semantic::Intrinsic::kTextureLoad:
       out << ".Load(";
       if (!texture_type->Is<type::StorageTexture>()) {
         pack_mip_in_coords = true;
       }
       break;
-    case ast::Intrinsic::kTextureStore:
+    case semantic::Intrinsic::kTextureStore:
       out << "[";
       break;
     default:
@@ -875,7 +876,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
     }
   }
 
-  if (ident->intrinsic() == ast::Intrinsic::kTextureStore) {
+  if (sem->intrinsic() == semantic::Intrinsic::kTextureStore) {
     out << "] = ";
     if (!EmitExpression(pre, out, params[pidx.value]))
       return false;
@@ -887,102 +888,102 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
 }  // namespace hlsl
 
 std::string GeneratorImpl::generate_builtin_name(
-    ast::IdentifierExpression* ident) {
+    const semantic::IntrinsicCall* call) {
   std::string out;
-  switch (ident->intrinsic()) {
+  switch (call->intrinsic()) {
-    case ast::Intrinsic::kAcos:
+    case semantic::Intrinsic::kAcos:
-    case ast::Intrinsic::kAny:
+    case semantic::Intrinsic::kAny:
-    case ast::Intrinsic::kAll:
+    case semantic::Intrinsic::kAll:
-    case ast::Intrinsic::kAsin:
+    case semantic::Intrinsic::kAsin:
-    case ast::Intrinsic::kAtan:
+    case semantic::Intrinsic::kAtan:
-    case ast::Intrinsic::kAtan2:
+    case semantic::Intrinsic::kAtan2:
-    case ast::Intrinsic::kCeil:
+    case semantic::Intrinsic::kCeil:
-    case ast::Intrinsic::kCos:
+    case semantic::Intrinsic::kCos:
-    case ast::Intrinsic::kCosh:
+    case semantic::Intrinsic::kCosh:
-    case ast::Intrinsic::kCross:
+    case semantic::Intrinsic::kCross:
-    case ast::Intrinsic::kDeterminant:
+    case semantic::Intrinsic::kDeterminant:
-    case ast::Intrinsic::kDistance:
+    case semantic::Intrinsic::kDistance:
-    case ast::Intrinsic::kDot:
+    case semantic::Intrinsic::kDot:
-    case ast::Intrinsic::kExp:
+    case semantic::Intrinsic::kExp:
-    case ast::Intrinsic::kExp2:
+    case semantic::Intrinsic::kExp2:
-    case ast::Intrinsic::kFloor:
+    case semantic::Intrinsic::kFloor:
-    case ast::Intrinsic::kFma:
+    case semantic::Intrinsic::kFma:
-    case ast::Intrinsic::kLdexp:
+    case semantic::Intrinsic::kLdexp:
-    case ast::Intrinsic::kLength:
+    case semantic::Intrinsic::kLength:
-    case ast::Intrinsic::kLog:
+    case semantic::Intrinsic::kLog:
-    case ast::Intrinsic::kLog2:
+    case semantic::Intrinsic::kLog2:
-    case ast::Intrinsic::kNormalize:
+    case semantic::Intrinsic::kNormalize:
-    case ast::Intrinsic::kPow:
+    case semantic::Intrinsic::kPow:
-    case ast::Intrinsic::kReflect:
+    case semantic::Intrinsic::kReflect:
-    case ast::Intrinsic::kRound:
+    case semantic::Intrinsic::kRound:
-    case ast::Intrinsic::kSin:
+    case semantic::Intrinsic::kSin:
-    case ast::Intrinsic::kSinh:
+    case semantic::Intrinsic::kSinh:
-    case ast::Intrinsic::kSqrt:
+    case semantic::Intrinsic::kSqrt:
-    case ast::Intrinsic::kStep:
+    case semantic::Intrinsic::kStep:
-    case ast::Intrinsic::kTan:
+    case semantic::Intrinsic::kTan:
-    case ast::Intrinsic::kTanh:
+    case semantic::Intrinsic::kTanh:
-    case ast::Intrinsic::kTrunc:
+    case semantic::Intrinsic::kTrunc:
-    case ast::Intrinsic::kMix:
+    case semantic::Intrinsic::kMix:
-    case ast::Intrinsic::kSign:
+    case semantic::Intrinsic::kSign:
-    case ast::Intrinsic::kAbs:
+    case semantic::Intrinsic::kAbs:
-    case ast::Intrinsic::kMax:
+    case semantic::Intrinsic::kMax:
-    case ast::Intrinsic::kMin:
+    case semantic::Intrinsic::kMin:
-    case ast::Intrinsic::kClamp:
+    case semantic::Intrinsic::kClamp:
-      out = builder_.Symbols().NameFor(ident->symbol());
+      out = semantic::intrinsic::str(call->intrinsic());
       break;
-    case ast::Intrinsic::kCountOneBits:
+    case semantic::Intrinsic::kCountOneBits:
       out = "countbits";
       break;
-    case ast::Intrinsic::kDpdx:
+    case semantic::Intrinsic::kDpdx:
       out = "ddx";
       break;
-    case ast::Intrinsic::kDpdxCoarse:
+    case semantic::Intrinsic::kDpdxCoarse:
       out = "ddx_coarse";
       break;
-    case ast::Intrinsic::kDpdxFine:
+    case semantic::Intrinsic::kDpdxFine:
       out = "ddx_fine";
       break;
-    case ast::Intrinsic::kDpdy:
+    case semantic::Intrinsic::kDpdy:
       out = "ddy";
       break;
-    case ast::Intrinsic::kDpdyCoarse:
+    case semantic::Intrinsic::kDpdyCoarse:
       out = "ddy_coarse";
       break;
-    case ast::Intrinsic::kDpdyFine:
+    case semantic::Intrinsic::kDpdyFine:
       out = "ddy_fine";
       break;
-    case ast::Intrinsic::kFaceForward:
+    case semantic::Intrinsic::kFaceForward:
       out = "faceforward";
       break;
-    case ast::Intrinsic::kFract:
+    case semantic::Intrinsic::kFract:
       out = "frac";
       break;
-    case ast::Intrinsic::kFwidth:
+    case semantic::Intrinsic::kFwidth:
-    case ast::Intrinsic::kFwidthCoarse:
+    case semantic::Intrinsic::kFwidthCoarse:
-    case ast::Intrinsic::kFwidthFine:
+    case semantic::Intrinsic::kFwidthFine:
       out = "fwidth";
       break;
-    case ast::Intrinsic::kInverseSqrt:
+    case semantic::Intrinsic::kInverseSqrt:
       out = "rsqrt";
       break;
-    case ast::Intrinsic::kIsFinite:
+    case semantic::Intrinsic::kIsFinite:
       out = "isfinite";
       break;
-    case ast::Intrinsic::kIsInf:
+    case semantic::Intrinsic::kIsInf:
       out = "isinf";
       break;
-    case ast::Intrinsic::kIsNan:
+    case semantic::Intrinsic::kIsNan:
       out = "isnan";
       break;
-    case ast::Intrinsic::kReverseBits:
+    case semantic::Intrinsic::kReverseBits:
       out = "reversebits";
       break;
-    case ast::Intrinsic::kSmoothStep:
+    case semantic::Intrinsic::kSmoothStep:
       out = "smoothstep";
       break;
     default:
       error_ = "Unknown builtin method: " +
-               builder_.Symbols().NameFor(ident->symbol());
+               std::string(semantic::intrinsic::str(call->intrinsic()));
       return "";
   }
 

src/writer/hlsl/generator_impl.h

@@ -30,7 +30,6 @@
 #include "src/ast/discard_statement.h"
 #include "src/ast/identifier_expression.h"
 #include "src/ast/if_statement.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/literal.h"
 #include "src/ast/loop_statement.h"
 #include "src/ast/member_accessor_expression.h"
@@ -41,10 +40,18 @@
 #include "src/ast/unary_op_expression.h"
 #include "src/program_builder.h"
 #include "src/scope_stack.h"
+#include "src/semantic/intrinsic.h"
 #include "src/type/struct_type.h"
 #include "src/writer/hlsl/namer.h"
 
 namespace tint {
+
+// Forward declarations
+namespace semantic {
+class TextureIntrinsicCall;
+class IntrinsicCall;
+}  // namespace semantic
+
 namespace writer {
 namespace hlsl {
 
@@ -146,10 +153,12 @@ class GeneratorImpl {
   /// @param pre the preamble for the expression stream
   /// @param out the output of the expression stream
   /// @param expr the call expression
+  /// @param sem the semantic information for the texture intrinsic call
   /// @returns true if the call expression is emitted
   bool EmitTextureCall(std::ostream& pre,
                        std::ostream& out,
-                       ast::CallExpression* expr);
+                       ast::CallExpression* expr,
+                       const semantic::TextureIntrinsicCall* sem);
   /// Handles a case statement
   /// @param out the output stream
   /// @param stmt the statement
@@ -346,9 +355,9 @@ class GeneratorImpl {
   std::string generate_storage_buffer_index_expression(std::ostream& pre,
                                                        ast::Expression* expr);
   /// Handles generating a builtin method name
-  /// @param expr the expression
+  /// @param call the semantic info for the intrinsic call
   /// @returns the name or "" if not valid
-  std::string generate_builtin_name(ast::IdentifierExpression* expr);
+  std::string generate_builtin_name(const semantic::IntrinsicCall* call);
   /// Converts a builtin to an attribute name
   /// @param builtin the builtin to convert
   /// @returns the string name of the builtin or blank on error

src/writer/hlsl/generator_impl_binary_test.cc

@@ -483,6 +483,10 @@ TEST_F(HlslGeneratorImplTest_Binary, Call_WithLogical) {
 
   Func("foo", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
+  Global("a", ast::StorageClass::kNone, ty.bool_());
+  Global("b", ast::StorageClass::kNone, ty.bool_());
+  Global("c", ast::StorageClass::kNone, ty.bool_());
+  Global("d", ast::StorageClass::kNone, ty.bool_());
 
   ast::ExpressionList params;
   params.push_back(create<ast::BinaryExpression>(ast::BinaryOp::kLogicalAnd,
@@ -497,6 +501,7 @@ TEST_F(HlslGeneratorImplTest_Binary, Call_WithLogical) {
                                     Expr("d"))));
 
   auto* expr = create<ast::CallStatement>(Call("foo", params));
+  WrapInFunction(expr);
 
   GeneratorImpl& gen = Build();
 

src/writer/hlsl/generator_impl_call_test.cc

@@ -30,11 +30,12 @@ namespace {
 using HlslGeneratorImplTest_Call = TestHelper;
 
 TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithoutParams) {
-  auto* call = Call("my_func");
-
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
 
+  auto* call = Call("my_func");
+  WrapInFunction(call);
+
   GeneratorImpl& gen = Build();
 
   ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error();
@@ -42,10 +43,13 @@ TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithoutParams) {
 }
 
 TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
-  auto* call = Call("my_func", "param1", "param2");
-
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
+  Global("param1", ast::StorageClass::kNone, ty.f32());
+  Global("param2", ast::StorageClass::kNone, ty.f32());
+
+  auto* call = Call("my_func", "param1", "param2");
+  WrapInFunction(call);
 
   GeneratorImpl& gen = Build();
 
@@ -54,10 +58,13 @@ TEST_F(HlslGeneratorImplTest_Call, EmitExpression_Call_WithParams) {
 }
 
 TEST_F(HlslGeneratorImplTest_Call, EmitStatement_Call) {
-  auto* call = create<ast::CallStatement>(Call("my_func", "param1", "param2"));
-
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
+  Global("param1", ast::StorageClass::kNone, ty.f32());
+  Global("param2", ast::StorageClass::kNone, ty.f32());
+
+  auto* call = create<ast::CallStatement>(Call("my_func", "param1", "param2"));
+  WrapInFunction(call);
 
   GeneratorImpl& gen = Build();
 

src/writer/hlsl/generator_impl_intrinsic_test.cc

@@ -17,6 +17,7 @@
 #include "src/ast/call_expression.h"
 #include "src/ast/identifier_expression.h"
 #include "src/program.h"
+#include "src/semantic/call.h"
 #include "src/type/f32_type.h"
 #include "src/type/vector_type.h"
 #include "src/type_determiner.h"
@@ -36,7 +37,7 @@ enum class ParamType {
 };
 
 struct IntrinsicData {
-  ast::Intrinsic intrinsic;
+  semantic::Intrinsic intrinsic;
   ParamType type;
   const char* hlsl_name;
 };
@@ -57,92 +58,92 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
   return out;
 }
 
-ast::CallExpression* GenerateCall(ast::Intrinsic intrinsic,
+ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
                                   ParamType type,
                                   ProgramBuilder* builder) {
   std::string name;
   std::ostringstream str(name);
   str << intrinsic;
   switch (intrinsic) {
-    case ast::Intrinsic::kAcos:
+    case semantic::Intrinsic::kAcos:
-    case ast::Intrinsic::kAsin:
+    case semantic::Intrinsic::kAsin:
-    case ast::Intrinsic::kAtan:
+    case semantic::Intrinsic::kAtan:
-    case ast::Intrinsic::kCeil:
+    case semantic::Intrinsic::kCeil:
-    case ast::Intrinsic::kCos:
+    case semantic::Intrinsic::kCos:
-    case ast::Intrinsic::kCosh:
+    case semantic::Intrinsic::kCosh:
-    case ast::Intrinsic::kDpdx:
+    case semantic::Intrinsic::kDpdx:
-    case ast::Intrinsic::kDpdxCoarse:
+    case semantic::Intrinsic::kDpdxCoarse:
-    case ast::Intrinsic::kDpdxFine:
+    case semantic::Intrinsic::kDpdxFine:
-    case ast::Intrinsic::kDpdy:
+    case semantic::Intrinsic::kDpdy:
-    case ast::Intrinsic::kDpdyCoarse:
+    case semantic::Intrinsic::kDpdyCoarse:
-    case ast::Intrinsic::kDpdyFine:
+    case semantic::Intrinsic::kDpdyFine:
-    case ast::Intrinsic::kExp:
+    case semantic::Intrinsic::kExp:
-    case ast::Intrinsic::kExp2:
+    case semantic::Intrinsic::kExp2:
-    case ast::Intrinsic::kFloor:
+    case semantic::Intrinsic::kFloor:
-    case ast::Intrinsic::kFract:
+    case semantic::Intrinsic::kFract:
-    case ast::Intrinsic::kFwidth:
+    case semantic::Intrinsic::kFwidth:
-    case ast::Intrinsic::kFwidthCoarse:
+    case semantic::Intrinsic::kFwidthCoarse:
-    case ast::Intrinsic::kFwidthFine:
+    case semantic::Intrinsic::kFwidthFine:
-    case ast::Intrinsic::kInverseSqrt:
+    case semantic::Intrinsic::kInverseSqrt:
-    case ast::Intrinsic::kIsFinite:
+    case semantic::Intrinsic::kIsFinite:
-    case ast::Intrinsic::kIsInf:
+    case semantic::Intrinsic::kIsInf:
-    case ast::Intrinsic::kIsNan:
+    case semantic::Intrinsic::kIsNan:
-    case ast::Intrinsic::kIsNormal:
+    case semantic::Intrinsic::kIsNormal:
-    case ast::Intrinsic::kLdexp:
+    case semantic::Intrinsic::kLdexp:
-    case ast::Intrinsic::kLength:
+    case semantic::Intrinsic::kLength:
-    case ast::Intrinsic::kLog:
+    case semantic::Intrinsic::kLog:
-    case ast::Intrinsic::kLog2:
+    case semantic::Intrinsic::kLog2:
-    case ast::Intrinsic::kNormalize:
+    case semantic::Intrinsic::kNormalize:
-    case ast::Intrinsic::kReflect:
+    case semantic::Intrinsic::kReflect:
-    case ast::Intrinsic::kRound:
+    case semantic::Intrinsic::kRound:
-    case ast::Intrinsic::kSin:
+    case semantic::Intrinsic::kSin:
-    case ast::Intrinsic::kSinh:
+    case semantic::Intrinsic::kSinh:
-    case ast::Intrinsic::kSqrt:
+    case semantic::Intrinsic::kSqrt:
-    case ast::Intrinsic::kTan:
+    case semantic::Intrinsic::kTan:
-    case ast::Intrinsic::kTanh:
+    case semantic::Intrinsic::kTanh:
-    case ast::Intrinsic::kTrunc:
+    case semantic::Intrinsic::kTrunc:
-    case ast::Intrinsic::kSign:
+    case semantic::Intrinsic::kSign:
       return builder->Call(str.str(), "f1");
-    case ast::Intrinsic::kAtan2:
+    case semantic::Intrinsic::kAtan2:
-    case ast::Intrinsic::kCross:
+    case semantic::Intrinsic::kCross:
-    case ast::Intrinsic::kDot:
+    case semantic::Intrinsic::kDot:
-    case ast::Intrinsic::kDistance:
+    case semantic::Intrinsic::kDistance:
-    case ast::Intrinsic::kPow:
+    case semantic::Intrinsic::kPow:
-    case ast::Intrinsic::kStep:
+    case semantic::Intrinsic::kStep:
       return builder->Call(str.str(), "f1", "f2");
-    case ast::Intrinsic::kFma:
+    case semantic::Intrinsic::kFma:
-    case ast::Intrinsic::kMix:
+    case semantic::Intrinsic::kMix:
-    case ast::Intrinsic::kFaceForward:
+    case semantic::Intrinsic::kFaceForward:
-    case ast::Intrinsic::kSmoothStep:
+    case semantic::Intrinsic::kSmoothStep:
       return builder->Call(str.str(), "f1", "f2", "f3");
-    case ast::Intrinsic::kAll:
+    case semantic::Intrinsic::kAll:
-    case ast::Intrinsic::kAny:
+    case semantic::Intrinsic::kAny:
       return builder->Call(str.str(), "b1");
-    case ast::Intrinsic::kAbs:
+    case semantic::Intrinsic::kAbs:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1");
       } else {
         return builder->Call(str.str(), "u1");
       }
-    case ast::Intrinsic::kCountOneBits:
+    case semantic::Intrinsic::kCountOneBits:
-    case ast::Intrinsic::kReverseBits:
+    case semantic::Intrinsic::kReverseBits:
       return builder->Call(str.str(), "u1");
-    case ast::Intrinsic::kMax:
+    case semantic::Intrinsic::kMax:
-    case ast::Intrinsic::kMin:
+    case semantic::Intrinsic::kMin:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1", "f2");
       } else {
         return builder->Call(str.str(), "u1", "u2");
       }
-    case ast::Intrinsic::kClamp:
+    case semantic::Intrinsic::kClamp:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1", "f2", "f3");
       } else {
         return builder->Call(str.str(), "u1", "u2", "u3");
       }
-    case ast::Intrinsic::kSelect:
+    case semantic::Intrinsic::kSelect:
       return builder->Call(str.str(), "f1", "f2", "b1");
-    case ast::Intrinsic::kDeterminant:
+    case semantic::Intrinsic::kDeterminant:
       return builder->Call(str.str(), "m1");
     default:
       break;
@@ -168,80 +169,92 @@ TEST_P(HlslIntrinsicTest, Emit) {
 
   GeneratorImpl& gen = Build();
 
-  EXPECT_EQ(
+  auto* sem = program->Sem().Get(call);
-      gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
+  ASSERT_NE(sem, nullptr);
-      param.hlsl_name);
+  auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+  ASSERT_NE(intrinsic, nullptr);
+
+  EXPECT_EQ(gen.generate_builtin_name(intrinsic), param.hlsl_name);
 }
 INSTANTIATE_TEST_SUITE_P(
     HlslGeneratorImplTest_Intrinsic,
     HlslIntrinsicTest,
     testing::Values(
-        IntrinsicData{ast::Intrinsic::kAbs, ParamType::kF32, "abs"},
+        IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32, "abs"},
-        IntrinsicData{ast::Intrinsic::kAbs, ParamType::kU32, "abs"},
+        IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "abs"},
-        IntrinsicData{ast::Intrinsic::kAcos, ParamType::kF32, "acos"},
+        IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32, "acos"},
-        IntrinsicData{ast::Intrinsic::kAll, ParamType::kBool, "all"},
+        IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool, "all"},
-        IntrinsicData{ast::Intrinsic::kAny, ParamType::kBool, "any"},
+        IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool, "any"},
-        IntrinsicData{ast::Intrinsic::kAsin, ParamType::kF32, "asin"},
+        IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32, "asin"},
-        IntrinsicData{ast::Intrinsic::kAtan, ParamType::kF32, "atan"},
+        IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32, "atan"},
-        IntrinsicData{ast::Intrinsic::kAtan2, ParamType::kF32, "atan2"},
+        IntrinsicData{semantic::Intrinsic::kAtan2, ParamType::kF32, "atan2"},
-        IntrinsicData{ast::Intrinsic::kCeil, ParamType::kF32, "ceil"},
+        IntrinsicData{semantic::Intrinsic::kCeil, ParamType::kF32, "ceil"},
-        IntrinsicData{ast::Intrinsic::kClamp, ParamType::kF32, "clamp"},
+        IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kF32, "clamp"},
-        IntrinsicData{ast::Intrinsic::kClamp, ParamType::kU32, "clamp"},
+        IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kU32, "clamp"},
-        IntrinsicData{ast::Intrinsic::kCos, ParamType::kF32, "cos"},
+        IntrinsicData{semantic::Intrinsic::kCos, ParamType::kF32, "cos"},
-        IntrinsicData{ast::Intrinsic::kCosh, ParamType::kF32, "cosh"},
+        IntrinsicData{semantic::Intrinsic::kCosh, ParamType::kF32, "cosh"},
-        IntrinsicData{ast::Intrinsic::kCountOneBits, ParamType::kU32,
+        IntrinsicData{semantic::Intrinsic::kCountOneBits, ParamType::kU32,
                       "countbits"},
-        IntrinsicData{ast::Intrinsic::kCross, ParamType::kF32, "cross"},
+        IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32, "cross"},
-        IntrinsicData{ast::Intrinsic::kDeterminant, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
                       "determinant"},
-        IntrinsicData{ast::Intrinsic::kDistance, ParamType::kF32, "distance"},
+        IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kDot, ParamType::kF32, "dot"},
+                      "distance"},
-        IntrinsicData{ast::Intrinsic::kDpdx, ParamType::kF32, "ddx"},
+        IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "dot"},
-        IntrinsicData{ast::Intrinsic::kDpdxCoarse, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32, "ddx"},
+        IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32,
                       "ddx_coarse"},
-        IntrinsicData{ast::Intrinsic::kDpdxFine, ParamType::kF32, "ddx_fine"},
+        IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kDpdy, ParamType::kF32, "ddy"},
+                      "ddx_fine"},
-        IntrinsicData{ast::Intrinsic::kDpdyCoarse, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32, "ddy"},
+        IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
                       "ddy_coarse"},
-        IntrinsicData{ast::Intrinsic::kDpdyFine, ParamType::kF32, "ddy_fine"},
+        IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kExp, ParamType::kF32, "exp"},
+                      "ddy_fine"},
-        IntrinsicData{ast::Intrinsic::kExp2, ParamType::kF32, "exp2"},
+        IntrinsicData{semantic::Intrinsic::kExp, ParamType::kF32, "exp"},
-        IntrinsicData{ast::Intrinsic::kFaceForward, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kExp2, ParamType::kF32, "exp2"},
+        IntrinsicData{semantic::Intrinsic::kFaceForward, ParamType::kF32,
                       "faceforward"},
-        IntrinsicData{ast::Intrinsic::kFloor, ParamType::kF32, "floor"},
+        IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32, "floor"},
-        IntrinsicData{ast::Intrinsic::kFma, ParamType::kF32, "fma"},
+        IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "fma"},
-        IntrinsicData{ast::Intrinsic::kFract, ParamType::kF32, "frac"},
+        IntrinsicData{semantic::Intrinsic::kFract, ParamType::kF32, "frac"},
-        IntrinsicData{ast::Intrinsic::kFwidth, ParamType::kF32, "fwidth"},
+        IntrinsicData{semantic::Intrinsic::kFwidth, ParamType::kF32, "fwidth"},
-        IntrinsicData{ast::Intrinsic::kFwidthCoarse, ParamType::kF32, "fwidth"},
+        IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kFwidthFine, ParamType::kF32, "fwidth"},
+                      "fwidth"},
-        IntrinsicData{ast::Intrinsic::kInverseSqrt, ParamType::kF32, "rsqrt"},
+        IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kIsFinite, ParamType::kF32, "isfinite"},
+                      "fwidth"},
-        IntrinsicData{ast::Intrinsic::kIsInf, ParamType::kF32, "isinf"},
+        IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kIsNan, ParamType::kF32, "isnan"},
+                      "rsqrt"},
-        IntrinsicData{ast::Intrinsic::kLdexp, ParamType::kF32, "ldexp"},
+        IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kLength, ParamType::kF32, "length"},
+                      "isfinite"},
-        IntrinsicData{ast::Intrinsic::kLog, ParamType::kF32, "log"},
+        IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32, "isinf"},
-        IntrinsicData{ast::Intrinsic::kLog2, ParamType::kF32, "log2"},
+        IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32, "isnan"},
-        IntrinsicData{ast::Intrinsic::kMax, ParamType::kF32, "max"},
+        IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32, "ldexp"},
-        IntrinsicData{ast::Intrinsic::kMax, ParamType::kU32, "max"},
+        IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32, "length"},
-        IntrinsicData{ast::Intrinsic::kMin, ParamType::kF32, "min"},
+        IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "log"},
-        IntrinsicData{ast::Intrinsic::kMin, ParamType::kU32, "min"},
+        IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32, "log2"},
-        IntrinsicData{ast::Intrinsic::kNormalize, ParamType::kF32, "normalize"},
+        IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32, "max"},
-        IntrinsicData{ast::Intrinsic::kPow, ParamType::kF32, "pow"},
+        IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "max"},
-        IntrinsicData{ast::Intrinsic::kReflect, ParamType::kF32, "reflect"},
+        IntrinsicData{semantic::Intrinsic::kMin, ParamType::kF32, "min"},
-        IntrinsicData{ast::Intrinsic::kReverseBits, ParamType::kU32,
+        IntrinsicData{semantic::Intrinsic::kMin, ParamType::kU32, "min"},
+        IntrinsicData{semantic::Intrinsic::kNormalize, ParamType::kF32,
+                      "normalize"},
+        IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "pow"},
+        IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32,
+                      "reflect"},
+        IntrinsicData{semantic::Intrinsic::kReverseBits, ParamType::kU32,
                       "reversebits"},
-        IntrinsicData{ast::Intrinsic::kRound, ParamType::kU32, "round"},
+        IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32, "round"},
-        IntrinsicData{ast::Intrinsic::kSign, ParamType::kF32, "sign"},
+        IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32, "sign"},
-        IntrinsicData{ast::Intrinsic::kSin, ParamType::kF32, "sin"},
+        IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "sin"},
-        IntrinsicData{ast::Intrinsic::kSinh, ParamType::kF32, "sinh"},
+        IntrinsicData{semantic::Intrinsic::kSinh, ParamType::kF32, "sinh"},
-        IntrinsicData{ast::Intrinsic::kSmoothStep, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kSmoothStep, ParamType::kF32,
                       "smoothstep"},
-        IntrinsicData{ast::Intrinsic::kSqrt, ParamType::kF32, "sqrt"},
+        IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32, "sqrt"},
-        IntrinsicData{ast::Intrinsic::kStep, ParamType::kF32, "step"},
+        IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32, "step"},
-        IntrinsicData{ast::Intrinsic::kTan, ParamType::kF32, "tan"},
+        IntrinsicData{semantic::Intrinsic::kTan, ParamType::kF32, "tan"},
-        IntrinsicData{ast::Intrinsic::kTanh, ParamType::kF32, "tanh"},
+        IntrinsicData{semantic::Intrinsic::kTanh, ParamType::kF32, "tanh"},
-        IntrinsicData{ast::Intrinsic::kTrunc, ParamType::kF32, "trunc"}));
+        IntrinsicData{semantic::Intrinsic::kTrunc, ParamType::kF32, "trunc"}));
 
 TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) {
   FAIL();

src/writer/msl/generator_impl.cc

@@ -50,6 +50,7 @@
 #include "src/ast/variable.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/program.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
 #include "src/semantic/variable.h"
@@ -442,11 +443,12 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
     return 0;
   }
 
-  if (ident->IsIntrinsic()) {
+  auto* call_sem = program_->Sem().Get(expr);
-    if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
+  if (auto* sem = call_sem->As<semantic::TextureIntrinsicCall>()) {
-      return EmitTextureCall(expr);
+    return EmitTextureCall(expr, sem);
   }
-    auto name = generate_builtin_name(ident);
+  if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
+    auto name = generate_builtin_name(sem);
     if (name.empty()) {
       return false;
     }
@@ -555,21 +557,20 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
   return true;
 }
 
-bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
+bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
+                                    const semantic::TextureIntrinsicCall* sem) {
   auto* ident = expr->func()->As<ast::IdentifierExpression>();
 
   auto params = expr->params();
-  auto* signature = static_cast<const ast::intrinsic::TextureSignature*>(
+  auto& pidx = sem->Params().idx;
-      ident->intrinsic_signature());
+  auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
-  auto& pidx = signature->params.idx;
-  auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
 
   assert(pidx.texture != kNotUsed);
   auto* texture_type =
       TypeOf(params[pidx.texture])->UnwrapAll()->As<type::Texture>();
 
-  switch (ident->intrinsic()) {
+  switch (sem->intrinsic()) {
-    case ast::Intrinsic::kTextureDimensions: {
+    case semantic::Intrinsic::kTextureDimensions: {
       std::vector<const char*> dims;
       switch (texture_type->dim()) {
         case type::TextureDimension::kNone:
@@ -623,7 +624,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
       }
       return true;
     }
-    case ast::Intrinsic::kTextureNumLayers: {
+    case semantic::Intrinsic::kTextureNumLayers: {
       out_ << "int(";
       if (!EmitExpression(params[pidx.texture])) {
         return false;
@@ -631,7 +632,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
       out_ << ".get_array_size())";
       return true;
     }
-    case ast::Intrinsic::kTextureNumLevels: {
+    case semantic::Intrinsic::kTextureNumLevels: {
       out_ << "int(";
       if (!EmitExpression(params[pidx.texture])) {
         return false;
@@ -639,7 +640,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
       out_ << ".get_num_mip_levels())";
       return true;
     }
-    case ast::Intrinsic::kTextureNumSamples: {
+    case semantic::Intrinsic::kTextureNumSamples: {
       out_ << "int(";
       if (!EmitExpression(params[pidx.texture])) {
         return false;
@@ -656,21 +657,21 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
 
   bool lod_param_is_named = true;
 
-  switch (ident->intrinsic()) {
+  switch (sem->intrinsic()) {
-    case ast::Intrinsic::kTextureSample:
+    case semantic::Intrinsic::kTextureSample:
-    case ast::Intrinsic::kTextureSampleBias:
+    case semantic::Intrinsic::kTextureSampleBias:
-    case ast::Intrinsic::kTextureSampleLevel:
+    case semantic::Intrinsic::kTextureSampleLevel:
-    case ast::Intrinsic::kTextureSampleGrad:
+    case semantic::Intrinsic::kTextureSampleGrad:
       out_ << ".sample(";
       break;
-    case ast::Intrinsic::kTextureSampleCompare:
+    case semantic::Intrinsic::kTextureSampleCompare:
       out_ << ".sample_compare(";
       break;
-    case ast::Intrinsic::kTextureLoad:
+    case semantic::Intrinsic::kTextureLoad:
       out_ << ".read(";
       lod_param_is_named = false;
       break;
-    case ast::Intrinsic::kTextureStore:
+    case semantic::Intrinsic::kTextureStore:
       out_ << ".write(";
       break;
     default:
@@ -766,115 +767,114 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr) {
 }
 
 std::string GeneratorImpl::generate_builtin_name(
-    ast::IdentifierExpression* ident) {
+    const semantic::IntrinsicCall* call) {
-  auto* type = TypeOf(ident);
   std::string out = "metal::";
-  switch (ident->intrinsic()) {
+  switch (call->intrinsic()) {
-    case ast::Intrinsic::kAcos:
+    case semantic::Intrinsic::kAcos:
-    case ast::Intrinsic::kAll:
+    case semantic::Intrinsic::kAll:
-    case ast::Intrinsic::kAny:
+    case semantic::Intrinsic::kAny:
-    case ast::Intrinsic::kAsin:
+    case semantic::Intrinsic::kAsin:
-    case ast::Intrinsic::kAtan:
+    case semantic::Intrinsic::kAtan:
-    case ast::Intrinsic::kAtan2:
+    case semantic::Intrinsic::kAtan2:
-    case ast::Intrinsic::kCeil:
+    case semantic::Intrinsic::kCeil:
-    case ast::Intrinsic::kCos:
+    case semantic::Intrinsic::kCos:
-    case ast::Intrinsic::kCosh:
+    case semantic::Intrinsic::kCosh:
-    case ast::Intrinsic::kCross:
+    case semantic::Intrinsic::kCross:
-    case ast::Intrinsic::kDeterminant:
+    case semantic::Intrinsic::kDeterminant:
-    case ast::Intrinsic::kDistance:
+    case semantic::Intrinsic::kDistance:
-    case ast::Intrinsic::kDot:
+    case semantic::Intrinsic::kDot:
-    case ast::Intrinsic::kExp:
+    case semantic::Intrinsic::kExp:
-    case ast::Intrinsic::kExp2:
+    case semantic::Intrinsic::kExp2:
-    case ast::Intrinsic::kFloor:
+    case semantic::Intrinsic::kFloor:
-    case ast::Intrinsic::kFma:
+    case semantic::Intrinsic::kFma:
-    case ast::Intrinsic::kFract:
+    case semantic::Intrinsic::kFract:
-    case ast::Intrinsic::kLength:
+    case semantic::Intrinsic::kLength:
-    case ast::Intrinsic::kLdexp:
+    case semantic::Intrinsic::kLdexp:
-    case ast::Intrinsic::kLog:
+    case semantic::Intrinsic::kLog:
-    case ast::Intrinsic::kLog2:
+    case semantic::Intrinsic::kLog2:
-    case ast::Intrinsic::kMix:
+    case semantic::Intrinsic::kMix:
-    case ast::Intrinsic::kNormalize:
+    case semantic::Intrinsic::kNormalize:
-    case ast::Intrinsic::kPow:
+    case semantic::Intrinsic::kPow:
-    case ast::Intrinsic::kReflect:
+    case semantic::Intrinsic::kReflect:
-    case ast::Intrinsic::kRound:
+    case semantic::Intrinsic::kRound:
-    case ast::Intrinsic::kSelect:
+    case semantic::Intrinsic::kSelect:
-    case ast::Intrinsic::kSin:
+    case semantic::Intrinsic::kSin:
-    case ast::Intrinsic::kSinh:
+    case semantic::Intrinsic::kSinh:
-    case ast::Intrinsic::kSqrt:
+    case semantic::Intrinsic::kSqrt:
-    case ast::Intrinsic::kStep:
+    case semantic::Intrinsic::kStep:
-    case ast::Intrinsic::kTan:
+    case semantic::Intrinsic::kTan:
-    case ast::Intrinsic::kTanh:
+    case semantic::Intrinsic::kTanh:
-    case ast::Intrinsic::kTrunc:
+    case semantic::Intrinsic::kTrunc:
-    case ast::Intrinsic::kSign:
+    case semantic::Intrinsic::kSign:
-    case ast::Intrinsic::kClamp:
+    case semantic::Intrinsic::kClamp:
-      out += program_->Symbols().NameFor(ident->symbol());
+      out += semantic::intrinsic::str(call->intrinsic());
       break;
-    case ast::Intrinsic::kAbs:
+    case semantic::Intrinsic::kAbs:
-      if (type->is_float_scalar_or_vector()) {
+      if (call->Type()->is_float_scalar_or_vector()) {
         out += "fabs";
       } else {
         out += "abs";
       }
       break;
-    case ast::Intrinsic::kCountOneBits:
+    case semantic::Intrinsic::kCountOneBits:
       out += "popcount";
       break;
-    case ast::Intrinsic::kDpdx:
+    case semantic::Intrinsic::kDpdx:
-    case ast::Intrinsic::kDpdxCoarse:
+    case semantic::Intrinsic::kDpdxCoarse:
-    case ast::Intrinsic::kDpdxFine:
+    case semantic::Intrinsic::kDpdxFine:
       out += "dfdx";
       break;
-    case ast::Intrinsic::kDpdy:
+    case semantic::Intrinsic::kDpdy:
-    case ast::Intrinsic::kDpdyCoarse:
+    case semantic::Intrinsic::kDpdyCoarse:
-    case ast::Intrinsic::kDpdyFine:
+    case semantic::Intrinsic::kDpdyFine:
       out += "dfdy";
       break;
-    case ast::Intrinsic::kFwidth:
+    case semantic::Intrinsic::kFwidth:
-    case ast::Intrinsic::kFwidthCoarse:
+    case semantic::Intrinsic::kFwidthCoarse:
-    case ast::Intrinsic::kFwidthFine:
+    case semantic::Intrinsic::kFwidthFine:
       out += "fwidth";
       break;
-    case ast::Intrinsic::kIsFinite:
+    case semantic::Intrinsic::kIsFinite:
       out += "isfinite";
       break;
-    case ast::Intrinsic::kIsInf:
+    case semantic::Intrinsic::kIsInf:
       out += "isinf";
       break;
-    case ast::Intrinsic::kIsNan:
+    case semantic::Intrinsic::kIsNan:
       out += "isnan";
       break;
-    case ast::Intrinsic::kIsNormal:
+    case semantic::Intrinsic::kIsNormal:
       out += "isnormal";
       break;
-    case ast::Intrinsic::kMax:
+    case semantic::Intrinsic::kMax:
-      if (type->is_float_scalar_or_vector()) {
+      if (call->Type()->is_float_scalar_or_vector()) {
         out += "fmax";
       } else {
         out += "max";
       }
       break;
-    case ast::Intrinsic::kMin:
+    case semantic::Intrinsic::kMin:
-      if (type->is_float_scalar_or_vector()) {
+      if (call->Type()->is_float_scalar_or_vector()) {
         out += "fmin";
       } else {
         out += "min";
       }
       break;
-    case ast::Intrinsic::kFaceForward:
+    case semantic::Intrinsic::kFaceForward:
       out += "faceforward";
       break;
-    case ast::Intrinsic::kReverseBits:
+    case semantic::Intrinsic::kReverseBits:
       out += "reverse_bits";
       break;
-    case ast::Intrinsic::kSmoothStep:
+    case semantic::Intrinsic::kSmoothStep:
       out += "smoothstep";
       break;
-    case ast::Intrinsic::kInverseSqrt:
+    case semantic::Intrinsic::kInverseSqrt:
       out += "rsqrt";
       break;
     default:
       error_ = "Unknown import method: " +
-               program_->Symbols().NameFor(ident->symbol());
+               std::string(semantic::intrinsic::str(call->intrinsic()));
       return "";
   }
   return out;

src/writer/msl/generator_impl.h

@@ -31,7 +31,6 @@
 #include "src/ast/else_statement.h"
 #include "src/ast/identifier_expression.h"
 #include "src/ast/if_statement.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/literal.h"
 #include "src/ast/loop_statement.h"
 #include "src/ast/member_accessor_expression.h"
@@ -42,11 +41,19 @@
 #include "src/ast/unary_op_expression.h"
 #include "src/program.h"
 #include "src/scope_stack.h"
+#include "src/semantic/intrinsic.h"
 #include "src/type/struct_type.h"
 #include "src/writer/msl/namer.h"
 #include "src/writer/text_generator.h"
 
 namespace tint {
+
+// Forward declarations
+namespace semantic {
+class TextureIntrinsicCall;
+class IntrinsicCall;
+}  // namespace semantic
+
 namespace writer {
 namespace msl {
 
@@ -114,8 +121,10 @@ class GeneratorImpl : public TextGenerator {
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param expr the call expression
+  /// @param sem the semantic information for the texture intrinsic call
   /// @returns true if the call expression is emitted
-  bool EmitTextureCall(ast::CallExpression* expr);
+  bool EmitTextureCall(ast::CallExpression* expr,
+                       const semantic::TextureIntrinsicCall* sem);
   /// Handles a case statement
   /// @param stmt the statement
   /// @returns true if the statement was emitted successfully
@@ -249,9 +258,9 @@ class GeneratorImpl : public TextGenerator {
   /// @returns the name
   std::string generate_name(const std::string& prefix);
   /// Handles generating a builtin name
-  /// @param ident the identifier to build the name from
+  /// @param call the semantic info for the intrinsic call
   /// @returns the name or "" if not valid
-  std::string generate_builtin_name(ast::IdentifierExpression* ident);
+  std::string generate_builtin_name(const semantic::IntrinsicCall* call);
 
   /// Checks if the global variable is in an input or output struct
   /// @param var the variable to check

src/writer/msl/generator_impl_call_test.cc

@@ -32,10 +32,12 @@ namespace {
 using MslGeneratorImplTest = TestHelper;
 
 TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithoutParams) {
-  auto* call = Call("my_func");
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
 
+  auto* call = Call("my_func");
+  WrapInFunction(call);
+
   GeneratorImpl& gen = Build();
 
   ASSERT_TRUE(gen.EmitExpression(call)) << gen.error();
@@ -43,9 +45,13 @@ TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithoutParams) {
 }
 
 TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithParams) {
-  auto* call = Call("my_func", "param1", "param2");
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
+  Global("param1", ast::StorageClass::kNone, ty.f32());
+  Global("param2", ast::StorageClass::kNone, ty.f32());
+
+  auto* call = Call("my_func", "param1", "param2");
+  WrapInFunction(call);
 
   GeneratorImpl& gen = Build();
 
@@ -54,16 +60,19 @@ TEST_F(MslGeneratorImplTest, EmitExpression_Call_WithParams) {
 }
 
 TEST_F(MslGeneratorImplTest, EmitStatement_Call) {
-  auto* call = Call("my_func", "param1", "param2");
   Func("my_func", ast::VariableList{}, ty.void_(), ast::StatementList{},
        ast::FunctionDecorationList{});
+  Global("param1", ast::StorageClass::kNone, ty.f32());
+  Global("param2", ast::StorageClass::kNone, ty.f32());
 
-  auto* expr = create<ast::CallStatement>(call);
+  auto* call = Call("my_func", "param1", "param2");
+  auto* stmt = create<ast::CallStatement>(call);
+  WrapInFunction(stmt);
 
   GeneratorImpl& gen = Build();
 
   gen.increment_indent();
-  ASSERT_TRUE(gen.EmitStatement(expr)) << gen.error();
+  ASSERT_TRUE(gen.EmitStatement(stmt)) << gen.error();
   EXPECT_EQ(gen.result(), "  my_func(param1, param2);\n");
 }
 

src/writer/msl/generator_impl_import_test.cc

@@ -24,6 +24,7 @@
 #include "src/ast/sint_literal.h"
 #include "src/ast/type_constructor_expression.h"
 #include "src/program.h"
+#include "src/semantic/call.h"
 #include "src/type/f32_type.h"
 #include "src/type/i32_type.h"
 #include "src/type/matrix_type.h"
@@ -57,8 +58,12 @@ TEST_P(MslImportData_SingleParamTest, FloatScalar) {
 
   GeneratorImpl& gen = Build();
 
-  ASSERT_EQ(
+  auto* sem = program->Sem().Get(call);
-      gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
+  ASSERT_NE(sem, nullptr);
+  auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+  ASSERT_NE(intrinsic, nullptr);
+
+  ASSERT_EQ(gen.generate_builtin_name(intrinsic),
             std::string("metal::") + param.msl_name);
 }
 INSTANTIATE_TEST_SUITE_P(MslGeneratorImplTest,

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -18,6 +18,7 @@
 #include "src/ast/call_expression.h"
 #include "src/ast/identifier_expression.h"
 #include "src/program.h"
+#include "src/semantic/call.h"
 #include "src/type/f32_type.h"
 #include "src/type/vector_type.h"
 #include "src/type_determiner.h"
@@ -38,7 +39,7 @@ enum class ParamType {
 };
 
 struct IntrinsicData {
-  ast::Intrinsic intrinsic;
+  semantic::Intrinsic intrinsic;
   ParamType type;
   const char* msl_name;
 };
@@ -59,92 +60,92 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
   return out;
 }
 
-ast::CallExpression* GenerateCall(ast::Intrinsic intrinsic,
+ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
                                   ParamType type,
                                   ProgramBuilder* builder) {
   std::string name;
   std::ostringstream str(name);
   str << intrinsic;
   switch (intrinsic) {
-    case ast::Intrinsic::kAcos:
+    case semantic::Intrinsic::kAcos:
-    case ast::Intrinsic::kAsin:
+    case semantic::Intrinsic::kAsin:
-    case ast::Intrinsic::kAtan:
+    case semantic::Intrinsic::kAtan:
-    case ast::Intrinsic::kCeil:
+    case semantic::Intrinsic::kCeil:
-    case ast::Intrinsic::kCos:
+    case semantic::Intrinsic::kCos:
-    case ast::Intrinsic::kCosh:
+    case semantic::Intrinsic::kCosh:
-    case ast::Intrinsic::kDpdx:
+    case semantic::Intrinsic::kDpdx:
-    case ast::Intrinsic::kDpdxCoarse:
+    case semantic::Intrinsic::kDpdxCoarse:
-    case ast::Intrinsic::kDpdxFine:
+    case semantic::Intrinsic::kDpdxFine:
-    case ast::Intrinsic::kDpdy:
+    case semantic::Intrinsic::kDpdy:
-    case ast::Intrinsic::kDpdyCoarse:
+    case semantic::Intrinsic::kDpdyCoarse:
-    case ast::Intrinsic::kDpdyFine:
+    case semantic::Intrinsic::kDpdyFine:
-    case ast::Intrinsic::kExp:
+    case semantic::Intrinsic::kExp:
-    case ast::Intrinsic::kExp2:
+    case semantic::Intrinsic::kExp2:
-    case ast::Intrinsic::kFloor:
+    case semantic::Intrinsic::kFloor:
-    case ast::Intrinsic::kFract:
+    case semantic::Intrinsic::kFract:
-    case ast::Intrinsic::kFwidth:
+    case semantic::Intrinsic::kFwidth:
-    case ast::Intrinsic::kFwidthCoarse:
+    case semantic::Intrinsic::kFwidthCoarse:
-    case ast::Intrinsic::kFwidthFine:
+    case semantic::Intrinsic::kFwidthFine:
-    case ast::Intrinsic::kInverseSqrt:
+    case semantic::Intrinsic::kInverseSqrt:
-    case ast::Intrinsic::kIsFinite:
+    case semantic::Intrinsic::kIsFinite:
-    case ast::Intrinsic::kIsInf:
+    case semantic::Intrinsic::kIsInf:
-    case ast::Intrinsic::kIsNan:
+    case semantic::Intrinsic::kIsNan:
-    case ast::Intrinsic::kIsNormal:
+    case semantic::Intrinsic::kIsNormal:
-    case ast::Intrinsic::kLdexp:
+    case semantic::Intrinsic::kLdexp:
-    case ast::Intrinsic::kLength:
+    case semantic::Intrinsic::kLength:
-    case ast::Intrinsic::kLog:
+    case semantic::Intrinsic::kLog:
-    case ast::Intrinsic::kLog2:
+    case semantic::Intrinsic::kLog2:
-    case ast::Intrinsic::kNormalize:
+    case semantic::Intrinsic::kNormalize:
-    case ast::Intrinsic::kReflect:
+    case semantic::Intrinsic::kReflect:
-    case ast::Intrinsic::kRound:
+    case semantic::Intrinsic::kRound:
-    case ast::Intrinsic::kSin:
+    case semantic::Intrinsic::kSin:
-    case ast::Intrinsic::kSinh:
+    case semantic::Intrinsic::kSinh:
-    case ast::Intrinsic::kSqrt:
+    case semantic::Intrinsic::kSqrt:
-    case ast::Intrinsic::kTan:
+    case semantic::Intrinsic::kTan:
-    case ast::Intrinsic::kTanh:
+    case semantic::Intrinsic::kTanh:
-    case ast::Intrinsic::kTrunc:
+    case semantic::Intrinsic::kTrunc:
-    case ast::Intrinsic::kSign:
+    case semantic::Intrinsic::kSign:
       return builder->Call(str.str(), "f1");
-    case ast::Intrinsic::kAtan2:
+    case semantic::Intrinsic::kAtan2:
-    case ast::Intrinsic::kCross:
+    case semantic::Intrinsic::kCross:
-    case ast::Intrinsic::kDot:
+    case semantic::Intrinsic::kDot:
-    case ast::Intrinsic::kDistance:
+    case semantic::Intrinsic::kDistance:
-    case ast::Intrinsic::kPow:
+    case semantic::Intrinsic::kPow:
-    case ast::Intrinsic::kStep:
+    case semantic::Intrinsic::kStep:
       return builder->Call(str.str(), "f1", "f2");
-    case ast::Intrinsic::kFma:
+    case semantic::Intrinsic::kFma:
-    case ast::Intrinsic::kMix:
+    case semantic::Intrinsic::kMix:
-    case ast::Intrinsic::kFaceForward:
+    case semantic::Intrinsic::kFaceForward:
-    case ast::Intrinsic::kSmoothStep:
+    case semantic::Intrinsic::kSmoothStep:
       return builder->Call(str.str(), "f1", "f2", "f3");
-    case ast::Intrinsic::kAll:
+    case semantic::Intrinsic::kAll:
-    case ast::Intrinsic::kAny:
+    case semantic::Intrinsic::kAny:
       return builder->Call(str.str(), "b1");
-    case ast::Intrinsic::kAbs:
+    case semantic::Intrinsic::kAbs:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1");
       } else {
         return builder->Call(str.str(), "u1");
       }
-    case ast::Intrinsic::kCountOneBits:
+    case semantic::Intrinsic::kCountOneBits:
-    case ast::Intrinsic::kReverseBits:
+    case semantic::Intrinsic::kReverseBits:
       return builder->Call(str.str(), "u1");
-    case ast::Intrinsic::kMax:
+    case semantic::Intrinsic::kMax:
-    case ast::Intrinsic::kMin:
+    case semantic::Intrinsic::kMin:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1", "f2");
       } else {
         return builder->Call(str.str(), "u1", "u2");
       }
-    case ast::Intrinsic::kClamp:
+    case semantic::Intrinsic::kClamp:
       if (type == ParamType::kF32) {
         return builder->Call(str.str(), "f1", "f2", "f3");
       } else {
         return builder->Call(str.str(), "u1", "u2", "u3");
       }
-    case ast::Intrinsic::kSelect:
+    case semantic::Intrinsic::kSelect:
       return builder->Call(str.str(), "f1", "f2", "b1");
-    case ast::Intrinsic::kDeterminant:
+    case semantic::Intrinsic::kDeterminant:
       return builder->Call(str.str(), "m1");
     default:
       break;
@@ -171,95 +172,127 @@ TEST_P(MslIntrinsicTest, Emit) {
 
   GeneratorImpl& gen = Build();
 
-  EXPECT_EQ(
+  auto* sem = program->Sem().Get(call);
-      gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
+  ASSERT_NE(sem, nullptr);
-      param.msl_name);
+  auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+  ASSERT_NE(intrinsic, nullptr);
+
+  EXPECT_EQ(gen.generate_builtin_name(intrinsic), param.msl_name);
 }
 INSTANTIATE_TEST_SUITE_P(
     MslGeneratorImplTest,
     MslIntrinsicTest,
     testing::Values(
-        IntrinsicData{ast::Intrinsic::kAbs, ParamType::kF32, "metal::fabs"},
+        IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kAbs, ParamType::kU32, "metal::abs"},
+                      "metal::fabs"},
-        IntrinsicData{ast::Intrinsic::kAcos, ParamType::kF32, "metal::acos"},
+        IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "metal::abs"},
-        IntrinsicData{ast::Intrinsic::kAll, ParamType::kBool, "metal::all"},
+        IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kAny, ParamType::kBool, "metal::any"},
+                      "metal::acos"},
-        IntrinsicData{ast::Intrinsic::kAsin, ParamType::kF32, "metal::asin"},
+        IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool,
-        IntrinsicData{ast::Intrinsic::kAtan, ParamType::kF32, "metal::atan"},
+                      "metal::all"},
-        IntrinsicData{ast::Intrinsic::kAtan2, ParamType::kF32, "metal::atan2"},
+        IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool,
-        IntrinsicData{ast::Intrinsic::kCeil, ParamType::kF32, "metal::ceil"},
+                      "metal::any"},
-        IntrinsicData{ast::Intrinsic::kClamp, ParamType::kF32, "metal::clamp"},
+        IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kClamp, ParamType::kU32, "metal::clamp"},
+                      "metal::asin"},
-        IntrinsicData{ast::Intrinsic::kCos, ParamType::kF32, "metal::cos"},
+        IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kCosh, ParamType::kF32, "metal::cosh"},
+                      "metal::atan"},
-        IntrinsicData{ast::Intrinsic::kCountOneBits, ParamType::kU32,
+        IntrinsicData{semantic::Intrinsic::kAtan2, ParamType::kF32,
+                      "metal::atan2"},
+        IntrinsicData{semantic::Intrinsic::kCeil, ParamType::kF32,
+                      "metal::ceil"},
+        IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kF32,
+                      "metal::clamp"},
+        IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kU32,
+                      "metal::clamp"},
+        IntrinsicData{semantic::Intrinsic::kCos, ParamType::kF32, "metal::cos"},
+        IntrinsicData{semantic::Intrinsic::kCosh, ParamType::kF32,
+                      "metal::cosh"},
+        IntrinsicData{semantic::Intrinsic::kCountOneBits, ParamType::kU32,
                       "metal::popcount"},
-        IntrinsicData{ast::Intrinsic::kCross, ParamType::kF32, "metal::cross"},
+        IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kDeterminant, ParamType::kF32,
+                      "metal::cross"},
+        IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
                       "metal::determinant"},
-        IntrinsicData{ast::Intrinsic::kDistance, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32,
                       "metal::distance"},
-        IntrinsicData{ast::Intrinsic::kDot, ParamType::kF32, "metal::dot"},
+        IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "metal::dot"},
-        IntrinsicData{ast::Intrinsic::kDpdx, ParamType::kF32, "metal::dfdx"},
+        IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kDpdxCoarse, ParamType::kF32,
                       "metal::dfdx"},
-        IntrinsicData{ast::Intrinsic::kDpdxFine, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32,
                       "metal::dfdx"},
-        IntrinsicData{ast::Intrinsic::kDpdy, ParamType::kF32, "metal::dfdy"},
+        IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kDpdyCoarse, ParamType::kF32,
+                      "metal::dfdx"},
+        IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32,
                       "metal::dfdy"},
-        IntrinsicData{ast::Intrinsic::kDpdyFine, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
                       "metal::dfdy"},
-        IntrinsicData{ast::Intrinsic::kExp, ParamType::kF32, "metal::exp"},
+        IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kExp2, ParamType::kF32, "metal::exp2"},
+                      "metal::dfdy"},
-        IntrinsicData{ast::Intrinsic::kFaceForward, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kExp, ParamType::kF32, "metal::exp"},
+        IntrinsicData{semantic::Intrinsic::kExp2, ParamType::kF32,
+                      "metal::exp2"},
+        IntrinsicData{semantic::Intrinsic::kFaceForward, ParamType::kF32,
                       "metal::faceforward"},
-        IntrinsicData{ast::Intrinsic::kFloor, ParamType::kF32, "metal::floor"},
+        IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kFma, ParamType::kF32, "metal::fma"},
+                      "metal::floor"},
-        IntrinsicData{ast::Intrinsic::kFract, ParamType::kF32, "metal::fract"},
+        IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "metal::fma"},
-        IntrinsicData{ast::Intrinsic::kFwidth, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kFract, ParamType::kF32,
+                      "metal::fract"},
+        IntrinsicData{semantic::Intrinsic::kFwidth, ParamType::kF32,
                       "metal::fwidth"},
-        IntrinsicData{ast::Intrinsic::kFwidthCoarse, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32,
                       "metal::fwidth"},
-        IntrinsicData{ast::Intrinsic::kFwidthFine, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32,
                       "metal::fwidth"},
-        IntrinsicData{ast::Intrinsic::kInverseSqrt, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32,
                       "metal::rsqrt"},
-        IntrinsicData{ast::Intrinsic::kIsFinite, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32,
                       "metal::isfinite"},
-        IntrinsicData{ast::Intrinsic::kIsInf, ParamType::kF32, "metal::isinf"},
+        IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kIsNan, ParamType::kF32, "metal::isnan"},
+                      "metal::isinf"},
-        IntrinsicData{ast::Intrinsic::kIsNormal, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32,
+                      "metal::isnan"},
+        IntrinsicData{semantic::Intrinsic::kIsNormal, ParamType::kF32,
                       "metal::isnormal"},
-        IntrinsicData{ast::Intrinsic::kLdexp, ParamType::kF32, "metal::ldexp"},
+        IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kLength, ParamType::kF32,
+                      "metal::ldexp"},
+        IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32,
                       "metal::length"},
-        IntrinsicData{ast::Intrinsic::kLog, ParamType::kF32, "metal::log"},
+        IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "metal::log"},
-        IntrinsicData{ast::Intrinsic::kLog2, ParamType::kF32, "metal::log2"},
+        IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kMax, ParamType::kF32, "metal::fmax"},
+                      "metal::log2"},
-        IntrinsicData{ast::Intrinsic::kMax, ParamType::kU32, "metal::max"},
+        IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kMin, ParamType::kF32, "metal::fmin"},
+                      "metal::fmax"},
-        IntrinsicData{ast::Intrinsic::kMin, ParamType::kU32, "metal::min"},
+        IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "metal::max"},
-        IntrinsicData{ast::Intrinsic::kNormalize, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kMin, ParamType::kF32,
+                      "metal::fmin"},
+        IntrinsicData{semantic::Intrinsic::kMin, ParamType::kU32, "metal::min"},
+        IntrinsicData{semantic::Intrinsic::kNormalize, ParamType::kF32,
                       "metal::normalize"},
-        IntrinsicData{ast::Intrinsic::kPow, ParamType::kF32, "metal::pow"},
+        IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "metal::pow"},
-        IntrinsicData{ast::Intrinsic::kReflect, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32,
                       "metal::reflect"},
-        IntrinsicData{ast::Intrinsic::kReverseBits, ParamType::kU32,
+        IntrinsicData{semantic::Intrinsic::kReverseBits, ParamType::kU32,
                       "metal::reverse_bits"},
-        IntrinsicData{ast::Intrinsic::kRound, ParamType::kU32, "metal::round"},
+        IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32,
-        IntrinsicData{ast::Intrinsic::kSelect, ParamType::kF32,
+                      "metal::round"},
+        IntrinsicData{semantic::Intrinsic::kSelect, ParamType::kF32,
                       "metal::select"},
-        IntrinsicData{ast::Intrinsic::kSign, ParamType::kF32, "metal::sign"},
+        IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kSin, ParamType::kF32, "metal::sin"},
+                      "metal::sign"},
-        IntrinsicData{ast::Intrinsic::kSinh, ParamType::kF32, "metal::sinh"},
+        IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "metal::sin"},
-        IntrinsicData{ast::Intrinsic::kSmoothStep, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kSinh, ParamType::kF32,
+                      "metal::sinh"},
+        IntrinsicData{semantic::Intrinsic::kSmoothStep, ParamType::kF32,
                       "metal::smoothstep"},
-        IntrinsicData{ast::Intrinsic::kSqrt, ParamType::kF32, "metal::sqrt"},
+        IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kStep, ParamType::kF32, "metal::step"},
+                      "metal::sqrt"},
-        IntrinsicData{ast::Intrinsic::kTan, ParamType::kF32, "metal::tan"},
+        IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32,
-        IntrinsicData{ast::Intrinsic::kTanh, ParamType::kF32, "metal::tanh"},
+                      "metal::step"},
-        IntrinsicData{ast::Intrinsic::kTrunc, ParamType::kF32,
+        IntrinsicData{semantic::Intrinsic::kTan, ParamType::kF32, "metal::tan"},
+        IntrinsicData{semantic::Intrinsic::kTanh, ParamType::kF32,
+                      "metal::tanh"},
+        IntrinsicData{semantic::Intrinsic::kTrunc, ParamType::kF32,
                       "metal::trunc"}));
 
 TEST_F(MslGeneratorImplTest, Intrinsic_Call) {

src/writer/spirv/builder.cc

@@ -40,7 +40,6 @@
 #include "src/ast/group_decoration.h"
 #include "src/ast/identifier_expression.h"
 #include "src/ast/if_statement.h"
-#include "src/ast/intrinsic.h"
 #include "src/ast/location_decoration.h"
 #include "src/ast/loop_statement.h"
 #include "src/ast/member_accessor_expression.h"
@@ -59,8 +58,10 @@
 #include "src/ast/variable.h"
 #include "src/ast/variable_decl_statement.h"
 #include "src/program.h"
+#include "src/semantic/call.h"
 #include "src/semantic/expression.h"
 #include "src/semantic/function.h"
+#include "src/semantic/intrinsic.h"
 #include "src/semantic/variable.h"
 #include "src/type/access_control_type.h"
 #include "src/type/alias_type.h"
@@ -165,25 +166,26 @@ type::Matrix* GetNestedMatrixType(type::Type* type) {
   return type->As<type::Matrix>();
 }
 
-uint32_t intrinsic_to_glsl_method(type::Type* type, ast::Intrinsic intrinsic) {
+uint32_t intrinsic_to_glsl_method(type::Type* type,
+                                  semantic::Intrinsic intrinsic) {
   switch (intrinsic) {
-    case ast::Intrinsic::kAbs:
+    case semantic::Intrinsic::kAbs:
       if (type->is_float_scalar_or_vector()) {
         return GLSLstd450FAbs;
       } else {
         return GLSLstd450SAbs;
       }
-    case ast::Intrinsic::kAcos:
+    case semantic::Intrinsic::kAcos:
       return GLSLstd450Acos;
-    case ast::Intrinsic::kAsin:
+    case semantic::Intrinsic::kAsin:
       return GLSLstd450Asin;
-    case ast::Intrinsic::kAtan:
+    case semantic::Intrinsic::kAtan:
       return GLSLstd450Atan;
-    case ast::Intrinsic::kAtan2:
+    case semantic::Intrinsic::kAtan2:
       return GLSLstd450Atan2;
-    case ast::Intrinsic::kCeil:
+    case semantic::Intrinsic::kCeil:
       return GLSLstd450Ceil;
-    case ast::Intrinsic::kClamp:
+    case semantic::Intrinsic::kClamp:
       if (type->is_float_scalar_or_vector()) {
         return GLSLstd450NClamp;
       } else if (type->is_unsigned_scalar_or_vector()) {
@@ -191,41 +193,41 @@ uint32_t intrinsic_to_glsl_method(type::Type* type, ast::Intrinsic intrinsic) {
       } else {
         return GLSLstd450SClamp;
       }
-    case ast::Intrinsic::kCos:
+    case semantic::Intrinsic::kCos:
       return GLSLstd450Cos;
-    case ast::Intrinsic::kCosh:
+    case semantic::Intrinsic::kCosh:
       return GLSLstd450Cosh;
-    case ast::Intrinsic::kCross:
+    case semantic::Intrinsic::kCross:
       return GLSLstd450Cross;
-    case ast::Intrinsic::kDeterminant:
+    case semantic::Intrinsic::kDeterminant:
       return GLSLstd450Determinant;
-    case ast::Intrinsic::kDistance:
+    case semantic::Intrinsic::kDistance:
       return GLSLstd450Distance;
-    case ast::Intrinsic::kExp:
+    case semantic::Intrinsic::kExp:
       return GLSLstd450Exp;
-    case ast::Intrinsic::kExp2:
+    case semantic::Intrinsic::kExp2:
       return GLSLstd450Exp2;
-    case ast::Intrinsic::kFaceForward:
+    case semantic::Intrinsic::kFaceForward:
       return GLSLstd450FaceForward;
-    case ast::Intrinsic::kFloor:
+    case semantic::Intrinsic::kFloor:
       return GLSLstd450Floor;
-    case ast::Intrinsic::kFma:
+    case semantic::Intrinsic::kFma:
       return GLSLstd450Fma;
-    case ast::Intrinsic::kFract:
+    case semantic::Intrinsic::kFract:
       return GLSLstd450Fract;
-    case ast::Intrinsic::kFrexp:
+    case semantic::Intrinsic::kFrexp:
       return GLSLstd450Frexp;
-    case ast::Intrinsic::kInverseSqrt:
+    case semantic::Intrinsic::kInverseSqrt:
       return GLSLstd450InverseSqrt;
-    case ast::Intrinsic::kLdexp:
+    case semantic::Intrinsic::kLdexp:
       return GLSLstd450Ldexp;
-    case ast::Intrinsic::kLength:
+    case semantic::Intrinsic::kLength:
       return GLSLstd450Length;
-    case ast::Intrinsic::kLog:
+    case semantic::Intrinsic::kLog:
       return GLSLstd450Log;
-    case ast::Intrinsic::kLog2:
+    case semantic::Intrinsic::kLog2:
       return GLSLstd450Log2;
-    case ast::Intrinsic::kMax:
+    case semantic::Intrinsic::kMax:
       if (type->is_float_scalar_or_vector()) {
         return GLSLstd450NMax;
       } else if (type->is_unsigned_scalar_or_vector()) {
@@ -233,7 +235,7 @@ uint32_t intrinsic_to_glsl_method(type::Type* type, ast::Intrinsic intrinsic) {
       } else {
         return GLSLstd450SMax;
       }
-    case ast::Intrinsic::kMin:
+    case semantic::Intrinsic::kMin:
       if (type->is_float_scalar_or_vector()) {
         return GLSLstd450NMin;
       } else if (type->is_unsigned_scalar_or_vector()) {
@@ -241,35 +243,35 @@ uint32_t intrinsic_to_glsl_method(type::Type* type, ast::Intrinsic intrinsic) {
       } else {
         return GLSLstd450SMin;
       }
-    case ast::Intrinsic::kMix:
+    case semantic::Intrinsic::kMix:
       return GLSLstd450FMix;
-    case ast::Intrinsic::kModf:
+    case semantic::Intrinsic::kModf:
       return GLSLstd450Modf;
-    case ast::Intrinsic::kNormalize:
+    case semantic::Intrinsic::kNormalize:
       return GLSLstd450Normalize;
-    case ast::Intrinsic::kPow:
+    case semantic::Intrinsic::kPow:
       return GLSLstd450Pow;
-    case ast::Intrinsic::kReflect:
+    case semantic::Intrinsic::kReflect:
       return GLSLstd450Reflect;
-    case ast::Intrinsic::kRound:
+    case semantic::Intrinsic::kRound:
       return GLSLstd450Round;
-    case ast::Intrinsic::kSign:
+    case semantic::Intrinsic::kSign:
       return GLSLstd450FSign;
-    case ast::Intrinsic::kSin:
+    case semantic::Intrinsic::kSin:
       return GLSLstd450Sin;
-    case ast::Intrinsic::kSinh:
+    case semantic::Intrinsic::kSinh:
       return GLSLstd450Sinh;
-    case ast::Intrinsic::kSmoothStep:
+    case semantic::Intrinsic::kSmoothStep:
       return GLSLstd450SmoothStep;
-    case ast::Intrinsic::kSqrt:
+    case semantic::Intrinsic::kSqrt:
       return GLSLstd450Sqrt;
-    case ast::Intrinsic::kStep:
+    case semantic::Intrinsic::kStep:
       return GLSLstd450Step;
-    case ast::Intrinsic::kTan:
+    case semantic::Intrinsic::kTan:
       return GLSLstd450Tan;
-    case ast::Intrinsic::kTanh:
+    case semantic::Intrinsic::kTanh:
       return GLSLstd450Tanh;
-    case ast::Intrinsic::kTrunc:
+    case semantic::Intrinsic::kTrunc:
       return GLSLstd450Trunc;
     default:
       break;
@@ -1813,11 +1815,12 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) {
     return 0;
   }
 
-  if (ident->IsIntrinsic()) {
+  auto* sem = builder_.Sem().Get(expr);
-    return GenerateIntrinsic(ident, expr);
+  if (auto* intrinsic = sem->As<semantic::IntrinsicCall>()) {
+    return GenerateIntrinsic(ident, expr, intrinsic);
   }
 
-  auto type_id = GenerateTypeIfNeeded(TypeOf(expr->func()));
+  auto type_id = GenerateTypeIfNeeded(sem->Type());
   if (type_id == 0) {
     return 0;
   }
@@ -1852,7 +1855,8 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) {
 }
 
 uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
-                                    ast::CallExpression* call) {
+                                    ast::CallExpression* call,
+                                    const semantic::IntrinsicCall* sem) {
   auto result = result_op();
   auto result_id = result.to_i();
 
@@ -1861,20 +1865,20 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
     return 0;
   }
 
-  auto intrinsic = ident->intrinsic();
+  auto intrinsic = sem->intrinsic();
 
-  if (ast::intrinsic::IsFineDerivative(intrinsic) ||
+  if (semantic::intrinsic::IsFineDerivative(intrinsic) ||
-      ast::intrinsic::IsCoarseDerivative(intrinsic)) {
+      semantic::intrinsic::IsCoarseDerivative(intrinsic)) {
     push_capability(SpvCapabilityDerivativeControl);
   }
 
-  if (ast::intrinsic::IsImageQueryIntrinsic(intrinsic)) {
+  if (semantic::intrinsic::IsImageQueryIntrinsic(intrinsic)) {
     push_capability(SpvCapabilityImageQuery);
   }
 
-  if (ast::intrinsic::IsTextureIntrinsic(intrinsic)) {
+  if (auto* tex_sem = sem->As<semantic::TextureIntrinsicCall>()) {
-    if (!GenerateTextureIntrinsic(ident, call, Operand::Int(result_type_id),
+    if (!GenerateTextureIntrinsic(ident, call, tex_sem,
-                                  result)) {
+                                  Operand::Int(result_type_id), result)) {
       return 0;
     }
     return result_id;
@@ -1883,11 +1887,11 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
   OperandList params = {Operand::Int(result_type_id), result};
 
   spv::Op op = spv::Op::OpNop;
-  if (intrinsic == ast::Intrinsic::kAny) {
+  if (intrinsic == semantic::Intrinsic::kAny) {
     op = spv::Op::OpAny;
-  } else if (intrinsic == ast::Intrinsic::kAll) {
+  } else if (intrinsic == semantic::Intrinsic::kAll) {
     op = spv::Op::OpAll;
-  } else if (intrinsic == ast::Intrinsic::kArrayLength) {
+  } else if (intrinsic == semantic::Intrinsic::kArrayLength) {
     if (call->params().empty()) {
       error_ = "missing param for runtime array length";
       return 0;
@@ -1920,35 +1924,35 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
       return 0;
     }
     return result_id;
-  } else if (intrinsic == ast::Intrinsic::kCountOneBits) {
+  } else if (intrinsic == semantic::Intrinsic::kCountOneBits) {
     op = spv::Op::OpBitCount;
-  } else if (intrinsic == ast::Intrinsic::kDot) {
+  } else if (intrinsic == semantic::Intrinsic::kDot) {
     op = spv::Op::OpDot;
-  } else if (intrinsic == ast::Intrinsic::kDpdx) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdx) {
     op = spv::Op::OpDPdx;
-  } else if (intrinsic == ast::Intrinsic::kDpdxCoarse) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdxCoarse) {
     op = spv::Op::OpDPdxCoarse;
-  } else if (intrinsic == ast::Intrinsic::kDpdxFine) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdxFine) {
     op = spv::Op::OpDPdxFine;
-  } else if (intrinsic == ast::Intrinsic::kDpdy) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdy) {
     op = spv::Op::OpDPdy;
-  } else if (intrinsic == ast::Intrinsic::kDpdyCoarse) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdyCoarse) {
     op = spv::Op::OpDPdyCoarse;
-  } else if (intrinsic == ast::Intrinsic::kDpdyFine) {
+  } else if (intrinsic == semantic::Intrinsic::kDpdyFine) {
     op = spv::Op::OpDPdyFine;
-  } else if (intrinsic == ast::Intrinsic::kFwidth) {
+  } else if (intrinsic == semantic::Intrinsic::kFwidth) {
     op = spv::Op::OpFwidth;
-  } else if (intrinsic == ast::Intrinsic::kFwidthCoarse) {
+  } else if (intrinsic == semantic::Intrinsic::kFwidthCoarse) {
     op = spv::Op::OpFwidthCoarse;
-  } else if (intrinsic == ast::Intrinsic::kFwidthFine) {
+  } else if (intrinsic == semantic::Intrinsic::kFwidthFine) {
     op = spv::Op::OpFwidthFine;
-  } else if (intrinsic == ast::Intrinsic::kIsInf) {
+  } else if (intrinsic == semantic::Intrinsic::kIsInf) {
     op = spv::Op::OpIsInf;
-  } else if (intrinsic == ast::Intrinsic::kIsNan) {
+  } else if (intrinsic == semantic::Intrinsic::kIsNan) {
     op = spv::Op::OpIsNan;
-  } else if (intrinsic == ast::Intrinsic::kReverseBits) {
+  } else if (intrinsic == semantic::Intrinsic::kReverseBits) {
     op = spv::Op::OpBitReverse;
-  } else if (intrinsic == ast::Intrinsic::kSelect) {
+  } else if (intrinsic == semantic::Intrinsic::kSelect) {
     op = spv::Op::OpSelect;
   } else {
     GenerateGLSLstd450Import();
@@ -1959,7 +1963,7 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
       return 0;
     }
     auto set_id = set_iter->second;
-    auto inst_id = intrinsic_to_glsl_method(TypeOf(ident), ident->intrinsic());
+    auto inst_id = intrinsic_to_glsl_method(sem->Type(), sem->intrinsic());
     if (inst_id == 0) {
       error_ = "unknown method " + builder_.Symbols().NameFor(ident->symbol());
       return 0;
@@ -1994,15 +1998,14 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
   return result_id;
 }
 
-bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
+bool Builder::GenerateTextureIntrinsic(
+    ast::IdentifierExpression* ident,
     ast::CallExpression* call,
+    const semantic::TextureIntrinsicCall* sem,
     Operand result_type,
     Operand result_id) {
-  auto* sig = static_cast<const ast::intrinsic::TextureSignature*>(
+  auto& pidx = sem->Params().idx;
-      ident->intrinsic_signature());
+  auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
-  assert(sig != nullptr);
-  auto& pidx = sig->params.idx;
-  auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
 
   assert(pidx.texture != kNotUsed);
   auto* texture_type =
@@ -2149,8 +2152,8 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
     return append_coords_to_spirv_params();
   };
 
-  switch (ident->intrinsic()) {
+  switch (sem->intrinsic()) {
-    case ast::Intrinsic::kTextureDimensions: {
+    case semantic::Intrinsic::kTextureDimensions: {
       // Number of returned elements from OpImageQuerySize[Lod] may not match
       // those of textureDimensions().
       // This might be due to an extra vector scalar describing the number of
@@ -2206,7 +2209,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
       }
       break;
     }
-    case ast::Intrinsic::kTextureNumLayers: {
+    case semantic::Intrinsic::kTextureNumLayers: {
       uint32_t spirv_dims = 0;
       switch (texture_type->dim()) {
         default:
@@ -2241,19 +2244,19 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
       }
       break;
     }
-    case ast::Intrinsic::kTextureNumLevels: {
+    case semantic::Intrinsic::kTextureNumLevels: {
       op = spv::Op::OpImageQueryLevels;
       append_result_type_and_id_to_spirv_params();
       spirv_params.emplace_back(gen_param(pidx.texture));
       break;
     }
-    case ast::Intrinsic::kTextureNumSamples: {
+    case semantic::Intrinsic::kTextureNumSamples: {
       op = spv::Op::OpImageQuerySamples;
       append_result_type_and_id_to_spirv_params();
       spirv_params.emplace_back(gen_param(pidx.texture));
       break;
     }
-    case ast::Intrinsic::kTextureLoad: {
+    case semantic::Intrinsic::kTextureLoad: {
       op = texture_type->Is<type::StorageTexture>() ? spv::Op::OpImageRead
                                                     : spv::Op::OpImageFetch;
       append_result_type_and_id_to_spirv_params_for_read();
@@ -2274,7 +2277,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
 
       break;
     }
-    case ast::Intrinsic::kTextureStore: {
+    case semantic::Intrinsic::kTextureStore: {
       op = spv::Op::OpImageWrite;
       spirv_params.emplace_back(gen_param(pidx.texture));
       if (!append_coords_to_spirv_params()) {
@@ -2283,7 +2286,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
       spirv_params.emplace_back(gen_param(pidx.value));
       break;
     }
-    case ast::Intrinsic::kTextureSample: {
+    case semantic::Intrinsic::kTextureSample: {
       op = spv::Op::OpImageSampleImplicitLod;
       append_result_type_and_id_to_spirv_params_for_read();
       if (!append_image_and_coords_to_spirv_params()) {
@@ -2291,7 +2294,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
       }
       break;
     }
-    case ast::Intrinsic::kTextureSampleBias: {
+    case semantic::Intrinsic::kTextureSampleBias: {
       op = spv::Op::OpImageSampleImplicitLod;
       append_result_type_and_id_to_spirv_params_for_read();
       if (!append_image_and_coords_to_spirv_params()) {
@@ -2302,7 +2305,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
           ImageOperand{SpvImageOperandsBiasMask, gen_param(pidx.bias)});
       break;
     }
-    case ast::Intrinsic::kTextureSampleLevel: {
+    case semantic::Intrinsic::kTextureSampleLevel: {
       op = spv::Op::OpImageSampleExplicitLod;
       append_result_type_and_id_to_spirv_params_for_read();
       if (!append_image_and_coords_to_spirv_params()) {
@@ -2326,7 +2329,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
       image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level});
       break;
     }
-    case ast::Intrinsic::kTextureSampleGrad: {
+    case semantic::Intrinsic::kTextureSampleGrad: {
       op = spv::Op::OpImageSampleExplicitLod;
       append_result_type_and_id_to_spirv_params_for_read();
       if (!append_image_and_coords_to_spirv_params()) {
@@ -2340,7 +2343,7 @@ bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
           ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddy)});
       break;
     }
-    case ast::Intrinsic::kTextureSampleCompare: {
+    case semantic::Intrinsic::kTextureSampleCompare: {
       op = spv::Op::OpImageSampleDrefExplicitLod;
       append_result_type_and_id_to_spirv_params();
       if (!append_image_and_coords_to_spirv_params()) {

src/writer/spirv/builder.h

@@ -57,6 +57,13 @@
 #include "src/writer/spirv/instruction.h"
 
 namespace tint {
+
+// Forward declarations
+namespace semantic {
+class TextureIntrinsicCall;
+class IntrinsicCall;
+}  // namespace semantic
+
 namespace writer {
 namespace spirv {
 
@@ -354,19 +361,23 @@ class Builder {
   /// Generates an intrinsic call
   /// @param ident the intrinsic expression
   /// @param call the call expression
+  /// @param sem the semantic information for the intrinsic call
   /// @returns the expression ID on success or 0 otherwise
   uint32_t GenerateIntrinsic(ast::IdentifierExpression* ident,
-                             ast::CallExpression* call);
+                             ast::CallExpression* call,
+                             const semantic::IntrinsicCall* sem);
   /// Generates a texture intrinsic call. Emits an error and returns false if
   /// we're currently outside a function.
   /// @param ident the texture intrinsic
   /// @param call the call expression
+  /// @param sem the semantic information for the texture intrinsic call
   /// @param result_type result type operand of the texture instruction
   /// @param result_id result identifier operand of the texture instruction
   /// parameters
   /// @returns true on success
   bool GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
                                 ast::CallExpression* call,
+                                const semantic::TextureIntrinsicCall* sem,
                                 spirv::Operand result_type,
                                 spirv::Operand result_id);
   /// Generates a sampled image